Kłopoty z uruchomieniem jąder CUDA kod inicjalizacji static

Question

Mam klasy, który wywołuje jądra w jej konstruktora, co następuje:

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

Jeśli uruchomię tę klasę na głównej (A a_object;), nie otrzymam żadnych błędów. Jednakże, jeśli utworzę go poza głównym, zaraz po jego zdefiniowaniu (class A {...} a_object;) otrzymuję komunikat o błędzie "nieprawidłowa funkcja urządzenia" po uruchomieniu jądra. Dlaczego tak się dzieje?

EDIT

kod bieżąco, aby zapewnić pełniejszy przykład.

EDIT 2

Po poradę w komentarzu przez Raxvan, chciałem powiedzieć, że mają dimensions zmienna użyta w konstruktorze ScalarField również zdefiniowany (w innej klasie) poza głównym, ale przed wszystkim innym. Czy to może być wyjaśnienie? Debugger miał jednak odpowiednią wartość dla dimensions.

Answer 1

Wersja krótka:

Podstawową przyczyną problemu, gdy class A jest tworzony poza głównym jest to, że zwłaszcza rutyna hak, który jest wymagany do zainicjowania biblioteki wykonawczego CUDA ze swoimi jądrami nie jest prowadzony przed konstruktor z class A jest wywoływany. Dzieje się tak, ponieważ nie ma gwarancji co do kolejności, w której obiekty statyczne są tworzone i inicjalizowane w modelu wykonawczym C++. Twoja globalna klasa zasięgu jest tworzona w instancji przed obiektami globalnego zasięgu, w których zainicjowano instalację CUDA. Twój kod jądra nigdy nie jest ładowany do kontekstu przed wywołaniem i pojawia się błąd wykonania.

Jak najlepiej rozumiem, jest to rzeczywiste ograniczenie interfejsu API środowiska wykonawczego CUDA, a nie coś łatwego do naprawienia w kodzie użytkownika. W swoim trywialnym przykładzie można zastąpić wywołanie jądra wywołaniem funkcji cudaMemset lub jedną z funkcji zestawu modułów wykonawczych API niezwiązanych z symbolami i to zadziała. Ten problem jest całkowicie ograniczony do jądra użytkownika lub symboli urządzeń załadowanych w środowisku wykonawczym za pośrednictwem interfejsu API środowiska wykonawczego. Z tego powodu pusty konstruktor domyślny również rozwiąże problem. Z punktu widzenia projektowania, byłbym bardzo wątpliwy wobec każdego wzorca, który wywołuje jądra w konstruktorze. Dodanie konkretnej metody instalacji/rozłączenia GPU klasy, która nie opiera się na domyślnym konstruktorze lub destruktorze, byłby znacznie czystszy i mniej podatny na błędy projekt, IMHO.

w szczegółach:

Jest wytworzone procedura (__cudaRegisterFatBinary), które muszą być prowadzone w celu załadowania i zarejestrować jądra, tekstur i statycznie określonych symboli urządzeń znajdujących się w polu danych fatbin każdego programu wykonawczego API z innymi Interfejs API sterownika CUDA przed jądrem można wywołać bez błędu. Jest to część "leniwej" funkcji inicjalizacji kontekstu API środowiska wykonawczego.Możesz to potwierdzić w następujący sposób:

Oto ślad gdb z poprawionego wysłanego przykładu. Uwaga wstawić przerwania do __cudaRegisterFatBinary, i że nie zostanie osiągnięty przed statycznym A konstruktor nazywa i uruchomienie jądra nie:

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] 

Oto ta sama procedura, tym razem z A instancji wewnątrz main (co jest gwarantowana zdarzyć po przedmioty, które wykonują leniwe konfiguracji zostały zainicjowane):

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] 

Jeśli jest to naprawdę wyniszczający dla Ciebie problem, chciałbym zaproponować kontaktując wsparcie programistów NVIDIA i podnoszenie raport o błędzie.