CUDA debug方法和一些问题总结（越界访问，debug与release结果不一致问题等）

yangpc_615

已于 2023-12-25 16:39:41 修改

阅读量2.1k

点赞数 2

分类专栏： C++/C CUDA 文章标签：开发语言深度学习

于 2022-12-02 17:48:52 首次发布

本文链接：https://blog.youkuaiyun.com/u014267567/article/details/128149249

版权

C++/C 同时被 2 个专栏收录

23 篇文章

订阅专栏

CUDA

1 篇文章

订阅专栏

现象

bt的堆栈信息

standard io上的错误输出

从报错信息上看是非法的内存访问，但是报错的位置不一定是真实的位置，因为GPU都是异步发起的，错误可能会被在后面的op捕捉。例如cudaEventDestory：

debug方式

思维方式

复现，解决问题一定要复现问题，不能复现的问题不能确定正真解决。所以首先要做的是复现。
定位，定位范围是逐渐缩小，优先排查自定义的代码。对于cuda-kernel，默认都是异步的操作，很多时候捕获错误的位置在抛出错误位置的后面。这种情况可以在可能出错的op后面加上下面代码：

 CUDA_RETURN_IF_ERROR(cudaStreamSynchronize(stream));

cuda-memcheck

命令行，使用cuda-memcheck的工具memcheck，当然它还有另外的三个工具，详情看文档。

cuda-memcheck --tool memcheck [application application-option]

结果返回发生越界的thread号和block号。

然后通过block和thread的坐标信息，借助cuda-gdb工具，定位到具体的位置。

cuda-gdb

在编译时，给nvcc加上-g -G标识(cmake中是CMAKE_CUDA_FLAGS)，就可以得到device的debug符号信息，其中的-G是标识开启device的debug编译。

set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -g -G")

Debug的操作指令：

# switch to the forced thread and block
cuda kernel 0 block (3,103,0) thread(0,0,0)

# condition breakpoint
b file:498 if blockIdx.y==103 && blockIdx.x==3

# set conditions for an existing breakpoint
condition 1 blockIdx.x == 0 && n > 3

找到了具体越界的位置。

cuda-gdb中打印数据

device上的地址或者指针前加上（@global type*），其他指令相同的方式使用。

x/249dg ((@global long *)0x7ff7ef111400)

在shared memory上打印数据的方法。

x/2048fh (@generic half * const) 0x7fff55004000

compute-sanitizer

对于debug和release版本结果不一致的问题，debug无法复现问题，这个时候需要借助一些工具，compute-sanitizer是一个非常有用的工具。其中有含有四个子工具，具体见文档。

debug与release code的差异：

1. warp的schedule方式不同，会有一些同步的差异；

2. 资源量不同，所以release版本可能有buffer溢出的问题；

3. 速度差异，release会更快，结合了更多的优化，这些优化会带来一些问题，同时更快的运行速度结合1.中的差异，也会带来脏读，脏写的问题。

结合具体问题使用下面code来排查：

compute-sanitizer --tool racecheck ./test

printf定位具体位置

对于代码比较简单，但是并发度很高的cuda-kernel-func，采用这个方式比较高效。但是需要注意：

ort里的cuda代码默认是开启了O3的优化等级的。如果注释掉一些代码，相关的一些代码可能被优化掉。解决方法：把删除的代码中的变量，用printf加上，可以屏蔽这个优化，当然也可以开启nvcc的debug模式。
cuda-kernel里的printf要注意数据的字节数，例如，int64的类型，在打印的时候需要加l（%ld），不然会引起错位，该字符串后面的输出，都会出现错误。