CUDA编程（六）：代码分析与调试

代码分析与调试方法

CUDA代码的运行时可能会遇到内存溢出、内存非法访问、核函数出错等一系列问题，需要进行Debug调试。在CUDA代码的分析与调试过程中，通常可以采用以下几种方法。

使用printf打印变量信息

在CUDA程序中，可以使用printf()函数打印变量信息。但是需要注意，由于CUDA是并行计算架构，因此在使用printf()函数时需要谨慎，避免对性能造成过大的影响。

使用CUDA的错误检查功能

在CUDA中，可以使用错误处理函数检查CUDA程序哪一个环节出现错误。在每次CUDA函数调用后，使用该函数进行检查，并打印错误信息以帮助诊断问题。另外，也可以使用cuda-memcheck工具来检查内存错误。

CUDA编程中错误处理相关的函数一共有4个：

// cudaGetErrorName函数接受一个错误码，返回错误名称
__host__ __device__ const char* cudaGetErrorName(cudaError_t error)
// cudaGetErrorString函数接受一个错误码，返回错误描述
__host__ __device__ const char* cudaGetErrorString(cudaError_t error)
// cudaGetLastError函数返回上一次CUDA runtime调用时的错误码，并将CUDA错误码置为cudaSuccess
__host__ __device__ cudaError_t cudaGetLastError(void)
// cudaPeekAtLastError函数与上面功能一样，但不会重置CUDA的错误信息
__host__ __device__ cudaError_t cudaPeekAtLastError(void)

使用CUDA-GDB进行调试

CUDA-GDB是一个基于命令行的调试工具，可用于调试CUDA C/C++应用程序。使用CUDA-GDB可以设置断点、单步调试、查看变量值等。需要在编译时使用-g、-G选项来生成可调试版本。CUDA-GDB使用方法如下：

# 编译程序的时候我们使用nvcc编译器，将其编译成可调式版本；-g 表示将CPU代码(host)编译成可调式版本，-G表示将GPU代码(kernel)编译成可调式版本。
nvcc -g -G XXX.cu -o XXX
# 进入调试器
cuda-gdb XXX
# run执行
run

使用Nsight进行调试

Nsight 开发工具套件提供深入的跟踪、调试、评测和其他分析，以优化跨 NVIDIA GPU 和 CPU 的复杂计算应用程序，包括 x86、ARM 和 Power 体系结构。一般在装CUDA时，在图形界面已安装好Nsight。用终端打开Nsight编译器，它可以直接对写好程序进行编译和Debug。关于如何使用Nsight调试程序，可以参考链接。

使用nvprof / nvvp工具

NVIDIA nvprof / nvvp工具是英伟达N卡GPU编程中用于观察的利器。全称是NVIDIA Visual Profiler，是由2008年起开始支持的性能分析器。它具有交互性好，利于使用的优点，并且可用于分析CUDA应用程序的性能瓶颈和优化方案，其在安装好CUDA toolkit后方可使用。

在CUDA程序调试过程中，Visual Profiler提供每个CUDA函数调用的时间分析，它还能给出如何调用内核函数以及存储器的使用情况等，有助于定位瓶颈可能出现的位置，并详细解释如何调用内核等。

记录运行日志时使用命令nvprof，可视化显示日志时使用命令nvvp。如果只想对某一段代码进行分析，在目标代码段前后加上 cudaProfilerStart() 和 cudaProfilerStop()。

# 使用 nvprof 在命令行查看分析结果
nvprof -o out.nvvp XXX
# 对于生成的 out.nvvp 文件，需要在使用 nvvp 进行查看。

以上是一些常用的CUDA代码调试的方法，希望对您有所帮助，欢迎在评论区留言。

参考文献

https://zhuanlan.zhihu.com/p/559682306