mac - cuda tutorial
¿Por qué no se ha implementado atomicAdd para dobles? (1)
¿Por qué no se ha implementado explícitamente atomicAdd() para dobles como parte de CUDA 4.0 o superior?
Desde el apéndice F, página 97 de la guía de programación 4.1 de CUDA, se implementaron las siguientes versiones de atomicAdd.
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)
La misma página continúa para dar una pequeña implementación de atomicAdd para dobles de la siguiente manera que acabo de comenzar a usar en mi proyecto.
__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 qué no definir el código anterior como parte de CUDA?
Edición: a partir de CUDA 8, la atomicAdd() doble precisión se implementa en CUDA con soporte de hardware en las GPU SM_6X (Pascal).
Actualmente, ningún dispositivo CUDA admite Como se señaló, puede implementarse en términos de atomicAdd para double en hardware.atomicCAS en atomicCAS de 64 bits, pero existe un costo de rendimiento no trivial para eso.
Por lo tanto, el equipo de software de CUDA decidió documentar una implementación correcta como una opción para los desarrolladores, en lugar de hacerlo parte de la biblioteca estándar de CUDA. De esta manera, los desarrolladores no están optando, sin saberlo, por un costo de rendimiento que no entienden.
Aparte: no creo que esta pregunta deba cerrarse como "no constructiva". Creo que es una pregunta perfectamente válida, +1.