Qual é a maneira canônica de verificar erros usando a API de tempo de execução CUDA?

Analisando as respostas e os comentários sobre as perguntas do CUDA, e no wiki da tag CUDA , vejo que muitas vezes é sugerido que o status de retorno de cada chamada da API seja verificado quanto a erros. A documentação da API contém funções como cudaGetLastError , cudaPeekAtLastError e cudaGetErrorString , mas qual é a melhor maneira de colocá-las juntas para capturar e relatar erros de forma confiável sem exigir muito código extra?

Provavelmente, a melhor maneira de verificar erros no código da API de tempo de execução é definir uma function de manipulador de estilo de declaração e uma macro de wrapper assim:

 #define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true) { if (code != cudaSuccess) { fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); if (abort) exit(code); } } 

Você pode então gpuErrchk cada chamada de API com a macro gpuErrchk , que processará o status de retorno da chamada de API que ela envolve, por exemplo:

 gpuErrchk( cudaMalloc(&a_d, size*sizeof(int)) ); 

Se houver um erro em uma chamada, uma mensagem de texto descrevendo o erro e o arquivo e a linha em seu código onde o erro ocorreu serão emitidos para stderr e o aplicativo será encerrado. Você pode modificar o gpuAssert para gerar uma exceção em vez de chamar exit() em um aplicativo mais sofisticado, se necessário.

Uma segunda questão relacionada é como verificar erros nos lançamentos do kernel, que não podem ser envolvidos diretamente em uma chamada de macro, como chamadas de API de tempo de execução padrão. Para kernels, algo assim:

 kernel< <<1,1>>>(a); gpuErrchk( cudaPeekAtLastError() ); gpuErrchk( cudaDeviceSynchronize() ); 

Em primeiro lugar, verificará se há argumento de boot inválido e, em seguida, forçará o host a aguardar até que o kernel pare e verifique se há um erro de execução. A synchronization pode ser eliminada se você tiver uma chamada de API de bloqueio subsequente assim:

 kernel< <<1,1>>>(a_d); gpuErrchk( cudaPeekAtLastError() ); gpuErrchk( cudaMemcpy(a_h, a_d, size * sizeof(int), cudaMemcpyDeviceToHost) ); 

cudaMemcpy caso, a chamada cudaMemcpy pode retornar os erros que ocorreram durante a execução do kernel ou os da própria cópia da memory. Isso pode ser confuso para o iniciante, e eu recomendaria usar a synchronization explícita após uma boot do kernel durante a debugging para facilitar a compreensão de onde os problemas podem estar surgindo.

A resposta de talonmies acima é uma ótima maneira de abortar uma aplicação de uma maneira estilo- assert .

Ocasionalmente, podemos desejar relatar e recuperar de uma condição de erro em um contexto C ++ como parte de um aplicativo maior.

Aqui está uma maneira razoavelmente concisa de fazer isso lançando uma exceção C ++ derivada de std::runtime_error usando thrust::system_error :

 #include  #include  #include  void throw_on_cuda_error(cudaError_t code, const char *file, int line) { if(code != cudaSuccess) { std::stringstream ss; ss < < file << "(" << line << ")"; std::string file_and_line; ss >> file_and_line; throw thrust::system_error(code, thrust::cuda_category(), file_and_line); } } 

Isso irá incorporar o nome do arquivo, o número da linha e uma descrição em inglês do cudaError_t no membro .what() da exceção lançada:

 #include  int main() { try { // do something crazy throw_on_cuda_error(cudaSetDevice(-1), __FILE__, __LINE__); } catch(thrust::system_error &e) { std::cerr < < "CUDA error after cudaSetDevice: " << e.what() << std::endl; // oops, recover cudaSetDevice(0); } return 0; } 

A saída:

 $ nvcc exception.cu -run CUDA error after cudaSetDevice: exception.cu(23): invalid device ordinal 

Um cliente de some_function pode distinguir erros CUDA de outros tipos de erros, se desejado:

 try { // call some_function which may throw something some_function(); } catch(thrust::system_error &e) { std::cerr < < "CUDA error during some_function: " << e.what() << std::endl; } catch(std::bad_alloc &e) { std::cerr << "Bad memory allocation during some_function: " << e.what() << std::endl; } catch(std::runtime_error &e) { std::cerr << "Runtime error during some_function: " << e.what() << std::endl; } catch(...) { std::cerr << "Some other kind of error during some_function" << std::endl; // no idea what to do, so just rethrow the exception throw; } 

Como thrust::system_error é std::runtime_error , podemos manipulá-lo alternativamente da mesma maneira que uma ampla class de erros se não precisarmos da precisão do exemplo anterior:

 try { // call some_function which may throw something some_function(); } catch(std::runtime_error &e) { std::cerr < < "Runtime error during some_function: " << e.what() << std::endl; } 

O C ++ – maneira canônica: Não verifique se há erros … use as ligações de C ++ que lançam exceções.

Eu costumava ficar incomodado com esse problema; e eu costumava ter uma solução macro-cum-wrapper, assim como nas respostas de Talonmies e Jared, mas, honestamente? Isso torna o uso da API de tempo de execução CUDA ainda mais feio e semelhante ao C.

Então eu me aproximei disso de uma maneira diferente e mais fundamental. Para uma amostra do resultado, aqui está parte da amostra vectorAdd CUDA – com verificação de erro completa de todas as chamadas da API de tempo de execução:

 // (... prepare host-side buffers here ...) auto current_device = cuda::device::current::get(); auto d_A = cuda::memory::device::make_unique(current_device, numElements); auto d_B = cuda::memory::device::make_unique(current_device, numElements); auto d_C = cuda::memory::device::make_unique(current_device, numElements); cuda::memory::copy(d_A.get(), h_A.get(), size); cuda::memory::copy(d_B.get(), h_B.get(), size); // (... prepare a launch configuration here... ) cuda::launch( vectorAdd, launch_config, d_A.get(), d_B.get(), d_C.get(), numElements ); cuda::memory::copy(h_C.get(), d_C.get(), size); // (... verify results here...) 

Novamente – todos os possíveis erros são verificados e relatados por meio de uma exceção gerada. Este código usa meu

Invólucros finos de C ++ modernos para a biblioteca da API de tempo de execução CUDA (Github)

Observe que as exceções transportam uma explicação de sequência e o código de status da API de tempo de execução CUDA após a chamada com falha

Alguns links para como os erros do CUDA são verificados automaticamente com esses wrappers:

  • Um programa de teste jogando e pegando um monte de exceções
  • Documentação para a funcionalidade relacionada a erros

A solução discutida aqui funcionou bem para mim. Esta solução utiliza funções corporativas integradas e é muito simples de implementar.

O código relevante é copiado abaixo:

 #include  #include  __global__ void foo(int *ptr) { *ptr = 7; } int main(void) { foo< <<1,1>>>(0); // make the host block until the device is finished with foo cudaDeviceSynchronize(); // check for error cudaError_t error = cudaGetLastError(); if(error != cudaSuccess) { // print the CUDA error message and exit printf("CUDA error: %s\n", cudaGetErrorString(error)); exit(-1); } return 0; }