CUDA atomicAdd für doppelte Definition Fehler

Question

In früheren Versionen von CUDA wurde atomicAdd nicht für doubles implementiert, so ist es üblich, dies wie here zu implementieren. Bei der neuen CUDA 8 RC stoße ich auf Probleme, wenn ich versuche, meinen Code zu kompilieren, der eine solche Funktion enthält. Ich denke, das liegt daran, dass mit Pascal und Compute Capability 6.0 eine native Doppelversion von atomicAdd hinzugefügt wurde, die aber für frühere Compute-Funktionen nicht korrekt ignoriert wird.

Der folgende Code verwendet zu kompilieren und mit früheren CUDA Versionen gut laufen, aber jetzt bekomme ich diese Kompilierungsfehler:

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

Aber wenn ich meine Implementierung entfernen, ich diesen Fehler stattdessen erhalten:

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

Ich sollte hinzufügen, dass ich das nur sehe, wenn ich mit -arch=sm_35 oder ähnlichem kompiliere. Wenn ich mit -arch=sm_60 kompiliere, erhalte ich das erwartete Verhalten, d. H. Nur den ersten Fehler, und eine erfolgreiche Kompilierung im zweiten Fall.

Edit: Auch ist es spezifisch für atomicAdd - wenn ich den Namen ändere, funktioniert es gut.

Es sieht wirklich wie ein Compiler Bug aus. Kann jemand anderes bestätigen, dass dies der Fall ist?

Beispielcode:

__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: Ich bekam eine Antwort von Nvidia, die dieses Problem erkennen, und hier ist das, was die Entwickler sagen über sie:

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

Das Aroma atomicAdd ist eine neue Methode, die für die Rechenleistung 6.0 eingeführt wurde. Sie können Ihre vorherige Durchführung anderer Rechenfähigkeiten halten es mit Makrodefinition Bewachung

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

Diese Makro mit dem Namen Architektur Identifizierung Makro here dokumentiert:

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.

Ich gehe davon aus NVIDIA nicht für vorherigen CC platziert haben zu Vermeiden Sie einen Konflikt für Benutzer, die ihn definieren und nicht zu Compute Capability> = 6.x wechseln. Ich würde es nicht als BUG, ​​sondern eher als Release-Delivery-Praxis betrachten.

EDIT: Makro Guard war unvollständig (behoben) - hier ein komplettes Beispiel.

#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 mit:

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

Befehlszeilen (beide erfolgreich):

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

Sie können feststellen, warum es mit der Datei enthalten funktioniert: sm_60_atomic_functions.h, wobei das Verfahren nicht deklariert wenn __CUDA_ARCH__ ist niedriger als 600.