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




url enlace (5)

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?


Answers

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

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

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.


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:


Para la GTX 970 hay 13 Multiprocesadores Streaming (SM) con 128 Cuda Cores cada uno. Los núcleos Cuda también se llaman Stream Processors (SP).

Puede definir las cuadrículas que asignan bloques a la GPU.

Puede definir bloques que mapeen hilos a Procesadores de flujo (los 128 Cuda Cores por SM).

Una urdimbre siempre está formada por 32 hilos y todos los hilos de una urdimbre se ejecutan de forma simultánea.

Para usar toda la potencia posible de una GPU necesita mucho más hilos por SM que SM tiene SP. Para cada capacidad de cómputo existe una cierta cantidad de subprocesos que pueden residir en un SM a la vez. Todos los bloques que defina están en cola y esperan que un SM tenga los recursos (número de SP libres), luego se carga. El SM comienza a ejecutar Warps. Como un Warp solo tiene 32 Threads y un SM tiene, por ejemplo, 128 SP, un SM puede ejecutar 4 Warps en un momento determinado. El problema es que si los subprocesos tienen acceso a la memoria, el hilo se bloqueará hasta que se satisfaga su solicitud de memoria. En números: un cálculo aritmético en el SP tiene una latencia de 18-22 ciclos, mientras que un acceso de memoria global no almacenado en caché puede tomar hasta 300-400 ciclos. Esto significa que si los hilos de un warp están esperando datos, solo un subconjunto de los 128 SP funcionaría. Por lo tanto, el planificador cambia para ejecutar otro warp, si está disponible. Y si este warp bloquea, ejecuta el siguiente y así sucesivamente. Este concepto se llama ocultación de latencia. El número de urdimbres y el tamaño del bloque determinan la ocupación (de cuántas deformaciones puede elegir ejecutar el SM). Si la ocupación es alta, es más improbable que no haya trabajo para los SP.

Su afirmación de que cada núcleo cuda ejecutará un bloque a la vez es incorrecta. Si habla de multiprocesadores de transmisión, puede ejecutar warps desde todos los hilos que residen en el SM. Si un bloque tiene un tamaño de 256 subprocesos y su GPU permite que 2048 subprocesos sean residentes por SM, cada SM tendría 8 bloques residentes desde los cuales el SM puede elegir distorsiones para ejecutar. Todos los hilos de los warps ejecutados se ejecutan en paralelo.

Aquí encontrará números para las diferentes Capacidades de cálculo y Arquitecturas de GPU: https://en.wikipedia.org/wiki/CUDA#Limitations

Puede descargar una hoja de cálculo de ocupación de la hoja de Cálculo de ocupación de Nvidia (por Nvidia) .