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

Question

¿Por qué ya no haya atomicAdd() en dobles sido implementados de manera explícita como parte de CUDA 4.0 o superior?

A partir del apéndice F Página 97 del CUDA programming guide 4.1 se han implementado 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 va a dar una pequeña aplicación de atomicAdd en dobles de la siguiente manera que acabo de comenzar a utilizar 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?

Answer 1

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

Actualmente, no hay dispositivos compatibles con CUDA atomicAdd para double en el hardware. Como se anotó, 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 CUDA eligió para documentar una aplicación correcta como una opción para los desarrolladores, en lugar de hacerlo parte de la biblioteca estándar de CUDA. De esta forma, los desarrolladores no están, sin saberlo, optando por un costo de rendimiento que no entienden.

Aparte: No creo que esta pregunta debe ser cerrado como "no constructiva". Creo que es una pregunta perfectamente válida, +1.