使用cuda-gdb调试cu程序

https://blog.csdn.net/FishSeeker/article/details/74178318

简要教程
仿佛对于.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 指令
1
2
3
使用方法
# 下面的语句可以获取当前所处的设备各种信息
(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;
1
2
3
4
5
6
7
8
9
10
11
12
获取所有的设备信息：

#　使用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
...
1
2
3
4
5
6
7
8
9
10
11
12
13
而且，竟然可以查看和操作寄存器：

 (cuda-gdb) info registers $R0 $R1 $R2
R0  0xf0        240
R1  0xfffc48    16776264
R2  0x7800  307
1
2
3
4
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