CUDA atomicAdd for doubles definition error
<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>