Problemas al iniciar kernels CUDA a partir del código de inicialización estático

Tengo una clase que llama un núcleo en su constructor, de la siguiente manera:

“ScalarField.h”

#include  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<<>>(v); ERROR_CHECK(cudaGetLastError(),"Kernel"); } }; 

“main.cu”

 #include "classA.h" A a_object; int main() { std::cout << "Main" << std::endl; return 0; } 

Si instalo esta clase en main ( A a_object; ) no obtengo ningún error. Sin embargo, si lo instalo fuera de main, justo después de definirlo ( class A {...} a_object; ) class A {...} a_object; un error de “función de dispositivo no válida” cuando se inicia el kernel. ¿Por qué sucede eso?

EDITAR

Código actualizado para proporcionar un ejemplo más completo.

EDIT 2

Siguiendo el consejo del comentario de Raxvan, quería decir que tengo la variable de dimensions utilizada en el constructor ScalarField también definida (en otra clase) fuera de main, pero antes que nada. ¿Podría ser esa la explicación? Sin embargo, el depurador mostraba el valor correcto para las dimensions .

La versión corta:

La razón subyacente del problema cuando se crea class A instancia de class A fuera de main es que no se está ejecutando una rutina de enlace particular que se requiere para inicializar la biblioteca de tiempo de ejecución de CUDA con sus núcleos antes de que se llame al constructor de class A Esto sucede porque no hay garantías sobre el orden en que los objetos estáticos se instancian e inicializan en el modelo de ejecución de C ++. Su clase de scope global se está instanciando antes de que se inicialicen los objetos de scope global que hacen la configuración de CUDA. El código del kernel nunca se carga en el contexto antes de que se llame y se produce un error de tiempo de ejecución.

Lo mejor que puedo decir es que esta es una limitación genuina de la API de tiempo de ejecución de CUDA y no es algo fácil de corregir en el código de usuario. En su ejemplo trivial, podría reemplazar la llamada del kernel con una llamada a cudaMemset o una de las funciones memset de la API en tiempo de ejecución no basadas en símbolos y funcionará. Este problema está completamente limitado a kernels de usuario o símbolos de dispositivos cargados en tiempo de ejecución a través de la API de tiempo de ejecución. Por esta razón, un constructor predeterminado vacío también resolvería su problema. Desde un punto de vista de diseño, dudaría mucho de cualquier patrón que llame kernels en el constructor. Agregar un método específico para la configuración / desassembly de GPU de clase que no dependa del constructor o destructor predeterminado sería un diseño mucho más limpio y menos propenso a errores, en mi humilde opinión.

En detalle:

Existe una rutina generada internamente ( __cudaRegisterFatBinary ) que debe ejecutarse para cargar y registrar kernels, texturas y símbolos de dispositivos definidos estáticamente en la carga útil de fatbin de cualquier API de tiempo de ejecución con la API del controlador CUDA antes de poder invocar el kernel sin error. Esto es parte de la característica de inicialización de contexto “vago” de la API de tiempo de ejecución. Puede confirmarlo usted mismo de la siguiente manera:

Aquí hay un rastro gdb del ejemplo revisado que publicó. Nota: inserto un punto de interrupción en __cudaRegisterFatBinary , y eso no se alcanza antes de que se llame al constructor de A estático y el lanzamiento del kernel falle:

 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  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: ... 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] 

Este es el mismo procedimiento, esta vez con A instanciación dentro de main (que se garantiza que ocurrirá después de que los objetos que realizan la configuración perezosa se hayan inicializado):

 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  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: ... 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] 

Si esto es realmente un problema paralizante para usted, le sugiero que se ponga en contacto con el soporte para desarrolladores de NVIDIA y genere un informe de errores.