programing - introduction to cuda

atomicAdd() para el doble en GPU

Question

Estoy haciendo un proyecto en GPU, y tengo que usar atomicAdd () para el doble, porque el cuda no es compatible con el doble, así que uso el siguiente código, que es NVIDIA.

__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);
}

Ahora quiero saber por qué el implemento requiere un bucle, mientras (¡se supone! = Viejo)

Answer 1

Básicamente porque la implementación requiere una carga, que no se puede realizar atómicamente. La operación de comparación e intercambio es una versión atómica de

(*address == assumed) ? (assumed + val) : *address

No hay garantía de que el valor de la address no cambie entre el ciclo en que se carga el valor de la address y el ciclo en que se atomicCAS llamada atomicCAS para almacenar el valor actualizado. Si eso sucede, el valor en la address no se actualizará. Por lo tanto, el ciclo garantiza que las dos operaciones se repitan hasta que no haya ningún cambio en el valor en la address entre la operación de lectura y la operación de comparación e intercambio, lo que implica que la actualización se llevó a cabo.