CUDA 工具与调试：从 Profiling 到 CUDA Graph

Q1：如何 Profile 一个 CUDA 程序？常用工具有哪些？

为什么要 Profile？

在 CUDA 程序优化中，"猜测瓶颈"几乎总是错的。先 Profile、后优化是正确顺序。Profile 能告诉你：

每个 Kernel 各耗时多少
内存带宽利用率是否饱和
SM 的 Occupancy（占用率）是否合理
CPU 与 GPU 之间是否在流水线重叠上出了问题
L1/L2 缓存命中率如何

主要工具一览

工具	用途	状态
Nsight Systems	系统级时间线（CPU+GPU 全局视角）	现役，推荐
Nsight Compute	Kernel 级性能分析（指标深挖）	现役，推荐
nvprof	命令行 Profiler（旧版）	已废弃（CUDA 11.x 后不再维护）
Visual Profiler	nvprof 的 GUI 版	已废弃
cuda-memcheck	内存错误检测	已被 compute-sanitizer 取代
compute-sanitizer	内存/竞争/初始化错误检测	现役，推荐

编译准备：为 Profiling 添加调试信息

直接 Profile 已编译的二进制也能工作，但加上以下编译选项能让工具显示源码行号：

# 为 Device 代码生成行号信息（不影响性能）
nvcc -lineinfo -o myapp myapp.cu

# 同时将源码嵌入 PTX（避免 -lineinfo 的路径依赖问题）
nvcc -lineinfo -src-in-ptx -o myapp myapp.cu

-lineinfo 不会降低性能，生产环境也可保留，方便随时 Profile。

Nsight Systems：系统级时间线

Nsight Systems 提供整个程序的时间线视图，适合回答以下问题：

“CPU 和 GPU 的工作是否重叠？”
“有没有不必要的同步导致 GPU 空闲？”
“数据拷贝是否成为瓶颈？”

命令行使用：

# 收集 Profile 数据
nsys profile --trace=cuda,nvtx -o myapp_profile ./myapp

# 生成报告（可在 Nsight Systems GUI 中打开）
nsys stats myapp_profile.nsys-rep

关键参数：

--trace=cuda：追踪 CUDA API 调用和 Kernel 执行
--trace=nvtx：追踪用户标注的 NVTX 范围（见下文）
--cuda-graph-trace=node：展示 CUDA Graph 中各节点在哪个 Green Context 执行

在代码中插入 NVTX 标注（让时间线更可读）：

#include <nvtx3/nvToolsExt.h>

nvtxRangePushA("矩阵乘法");
matmul<<<grid, block>>>(A, B, C);
nvtxRangePop();

时间线上会出现带颜色的区间，一眼看出各阶段耗时。

Nsight Compute：Kernel 级深度分析

Nsight Compute 针对单个 Kernel，收集数百项硬件性能指标，适合回答：

“这个 Kernel 是计算瓶颈还是内存瓶颈？”
“Warp Stall 主要因为什么？”
“Tensor Core 利用率是多少？”

命令行使用：

# 收集 Kernel 性能报告（默认采集全套 section）
ncu -o kernel_report ./myapp

# 只分析特定 Kernel
ncu --kernel-name myKernel -o report ./myapp

# 直接打印到终端（不保存文件）
ncu --print-summary per-kernel ./myapp

关键指标：

指标	含义	目标
SM Active Cycles	SM 实际有效工作的时钟周期比例	越高越好
Achieved Occupancy	实际 Warp 占用率	≥ 50% 通常足够
Memory Throughput	全局内存实际带宽利用率	接近硬件峰值
L2 Hit Rate	L2 缓存命中率	越高越好
Warp State: Stall (reason)	Warp 停滞原因分解	找主要原因优化
Tensor Active Cycles	Tensor Core 利用率	GEMM 应接近 80%+

旧工具：nvprof（了解即可）

nvprof 是早期的命令行 Profiler，已在 CUDA 12 中被完全移除：

1
2
3

# 仅供了解，现代 CUDA 不再支持
nvprof ./myapp
nvprof --metrics achieved_occupancy,gld_efficiency ./myapp

迁移：将 nvprof 的 --metrics 替换为 Nsight Compute 的 --metrics；将 nvprof 的系统时间线替换为 Nsight Systems。

Q2：如何使用 Nsight Compute 分析 Kernel 的瓶颈？

工作流程

1. 运行 ncu 收集报告
2. 打开 Nsight Compute GUI / 查看 CLI 输出
3. 看 Speed Of Light（SOL）总览 → 确定是计算还是内存瓶颈
4. 根据瓶颈类型，深入对应的 Section
5. 定位问题 → 修改代码 → 对比前后数据

第一步：Speed Of Light（SOL）总览

SOL Section 给出最直观的性能画像：

1 2	Compute Throughput: 42.3% ← SM 计算单元利用率 Memory Throughput: 87.6% ← 内存系统利用率

内存瓶颈（Memory >> Compute）：优化内存访问模式（合并访问、共享内存缓存、向量化读写）
计算瓶颈（Compute >> Memory）：提高 ILP（每线程更多计算、Tensor Core）
两者都低：Occupancy 太低，或者存在严重的同步/等待

第二步：Warp State Statistics（分析停滞原因）

Warp 停滞（Stall）是 GPU 效率损失的主要来源。Nsight Compute 将停滞分类：

停滞类型	含义	优化方向
`Stall Long Scoreboard`	等待全局内存返回数据	提高 Occupancy、使用共享内存、异步预取
`Stall Short Scoreboard`	等待共享内存或寄存器依赖	增加 ILP，避免连续依赖链
`Stall Barrier`	等待 `__syncthreads()`	减少不必要的同步
`Stall Not Selected`	Warp 就绪但 Warp Scheduler 未选中	正常现象，但若比例极高说明 Occupancy 过高
`Stall Branch Resolving`	等待分支判断结果	减少 Warp 发散
`Stall Membar`	等待内存屏障	减少不必要的内存栅栏

典型场景：若 Stall Long Scoreboard 占比超过 40%，说明程序是强内存延迟受限，需要通过提升 Occupancy 隐藏延迟或使用异步内存拷贝（__pipeline_memcpy_async）。

第三步：Memory Workload Analysis

Global Memory Load  Throughput: 312 GB/s  (peak: 2039 GB/s)  → 15%
Global Memory Store Throughput: 156 GB/s
L2 Hit Rate: 23.4%   ← 低，说明缓存压力大
L1/TEX Hit Rate: 67.1%

重点关注：

Global Load/Store Efficiency：若远低于峰值，检查是否内存不合并
L2 Hit Rate：若很低，考虑数据重用策略（共享内存、L2 Persistence API）
Shared Memory Bank Conflicts：若有大量 Bank Conflict，检查访问步长

第四步：Occupancy Analysis

1 2	Theoretical Occupancy: 75% ← 理论上限（受寄存器/共享内存限制） Achieved Occupancy: 71% ← 实际测量值

Nsight Compute 同时给出限制 Occupancy 的原因：

1
2
3

Limiting Factors:
  Registers per Thread: 32   (limit: 255)
  Shared Memory per Block: 16384 B  ← 这限制了每 SM 只能跑 4 个 Block

若 Occupancy 因共享内存过大而受限，可以考虑：

减少 Tile 大小
将静态共享内存改为动态分配（在运行时根据 Occupancy API 调整）

第五步：Source Counter（热点行定位）

配合 -lineinfo 编译，Nsight Compute 可以将指标映射到源码行：

Line 47:  smem[threadIdx.x][threadIdx.y] = ...
  Warp Stall (Long Scoreboard): 8.3 cycles/inst  ← 这行是主要延迟来源
Line 52:  C[row * N + col] = ...
  Memory Store: 128 bytes/transaction  ← 合并写，很好

常见瓶颈诊断速查表

SOL 状态	Warp Stall	诊断结论	优化方向
Memory 高，Compute 低	Long Scoreboard	内存延迟受限	提 Occupancy / 共享内存
Memory 高，Compute 低	Barrier	过度同步	合并 `__syncthreads`
Compute 高，Memory 低	Short Scoreboard	计算延迟受限	ILP、更多 ILP 循环
两者均低	Not Selected	Occupancy 问题	调整 Block 大小/寄存器
两者均低	Branch	Warp 发散	消除 if-else 发散

Q3：CUDA 中常见的错误有哪些？如何调试？

一、CUDA 错误检查机制

所有 CUDA Runtime API 都返回 cudaError_t，但大量示例代码为了简洁而省略了错误检查，这在生产代码中是绝对不允许的。

推荐的错误检查宏：

#define CUDA_CHECK(expr) do {                                           \
    cudaError_t result = (expr);                                        \
    if (result != cudaSuccess) {                                        \
        fprintf(stderr, "CUDA Error: %s:%d: %s\n",                    \
                __FILE__, __LINE__, cudaGetErrorString(result));        \
        exit(1);                                                        \
    }                                                                   \
} while(0)

// 使用示例
CUDA_CHECK(cudaMalloc(&d_ptr, size));
CUDA_CHECK(cudaMemcpy(d_ptr, h_ptr, size, cudaMemcpyHostToDevice));

Kernel 启动不返回错误值，需要在启动后显式检查：

1
2
3

myKernel<<<grid, block>>>(args...);
CUDA_CHECK(cudaGetLastError());      // 检查 Launch 参数错误
CUDA_CHECK(cudaDeviceSynchronize()); // 等待执行完毕，捕获运行时错误

cudaGetLastError() 会清空错误状态；cudaPeekLastError() 查看但不清空。

二、异步错误的陷阱

CUDA 大量操作是异步的，错误往往不在发生的那一行被捕获，而是在下一次检查时才冒出来：

// 错误代码：内核崩溃，但下面几行看起来"成功"
badKernel<<<1, 1>>>();
// cudaGetLastError() 在此处可能返回 cudaSuccess（崩溃还没被检测到）

// 第二个 API 调用返回了前一个内核的错误
cudaError_t e = cudaMemcpy(...);  // 这里可能返回前一个内核的错误！

最佳实践：

// 每个 Kernel 后加双重检查
myKernel<<<grid, block>>>(args...);
CUDA_CHECK(cudaGetLastError());         // 捕获 Launch 配置错误
CUDA_CHECK(cudaDeviceSynchronize());    // 等待 + 捕获执行错误

调试时可以设置环境变量，让 CUDA 将错误信息写入文件：

1
2
3

CUDA_LOG_FILE=cuda_errors.log ./myapp
# 或直接打到终端
CUDA_LOG_FILE=stderr ./myapp

日志会包含比 API 返回值更详细的错误原因，例如具体是哪个 cuLaunchKernel 失败及其参数。

三、常见错误类型与排查

1. 越界访问（Out-of-Bounds）

// 典型错误：忘记边界检查
__global__ void badKernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    data[idx] = 1.0f;  // 若 gridDim 设置过大，idx >= n 时越界
}

// 正确：加边界检查
__global__ void goodKernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) data[idx] = 1.0f;
}

症状：cudaErrorIllegalAddress，程序崩溃或结果错误。

2. 线程配置错误

// 错误：Block 大小超过 1024
myKernel<<<1, 2048>>>();
// cudaGetLastError() → cudaErrorInvalidValue: "invalid argument"

// 常见 Block 大小：128 / 256 / 512，不超过 1024

CUDA 的 Block 最大线程数为 1024（三维乘积），Grid 最大维度为 (2^31-1, 65535, 65535)。

3. 共享内存竞态（Race Condition）

// 错误：缺少 __syncthreads()
__global__ void raceKernel(float* data) {
    __shared__ float smem[256];
    smem[threadIdx.x] = data[threadIdx.x];
    // 缺少 __syncthreads()！其他线程可能还没写完
    float val = smem[(threadIdx.x + 1) % 256];  // 读到未定义数据
}

症状：结果非确定性，时对时错，换机器/重新编译结果变化。

4. 死锁（`__syncthreads` 在分支内）

// 危险：不同线程可能执行不同的 __syncthreads 路径
__global__ void deadlockKernel(int* flag) {
    if (threadIdx.x < 16) {
        __syncthreads();  // 只有前 16 个线程到达这里
    } else {
        __syncthreads();  // 后 16 个线程到达另一个 __syncthreads
    }
    // 两者永远无法匹配 → 死锁
}

// 正确：__syncthreads() 必须所有线程无条件都执行到

5. 主机与设备指针混用

float* h_data = new float[N];
float* d_data;
cudaMalloc(&d_data, N * sizeof(float));

// 错误：在 Host 上读 Device 指针
float val = d_data[0];    // 未定义行为，通常段错误

// 错误：在 Kernel 中写 Host 指针
myKernel<<<1, 1>>>(h_data);  // Kernel 内访问 Host 内存 → 越界

症状：cudaErrorInvalidDevicePointer 或直接 Segfault。

6. Kernel 未定义符号

1
2
3

// 编译多文件项目时，若 Device 函数没有正确链接
// 错误信息：undefined reference to `__cudaRegisterLinkedBinary_...`
// 解决：确保 nvcc 编译所有包含 __device__ 的文件，或使用 -rdc=true（relocatable device code）

四、调试工具链

问题类型	推荐工具
内存越界/未初始化读	`compute-sanitizer --tool memcheck`
竞态条件	`compute-sanitizer --tool racecheck`
非法指令/崩溃	`cuda-gdb`
数值错误（NaN/Inf）	`compute-sanitizer --tool initcheck` + 检查 overflow
性能问题	`ncu` + `nsys`

Q4：如何用 compute-sanitizer 检测内存错误？

compute-sanitizer 简介

compute-sanitizer 是 CUDA 官方的运行时错误检测工具，取代了旧版的 cuda-memcheck。它能在 不修改源码 的情况下，以慢速运行程序并检测各类错误。

# 旧版（已废弃，不再维护）
cuda-memcheck ./myapp

# 新版（推荐）
compute-sanitizer ./myapp
compute-sanitizer --tool memcheck ./myapp  # 默认即 memcheck

四种检测工具

1. memcheck：内存访问错误

检测越界访问、访问已释放内存、Misaligned 访问等：

1	compute-sanitizer --tool memcheck ./myapp

示例输出：

========= Invalid __global__ write of size 4 bytes
=========     at 0x000001e0 in badKernel(float*, int)
=========     by thread (17,0,0) in block (0,0,0)
=========     Address 0x7f8b40004004 is out of bounds
=========     and is 4 bytes after the allocated region of size 4096 bytes

信息包含：操作类型（read/write）、大小、Kernel 名、线程 ID（方便定位）、越界偏移量。

配合 -lineinfo 编译后，还会显示源码行号：

1
2
3

nvcc -lineinfo -o myapp myapp.cu
compute-sanitizer --tool memcheck ./myapp
# 输出中会包含类似：main.cu:47

常见检测场景：

// 场景 1：越界写
__global__ void test(float* a, int n) {
    a[blockDim.x * blockIdx.x + threadIdx.x] = 1.0f;  // Grid 设置过大时越界
}

// 场景 2：访问已释放内存
float* d_ptr;
cudaMalloc(&d_ptr, 100 * sizeof(float));
cudaFree(d_ptr);
myKernel<<<1, 1>>>(d_ptr);  // 使用悬空指针 → memcheck 报告

// 场景 3：未对齐访问（float4 需要 16 字节对齐）
float4* ptr = (float4*)(base + 1);  // 地址偏移 1 字节，未对齐

2. racecheck：共享内存竞态

检测共享内存中的读写竞争（缺少 __syncthreads()）：

1	compute-sanitizer --tool racecheck ./myapp

示例输出：

========= ERROR: Race condition at shared 0x0000000000000004
=========     Write Thread (1,0,0) in block (0,0,0) wrote 4 bytes
=========     Read  Thread (0,0,0) in block (0,0,0) read  4 bytes
=========     Both accesses occurred without synchronization

3. initcheck：未初始化内存读取

检测读取未初始化的设备全局内存（比 memcheck 更严格）：

1	compute-sanitizer --tool initcheck ./myapp

1 2	========= Uninitialized __global__ read of size 4 bytes ========= at main.cu:23 in myKernel

这类错误在结果验证中很难发现（数据恰好"看起来对"），initcheck 能提前暴露。

4. synccheck：非法同步

检测 __syncthreads() 的非法使用（如在发散分支中）：

1	compute-sanitizer --tool synccheck ./myapp

1 2	========= ERROR: Barrier error detected. Divergent thread(s) in block ========= at main.cu:56 in badKernel

实用技巧

只检测特定 Kernel：

1	compute-sanitizer --kernel-name myKernel ./myapp

组合多种检测（但会更慢）：

1	compute-sanitizer --tool memcheck --tool racecheck ./myapp

性能影响：compute-sanitizer 会使程序变慢 10x ~ 100x，仅在调试阶段使用。

与 ASAN 配合（同时检测 CPU 侧内存错误）：

1
2
3

# 编译时加 AddressSanitizer
nvcc -Xcompiler "-fsanitize=address" -o myapp myapp.cu
compute-sanitizer ./myapp

Q5：什么是 CUDA Graph？它能带来哪些性能提升？

为什么需要 CUDA Graph？

每次调用 kernel<<<grid, block, smem, stream>>>(args...) 时，CUDA Runtime 在幕后执行一系列工作：

验证 Kernel 参数
设置 GPU 寄存器（Push 执行配置）
将 Kernel 加入 Stream 队列
驱动层将命令打包发送给 GPU

这些 CPU 侧开销（Launch Overhead） 通常在 5 ~ 20 μs 每次。对于执行时间本身只有几十微秒的短 Kernel，启动开销可能占总时间的 50% 以上。

当同样的一批 Kernel 需要反复执行（如推理服务的每次请求、训练的每个迭代），重复支付这些 CPU 开销就非常浪费。CUDA Graph 的核心思想：把操作图定义一次、执行多次。

CUDA Graph 的工作原理

CUDA Graph 将工作分为三个阶段：

1
2
3

① 定义（Define）：一次性描述操作图（节点 + 依赖关系）
② 实例化（Instantiate）：验证图、预构建运行时结构、最小化后续启动开销
③ 执行（Launch）：反复高效执行实例化后的图

传统 Stream 方式（每次迭代）：
  [CPU 开销] Kernel A → [CPU 开销] Kernel B → [CPU 开销] Kernel C → ...

CUDA Graph 方式：
  首次：[较高 CPU 开销] 定义 + 实例化
  后续：[极低 CPU 开销] 仅一次图启动 → GPU 自主执行 A, B, C, ...

图的节点类型

CUDA Graph 的操作被表示为图中的节点（Node），支持的节点类型包括：

Kernel 节点（最常用）
cudaMemcpy 节点（设备间/主机设备拷贝）
cudaMemset 节点
CPU 函数调用节点（cudaLaunchHostFunc）
等待/记录 CUDA Event 节点
子图节点（嵌套图）
条件节点（if/while 逻辑，CUDA 12+）

节点之间的依赖边定义执行顺序，无依赖关系的节点可以并行执行。

创建方式一：Stream Capture（推荐）

Stream Capture 是最简单的图创建方式，几乎不需要修改现有代码：

cudaGraph_t    graph;
cudaGraphExec_t graphExec;
bool graphCreated = false;

for (int step = 0; step < NSTEP; step++) {
    if (!graphCreated) {
        // ===== 开始捕获 =====
        cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

        // 把原来的操作照搬进来
        cudaMemcpyAsync(d_input, h_input, size, cudaMemcpyHostToDevice, stream);
        preprocessKernel<<<grid, block, 0, stream>>>(d_input, d_temp);
        mainKernel<<<grid, block, 0, stream>>>(d_temp, d_output);
        postprocessKernel<<<grid, block, 0, stream>>>(d_output);

        // ===== 结束捕获，得到图对象 =====
        cudaStreamEndCapture(stream, &graph);

        // ===== 实例化图 =====
        cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);
        cudaGraphDestroy(graph);  // 图模板已不再需要

        graphCreated = true;
    }

    // ===== 每次迭代：只需这一行启动整个操作链 =====
    cudaGraphLaunch(graphExec, stream);
    cudaStreamSynchronize(stream);
}

cudaGraphExecDestroy(graphExec);

捕获期间，提交到 stream 的操作不会立即执行，而是被记录进内部图结构。

注意：以下操作不能在捕获期间使用：

cudaMemcpy()（同步版本，因为它会同步 Legacy Stream）
任何会隐式同步的 API
对 Legacy Stream（NULL stream）的任何操作

创建方式二：Graph API（显式构建）

适合图结构复杂或需要精细控制的场景：

cudaGraph_t graph;
cudaGraphCreate(&graph, 0);

// 添加节点和依赖
cudaGraphNode_t memcpyNode, kernelNode;
cudaMemcpy3DParms memParams = {0};
// ... 填充 memParams ...
cudaGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &memParams);

cudaKernelNodeParams kParams = {0};
kParams.func = (void*)myKernel;
kParams.gridDim = dim3(gridX);
kParams.blockDim = dim3(blockX);
// ...
// kernelNode 依赖于 memcpyNode
cudaGraphAddKernelNode(&kernelNode, graph, &memcpyNode, 1, &kParams);

// 实例化并执行
cudaGraphExec_t graphExec;
cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);
cudaGraphLaunch(graphExec, stream);

处理参数变化：Graph Update

当每次迭代图结构不变、但参数变化（如输入指针、数组大小），无需重新实例化，只需更新：

cudaGraphExec_t graphExec = NULL;

for (int i = 0; i < ITERATIONS; i++) {
    cudaGraph_t graph;
    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
    do_cuda_work(stream);  // 每次可能有不同参数
    cudaStreamEndCapture(stream, &graph);

    if (graphExec != NULL) {
        // 尝试用新图更新已有实例（成本远低于重新实例化）
        cudaGraphExecUpdateResultInfo updateResult;
        cudaGraphExecUpdate(graphExec, graph, &updateResult);

        if (updateResult.result != cudaGraphExecUpdateSuccess) {
            // 图结构变化，必须重新实例化
            cudaGraphExecDestroy(graphExec);
            graphExec = NULL;
        }
    }

    if (graphExec == NULL) {
        cudaGraphInstantiate(&graphExec, graph, 0);
    }

    cudaGraphDestroy(graph);
    cudaGraphLaunch(graphExec, stream);
    cudaStreamSynchronize(stream);
}

跨流依赖的捕获（DAG 图）

Stream Capture 支持通过 Event 表达跨流依赖，创建真正的 DAG：

cudaStreamBeginCapture(stream1, cudaStreamCaptureModeGlobal);

kernel_A<<<..., stream1>>>();

// Fork：stream2 和 stream3 依赖 stream1 的进度
cudaEventRecord(event1, stream1);
cudaStreamWaitEvent(stream2, event1, 0);
cudaStreamWaitEvent(stream3, event1, 0);

kernel_B<<<..., stream1>>>();  // 与 C、D 并行
kernel_C<<<..., stream2>>>();
kernel_D<<<..., stream3>>>();

// Join：stream1 等待 stream2 和 stream3 完成
cudaEventRecord(event2, stream2);
cudaEventRecord(event3, stream3);
cudaStreamWaitEvent(stream1, event2, 0);
cudaStreamWaitEvent(stream1, event3, 0);

kernel_E<<<..., stream1>>>();

cudaStreamEndCapture(stream1, &graph);
// 结果是一个 A→(B,C,D)→E 的 DAG 图

性能提升实测数据

CUDA Graph 的收益在以下场景中最显著：

场景	传统方式	CUDA Graph	加速比
100 个短 Kernel（各 10μs）	~2ms（含 Launch）	~1ms	~2x
推理服务（小 batch）	高延迟	显著降低延迟	2x ~ 5x
多 Stream 复杂依赖	手动 Event 同步	图自动调度	代码简化
相同图执行 1000+ 次	重复 CPU 开销	只付一次实例化	随次数增加收益增大

适用场景与限制

适合使用 CUDA Graph 的场景：

推理服务（每个请求执行相同操作图）
训练迭代（每步操作图相同，只有数据变化）
任何需要重复执行相同操作序列的场景
Kernel 执行时间短（<100μs），Launch 开销占比高

不适合或注意事项：

图的操作数量太少（单个 Kernel）：收益微弱
每次迭代图结构需要大幅改变：重新实例化代价大
捕获期间不能使用同步 API（需要将 cudaMemcpy 替换为 cudaMemcpyAsync）
分析 CUDA Graph 时，Nsight Systems 需要指定 --cuda-graph-trace=node 才能看到各节点的实际执行 Green Context

与 Persistent Kernel 的对比

另一种减少 Launch 开销的方式是持久 Kernel（Persistent Kernel）：让一个 Kernel 一直运行，通过全局内存中的信号量等待新工作。

方式	优点	缺点
CUDA Graph	代码简单，无需修改 Kernel	不支持动态决策
Persistent Kernel	支持动态工作分发	代码复杂，调试难

大多数情况下，CUDA Graph 是更推荐的方式。

总结

工具/技术	主要用途	何时使用
Nsight Systems	全局时间线，CPU/GPU 重叠分析	系统级性能诊断
Nsight Compute	Kernel 指标深挖（SOL/Warp Stall）	确定单个 Kernel 瓶颈
compute-sanitizer memcheck	内存越界/悬空指针	所有调试阶段
compute-sanitizer racecheck	共享内存竞态	怀疑结果非确定性时
compute-sanitizer initcheck	未初始化内存读取	结果莫名其妙时
CUDA_CHECK 宏	API 返回值检查	所有生产代码
CUDA_LOG_FILE	驱动层详细错误日志	追踪异步错误
CUDA Graph	减少重复 Kernel 的 Launch 开销	推理/训练循环