AnsweredAssumed Answered

atomicadd opencl

Question asked by shailu1995 on Jul 31, 2016
Latest reply on Aug 4, 2016 by shailu1995


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