title: Why has atomicAdd not been implemented for doubles?
body: <p>Why hasnt <code>atomicAdd()</code> for doubles been implemented explicitly as a part of CUDA 4.0 or higher?</p>
  
  <p>From the appendix F Page 97 of the <a href="http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/CUDA_C_Programming_Guide.pdf">CUDA programming guide 4.1</a> the following versions of
  atomicAdd have been implemented.</p>
  
  <pre><code>int atomicAdd(int* address, int val);
  unsigned int atomicAdd(unsigned int* address,
                         unsigned int val);
  unsigned long long int atomicAdd(unsigned long long int* address,
                                   unsigned long long int val);
  float atomicAdd(float* address, float val)
  </code></pre>
  
  <p>The same page goes on to give a small implementation of atomicAdd for doubles as follows
  which I have just started using in my project. </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);
  }
  </code></pre>
  
  <p>Why not define the above code as a part of CUDA ? </p>