shingoxlf

can I do atomic add on float using compare and swap

Discussion created by shingoxlf on Oct 11, 2011
Latest reply on Oct 12, 2011 by MicahVillmow

In the cuda manual, attomic add can be implmented using compare and swap like this:

 

 

__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);
}
So I tried to implement atomic add in Opencl like this:
double atomicDoubleAdd(__global double* address, double val)
{
ulong * address_as_ull =(ulong *)address;
ulong old = *address_as_ull, assumed;
do {
assumed = old;
old = atom_cmpxchg(address_as_ull, assumed,as_ulong(val +as_double(assumed)));
} while (assumed != old);
return as_double(old);
}
however, it fails when I tried to build the program, so if anybody knows how to do this?


Outcomes