¿Por qué atomicAdd no se ha implementado para dobles?

Question

¿Por qué atomicAdd no se ha implementado para dobles?

¿Por qué no se ha implementado atomicAdd() para los dobles explícitamente como parte de CUDA 4.0 o superior?

Del apéndice F Página 97 de la guía de programación CUDA 4.1 las siguientes versiones de atomicAdd se han implementado.

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 empezar 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 ?

31

cuda

Author: harrism, 2012-09-27

Source

1 answers

score 34 · Accepted Answer

Editar: A partir de CUDA 8, la doble precisión atomicAdd() se implementa en CUDA con soporte de hardware en GPU SM_6X (Pascal).

~~Actualmente, ningún dispositivo CUDA soporta atomicAdd para double en hardware.~~ Como ha señalado, se puede implementar en términos de atomicCAS en enteros de 64 bits, pero hay un costo de rendimiento no trivial para eso.

Por lo tanto, el equipo de software de CUDA eligió documentar una implementación correcta como una opción para los desarrolladores, en lugar de hacerlo parte del estándar CUDA biblioteca. 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.