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

url canonica blogger

Question

Mirando las respuestas y los comentarios sobre las preguntas de CUDA, y en el wiki de la etiqueta de CUDA , veo que a menudo se sugiere que el estado de retorno de cada llamada a la API debería revisarse en busca de errores. La documentación de la API contiene funciones como cudaGetLastError , cudaPeekAtLastError y cudaGetErrorString , pero ¿cuál es la mejor manera de cudaGetErrorString para capturar e informar errores de manera confiable sin requerir mucho código adicional?

Answer 1

La forma canónica de C ++: No verifique los errores ... use los enlaces de C ++ que generan excepciones.

Solía ​​estar molesto por este problema; y solía tener una solución de función macro-cum-wrapper como en las respuestas de Talonmies y Jared, pero, sinceramente? Hace que el uso de CUDA Runtime API sea aún más feo y parecido a C.

Así que he abordado esto de una manera diferente y más fundamental. Para una muestra del resultado, aquí hay parte de la muestra vectorAdd CUDA, con una completa comprobación de errores en cada llamada a la API en 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...)

De nuevo, todos los errores potenciales se verifican y se informan mediante una excepción lanzada. Este código usa mi

Thin Modern-C ++ wrappers para la biblioteca de API de tiempo de ejecución CUDA (Github)

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

Algunos enlaces a cómo se comprueban automáticamente los errores de CUDA con estos envoltorios:

• Un programa de prueba lanzando y atrapando un montón de excepciones.
• Documentación para la funcionalidad relacionada con el error.

Answer 2

La respuesta anterior de Talonmies es una buena manera de abortar una aplicación de manera asertiva.

En ocasiones, es posible que deseamos informar y recuperar una condición de error en un contexto de C ++ como parte de una aplicación más grande.

Aquí hay una manera razonablemente concisa de hacerlo lanzando una excepción de C ++ derivada de std::runtime_error usando 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 en idioma inglés de cudaError_t en el miembro cudaError_t .what() 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_function puede 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;
}

Debido a que thrust::system_error es un std::runtime_error , alternativamente podemos manejarlo 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;
}

Answer 3

La solución discutida here funcionó bien para mí. Esta solución utiliza funciones de cuda integradas 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;
}

Answer 4

Probablemente, la mejor manera de verificar errores en el código de la API en tiempo de ejecución es definir una función de manejador de estilo de afirmació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 macro gpuErrchk , que procesará el estado de retorno de la llamada a la API, 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 a stderr y la aplicación se cerrará. gpuAssert podría modificar gpuAssert para generar una excepción en lugar de llamar a exit() en una aplicación más sofisticada si fuera necesario.

Una segunda pregunta relacionada es cómo verificar errores en los lanzamientos del kernel, que no pueden ser envueltos directamente en una llamada macro como las llamadas API estándar de tiempo de ejecución. Para los granos, 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 kernel se detenga y verifique un error de ejecución. La sincronización puede eliminarse 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 llamada cudaMemcpy puede devolver los errores que se produjeron durante la ejecución del kernel o los de la memoria en sí. Esto puede ser confuso para el principiante, y recomendaría el uso de la sincronización explícita después de un inicio del kernel durante la depuración para que sea más fácil comprender dónde pueden surgir los problemas.