cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

shingoxlf
Journeyman III

can I do atomic add on float using compare and swap

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?


0 Likes
7 Replies

This is not OpenCL code, please use OpenCL and not cuda.
0 Likes

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);
}
this is opencl code, right?


0 Likes

Ahh sorry about that, the text and the code was hard to distinguish, I thought it was two functions. It would be clearer if you use the 'attach code' button for code.

Anyways, does the device you target support 64bit atomics? If not, this won't work.
0 Likes

Hi Micah, thank you for your reply.

 

I was trying this code on a Nvidia Fermi card C2050.

 

I can use atomicdoubleadd from cuda on this card, so I guess it supports 64-bit atomic.

However, I cannot build the program successfully due to this function, so could you please me to see what's wrong with my code?

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); }

0 Likes

shingoxlf,
There are three issues here:
1) We don't support NVidia products.
2) None of our GPU's support 64bit atomics.
3) You are not correctly enabling an extension before using it(double's).
0 Likes

Anyway, thanks for your reply.

So will there be support for double precision atomic operation in the near furture?

0 Likes

No, there will not.
0 Likes