atomicAdd with signed long long not working


I am trying to add an integer to a (signed) long long global variable using the atomicAdd function using Cuda 3.2. Unfortunately atomicAdd takes unsigned long long types and no signed long long types. The Cuda C Programming Guide says that any atomic operation can be implemented based on atomicCAS(). Which is what I am trying, but I am not succesfull. Below is my code:

__device__ long long atomicAdd(long long* address, int 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, val + assumed);

	} while (assumed != old); 


	return old;


I am also unable to print long long variables to my console. I have tried cuPrintf("%lld", long_value), cuPrintf("%d", long_value) and cuPrintf("%1f", long_value) but none of them are working.

Can anyone suggest a function that does work or another way to add an integer to a signed long long variable in global memory in an atomic fashion. Thank you very much in advance.

Kind regards,


There is no need to use atomicCAS(). Just use atomicAdd(), casting its arguments to (unsigned long long*), (unsigned long long).

Signed and unsigned add give bit-identical results. The difference is only how the result is interpreted.

Thank you very much for the help. The casting worked perfectly. I’m glad it was this easy. Thanks!