CUDA atomicAdd per errore di definizione doppi

Question

Nelle versioni precedenti di CUDA, atomicAdd non è stato implementato per i doppi, quindi è normale implementarlo come here. Con il nuovo CUDA 8 RC, mi imbatto in problemi quando provo a compilare il mio codice che include tale funzione. Immagino che ciò sia dovuto al fatto che con Pascal e Compute Capability 6.0 è stata aggiunta una versione nativa doppia di atomicAdd, ma in qualche modo non è correttamente ignorata per le precedenti capacità di calcolo.

Il codice di seguito utilizzato per compilare ed eseguire bene con le versioni precedenti CUDA, ma ora ottenere questo errore di compilazione:

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

Ma se rimuovo la mia applicazione, ho invece ottenere questo errore:

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

Devo aggiungere che vedo solo questo se compilo con -arch=sm_35 o simile. Se compilo con -arch=sm_60 ottengo il comportamento previsto, cioè solo il primo errore e la compilazione riuscita nel secondo caso.

Modifica: Inoltre, è specifico per atomicAdd - se cambio il nome, funziona bene.

Sembra davvero un bug del compilatore. Qualcun altro può confermare che questo è il caso?

codice Esempio:

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

---

Edit: ho ricevuto una risposta da Nvidia che riconoscono questo problema, e qui è ciò che gli sviluppatori dicono su di esso:

The sm_60 architecture, that is newly supported in CUDA 8.0, has native fp64 atomicAdd function. Because of the limitations of our toolchain and CUDA language, the declaration of this function needs to be present even when the code is not being specifically compiled for sm_60. This causes a problem in your code because you also define a fp64 atomicAdd function.

CUDA builtin functions such as atomicAdd are implementation-defined and can be changed between CUDA releases. Users should not define functions with the same names as any CUDA builtin functions. We would suggest you to rename your atomicAdd function to one that is not the same as any CUDA builtin functions.

Answer 1

quel sapore di atomicAdd è un nuovo metodo introdotto per la capacità di calcolo 6.0. Si può mantenere la precedente implementazione di altre funzionalità di calcolo guardia utilizzando macro definizione

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

Questa macro denominata architettura macro di identificazione è documentata here:

5.7.4. Virtual Architecture Identification Macro

The architecture identification macro __CUDA_ARCH__ is assigned a three-digit value string xy0 (ending in a literal 0) during each nvcc compilation stage 1 that compiles for compute_xy.

This macro can be used in the implementation of GPU functions for determining the virtual architecture for which it is currently being compiled. The host code (the non-GPU code) must not depend on it.

Presumo NVIDIA non ha posto per precedente CC evitare conflitti per gli utenti che lo definiscono e non si spostano in Capacità di calcolo> = 6.x. Non lo considererei un BUG, ​​piuttosto una pratica di rilascio del rilascio.

EDIT: la protezione macro era incompleta (corretta) - qui un esempio 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() ; 
} 

Compilation 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 

righe di comando (sia successo):

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

Si possono trovare perché funziona con il file include: sm_60_atomic_functions.h, in cui il metodo non è dichiarato se __CUDA_ARCH__ è inferiore a 600.