Porque noatomicAdd() ¿Se ha implementado explícitamente el doble como parte de CUDA 4.0 o superior?

Del apéndice F Página 97 de laGuía de programación CUDA 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 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?