使用cuda-gdb调试cu程序

简要教程

仿佛对于.cu程序的调试是需要关闭x的，所以学习了一下如何用cuda-gdb在tty下调试.cu程序.
1. 对于.cu文件是需要编译好才能调试的，首先cd到想要编译的.cu目录下，使用以下命令进行编译：
nvcc -g -G filename.cu -arch sm_50 -o outputfilename
这样就生成了一个可执行文件，用于调试。
2. 关闭x
sudo /etc/init.d/lightdm stop
3. 直接在命令行里输入cuda-gdb就可以进入cuda-gdb了
4. 使用gdb调试程序：

• 首先使用file filename来打开刚才编译好的可执行文件
• l：输入一个l然后回车可以浏览所有的源代码，看到代码对应的行数
• b line_index：给标号为line_index的行设置断点
• r：run 运行程序，并在断点停止
• n：next 下一步执行
• p param：print 打印参数param的值
• delete breakpoint line_index:删除line_index行的断点

进阶教程

在运行到断点时，使用cuda-gdb能够看到很多信息。

获取帮助

使用下面的语句可以获得相应指令的帮助

(cuda-gdb) help cuda cuda指令
(cuda-gdb) help set cuda cuda设置
(cuda-gdb) help info cuda info-cuda 指令

使用方法

# 下面的语句可以获取当前所处的设备各种信息
(cuda-gdb) cuda device sm warp lane block thread
block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0
(cuda-gdb) cuda kernel block thread
kernel 1, block (0,0,0), thread (0,0,0)
(cuda-gdb) cuda kernel
kernel 1
# 使用下面的语句可以讲关注点移动到另外的单元上去（这里的单元包括device,sm,block,warp等）
(cuda-gdb) cuda device 0 sm 1 warp 2 lane 3
[Switching focus to CUDA kernel 1, grid 2, block (8,0,0), thread
(67,0,0), device 0, sm 1, warp 2, lane 3]
374 int totalThreads = gridDim.x * blockDim.x;

获取所有的设备信息：

#　使用info cuda **的命令可以获取当前所有的单元信息，而且星花所指向的当前所处(focus on)的单元：
(cuda-gdb) info cuda sms SM
Active Warps Mask Device 0
* 0 0xffffffffffffffff
1 0xffffffffffffffff
2 0xffffffffffffffff
3 0xffffffffffffffff
4 0xffffffffffffffff
5 0xffffffffffffffff
6 0xffffffffffffffff
7 0xffffffffffffffff
8 0xffffffffffffffff
...

而且，竟然可以查看和操作寄存器：

 (cuda-gdb) info registers $R0 $R1 $R2
R0  0xf0        240
R1  0xfffc48    16776264
R2  0x7800  307

ps:
使用下面的命令可以编译出ptx代码：
nvcc -g -G --ptx print_test.cu -arch sm_50 -o print_ptx2
　
下面这个可以生成许多运行时候的信息(诸如ipc之类)
nvprof --metrics ipc ./print_test 或 nvprof --metrics all ./print_test