Archives Discussions

[shailu1995]

Hi,

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;

do

{ 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

[dipak · Accepted solution]

Please add "__attribute__((overloadable))".

E.g.

double __attribute__((overloadable)) atomic_add(__global double *valq,double delta)

{ ... }

[shailu1995]

Hi dipak,

It worked.

Thanks

Shailesh