CUDA进阶第一篇：CUDA调试

0. 写在前面
   ==

“初学CUDA，好不容易自己写完一段cuda代码，一运行，满屏的语法bug，语法bug还好说，竟然还有逻辑bug，逻辑bug怎么改啊，wtf！！”

“从别人手里接到一段CUDA代码，WTF，为什么还有bug！！还没有注释！！没有文档！！写代码的人怎么不去死啊！！”

同事的代码出bug了，找你调bug，内心独白：“tmd这写的都是啥”，“tmd这不是我写的代码，出bug为啥要找我”，“tmd这神马破代码，代码注释都没有”

由于CUDA调试工具的不完善、CUDA调试工具上手难度较高，并行思想本身就难调试等因素，CUDA调试一直都是一件很蛋疼的事情。写CUDA也有三四年了，前段时间在群里见别人问CUDA调试的问题，突然有想法写个CUDA调试的博客。自己经验尚浅，希望各位大大看过后能够在评论里指点一二，共同完善这篇博客。
本博客只针对逻辑bug。

1 定位bug

出现bug的第一想法自然是定位bug。cuda比较奇特的地方在于，有时报错bug在500行，但500行出的代码没有错误，而是在1000行的地方逻辑错了，十分头疼。
下面介绍三种我总结的定位bug方法：

1.1 二分法

一半一半的注释代码，定位bug。比较笨拙和麻烦，但是十分好用。

1.2 输出定位法

将整体代码分为几个模块，正常的CUDA代码大概可以分为数据初始化，内存申请，内存拷贝，核函数执行，结果拷贝等模块。在每个模块结束后输出标志，示例如图1。这样在调试时就可以根据输出快速定位bug大约在什么位置。如下图：

1.3 调试工具

cuda-gdb
对于部分bug，可以用调试工具更快速的定位。
在linux下，对于API调用出错等问题，cuda gdb可以直接定位在崩溃那一行。
win下是Nsight，我不熟悉nsight，求大神补充。

cuda-memcheck
而对于核函数内访存类bug，有时候明明越界了，但是运行的时候却没有报错，造成结果结果，而cuda-memcheck可以直接定位这类bug。
建议在写完代码后，无论有没有bug，先用cuda-memcheck跑一遍。

2 解决bug

比较简单的bug，定位后基本看一眼就能解决。但对于复杂的bug，还是比较费劲的。

2.1 调试工具

单步调试，打断点。无论是cuda gdb还是Nsight，都可以定位到某一个线程上进行调试，可以说是非常强大。cuda gdb和nsight都有英文官方文档，建议大家都学一学，熟练后调试事半功倍。
但因为大量线程随机并行执行，有时并不知道该定位到哪个线程上；线程调试不容易控制；定位到单线程调试比较费劲，费时间；教程少（虽然有官方文档），上手难度较大一些。这些都是CUDA调试工具没有被广泛接受的原因。

2.2 缩小数据量或线程数并在核函数中打印

大量线程并行是导致CUDA调试难度大的最大原因，尽量的减少并行量是一个非常好的降低调试难度的办法。“小并行”甚至“串行”能够大大方便调试。

在Fermi以后的架构中，可以在核函数中使用printf。
在合理范围内缩小数据量，进而减少线程数，比如输入图像大小改为16*16。或者修改线程为<<<1,1>>>，printf打印，看是否与预期结果相同。
祝：如果核函数运行不打印输出结果，加一个设备同步函数cudaDeviceSynchronize()即可。

3 预防bug

每一个写CUDAer大概都有花几个小时甚至几天调一个bug的经历。既然bug这么难调，那么预防bug就显得尤其重要了。

3.1 写代码前一定要完全构思好架构

社会快速发展，人的心也变得着急了。写CUDA代码之前，一定要沉得住气，多花点时间在纸上构思代码，将代码模块化，哪里容易出问题，哪里该写输出，哪里该检查。

3.2 函数返回结果检查

函数返回结果检查能够非常好的定位bug，是基本的编程意识。
虽然代码可能看起来会比较冗余，示例如下，也可以参考cuda sample里的代码。

3.3 函数输入检查

在调用比较重要的函数时，建议用assert检查输入参数与预期值是否相同。

3.4 核函数内检查

举个例子，遇到的情况是在拷贝shred memory时，拷贝逻辑比较复杂。此时可以写一个检查函数，以保证拷贝的正确性。
核函数内代码如下：

3.5 练习肉眼看代码

2年后重新完善这篇博客，想了想还是加一条建议：练习肉眼看代码的能力。cuda调试是一件非常耗时间和精力的事情，所以在出错调试之前，不妨先肉眼看几遍代码，在脑子里过几遍，很有可能就能发现代码的bug，并且随着经验的积累，该办法的效率越高。

写在后面

OpenCUDA，CUDA图像算法开源项目，算法内都有详细的注释，大家一起学习。

---