CUDA-GDB 入门与原理剖析

1
(cuda-gdb) break main
2
Breakpoint 1 at 0x18e1: file bitreverse.cu, line 26.
3
(cuda-gdb) break bitreverse
4
Breakpoint 2 at 0x18a1: file bitreverse.cu, line 8.
5
(cuda-gdb) break 23
6
Breakpoint 3 at 0x18ac: file bitreverse.cu, line 23.

1
(cuda-gdb) run
2
Starting program: /path/to/bitreverse
3
[Thread debugging using libthread_db enabled]
4
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
5

6
Breakpoint 1, main () at /path/to/bitreverse.cu:26
7
24      int main(void) {

1
(cuda-gdb) continue
2
Continuing.
3
[New Thread 0x7ffff1dff000 (LWP 3111508)]
4
[New Thread 0x7ffff09ff000 (LWP 3111509)]
5
[Detaching after fork from child process 3111510]
6
[New Thread 0x7fffe992a000 (LWP 3111532)]
7
[New Thread 0x7fffe9129000 (LWP 3111533)]
8
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
9

10
CUDA thread hit Breakpoint 2.2, bitreverse<<<(1,1,1),(256,1,1)>>> (data=0x7fffd1e00000) at bitreverse.cu:12
11
12        array[threadIdx.x] = idata[threadIdx.x];

1
(cuda-gdb) info cuda threads
2
  BlockIdx ThreadIdx To BlockIdx To ThreadIdx Count                 PC      Filename  Line
3
Kernel 0
4
*  (0,0,0)   (0,0,0)     (0,0,0)    (255,0,0)   256 0x00007fffd72738b0 bitreverse.cu    12

1
(cuda-gdb) print blockIdx
2
$1 = {x = 0, y = 0, z = 0}
3
(cuda-gdb) print threadIdx
4
$2 = {x = 0, y = 0, z = 0}

1
(cuda-gdb) print gridDim
2
$3 = {x = 1, y = 1, z = 1}
3
(cuda-gdb) print blockDim
4
$4 = {x = 256, y = 1, z = 1}

1
(cuda-gdb) next
2
12       array[threadIdx.x] = idata[threadIdx.x];
3
(cuda-gdb) next
4
14       array[threadIdx.x] = ((0xf0f0f0f0 & array[threadIdx.x]) >> 4) |
5
(cuda-gdb) next
6
16       array[threadIdx.x] = ((0xcccccccc & array[threadIdx.x]) >> 2) |
7
(cuda-gdb) next
8
18       array[threadIdx.x] = ((0xaaaaaaaa & array[threadIdx.x]) >> 1) |
9
(cuda-gdb) next
10

11
Breakpoint 3, bitreverse <<<(1,1),(256,1,1)>>> (data=0x100000) at bitreverse.cu:23
12
21             idata[threadIdx.x] = array[threadIdx.x];
13
(cuda-gdb) print array[0]@12
14
$7 = {0, 128, 64, 192, 32, 160, 96, 224, 16, 144, 80, 208}

1
(cuda-gdb) print/x array[0]@12
2
$8 = {0x0, 0x80, 0x40, 0xc0, 0x20, 0xa0, 0x60, 0xe0, 0x10, 0x90, 0x50,
3
0xd0}
4

5
(cuda-gdb) print &data
6
$9 = (@global void * @parameter *) 0x10
7
(cuda-gdb) print *(@global void * @parameter *) 0x10
8
$10 = (@global void * @parameter) 0x100000

1
(cuda-gdb) cuda thread 170
2
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread
3
(170,0,0), device 0, sm 0, warp 5, lane 10]

1
(cuda-gdb) delete breakpoints
2
Delete all breakpoints? (y or n) y
3
(cuda-gdb) continue
4
Continuing.
5

6
Program exited normally.
7
(cuda-gdb)

1
while (1) {
2
  // 非阻塞 waitpid 检查事件
3
  pid = waitpid(-1, &status, WNOHANG | __WALL);
4

5
  if (pid > 0) {
6
    // 有事件发生
7
    handle_event(pid, status);
8
  } else {
9
    // 无事件, 等待 SIGCHLD
10
    sigsuspend(&blocked_mask);
11
  }
12
}

1
// PTRACE_SINGLESTEP 让内核设置 TF (Trap Flag)
2
ptrace(PTRACE_SINGLESTEP, pid, 0, 0);
3
// CPU 执行一条指令后，因 TF 触发 #DB 异常
4
// 内核将进程停止，GDB 通过 waitpid 检测到

1
// 在下一条指令地址插入临时断点
2
insert_single_step_breakpoint(gdbarch, aspace, next_pc);
3
// 然后用 PTRACE_CONT 执行
4
ptrace(PTRACE_CONT, pid, 0, 0);

1
struct lwp_info {
2
  ptid_t ptid;              // 线程 ID
3
  int stopped;              // 是否停止
4
  int resumed;              // 是否已恢复
5
  int step;                 // 是否单步
6
  int last_resume_kind;     // 最后恢复类型
7
  target_waitstatus status; // 等待状态
8
  CORE_ADDR stop_pc;        // 停止时的 PC
9
  };

1
// 发送 SIGSTOP 给所有线程
2
for (each LWP) {
3
  if (!lp->stopped) {
4
    ptrace(PTRACE_CONT, pid, 0, SIGSTOP);
5
    lp->stop_requested = true;
6
  }
7
}
8
// 等待所有线程停止
9
while (threads_not_stopped) {
10
  waitpid(-1, &status, __WALL);
11
}