Problema ao iniciar os kernels CUDA a partir do código de boot estática

Eu tenho uma class que chama um kernel em seu construtor, da seguinte maneira:

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

Se eu instancio esta class no main ( A a_object; ) não recebo nenhum erro. No entanto, se eu instanciá-lo fora principal, logo após defini-lo ( class A {...} a_object; ) recebo um erro “function de dispositivo inválido” quando o kernel é iniciado. Por que isso acontece?

EDITAR

Código atualizado para fornecer um exemplo mais completo.

EDIT 2

Seguindo o conselho no comentário de Raxvan, eu queria dizer que eu tenho a variável de dimensions usada no construtor ScalarField também definida (em outra class) fora do main, mas antes de tudo. Essa poderia ser a explicação? O depurador estava mostrando o valor correto para as dimensions .

A versão curta:

O motivo subjacente para o problema quando a class A é instanciada fora do main é que uma rotina de gancho específica que é necessária para inicializar a biblioteca de tempo de execução CUDA com seus kernels não está sendo executada antes do construtor da class A ser chamado. Isso acontece porque não há garantias sobre a ordem na qual os objects estáticos são instanciados e inicializados no modelo de execução do C ++. Sua class de escopo global está sendo instanciada antes dos objects de escopo global que fazem a configuração CUDA serem inicializados. Seu código do kernel nunca está sendo carregado no contexto antes de ser chamado e ocorre um erro de tempo de execução.

Tanto quanto eu posso dizer, esta é uma limitação genuína da API de tempo de execução CUDA e não algo facilmente corrigido no código do usuário. Em seu exemplo trivial, você poderia replace a chamada do kernel por uma chamada para cudaMemset ou uma das funções memset da API de tempo de execução baseadas em não-símbolos e ela funcionará. Esse problema é completamente limitado a kernels de usuários ou símbolos de dispositivos carregados em tempo de execução por meio da API de tempo de execução. Por esse motivo, um construtor padrão vazio também resolveria seu problema. Do ponto de vista do design, eu seria muito duvidoso de qualquer padrão que chama kernels no construtor. Adicionar um método específico para configuração / desassembly de GPU de class que não dependa do construtor ou destruidor padrão seria um projeto muito mais limpo e menos propenso a erros, o IMHO.

Em detalhe:

Existe uma rotina gerada internamente ( __cudaRegisterFatBinary ) que deve ser executada para carregar e registrar kernels, texturas e símbolos de dispositivos estaticamente definidos contidos na carga útil de qualquer programa API de runtime com a API do driver CUDA antes que o kernel possa ser chamado sem erro. Essa é uma parte do recurso de boot de contexto “preguiçoso” da API de tempo de execução. Você pode confirmar isso por si mesmo da seguinte maneira:

Aqui está um rastreamento do gdb do exemplo revisado que você postou. Observe que insiro um ponto de interrupção em __cudaRegisterFatBinary e isso não é alcançado antes que seu construtor estático A seja chamado e a boot do kernel falhe:

 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] 

Aqui está o mesmo procedimento, desta vez com A instanciação dentro do main (que é garantido que acontecerá depois que os objects que executam a configuração lenta foram inicializados):

 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] 

Se esse for realmente um problema incapacitante para você, sugiro entrar em contato com o suporte ao desenvolvedor da NVIDIA e criar um relatório de bug.

    Intereting Posts