2 Replies Latest reply on Aug 4, 2016 1:23 PM by shailu1995

    atomicadd opencl



      The CUDA implementation of overloaded atomicAdd() operation for double data type is:

      __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;


      { assumed = old;

      old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));

      // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN) }

      while (assumed != old);

      return __longlong_as_double(old);



      Is there a similar OpenCL implementation of atomic_add() for double data type. I tried using this:

      #pragma OPENCL EXTENSION cl_khr_fp64: enable

      #pragma OPENCL EXTENSION cl_khr_int64_base_atomics: enable

      double atomic_add(__global double *valq,double delta) {

         union {

           double f;

           unsigned long  i;

         } old;

         union {

           double f;

           unsigned long  i;

         } new1;

        do {

           old.f = *valq;

           new1.f = old.f + delta;

         } while (atom_cmpxchg((volatile __global unsigned long *)valq, old.i, new1.i) != old.i);

         return old.f;


      This works fine when I change the name to something else but with this function name it gives the error:

      <kernel>:11:9: error: overloaded function 'atom_add' must have the 'overloadable' attribute

      double atomic_add(__global double *valq,double delta) {


      cl_kernel.h:4191:20: note: expanded from macro 'atomic_add'

      #define atomic_add atom_add


      cl_kernel.h:4094:24: note: previous overload of function is here

      ulong __OVERLOADABLE__ atom_add(__local volatile ulong *ptr, ulong val);



      Is there something I need to add to overload the function?


      Thanks in advance

      Shailesh Tripathi