《CUDA Fortran Programming Guide and Reference》阅读笔记之 错误处理

在进行CUDA Fortran程序的测试时，发现结果不对，但是又不存在语法和逻辑错误，所以十分需要获取程序错误信息，但是Nvidia和CUDA提供了很多的可视化(Nsight、NVVP)和命令行（cuda-gdb）的形式，但是这些都支持CUDA C。对于CUDA Fortran的调试，在官方文档中提供了Error Handling的方式，也即一些cudaError函数。本文结合《CUDA Fortran Programming Guide and Reference》和《CUDA C Programming Guide and Reference》两个文档。

运行时为每个初始化为cudaSuccess的主机线程维护一个错误变量，并在每次发生错误时由错误代码覆盖(无论是参数验证错误还是异步错误)。

1、cudaGetErrorString()

cudaGetErrorString返回与给定错误代码相关联的消息字符串。

function cudaGetErrorString( errcode )
    integer, intent(in) :: errcode
    character*(*) :: cudaGetErrorString

可以看出cudaGetErrorString()的输入是error code，其输出的cudaSuccess或程序的错误信息

2、cudaGetLastError()

cudaGetLastError返回最近从该主机线程的任何运行时调用中返回的错误代码。

integer function cudaGetLastError()

可以看出cudaGetLastError()没有输入参数，在CUDA Fortran中该函数是integer型，说明其使用时需要返回给一个integer的变量，返回值为程序的错误信息，也就是error code，并将其重置为cudaSuccess

3、cudaPeekAtLastError()

cudaPeekAtLastError返回CUDA运行时生成的最后一个错误代码，而不需要像cudaGetLastError那样将错误代码重置为cudaSuccess。

integer function cudaPeekAtLastError()

可以看出cudaPeekAtLastError()没有输入参数，在CUDA Fortran中该函数是integer型，说明其使用时需要返回给一个integer的变量。返回值为程序的错误信息，也就是error code，但不会重置该变量。

所有运行时函数都返回一个错误代码，在使用时要确保error code设置为cudaSuccess。但对于异步函数，此错误代码不能报告设备上可能发生的任何异步错误，因为该函数在设备完成任务之前已经返回主机执行后续程序。

内核启动是异步的，因此为了检查异步错误，应用程序必须在内核启动和调用cudaPeekAtLastError()或cudaGetLastError()之间进行同步。同步方法: call cudaDeviceSynchronize() (或使用其他同步机制) 。确保任何错误返回cudaPeekAtLastError()或cudaGetLastError()并不源自调用内核启动前。但是进行错误检查也是有代价的，会降低程序的性能，特别是对cudaDeviceSynchronize()的调用，会阻塞线程直到kernel 函数执行完毕。

使用举例

integer :: istat1, istat2
call addKernel<<dimGrid, dimBlock>>>(a)
istat1 = cudaGetLastError()
istat2 = cudaDeviceSynchronize() 
if (istat1 .ne. cudaSuccess) print *, 'Sync kernel error:', cudaGetErrorString(istat1)
if (istat2 .ne. cudaSuccess) print *, 'ASync kernel error:', cudaGetErrorString(istat2)

第二行调用kernel函数，第三行提供了kernel函数启动时可能出现的错误，第四行提供了kernel函数执行时可能出现的错误。

参考文档

[1]: 《CUDA Fortran Programming Guide and Reference》
[2]: 《CUDA C Programming Guide and Reference》
[3]: 《CUDA for Engineers》