CUDA 异步报错：为什么调用栈总是不在真正出错处

Posted on In , Views: Word count in article: 6.3k Reading time ≈ 11 mins.

CUDA 报错最迷惑人的地方，不是它报错，而是它经常报在“后面那一行”。

📚 基本概念速读

名称 定义 省流
Host CPU 端代码，也就是发起 CUDA 调用的一侧。 下命令的人
Device GPU 端代码，真正执行 kernel 的一侧。 干活的人
Kernel 在 GPU 上并行执行的函数。 GPU 任务
Stream CUDA 的任务队列，同一个 stream 内通常按提交顺序执行。 GPU 队列
Synchronize CPU 停下来等 GPU 做完前面的任务。 等结果
CUDA error code CUDA runtime/driver 返回给 CPU 的错误状态。 一句错误摘要

一句话省流：

CPU 端调用栈通常表示“CPU 在哪里发现 GPU 出错”，不一定表示“GPU 真正在哪里出错”。

🧠 先建立一个心智模型

CUDA 程序里通常有两条时间线：

1
2
3
4
5
6
CPU 时间线:
  launch A  ->  launch B  ->  print  ->  synchronize  ->  报错

GPU 时间线:
                 execute A  ->  A 内部出错
                              execute B 或被错误状态影响

CPU 发起 kernel 时，很多情况下只是把任务提交给 GPU 队列，然后自己继续往下跑。

也就是说：

1
myKernel<<<grid, block>>>(args);

这行代码返回成功，通常只能说明：

kernel launch 请求成功提交了。

它不等价于：

kernel 已经在 GPU 上执行成功了。

这就是 CUDA 异步报错的根源。

🔁 CPU 和 GPU 怎么协作

一个简化版流程如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
+-------------------+
| CPU Host Thread   |
+---------+---------+
          |
          | 1. 准备 kernel 参数
          | 2. 调用 CUDA runtime/driver
          | 3. 把命令放入 stream
          v
+-------------------+
| CUDA Driver       |
+---------+---------+
          |
          | 4. 提交到 GPU command queue
          v
+-------------------+
| GPU Device        |
+---------+---------+
          |
          | 5. 按 stream 顺序执行 kernel / memcpy / event
          v
+-------------------+
| Device Memory     |
+-------------------+

关键点是：提交任务和任务执行完成不是同一件事。

普通 CPU 函数调用更像这样：

1
调用函数 -> 函数执行 -> 返回结果/异常

CUDA kernel launch 更像这样：

1
提交任务 -> 立即返回 -> GPU 稍后执行 -> CPU 之后某个点才知道结果

所以 CPU 端代码可能已经走到很后面，GPU 才刚执行到前面提交的 kernel。

🧩 Stream 是什么

可以把 CUDA stream 理解成 GPU 任务队列：

1
2
3
4
5
stream 0:
  kernel A
  kernel B
  cudaMemcpy
  kernel C

同一个 stream 里的任务通常按顺序执行。但 CPU 提交任务时不一定等待每个任务结束：

1
2
3
4
5
时间 --->

CPU:  launch A | launch B | launch C | print("done") | synchronize

GPU:             |---- A ----|---- B ----|---- C ----|

所以 print("done") 出现，不代表 GPU kernel 已经 done。

这在 PyTorch 里很常见：

1
2
3
4
logits = model(x)
loss = criterion(logits, target)
loss.backward()
print(loss.item())

前面几行可能只是不断向 GPU 提交算子。到了 loss.item()，CPU 需要把 GPU 上的标量拿回来，这时必须等待 GPU，于是错误可能在这里才爆出来。

💥 GPU 出错时发生了什么

假设某个 kernel 里有越界写：

1
2
3
4
__global__ void badKernel(float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    out[i] = 1.0f;  // 如果 i >= n，就是越界写
}

CPU 端这样调用：

1
2
3
badKernel<<<grid, block>>>(out, n);
printf("after launch\n");
cudaMemcpy(host, out, bytes, cudaMemcpyDeviceToHost);

可能的结果是：

1
2
after launch
CUDA error: an illegal memory access was encountered

报错位置看起来在：

1
cudaMemcpy(...)

但真正的事故现场可能是前面的：

1
badKernel<<<grid, block>>>(out, n);

原因是 cudaMemcpy(...DeviceToHost...) 需要等待前面的 GPU 工作完成，CPU 到这里才第一次认真问 GPU：

前面的任务都做完了吗？我要把结果拿回来了。

GPU/driver 这时返回：

前面那个 kernel 已经出错了。

📡 GPU 怎么通知 CPU

默认情况下，GPU 端错误通常不是像 CPU 异常一样立刻带着完整调用栈抛回来。

更接近下面这个过程：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
+-------------------------------+
| GPU 执行 kernel               |
+---------------+---------------+
                |
                | 出现非法访问 / assert / launch failure
                v
+-------------------------------+
| Driver/context 记录错误状态   |
+---------------+---------------+
                |
                | CPU 后续调用同步、拷贝、查询、相关 CUDA API
                v
+-------------------------------+
| CUDA runtime 返回错误码       |
+---------------+---------------+
                |
                | PyTorch/C++ runtime 包装成异常或返回值
                v
+-------------------------------+
| 打印 CPU/Python/C++ 调用栈    |
+-------------------------------+

所以你看到的经常是：

1
CPU 调用栈 + 一句 GPU 错误摘要

例如：

1
RuntimeError: CUDA error: device-side assert triggered

或者：

1
CUDA error: an illegal memory access was encountered

这类信息说明 GPU 端发生了什么类型的错误，但它默认不一定告诉你：

  • 哪个 kernel 的哪一行出错；
  • 哪个 block/thread 出错；
  • GPU 上完整的调用链是什么；
  • Python 栈中哪一行才是真正提交问题 kernel 的位置。

🧾 CPU 调用栈到底说明什么

当 PyTorch 报错栈指向：

1
2
3
4
loss.item()
tensor.cpu()
torch.cuda.synchronize()
loss.backward()

它通常表示：

CPU 在这个 API 调用处发现 CUDA 已经处于错误状态。

它不必然表示：

这个 API 自己就是根因。

更准确地说，调用栈里的位置是通知点，不一定是事故现场。

常见通知点包括：

操作 为什么会暴露错误
cudaDeviceSynchronize() 明确等待 GPU 之前的任务完成。
cudaStreamSynchronize(stream) 等某个 stream 完成。
cudaMemcpy(...DeviceToHost...) 要从 GPU 拿数据回 CPU，必须等结果可用。
tensor.cpu() PyTorch 需要把 GPU Tensor 拷回 CPU。
loss.item() 需要把 GPU 标量读到 Python 数值。
后续 CUDA API 发现当前 CUDA context 已经有错误状态。

🔍 为什么 GPU 错误默认只有一句话

CPU 异常容易打印调用栈，是因为 CPU 线程有清晰的栈帧和同步执行路径。

GPU kernel 不一样：

  • 一个 kernel 可能有成千上万个线程同时执行；
  • 每个线程都有自己的 threadIdx/blockIdx；
  • 错误可能只发生在某个 thread；
  • kernel launch 和 kernel 执行在时间线上分离；
  • 上层框架默认更关心“这个 CUDA 操作失败了”，而不是自动进入 GPU 调试模式。

所以默认报错往往是一个 CUDA error code 加一段短描述：

1
2
3
4
5
illegal memory access
device-side assert triggered
misaligned address
unspecified launch failure
an illegal instruction was encountered

这些描述很有用，但它们不是完整定位结果。

🛠️ CUDA_LAUNCH_BLOCKING=1 做了什么

调试 PyTorch CUDA 错误时，最常用的第一步是：

1
CUDA_LAUNCH_BLOCKING=1 python train.py

它的作用可以简化理解为：让很多 CUDA launch 变得更同步。

默认情况：

1
2
CPU: launch A -> launch B -> launch C -> 后面某处报错
GPU:      A 执行并出错

开启后：

1
2
CPU: launch A -> 等 A 完成 -> 如果 A 出错，就更早报错
GPU:      A 执行并出错

这不会修复 bug，但能让 Python 调用栈更接近真正提交错误 kernel 的位置。

代价是程序会明显变慢，因为 CPU/GPU 原本可以重叠执行的部分被压平了。

🧪 一个 PyTorch 里的典型例子

分类任务里，target 必须在 [0, num_classes - 1] 范围内。

例如模型输出 10 类：

1
logits.shape  # [batch_size, 10]

但 target 里出现了非法类别：

1
target.max()  # 12

这时：

1
loss = torch.nn.functional.cross_entropy(logits, target)

GPU kernel 可能触发 device-side assert。你看到的报错位置却可能在：

1
loss.backward()

或者：

1
loss.item()

真正要查的是：

1
print(target.min(), target.max(), logits.shape)

如果是 cross_entropy，优先确认：

1
2
3
0 <= target.min()
target.max() < logits.shape[1]
target.dtype == torch.long

🧰 C++/CUDA 里怎么抓得更准

自定义 CUDA kernel 调试时，可以在 kernel 后面临时加两类检查：

1
2
3
4
5
6
7
8
9
10
11
myKernel<<<grid, block>>>(args);

cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("launch error: %s\n", cudaGetErrorString(err));
}

err = cudaDeviceSynchronize();
if (err != cudaSuccess) {
    printf("runtime error: %s\n", cudaGetErrorString(err));
}

区别是：

API 主要检查什么
cudaGetLastError() kernel launch 是否成功，比如配置、参数、符号等 launch 阶段错误。
cudaDeviceSynchronize() 等 GPU 实际执行完，能暴露 kernel 执行期间的错误。

如果要进一步定位到源码行，可以用更重的工具：

1
compute-sanitizer --tool memcheck ./your_program

典型报告可能会给出：

1
2
3
Invalid __global__ write of size 4
    at kernel.cu:123
    by thread (31,0,0) in block (12,0,0)

这比默认的一句 illegal memory access 更接近真正的 GPU 事故现场。

⚠️ 常见误区

误区 正解
报错栈指向 loss.item()，所以 loss.item() 写错了。 它可能只是第一个同步点，真正错误在更早的 GPU kernel。
kernel launch 没报错，说明 kernel 执行没问题。 launch 成功只代表提交成功，不代表执行成功。
CUDA_LAUNCH_BLOCKING=1 能修复 CUDA bug。 它只是让报错更同步、更容易定位。
GPU 报错只有一句话，所以没法查。 默认信息少，但可以用同步、断言、Compute Sanitizer、cuda-gdb 等工具补充。
后续所有 CUDA 调用都报错，说明每一行都有 bug。 一个早期 GPU 错误可能污染后续 CUDA context，后面只是连续发现错误状态。

✅ 实战排查顺序

遇到 CUDA 异步报错，我一般按这个顺序做：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
1. 开启同步定位
   CUDA_LAUNCH_BLOCKING=1 python train.py

2. 看报错更靠近哪一行
   model forward / loss / backward / tensor copy

3. 查最常见数据问题
   shape / dtype / device / index range / NaN / Inf

4. 如果是分类 loss
   检查 target.min()、target.max()、target.dtype

5. 如果是自定义 CUDA kernel
   kernel 后加 cudaGetLastError() + cudaDeviceSynchronize()

6. 仍然不清楚
   compute-sanitizer --tool memcheck ...

尤其是 PyTorch 里，下面这些检查非常便宜：

1
2
3
4
print('logits:', logits.shape, logits.dtype, logits.device)
print('target:', target.shape, target.dtype, target.device)
print('target range:', target.min().item(), target.max().item())
print('finite:', torch.isfinite(logits).all().item())

先查这些，往往比盯着最后一行报错栈更有效。

✅ 总结

CUDA 异步报错里，CPU 调用栈通常是“发现错误的位置”，GPU 错误摘要才说明“发生了哪类 GPU 错误”，真正的出错 kernel 往往在更早提交的任务里。

调试时不要把报错栈当成唯一事实。先把异步变同步，再检查数据范围、shape、dtype 和自定义 kernel 的边界访问；需要更细粒度定位时，再上 Compute Sanitizer。

Happy Hacking! 🎉