Try to implement a AtomicAddDouble function

Accelerated Computing CUDA CUDA Programming and Performance
1

I try to implement a Cuda atomicAddDouble function on a Tesla C2050. But, I get that problem:

Description	Resource	Path	Location	Type

identifier "__double_as_longlong" is undefined	MatchedFilter.cu	/MatchedFilter/Source	line 254	C/C++ Problem

identifier "__longlong_as_double" is undefined	MatchedFilter.cu	/MatchedFilter/Source	line 252	C/C++ Problem

identifier "atomicCAS" is undefined	MatchedFilter.cu	/MatchedFilter/Source	line 253	C/C++ Problem

make: *** [Source/cu_MatchedFilter.o] Error 2	MatchedFilter		line 0	C/C++ Problem

My flag setting are

-g -G -gencode arch=compute_20,code=sm_20 -gencode arch=compute_10,code=compute_10

It’s the function.

__device__ double atomicAddDouble(double* address, double val)

{

double old = *address, assumed;

do {

assumed = old;

old =

__longlong_as_double(

atomicCAS((unsigned long long int*)address,

__double_as_longlong(assumed),

__double_as_longlong(val + assumed)));

} while (assumed != old);

return old;

}
2

These function flagged are not supported for sm_10. If you leave off “-gencode arch=compute_10,code=compute_10” these errors should go away. Did you mean sm_13 by any chance (that’s the minimum compute capability that includes double-precision support) ?

3

Your function won’t work. You have a race in the initialization of old. You first evaluate old=*address, but the value may have changed before you get into your loop.

4

Good spotting. It’s not a race condition though, the problem rather is that

double old = *address, assumed;

while looking like self-explanatory code will actually set [font=“Courier New”]old[/font] to the uninitialized value of the variable [font=“Courier New”]assumed[/font].

5

Tera, good double-spotting!

6

Triple-spotting: Forget about my double-spotting, the code is actually going to work. Somehow for a moment I parsed the expression with the relative precedence of “=” and “,” inverted.

7

You guys really had me going. So much, that thought for a moment that:

double old = *address, assumed;

was a new kind of atomic syntax instead of a declaration of an iniatialized and an uninitialized double variable.

The code is straight from the programming guide, section B11.

Thanks for this forum, it forces me to read stuff in the manuals!

8

thanks