Problemi lanciare i kernel CUDA dal codice di inizializzazione statica

Question

Ho una classe che chiama un kernel nel suo costruttore, come segue:

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

Se istanzio questa classe su main (A a_object;) non ottengo errori. Tuttavia, se lo istanziamo all'esterno del main, appena dopo averlo definito (class A {...} a_object;), ricevo un errore di "funzione dispositivo non valida" all'avvio del kernel. Perché succede?

EDIT

Aggiornato codice per fornire un esempio più completo.

EDIT 2

Seguendo i consigli nel commento di Raxvan, volevo dire che ho la variabile dimensions utilizzato nel costruttore ScalarField anche definito (in un'altra classe) al di fuori principale, ma prima di tutto il resto. Potrebbe essere questa la spiegazione? Il debugger stava mostrando il giusto valore per dimensions.

Answer 1

La versione corta:

La ragione per il problema quando class A viene istanziata fuori della principale è che una routine particolare gancio che è richiesto per inizializzare la libreria di runtime CUDA con le kernel non viene eseguito prima viene chiamato il costruttore di class A. Questo accade perché non ci sono garanzie sull'ordine in cui gli oggetti statici vengono istanziati e inizializzati nel modello di esecuzione C++. La classe di ambito globale viene creata un'istanza prima che gli oggetti ambito globale che eseguono l'installazione CUDA vengano inizializzati. Il codice del kernel non viene mai caricato nel contesto prima della chiamata e viene generato un errore di runtime.

Per quello che posso dire, si tratta di una vera e propria limitazione dell'API di runtime CUDA e non di un problema facilmente risolto nel codice utente. Nel tuo banale esempio, potresti sostituire la chiamata del kernel con una chiamata a cudaMemset o una delle funzioni di memset API di runtime non basate su simboli e funzionerà. Questo problema è completamente limitato ai kernel utente o ai simboli del dispositivo caricati in fase di runtime tramite l'API runtime. Per questo motivo, un costruttore predefinito vuoto risolverebbe anche il tuo problema. Dal punto di vista del design, sarei molto dubbioso su qualsiasi modello che chiama kernel nel costruttore. L'aggiunta di un metodo specifico per l'installazione/rimozione della GPU di classe che non si basa sul costruttore o sul distruttore di default sarebbe un progetto molto più pulito e meno incline agli errori, IMHO.

In dettaglio:

C'è una routine generato internamente (__cudaRegisterFatBinary) che deve essere eseguito per caricare e registrare kernel, texture e simboli di unità definite staticamente contenute nel payload fatbin di qualsiasi programma runtime API con il API del driver CUDA prima che il kernel possa essere chiamato senza errori. Questa è una parte della funzione di inizializzazione del contesto "lazy" dell'API runtime.Puoi confermare questo per te come segue:

Ecco una traccia gdb dell'esempio revisionato che hai postato. Nota inserisco un punto di interruzione in __cudaRegisterFatBinary, e che non viene raggiunto prima della A costruttore statico viene chiamato e il lancio del kernel fallisce:

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] 

Ecco la stessa procedura, questa volta con A instanziazione all'interno main (che è garantito per accadere dopo che gli oggetti che eseguono l'installazione pigri sono state inizializzato):

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] 

Se questo è davvero un problema paralizzante per te, vorrei suggerire di contattare l'assistenza degli sviluppatori di NVIDIA e sollevando un bug report.