调试 CUDA Kernel 并进入 __device__ 函数是 CUDA 开发中一项非常重要的技能。这主要依赖于 NVIDIA 的官方调试器 NVIDIA Nsight Systems (用于系统级分析) 和 NVIDIA Nsight Compute (用于内核级分析) 以及经典的 cuda-gdb (命令行调试器)。

这里将重点介绍两种最常用和强大的方法：使用 Nsight Visual Studio Code Edition (图形化界面，推荐) 和 cuda-gdb (命令行)。

方法一：使用 Nsight VSCode Edition (图形化界面，最推荐)

Nsight for VSCode 提供了类似于调试 CPU 代码的直观体验，是当前调试 CUDA 的首选工具。

前提条件

安装 NVIDIA 驱动和 CUDA Toolkit: 确保你的系统安装了正确版本的驱动和 CUDA。
安装 VSCode: 从 VSCode 官网下载并安装。
安装 Nsight VSCode 插件: 在 VSCode 的扩展市场中搜索 “NVIDIA Nsight” 并安装。
确保你的代码是可调试的: 在编译你的 CUDA 代码时，必须使用 -G 或 -lineinfo 标志来生成调试信息。
- -G: 生成完整的调试信息，但会严重禁用所有编译器优化，极大影响性能，仅用于调试。
- -lineinfo (--generate-line-info): 生成行号信息，允许调试和性能分析，但对性能影响较小，是 Nsight Compute 性能分析的首选。对于调试 __device__ 函数，-lineinfo 通常就足够了。
示例编译命令:
```
nvcc -g -G -o my_program my_program.cu # 使用 -g -G 进行完整调试
# 或者
nvcc --device-debug -o my_program my_program.cu # CUDA 11.2+ 推荐方式，等同于 -g -G
# 或者 (如果主要为了分析，附带一些调试能力)
nvcc -lineinfo -o my_program my_program.cu
```

调试步骤

打开项目: 在 VSCode 中打开你的 CUDA 项目文件夹。
创建调试配置文件:
- 点击侧边栏的 “运行和调试” 图标 (或 Ctrl+Shift+D)。
- 点击 “create a launch.json file”。
- 在弹出的环境选择中，选择 CUDA C++。
- 这会在项目下生成一个 .vscode/launch.json 文件。

配置 launch.json:

一个基本的配置如下。关键是指定正确的程序路径和 cuda-gdb 路径。

{"version": "0.2.0","configurations": [{"name": "CUDA C++: Launch","type": "cuda-gdb","request": "launch","program": "${workspaceFolder}/my_program", // 你的可执行文件路径"stopAtEntry": false, // 设为 true 会在 main 函数入口处暂停"cwd": "${workspaceFolder}","args": [], // 传递给程序的命令行参数"environment": [], // 环境变量"externalConsole": false}]
}

设置断点:
- 在你的 CUDA 源码 (.cu 或 .cuh 文件) 中，在你想要中断的行的左侧空白处点击。
- 你可以在 main 函数、<<<...>>> 调用的 kernel 函数、以及任何 __device__ 函数中设置断点。
开始调试:
- 选择刚刚创建的 “CUDA C++: Launch” 配置。
- 按 F5 或点击绿色的 “开始调试” 按钮。
- 程序开始运行，并在遇到你设置的断点时暂停。
步入 __device__ 函数:
- 当程序在 kernel 的某一行暂停时，你可以使用调试控制栏的按钮：
  - Step Into (F11): 如果当前行调用了某个 __device__ 函数，按 F11 会进入该 __device__ 函数的函数体。
  - Step Over (F10): 执行当前行，但不进入函数内部。
  - Step Out (Shift+F11): 执行完当前函数，返回到调用它的地方。
查看变量和调用堆栈:
- 在调试过程中，你可以在 VSCode 的左侧面板查看变量的当前值。
- 调用堆栈 面板会显示你当前的执行位置，从 main 到 kernel 再到 __device__ 函数，清晰明了。
- 你还可以将鼠标悬停在源码中的变量上来查看其值。

方法二：使用 cuda-gdb (命令行)

但其实我们更习惯命令行操作，或者在没有图形界面的远程服务器上工作，cuda-gdb 是强大的选择。

前提条件

同样，编译时必须使用 -G 或 -lineinfo 标志。

nvcc -g -G -o my_program my_program.cu

例如： nvcc -g -G --gpu-architecture=sm_120 sgemm_1.cu -o sgemm_1 -I ../cutlass/include/ -I ../cutlass/tools/util/include -I /usr/local/cuda/include

调试步骤

启动调试器:
```
cuda-gdb ./my_program
```
设置断点:
- 在 kernel 函数处设置断点：
```
(cuda-gdb) break my_kernel_function
```
- 在 __device__ 函数处设置断点：
```
(cuda-gdb) break my_device_function
```
- 在特定行设置断点：
```
(cuda-gdb) break file.cu:123
```
运行程序:
```
(cuda-gdb) run
```
程序会开始执行，并在第一个断点处停止。
控制执行和步入 __device__ 函数:
- next (n): Step Over，执行下一行。
- step (s): Step Into，进入函数。如果下一行是 __device__ 函数调用，这会进入该 __device__ 函数体。
- continue (c): 继续运行直到下一个断点。
- finish: Step Out，运行到当前函数返回。

检查线程和变量:

CUDA 调试的核心是理解线程。你可以切换当前关注的线程：

(cuda-gdb) cuda thread 1  # 切换到 block 0, thread 0
(cuda-gdb) cuda thread (1, 2, 3)  # 切换到 blockIdx(1,2), threadIdx(3)

打印变量值：
```
(cuda-gdb) print variable_name
```

退出:
```
(cuda-gdb) quit
```

重要提示和常见问题

硬件兼容性: 并非所有 NVIDIA GPU 都支持调试功能。请查阅 NVIDIA 文档确认你的 GPU 是否支持 “TCC 模式” 或 “Compute Preemption”，这对于调试至关重要。消费级显卡 (GeForce) 的调试支持可能不如专业卡 (Tesla, Quadro) 完善。
性能影响: 使用 -G 编译会极大降低 kernel 运行速度，并且调试本身也会引入开销。这是正常的，目的是为了获得精确的调试状态。
焦点线程: 在 CUDA 调试中，任何时候你都只在一个特定的 GPU 线程上下文中查看变量和执行代码。确保你正在检查你感兴趣的线程。
Nsight vs cuda-gdb: Nsight VSCode 底层调用的也是 cuda-gdb，但它提供了无比友好的图形界面，极大地简化了操作，强烈推荐初学者和绝大多数开发者使用。