In früheren CUDA-Versionen wurde atomicAdd nicht für Doubles implementiert, daher ist es üblich, dies wie @ zu implementiereHie. Mit der neuen CUDA 8 RC habe ich Probleme, wenn ich versuche, meinen Code zu kompilieren, der eine solche Funktion enthält. Ich vermute, das liegt an der Tatsache, dass mit Pascal und Compute Capability 6.0 eine native Doppelversion von atomicAdd hinzugefügt wurde, die jedoch für frühere Compute Capabilities nicht richtig ignoriert wird.

Der folgende Code wurde zum Kompilieren und Ausführen mit früheren CUDA-Versionen verwendet, aber jetzt erhalte ich den folgenden Kompilierungsfehler:

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

Aber wenn ich meine Implementierung entferne, erhalte ich stattdessen den folgenden Fehler:

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 @ kompilie-arch=sm_35 o.ä. Wenn ich mit @ kompilie-arch=sm_60 Ich erhalte das erwartete Verhalten, d. H. Nur den ersten Fehler, und im zweiten Fall die erfolgreiche Kompilierung.

Bearbeiten: Auch ist es spezifisch füratomicAdd - wenn ich den Namen ändere, funktioniert es gut.

Es sieht wirklich aus wie ein Compiler-Fehler. 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 habe eine Antwort von Nvidia erhalten, die dieses Problem erkannt hat, und hier ist, was die Entwickler dazu sagen:

Die sm_60-Architektur, die in CUDA 8.0 neu unterstützt wird, verfügt über die native Funktion fp64 atomicAdd. Aufgrund der Einschränkungen unserer Toolchain und der CUDA-Sprache muss die Deklaration dieser Funktion auch dann vorhanden sein, wenn der Code nicht speziell für sm_60 kompiliert wird. Dies verursacht ein Problem in Ihrem Code, da Sie auch eine atomicAdd-Funktion für fp64 definieren.

CUDA-integrierte Funktionen wie atomicAdd sind implementierungsdefiniert und können zwischen CUDA-Releases geändert werden. Benutzer sollten keine Funktionen mit denselben Namen wie in CUDA integrierte Funktionen definieren. Wir empfehlen Ihnen, Ihre atomicAdd-Funktion in eine Funktion umzubenennen, die nicht mit den in CUDA integrierten Funktionen identisch ist.