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

Examinando as respostas e os comentários sobre as perguntas da CUDA, e no wiki da tag CUDA , vejo que é frequentemente sugerido que o status de retorno de cada chamada de API seja verificado quanto a erros. A documentação da API contém funções como cudaGetLastError, cudaPeekAtLastErrore cudaGetErrorString, mas o que é a melhor maneira de colocá-los juntos para erros de captura e relatório de forma confiável sem a necessidade de um monte de código extra?

Provavelmente, a melhor maneira de verificar se há erros no código da API de tempo de execução é definir uma função de manipulador de estilo assert e uma macro de wrapper como esta:

#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);
   }
}

Em seguida, você pode agrupar cada chamada de API com a gpuErrchkmacro, que processará o status de retorno da API chamada de quebra, por exemplo:

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

Se houver um erro em uma chamada, será emitida uma mensagem de texto descrevendo o erro e o arquivo e linha no seu código em que o erro ocorreu stderre o aplicativo será encerrado. É possível modificar gpuAssertpara gerar uma exceção, em vez de chamar exit()um aplicativo mais sofisticado, se necessário.

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

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

primeiro verificará o argumento de inicialização inválido e forçará o host a esperar até que o kernel pare e verifique se há um erro de execução. A sincronização pode ser eliminada se você tiver uma chamada subsequente à API de bloqueio como esta:

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

nesse caso, a cudaMemcpychamada pode retornar os erros que ocorreram durante a execução do kernel ou os da própria cópia de memória. Isso pode ser confuso para o iniciante, e eu recomendaria o uso da sincronização explícita após o lançamento do kernel durante a depuração para facilitar a compreensão de onde os problemas podem estar surgindo.

Observe que, ao usar o CUDA Dynamic Parallelism , uma metodologia muito semelhante pode e deve ser aplicada a qualquer uso da API de tempo de execução CUDA nos kernels do dispositivo, bem como após o lançamento de qualquer kernel do dispositivo:

#include <assert.h>
#define cdpErrchk(ans) { cdpAssert((ans), __FILE__, __LINE__); }
__device__ void cdpAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      printf("GPU kernel assert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) assert(0);
   }
}

A resposta das garras acima é uma ótima maneira de abortar um aplicativo de assertmaneira estilo.

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 do std::runtime_erroruso thrust::system_error:

#include <thrust/system_error.h>
#include <thrust/system/cuda/error.h>
#include <sstream>

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 incorporará o nome do arquivo, o número da linha e uma descrição do idioma inglês cudaError_tno .what()membro da exceção lançada :

#include <iostream>

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_functionpode distinguir erros CUDA de outros tipos de erros, se desejar:

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é a std::runtime_error, podemos alternativamente lidar com isso da mesma maneira que uma ampla classe de erros se não exigirmos a 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;
}

A maneira canônica de C ++: Não verifique se há erros ... use as ligações C ++ que geram exceções.

Eu costumava ficar irritado com esse problema; e eu costumava ter uma solução de função macro-cum-wrapper, como nas respostas de Talonmies e Jared, mas, honestamente? Torna o uso da API do CUDA Runtime ainda mais feio e parecido com o C.

Então, eu abordei isso de uma maneira diferente e mais fundamental. Para uma amostra do resultado, aqui está parte da vectorAddamostra CUDA - com verificação completa de erros de cada chamada da API em tempo de execução:

// (... prepare host-side buffers here ...)

auto current_device = cuda::device::current::get();
auto d_A = cuda::memory::device::make_unique<float[]>(current_device, numElements);
auto d_B = cuda::memory::device::make_unique<float[]>(current_device, numElements);
auto d_C = cuda::memory::device::make_unique<float[]>(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 erros em potencial são verificados e uma exceção se ocorrer um erro (ressalva: se o kernel causou algum erro após o lançamento, ele será capturado após a tentativa de copiar o resultado, não antes; para garantir que o kernel tenha êxito, precisa verificar se há um erro entre o lançamento e a cópia com um cuda::outstanding_error::ensure_none()comando).

O código acima usa meu

Wrappers Modern-C ++ finos para a biblioteca de API de tempo de execução CUDA (Github)

Observe que as exceções contêm uma explicação de cadeia 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 CUDA são verificados automaticamente com estes wrappers:

• Um programa de teste lançando e capturando várias exceções
• Documentação para a funcionalidade relacionada a erros

A solução discutida aqui funcionou bem para mim. Esta solução usa funções cuda embutidas e é muito simples de implementar.

O código relevante é copiado abaixo:

#include <stdio.h>
#include <stdlib.h>

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