Por que o atomicAdd não foi implementado para duplas?
Por que nãoatomicAdd()
para duplas foram implementadas explicitamente como parte do CUDA 4.0 ou superior?
Do apêndice F página 97 daGuia de programação CUDA 4.1 as seguintes versões do atomicAdd foram implementadas.
int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,
unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,
unsigned long long int val);
float atomicAdd(float* address, float val)
A mesma página continua dando uma pequena implementação de atomicAdd para duplas da seguinte forma que eu comecei a usar no meu projeto.
__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);
}
Por que não definir o código acima como parte do CUDA?