¿Cuál es la forma canónica de verificar errores usando la API de tiempo de ejecución CUDA?

Al revisar las respuestas y los comentarios sobre las preguntas de CUDA, y en el wiki de etiquetas de CUDA , veo que a menudo se sugiere que el estado de devolución de cada llamada a la API debe verificarse en busca de errores. La documentación de la API contiene funciones como cudaGetLastError, cudaPeekAtLastErrory cudaGetErrorString, pero ¿cuál es la mejor manera de poner éstos juntos para detectar los errores e informe de forma fiable sin requerir una gran cantidad de código extra?

Probablemente la mejor manera de verificar si hay errores en el código API de tiempo de ejecución es definir una función de controlador de estilo de aserción y una macro de envoltura 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);
   }
}

Luego, puede ajustar cada llamada a la API con la gpuErrchkmacro, que procesará el estado de retorno de la llamada a la API que se ajusta, por ejemplo:

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

Si hay un error en una llamada, se emitirá un mensaje de texto que describe el error y el archivo y la línea en su código donde se produjo el error stderry la aplicación se cerrará. Podría modificarse gpuAssertpara generar una excepción en lugar de llamar exit()a una aplicación más sofisticada si fuera necesario.

Una segunda pregunta relacionada es cómo verificar si hay errores en los lanzamientos de kernel, que no se pueden incluir directamente en una llamada de macro como las llamadas API estándar de tiempo de ejecución. Para los núcleos, algo como esto:

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

primero verificará si hay un argumento de inicio no válido, luego forzará al host a esperar hasta que el núcleo se detenga y verifique si hay un error de ejecución. La sincronización se puede eliminar si tiene una llamada de API de bloqueo posterior como esta:

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

en cuyo caso la cudaMemcpyllamada puede devolver errores que ocurrieron durante la ejecución del kernel o los de la copia de memoria. Esto puede ser confuso para el principiante, y recomendaría usar la sincronización explícita después del inicio del kernel durante la depuración para que sea más fácil entender dónde pueden surgir problemas.

Tenga en cuenta que cuando se utiliza el paralelismo dinámico de CUDA , una metodología muy similar puede y debe aplicarse a cualquier uso de la API de tiempo de ejecución de CUDA en los núcleos de dispositivos, así como después de que se inicie cualquier núcleo de dispositivos:

#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 respuesta anterior de talonmies es una buena manera de abortar una aplicación de una assertmanera estilo.

Ocasionalmente, podemos informar y recuperar una condición de error en un contexto C ++ como parte de una aplicación más grande.

Aquí hay una manera razonablemente breve de hacerlo lanzando una excepción C ++ derivada del 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);
  }
}

Esto incorporará el nombre de archivo, el número de línea y una descripción cudaError_ten inglés del .what()miembro de la excepción lanzada :

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

La salida:

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

Un cliente de some_functionpuede distinguir los errores de CUDA de otros tipos de errores si lo desea:

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_errores un std::runtime_error, podemos manejarlo alternativamente de la misma manera que una amplia clase de errores si no requerimos la precisión del ejemplo 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;
}

La forma canónica de C ++: no compruebe si hay errores ... use los enlaces de C ++ que arrojan excepciones.

Solía ​​molestarme este problema; y solía tener una solución de función macro-cum-wrapper como en las respuestas de Talonmies y Jared, pero, ¿honestamente? Hace que usar la API CUDA Runtime sea aún más feo y similar a C.

Así que he abordado esto de una manera diferente y más fundamental. Para obtener una muestra del resultado, aquí hay parte de la vectorAddmuestra de CUDA , con una verificación completa de errores de cada llamada a la API de tiempo de ejecución:

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

Una vez más, se verifican todos los posibles errores, y una excepción si se produce un error (advertencia: si el kernel causó algún error después del lanzamiento, se detectará después del intento de copiar el resultado, no antes; para asegurarse de que el kernel tuvo éxito, lo haría necesita verificar el error entre el inicio y la copia con un cuda::outstanding_error::ensure_none()comando).

El código anterior usa mi

Thin Modern-C ++ wrappers para la biblioteca de API CUDA Runtime (Github)

Tenga en cuenta que las excepciones llevan una explicación de cadena y el código de estado de API de tiempo de ejecución CUDA después de la llamada fallida.

Algunos enlaces a cómo los errores CUDA se verifican automáticamente con estos contenedores:

• Un programa de prueba que arroja y atrapa un montón de excepciones
• Documentación para la funcionalidad relacionada con errores

La solución discutida aquí funcionó bien para mí. Esta solución utiliza funciones integradas de cuda y es muy sencilla de implementar.

El código relevante se copia a continuación:

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