programming - CUDA atomicAdd para el error de definición de dobles

cuda tutorial (1)

En versiones anteriores de CUDA, atomicAdd no se implementó para los dobles, por lo que es común implementar esto como aquí . Con el nuevo CUDA 8 RC, me encuentro con problemas cuando trato de compilar mi código que incluye dicha función. Supongo que esto se debe al hecho de que con Pascal y Compute Capability 6.0, se ha agregado una versión doble nativa de atomicAdd, pero de alguna manera eso no se ignora adecuadamente para las Capacidades Compute previas.

El siguiente código se usa para compilar y ejecutar bien con versiones anteriores de CUDA, pero ahora obtengo este error de compilación:

test.cu(3): error: function "atomicAdd(double *, double)" has already been defined

Pero si elimino mi implementación, en cambio, obtengo este error:

test.cu(33): error: no instance of overloaded function "atomicAdd" matches the argument list argument types are: (double *, double)

Debo añadir que solo veo esto si compilo con -arch=sm_35 o similar. Si compilo con -arch=sm_60 obtengo el comportamiento esperado, es decir, solo el primer error y la compilación exitosa en el segundo caso.

Editar: Además, es específico para atomicAdd : si cambio el nombre, funciona bien.

Realmente parece un error de compilación. ¿Alguien más puede confirmar que este es el caso?

Código de ejemplo:

__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); } __global__ void kernel(double *a) { double b=1.3; atomicAdd(a,b); } int main(int argc, char **argv) { double *a; cudaMalloc(&a,sizeof(double)); kernel<<<1,1>>>(a); cudaFree(a); return 0; }

Editar: recibí una respuesta de Nvidia que reconoce este problema, y esto es lo que los desarrolladores dicen al respecto:

La arquitectura sm_60, que es recientemente compatible con CUDA 8.0, tiene la función nativa fp64 atomicAdd. Debido a las limitaciones de nuestra cadena de herramientas y del lenguaje CUDA, la declaración de esta función debe estar presente incluso cuando el código no se compile específicamente para sm_60. Esto causa un problema en su código porque también define una función fp64 atomicAdd.
Las funciones integradas de CUDA, como atomicAdd, están definidas por la implementación y se pueden cambiar entre las versiones de CUDA. Los usuarios no deben definir funciones con los mismos nombres que cualquier función incorporada de CUDA. Le sugerimos que cambie el nombre de su función atomicAdd a una que no sea la misma que cualquier función incorporada de CUDA.

Ese sabor de atomicAdd es un nuevo método introducido para la capacidad de cálculo 6.0. Puede mantener su implementación anterior de otras capacidades de cómputo protegiéndola usando definición de macro

#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600 #else <... place here your own pre-pascal atomicAdd definition ...> #endif

Esta macro llamada macro de identificación de arquitectura está documentada aquí :

5.7.4. Macro de identificación de arquitectura virtual
A la macro de identificación de arquitectura __CUDA_ARCH__ se le asigna una cadena de valores de tres dígitos xy0 (que termina en un literal 0) durante cada etapa de compilación nvcc 1 que se compila para compute_xy.
Esta macro se puede usar en la implementación de funciones de GPU para determinar la arquitectura virtual para la cual se está compilando actualmente. El código de host (el código que no es GPU) no debe depender de él.

Supongo que NVIDIA no lo colocó para el CC anterior para evitar conflictos para los usuarios que lo definen y no se mueve a Capacidad de cálculo> = 6.x. Sin embargo, no lo consideraría un ERROR, sino una práctica de liberación de entregas.

EDIT : macro guard estaba incompleto (fijo) - aquí un ejemplo completo.

#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600 #else __device__ double atomicAdd(double* a, double b) { return b; } #endif __device__ double s_global ; __global__ void kernel () { atomicAdd (&s_global, 1.0) ; } int main (int argc, char* argv[]) { kernel<<<1,1>>> () ; return ::cudaDeviceSynchronize () ; }

Compilación con:

Líneas de comando (ambas exitosas):

$> nvcc main.cu -arch=sm_60 $> nvcc main.cu -arch=sm_35

Puede encontrar por qué funciona con el archivo de inclusión: sm_60_atomic_functions.h , donde el método no se declara si __CUDA_ARCH__ es menor que 600.