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

Repasando las respuestas y los comentarios sobre las preguntas de CUDA, y en la wiki de la etiqueta de CUDA , veo que a menudo se sugiere que el estado de devolución de cada llamada API se revise para detectar errores. La documentación de la API contiene funciones como cudaGetLastError , cudaPeekAtLastError y cudaGetErrorString , pero ¿cuál es la mejor manera de juntarlas para detectar e informar errores de manera confiable sin requerir muchos códigos adicionales?

Probablemente, la mejor forma de comprobar si hay errores en el código API de tiempo de ejecución es definir una función de controlador 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); } } 

A continuación, puede ajustar cada llamada a la API con la macro gpuErrchk , que procesará el estado de devolución de la llamada API que envuelve, por ejemplo:

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

Si hay un error en una llamada, se enviará a stderr un mensaje de texto que describe el error y el archivo y la línea en su código donde ocurrió el error 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 si hay errores en los inicios del kernel, que no pueden ser envueltos directamente en una llamada de macro como las llamadas API de tiempo de ejecución estándar. Para kernels, algo como esto:

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

primero verificará si hay un argumento de inicio no válido, y 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 puede eliminarse si tiene una llamada API de locking 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 ocurrieron durante la ejecución del kernel o los de la propia copia de la memoria. Esto puede ser confuso para el principiante, y recomendaría usar la sincronización explícita después de un lanzamiento del kernel durante la depuración para que sea más fácil entender dónde podrían surgir problemas.

La respuesta de talonmies anterior es una buena forma de abortar una aplicación de una manera de estilo assert .

Ocasionalmente, es posible que deseemos informarnos y recuperarnos de una condición de error en un contexto de C ++ como parte de una aplicación más grande.

Esta es una forma razonable de hacerlo lanzando una excepción de C ++ derivada de std::runtime_error utilizando thrust::system_error :

 #include  #include  #include  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  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 std::runtime_error , alternativamente podemos manejarlo de la misma manera que una clase amplia 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; } 

El C ++ – forma canónica: no verifique los errores … use los enlaces C ++ que generan excepciones.

Antes me molestaba este problema; y solía tener una solución macro-cum-wrapper-function como en Talonmies y las respuestas de Jared, pero, ¿sinceramente? Hace que utilizar CUDA Runtime API sea aún más feo y similar a C.

Así que me he acercado a esto de una manera diferente y más fundamental. Para obtener una muestra del resultado, aquí hay una parte de la muestra CUDA vectorAdd , con la comprobación completa de errores de cada llamada 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(current_device, numElements); auto d_B = cuda::memory::device::make_unique(current_device, numElements); auto d_C = cuda::memory::device::make_unique(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 posibles errores se verifican y se informan a través de una excepción lanzada. Este código usa mi

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

Tenga en cuenta que las excepciones incluyen tanto una explicación de cadena como 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 de CUDA se verifican automágicamente con estos envoltorios:

  • Un progtwig de prueba arrojando y atrapando 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 usa funciones de cuda integradas y es muy simple de implementar.

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

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