Ärger starten CUDA Kerne aus statischen Initialisierungscode

Question

Ich habe eine Klasse, die einen Kern in seinem Konstruktor aufruft, wie folgt:

"ScalarField.h"

#include <iostream> 

    void ERROR_CHECK(cudaError_t err,const char * msg) { 
     if(err!=cudaSuccess) { 
      std::cout << msg << " : " << cudaGetErrorString(err) << std::endl; 
      std::exit(-1); 
     } 
    } 

    class ScalarField { 
    public: 
     float* array; 
     int dimension; 

     ScalarField(int dim): dimension(dim) { 
      std::cout << "Scalar Field" << std::endl; 
      ERROR_CHECK(cudaMalloc(&array, dim*sizeof(float)),"cudaMalloc"); 
     } 
    }; 

"classA.h"

#include "ScalarField.h" 


static __global__ void KernelSetScalarField(ScalarField v) { 
    int index = threadIdx.x + blockIdx.x * blockDim.x; 
    if (index < v.dimension) v.array[index] = 0.0f; 
} 

class A { 
public: 
    ScalarField v; 

    A(): v(ScalarField(3)) { 
     std::cout << "Class A" << std::endl; 
     KernelSetScalarField<<<1, 32>>>(v); 
     ERROR_CHECK(cudaGetLastError(),"Kernel"); 
    } 
}; 

"main.cu"

#include "classA.h" 

A a_object; 

int main() { 
    std::cout << "Main" << std::endl; 
    return 0; 
} 

Wenn ich diese Klasse auf Haupt (A a_object;) instanziiere, erhalte ich keine Fehler. Wenn ich es jedoch außerhalb von main initiiere, gleich nachdem ich es definiert habe (class A {...} a_object;), erhalte ich beim Start des Kernels den Fehler "ungültige Gerätefunktion". Warum passiert das?

EDIT

Aktualisiert Code ein vollständigeres Beispiel zu liefern.

EDIT 2

Im Anschluss an die Beratung im Kommentar von Raxvan wollte ich ich die auch in ScalarField Konstruktor verwendet dimensions Variable muss sagen, definiert (in einer anderen Klasse), die außerhalb Haupt, aber vor allem anderen. Könnte das die Erklärung sein? Der Debugger zeigte jedoch den richtigen Wert für dimensions.

Answer 1

Die kurze Version:

Der eigentliche Grund für das Problem, wenn class A außerhalb der Haupt-instanziiert ist, dass eine bestimmte Hook-Routine, die die CUDA-Laufzeitbibliothek mit Ihrem Kernel initialisieren erforderlich ist, wird nicht vor laufen wird Der Konstruktor von class A wird aufgerufen. Dies geschieht, weil es keine Garantien für die Reihenfolge gibt, in der statische Objekte im C++ - Ausführungsmodell instanziiert und initialisiert werden. Ihre globale Bereichsklasse wird instanziiert, bevor die globalen Bereichsobjekte initialisiert werden, die das CUDA-Setup initialisieren. Ihr Kernel-Code wird niemals vor dem Aufruf in den Kontext geladen, und es tritt ein Laufzeitfehler auf.

So gut ich kann sagen, dies ist eine echte Einschränkung der CUDA-Laufzeit-API und nicht etwas leicht im Benutzercode behoben. In Ihrem trivialen Beispiel könnten Sie den Kernel-Aufruf durch einen Aufruf an cudaMemset oder eine der nicht symbolbasierten Runtime-API-memset-Funktionen ersetzen und es wird funktionieren. Dieses Problem ist vollständig auf Benutzerkerne oder Gerätesymbole beschränkt, die zur Laufzeit über die Laufzeit-API geladen werden. Aus diesem Grund würde ein leerer Standardkonstruktor auch Ihr Problem lösen. Aus der Sicht des Designs wäre ich sehr fragwürdig, ob ein Muster Kernel im Konstruktor aufruft. Das Hinzufügen einer spezifischen Methode für das GPU-Setup/-Teardown der Klasse, die nicht auf dem Standardkonstruktor oder Destruktor beruht, wäre ein viel saubereres und weniger fehleranfälliges Design, IMHO.

Im Detail:

Es ist eine intern erzeugte Routine (__cudaRegisterFatBinary), die ausgeführt werden muss, um Kerne, Texturen zu laden und registrieren und statisch definierten Vorrichtung Symbole in der fatbin Nutzlast von jeder Laufzeit-API-Programm enthalten ist, mit der CUDA-Treiber API vor dem Kernel kann ohne Fehler aufgerufen werden. Dies ist ein Teil der "faulen" Kontextinitialisierungsfunktion der Laufzeit-API.Sie können dies für sich selbst wie folgt bestätigen:

Hier ist ein gdb-Trace des überarbeiteten Beispiels, das Sie gepostet haben. Hinweis füge ich einen Haltepunkt in __cudaRegisterFatBinary, und das nicht erreicht wird, bevor Sie Ihr statischer A Konstruktor aufgerufen wird und der Kernel Start fehlschlägt:

talonmies@box:~$ gdb a.out 
GNU gdb (Ubuntu/Linaro 7.4-2012.04-0ubuntu2.1) 7.4-2012.04 
Copyright (C) 2012 Free Software Foundation, Inc. 
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html> 
This is free software: you are free to change and redistribute it. 
There is NO WARRANTY, to the extent permitted by law. Type "show copying" 
and "show warranty" for details. 
This GDB was configured as "x86_64-linux-gnu". 
For bug reporting instructions, please see: 
<http://bugs.launchpad.net/gdb-linaro/>... 
Reading symbols from /home/talonmies/a.out...done. 
(gdb) break '__cudaRegisterFatBinary' 
Breakpoint 1 at 0x403180 
(gdb) run 
Starting program: /home/talonmies/a.out 
[Thread debugging using libthread_db enabled] 
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1". 
Scalar Field 
[New Thread 0x7ffff5a63700 (LWP 10774)] 
Class A 
Kernel : invalid device function 
[Thread 0x7ffff5a63700 (LWP 10774) exited] 
[Inferior 1 (process 10771) exited with code 0377] 

Hier wird das gleiche Verfahren ist, diesmal mit A Instanziierung innerhalb main (was garantierter geschehen, nachdem die Objekte, die faul Setup ausführen haben initialisiert):

talonmies@box:~$ cat main.cu 
#include "classA.h" 


int main() { 
    A a_object; 
    std::cout << "Main" << std::endl; 
    return 0; 
} 

talonmies@box:~$ nvcc --keep -arch=sm_30 -g main.cu 
talonmies@box:~$ gdb a.out 
GNU gdb (Ubuntu/Linaro 7.4-2012.04-0ubuntu2.1) 7.4-2012.04 
Copyright (C) 2012 Free Software Foundation, Inc. 
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html> 
This is free software: you are free to change and redistribute it. 
There is NO WARRANTY, to the extent permitted by law. Type "show copying" 
and "show warranty" for details. 
This GDB was configured as "x86_64-linux-gnu". 
For bug reporting instructions, please see: 
<http://bugs.launchpad.net/gdb-linaro/>... 
Reading symbols from /home/talonmies/a.out...done. 
(gdb) break '__cudaRegisterFatBinary' 
Breakpoint 1 at 0x403180 
(gdb) run 
Starting program: /home/talonmies/a.out 
[Thread debugging using libthread_db enabled] 
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1". 

Breakpoint 1, 0x0000000000403180 in __cudaRegisterFatBinary() 
(gdb) cont 
Continuing. 
Scalar Field 
[New Thread 0x7ffff5a63700 (LWP 11084)] 
Class A 
Main 
[Thread 0x7ffff5a63700 (LWP 11084) exited] 
[Inferior 1 (process 11081) exited normally] 

Wenn dies für Sie wirklich eine lähmende Problem ist, würde ich vorschlagen, NVIDIA Entwickler Support kontaktieren und einen Fehlerbericht zu erhöhen.