Quelle est la manière canonique de vérifier les erreurs à l'aide de l'API d'exécution CUDA?

En parcourant les réponses et les commentaires sur les questions CUDA, et dans le wiki des balises CUDA , je constate qu'il est souvent suggéré que l'état de retour de chaque appel d'API devrait vérifier les erreurs. La documentation de l'API contient des fonctions commecudaGetLastError que , cudaPeekAtLastErroret cudaGetErrorString, mais quelle est la meilleure façon de les mettre ensemble pour les prises de manière fiable et des erreurs rapport sans nécessiter beaucoup de code supplémentaire?

La meilleure façon de vérifier les erreurs dans le code API d'exécution est probablement de définir une fonction de gestionnaire de style d'assertion et une macro d'encapsuleur comme celle-ci:

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

Vous pouvez ensuite encapsuler chaque appel d'API avec la gpuErrchkmacro, qui traitera le statut de retour de l'appel d'API qu'il encapsule, par exemple:

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

S'il y a une erreur dans un appel, un message textuel décrivant l'erreur et le fichier et la ligne dans votre code où l'erreur s'est produite seront émis stderret l'application se fermera. Vous pouvez éventuellement modifier gpuAssertpour déclencher une exception plutôt que d'appelerexit() une application plus sophistiquée si cela était nécessaire.

Une deuxième question connexe est de savoir comment vérifier les erreurs dans les lancements du noyau, qui ne peuvent pas être directement encapsulées dans un appel de macro comme les appels d'API d'exécution standard. Pour les noyaux, quelque chose comme ceci:

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

vérifiera d'abord un argument de lancement non valide, puis forcera l'hôte à attendre jusqu'à ce que le noyau s'arrête et recherche une erreur d'exécution. La synchronisation peut être éliminée si vous avez un appel d'API de blocage ultérieur comme celui-ci:

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

dans ce cas, l' cudaMemcpyappel peut renvoyer soit des erreurs survenues lors de l'exécution du noyau, soit celles de la copie mémoire elle-même. Cela peut être déroutant pour le débutant, et je recommanderais d'utiliser une synchronisation explicite après un lancement du noyau pendant le débogage pour faciliter la compréhension des problèmes potentiels.

Notez que lorsque vous utilisez CUDA Dynamic Parallelism , une méthodologie très similaire peut et doit être appliquée à toute utilisation de l'API d'exécution CUDA dans les noyaux de périphériques, ainsi qu'après tout lancement de noyau de périphériques:

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

La réponse de talonmies ci-dessus est un bon moyen d'interrompre une application de assertmanière -style.

Parfois, nous pouvons souhaiter signaler et récupérer une condition d'erreur dans un contexte C ++ dans le cadre d'une application plus grande.

Voici une façon raisonnablement concise de le faire en lançant une exception C ++ dérivée de l' std::runtime_errorutilisation 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);
  }
}

Cela incorporera le nom de fichier, le numéro de ligne et une description en anglais du cudaError_tdans le .what()membre de l'exception levée :

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

Le résultat:

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

Un client de some_functionpeut distinguer les erreurs CUDA des autres types d'erreurs s'il le souhaite:

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

Parce que thrust::system_errorest un std::runtime_error, nous pouvons le gérer de la même manière d'une large classe d'erreurs si nous n'avons pas besoin de la précision de l'exemple précédent:

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

La méthode canonique C ++: ne vérifiez pas les erreurs ... utilisez les liaisons C ++ qui lèvent des exceptions.

J'étais irrité par ce problème; et j'avais l'habitude d'avoir une solution de fonction macro-wrapper comme dans les réponses de Talonmies et Jared, mais, honnêtement? Cela rend l'utilisation de l'API CUDA Runtime encore plus laide et semblable à C.

J'ai donc abordé cela d'une manière différente et plus fondamentale. Pour un exemple du résultat, voici une partie de l' vectorAddexemple CUDA - avec une vérification complète des erreurs de chaque appel d'API d'exécution:

// (... 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...)

Encore une fois - toutes les erreurs potentielles sont vérifiées, et une exception en cas d'erreur (mise en garde: si le noyau a causé une erreur après le lancement, il sera détecté après la tentative de copie du résultat, pas avant; pour vous assurer que le noyau a réussi, vous le feriez besoin de vérifier les erreurs entre le lancement et la copie avec uncuda::outstanding_error::ensure_none() commande).

Le code ci-dessus utilise mon

Enveloppeurs Thin Modern-C ++ pour la bibliothèque API CUDA Runtime (Github)

Notez que les exceptions portent à la fois une explication de chaîne et le code d'état de l'API d'exécution CUDA après l'appel ayant échoué.

Quelques liens vers la façon dont les erreurs CUDA sont vérifiées automatiquement avec ces wrappers:

• Un programme de test lançant et interceptant un tas d'exceptions
• Documentation pour la fonctionnalité liée aux erreurs

La solution discutée ici a bien fonctionné pour moi. Cette solution utilise des fonctions cuda intégrées et est très simple à implémenter.

Le code correspondant est copié ci-dessous:

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