title: CUDA atomicAdd for doubles definition error
body: <p>In previous versions of CUDA, atomicAdd was not implemented for doubles, so it is common to implement this like <a href="https://stackoverflow.com/questions/12626096/why-has-atomicadd-not-been-implemented-for-doubles">here</a>. With the new CUDA 8 RC, I run into troubles when I try to compile my code which includes such a function. I guess this is due to the fact that with Pascal and Compute Capability 6.0, a native double version of atomicAdd has been added, but somehow that is not properly ignored for previous Compute Capabilities.</p>
  
  <p>The code below used to compile and run fine with previous CUDA versions, but now I get this compilation error:</p>
  
  <pre><code>test.cu(3): error: function "atomicAdd(double *, double)" has already been defined
  </code></pre>
  
  <p>But if I remove my implementation, I instead get this error:</p>
  
  <pre><code>test.cu(33): error: no instance of overloaded function "atomicAdd" matches the argument list
              argument types are: (double *, double)
  </code></pre>
  
  <p>I should add that I only see this if I compile with <code>-arch=sm_35</code> or similar. If I compile with <code>-arch=sm_60</code> I get the expected behavior, i.e. only the first error, and successful compilation in the second case.</p>
  
  <p>Edit: Also, it is specific for <code>atomicAdd</code> -- if I change the name, it works well.</p>
  
  <p>It really looks like a compiler bug. Can someone else confirm that this is the case?</p>
  
  <p>Example code:</p>
  
  <pre><code>__device__ double atomicAdd(double* address, double val)
  {
      unsigned long long int* address_as_ull = (unsigned long long int*)address;
      unsigned long long int old = *address_as_ull, assumed;
      do {
          assumed = old;
          old = atomicCAS(address_as_ull, assumed,
                  __double_as_longlong(val + __longlong_as_double(assumed)));
      } while (assumed != old);
      return __longlong_as_double(old);
  }
  
  __global__ void kernel(double *a)
  {
      double b=1.3;
      atomicAdd(a,b);
  }
  
  int main(int argc, char **argv)
  {
      double *a;
      cudaMalloc(&amp;a,sizeof(double));
  
      kernel&lt;&lt;&lt;1,1&gt;&gt;&gt;(a);
  
      cudaFree(a);
      return 0;
  }
  </code></pre>
  
  <hr>
  
  <p>Edit: I got an answer from Nvidia who recognize this problem, and here is what the developers say about it:</p>
  
  <blockquote>
    <p>The sm_60 architecture, that is newly supported in CUDA 8.0, has
    native fp64 atomicAdd function. Because of the limitations of our
    toolchain and CUDA language, the declaration of this function needs to
    be present even when the code is not being specifically compiled for
    sm_60. This causes a problem in your code because you also define a
    fp64 atomicAdd function.</p>
    
    <p>CUDA builtin functions such as atomicAdd are implementation-defined
    and can be changed between CUDA releases. Users should not define
    functions with the same names as any CUDA builtin functions. We would
    suggest you to rename your atomicAdd function to one that is not the
    same as any CUDA builtin functions.</p>
  </blockquote>