it-swarm-es.tech

CUDA atomicAdd para error de definición de dobles

En versiones anteriores de CUDA, atomicAdd no se implementó para dobles, por lo que es común implementar esto como aquí . Con el nuevo CUDA 8 RC, me encuentro con problemas cuando intento 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 de cómputo anteriores.

El siguiente código solía compilar y funcionar bien con versiones anteriores de CUDA, pero ahora recibo este error de compilación:

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

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

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

Debo agregar 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 una 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 dicen los desarrolladores al respecto:

La arquitectura sm_60, que se admite recientemente en CUDA 8.0, tiene la función nativa fp64 atomicAdd. Debido a las limitaciones de nuestra cadena de herramientas y el lenguaje CUDA, la declaración de esta función debe estar presente incluso cuando el código no se está compilando 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 versiones de CUDA. Los usuarios no deben definir funciones con los mismos nombres que las funciones integradas 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.

13
kalj

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 la 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

La macro de identificación de arquitectura __CUDA_Arch__ se le asigna una cadena de valor de tres dígitos xy0 (que termina en un literal 0) durante cada etapa de compilación nvcc 1 que compila para compute_xy.

Esta macro se puede usar en la implementación de funciones de GPU para determinar la arquitectura virtual para la que 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 CC anterior para evitar conflictos para los usuarios que lo definieron y no se movieron a Capacidad de cómputo> = 6.x. Sin embargo, no lo consideraría un ERROR, sino una práctica de entrega de lanzamientos.

[~ # ~] editar [~ # ~] : el protector de macro 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:

$> nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Wed_May__4_21:01:56_CDT_2016
Cuda compilation tools, release 8.0, V8.0.26

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 inferior a 600.

12
Florent DUGUET