CUDA Core Dump：调试内存访问问题及其他问题的有效工具
TL;DR：如果你遇到了 an illegal memory access was encountered 错误，可以启用 CUDA 核心转储来调试该问题。只需设置以下环境变量并重新运行程序以收集核心转储文件，随后即可使用 cuda-gdb 进行调试。
CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 \
CUDA_COREDUMP_SHOW_PROGRESS=1 \
CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory' \
CUDA_COREDUMP_FILE="/tmp/cuda_coredump_%h.%p.%t"简介
在开发 CUDA 内核时，你是否经常遇到非法内存访问（简称 IMA）问题却束手无策？我们在开发高性能 LLM 推理引擎 vLLM 时，也曾反复经历这种痛苦。
如果你也面临过此类问题，这篇博客正是为你准备的！我们将揭示一些高级调试技术，帮助用户在 vLLM 中排查复杂的 IMA 等问题。
例如，以下是来自 PyTorch 的一个错误：
RuntimeError: CUDA error: an illegal memory access was encountered
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.
For debugging consider passing CUDA_LAUNCH_BLOCKING=1
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.这里最具挑战性的点在于：CUDA 内核错误可能是异步报告的，通常在后续的其他 API 调用中才会触发，因此下方的堆栈跟踪信息往往不准确。根据我们的经验，这类异常的 Python 堆栈跟踪基本上永远是不准确的，且毫无价值。为了解决这个问题，错误消息建议在运行代码时添加 CUDA_LAUNCH_BLOCKING=1。然而，依然存在两个问题：
- 许多人使用
kernel<<<>>>语法启动 CUDA 内核时没有添加错误检查（例如这段 代码）。在这种情况下，即使设置了CUDA_LAUNCH_BLOCKING=1，也无法定位出错的内核。 - 如果非法内存访问发生在 CUDA 图内的内核中，即使启用了
CUDA_LAUNCH_BLOCKING=1，我们也只能在启动 CUDA 图时发现存在问题，但依然无法定位具体的出错内核。
为了精确查明此类问题，我们需要在非法内存访问发生的瞬间做出反应。当然，这不仅仅是用户可以做到的——必须得到 CUDA 驱动程序本身的支持。
CUDA 核心转储功能正是为此设计。它允许 CUDA 驱动程序在发生非法内存访问时转储 GPU 状态，以便用户后续分析 GPU 状态，从而找出哪个内核引发了错误以及非法内存访问的具体内容。
什么是核心转储（Core Dump）？
GPU 本质上是一个大规模并行处理器，它的许多概念在 CPU 中都能找到对应点。
核心转储（Core dump）是由 CPU 和操作系统共同提供的功能。当程序在执行过程中崩溃时，操作系统可以记录程序的内存数据、运行时状态及其他信息，以供后续分析和调试。程序崩溃是一个硬件层面的概念。当 CPU 在执行某些指令时遇到错误，它会进入 trap（陷阱）状态。此时，操作系统接管程序并执行相应的异常处理程序（默认情况下会终止程序，但可以通过配置生成核心转储以便分析。例如，ulimit -c 1 可以启用核心转储，echo "core.%e.%p" > /proc/sys/kernel/core_pattern 可以指定核心转储文件的路径）。
类比之下，GPU 上的核心转储功能需要 GPU 硬件与 GPU 驱动程序的协作。当 GPU 上的线程在执行过程中崩溃时，GPU 硬件需要触发异常并将其传递给 GPU 驱动程序，由驱动程序立即处理。然而，根据 论坛讨论，GPU 驱动程序在处理异常时的默认行为是将当前的 CUDA 上下文标记为不可用，而不是直接终止程序。
如何启用 CUDA 核心转储
启用 CUDA 核心转储非常简单；只需设置 CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 环境变量。但为了获得更好的体验，建议同时设置一些额外的环境变量。
- 默认情况下，CUDA 核心转储会将文件保存在当前目录且不打印文件路径。你可以启用
CUDA_COREDUMP_SHOW_PROGRESS=1环境变量来显示转储过程的进度和详细信息。最重要的是，它会在过程完成后显示核心转储文件的路径，从而方便后续的调试和分析。 - 许多任务运行在容器中，一旦任务失败，容器就会被销毁，导致无法保留核心转储文件。在这种情况下，你可以使用
CUDA_COREDUMP_FILE环境变量来指定核心转储文件的路径模板。例如，你可以将文件存储在持久化存储目录中：CUDA_COREDUMP_FILE="/persistent_dir/cuda_coredump_%h.%p.%t"，其中%h是主机名，%p是进程 ID，%t是核心转储的时间戳。 - 默认情况下，核心转储过程会保存整个 GPU 上下文。对于像大模型推理这类占用几乎全部 GPU 显存的任务，全量核心转储是不切实际的（数据量达数百 GiB）。你可以使用
CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory'环境变量来跳过保存 GPU 全局内存、共享内存和本地内存，从而减小核心转储文件的大小。skip_constbank_memory标志虽然未出现在官方文档中，但确实是 CUDA 核心转储功能支持的，在许多 GPU 线程同时出错时有时是必要的。
文档还提到，在 CUDA_COREDUMP_GENERATION_FLAGS 中添加 skip_abort 可以防止 CPU 进程在核心转储完成后中止。这允许 CPU 进程添加自己的错误跟踪信息，从而提供更多的调试信息。然而，实验表明该功能存在重大 缺陷，可能导致 GPU 上的非法内存访问错误被忽略。这种情况下，后续代码可能会继续运行，但程序的内存数据可能早已损坏。这对于训练任务来说是不可接受的，对于推理任务也是不希望发生的。因此，该功能通常不可靠，不建议使用。
此外，文档指出启用 CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 不仅会启用 CUDA 核心转储，默认还会生成 CPU 核心转储。但在实际使用中，我们发现 CPU 核心转储几乎不包含有用的信息，且难以分析。
如果你想要用于调试的实时数据，也可以启用 CUDA_DEVICE_WAITS_ON_EXCEPTION=1 环境变量。该选项不使用 CUDA 核心转储，而是在异常发生时立即挂起 GPU 执行，等待用户连接调试器（如 cuda-gdb）来检查 GPU 状态（此时完整的 GPU 显存仍然完好）。不过，这种方法自动化程度较低，需要更多的人工干预。
总之，使用 CUDA 核心转储功能时，建议采用以下环境变量组合：
CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 CUDA_COREDUMP_SHOW_PROGRESS=1 CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory' CUDA_COREDUMP_FILE="/persistent_dir/cuda_coredump_%h.%p.%t"
CUDA 核心转储的使用示例
让我们通过一些代码来验证 CUDA 核心转储的有效性。
调试不当的内核启动
// test.cu
#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>
// CUDA error checking macro
#define cuda_check(call) do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
printf("CUDA Error at %s:%d - %s: %s\n", __FILE__, __LINE__, #call, cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
} while(0)
// Kernel with illegal memory access - accesses memory beyond allocated bounds
__global__ void illegalMemoryAccessKernel(int* data, int size) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// This will cause illegal memory access - accessing beyond allocated memory
// We allocate 'size' elements but access up to size * 2
if (idx < size * 2) { // Access twice the allocated size
for (int i = 0; i < 10000; i++) {
data[idx - 1000000000 + i] = idx; // This will cause illegal access for idx == 0
}
}
}
// Simple kernel with no errors
__global__ void normalKernel(int* data, int size) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
data[idx] = idx;
}
}
int main() {
printf("CUDA Illegal Memory Access Test\n");
printf("===============================\n\n");
int size = 100;
int* h_data = (int*)malloc(size * sizeof(int));
int* d_data;
// Initialize host memory
for (int i = 0; i < size; i++) {
h_data[i] = 0;
}
// Allocate device memory
cuda_check(cudaMalloc(&d_data, (unsigned long long)(size) * sizeof(int)));
cuda_check(cudaMemcpy(d_data, h_data, size * sizeof(int), cudaMemcpyHostToDevice));
// Launch kernel with illegal memory access
int blockSize = 256;
int numBlocks = (size + blockSize - 1) / blockSize;
printf("Launching kernel with out-of-bounds access...\n");
illegalMemoryAccessKernel<<<numBlocks, blockSize>>>(d_data, size);
normalKernel<<<numBlocks, blockSize>>>(d_data, size);
cuda_check(cudaMemcpy(h_data, d_data, size * sizeof(int), cudaMemcpyDeviceToHost));
for (int i = 0; i < 5; i++) {
printf("%d ", h_data[i]);
}
printf("\n");
// Synchronize to catch any runtime errors
cuda_check(cudaDeviceSynchronize());
printf("Test completed.\n");
// Cleanup
cuda_check(cudaFree(d_data));
free(h_data);
return 0;
}这段代码连续启动了两个内核（illegalMemoryAccessKernel 和 normalKernel）。在执行过程中，你会遇到如下错误消息：CUDA Error at test.cu:62 - cudaMemcpy(h_data, d_data, size * sizeof(int), cudaMemcpyDeviceToHost): an illegal memory access was encountered，且该错误仅在 cudaMemcpy 的返回值中被检测到。即使设置了 CUDA_LAUNCH_BLOCKING=1，也无法识别出导致错误的特定内核。
添加 CUDA 核心转储相关的环境变量后，我们可以观察到：
[06:43:15.209195] coredump: Detected an exception of type CUDBG_EXCEPTION_WARP_ILLEGAL_ADDRESS (14)
[06:43:15.209202] coredump: - Device: 0
[06:43:15.209206] coredump: - SM: 124
[06:43:15.209208] coredump: - Warp: 0
[06:43:15.209210] coredump: - PC 0x7462c3bac310
[06:43:15.209477] coredump: Stack trace (lane masks: active 0xFFFFFFFF, valid 0xFFFFFFFF):
[06:43:15.209486] coredump: #0 0x7462c3bac620 _Z25illegalMemoryAccessKernelPii
[00:40:46.806153] coredump: Writing ELF file to /tmp/cuda_coredump_xxx.1799919.1754898045
[1] 1799919 IOT instruction (core dumped) CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 CUDA_COREDUMP_SHOW_PROGRESS=1 = = ./test3当 GPU 线程触发非法内存访问后，CPU 会立即生成一个核心转储文件，随后触发 CPU 异常，直接终止程序。此时，我们获得了一个核心转储文件 /tmp/cuda_coredump_xxx.1799919.1754898045。我们可以使用 cuda-gdb 打开它（命令：target cudacore /path/to/coredump_file，其中 cudacore 指的是 CUDA 核心转储）。
$ cuda-gdb
(cuda-gdb) target cudacore /tmp/cuda_coredump_xxx.1799919.1754898045
Opening GPU coredump: /tmp/cuda_coredump_xxx.1799919.1754898045
CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x7f31abb9f6d0 illegalMemoryAccessKernel(int*, int)
[Current focus set to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 124, warp 0, lane 0]
#0 0x00007f31abb9f6e0 in illegalMemoryAccessKernel(int*, int)<<<(1,1,1),(256,1,1)>>> ()我们可以清楚地看到，该异常是由 illegalMemoryAccessKernel 在 kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 124, warp 0, lane 0 处引起的。
调试 CUDA 图（CUDA Graphs）中的内核异常
这里有一个更复杂的例子：将一个会导致非法内存访问的内核插入到 CUDA 图中：
# core_dump.py
import torch
import torch.nn as nn
from dataclasses import dataclass
@dataclass
class CupyWrapper:
data_ptr: int
size_in_bytes: int
@property
def __cuda_array_interface__(self):
return {
"shape": (self.size_in_bytes,),
"typestr": '|u1',
"data": (self.data_ptr, False),
"version": 3,
}
def from_buffer(data_ptr: int, size_in_bytes: int) -> torch.Tensor:
out = torch.as_tensor(CupyWrapper(data_ptr, size_in_bytes))
assert data_ptr == out.data_ptr(), "not zero-copy convert, something must be wrong!"
return out
class NeuralNetwork(nn.Module):
def __init__(self):
super(NeuralNetwork, self).__init__()
# First layer: [B, 10] -> [B, 20] with ReLU activation
self.layer1 = nn.Linear(10, 20)
self.relu = nn.ReLU()
# Second layer: [B, 20] -> [B, 30]
self.layer2 = nn.Linear(20, 30)
self.num_called = 0
def forward(self, x):
# Input shape: [B, 10]
x = self.layer1(x) # [B, 20]
x = self.relu(x) # [B, 20] with ReLU activation
self.num_called += 1
if self.num_called > 1:
y = from_buffer(x.data_ptr(), x.numel() * 1024 * 1024)
# will trigger illegal memory access
y.fill_(1)
x = self.layer2(x) # [B, 30]
return x
# Example usage
if __name__ == "__main__":
# Check if CUDA is available
device = torch.device('cuda' if torch.cuda.is_available() else 'cpu')
print(f"Using device: {device}")
# Create the model and move to CUDA
model = NeuralNetwork().to(device)
# Create sample input with batch size B=4 and move to CUDA
batch_size = 4
input_tensor = torch.randn(batch_size, 10).to(device)
print(f"Input shape: {input_tensor.shape}")
print(f"Input device: {input_tensor.device}")
# Forward pass
with torch.no_grad():
# warmup
output = model(input_tensor)
# capture graph
g = torch.cuda.CUDAGraph()
with torch.cuda.graph(g):
output = model(input_tensor)
# replay graph
g.replay()
print(f"Output shape: {output.shape}")
print(f"Output device: {output.device}")
print(f"Output: {output.sum()}")
# Print model summary
print("\nModel architecture:")
print(model)
# Print number of parameters
total_params = sum(p.numel() for p in model.parameters())
print(f"\nTotal parameters: {total_params}")
# Verify model is on CUDA
print(f"Model device: {next(model.parameters()).device}")直接执行会导致以下错误：
Using device: cuda
Input shape: torch.Size([4, 10])
Input device: cuda:0
Output shape: torch.Size([4, 30])
Output device: cuda:0
Traceback (most recent call last):
File "core_dump.py", line 76, in <module>
print(f"Output: {output.sum()}")
RuntimeError: CUDA error: an illegal memory access was encountered
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.
For debugging consider passing CUDA_LAUNCH_BLOCKING=1
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.在 output.sum() 触发设备同步并揭示非法内存访问之前，错误不会被打印出来。然而，由于 CUDA 内核是异步执行的，我们无法得知是哪个内核引发了非法访问。
添加 CUDA_LAUNCH_BLOCKING=1 后，错误消息变为：
Using device: cuda
Input shape: torch.Size([4, 10])
Input device: cuda:0
Traceback (most recent call last):
File "core_dump.py", line 71, in <module>
g.replay()
File "/uv_envs/py310/lib/python3.10/site-packages/torch/cuda/graphs.py", line 88, in replay
super().replay()
RuntimeError: CUDA error: an illegal memory access was encountered
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.可以推断出异常发生在 CUDA 图内的一个内核中。但常规方法只能提供到此为止的信息。
添加环境变量 CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 CUDA_COREDUMP_SHOW_PROGRESS=1 CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory' CUDA_COREDUMP_FILE="/tmp/cuda_coredump_%h.%p.%t" 后，我们可以清楚地识别出导致错误的内核：
(cuda-gdb) target cudacore /tmp/cuda_coredump_flow-matic.1929094.1754901120
Opening GPU coredump: /tmp/cuda_coredump_flow-matic.1929094.1754901120
CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x7fc2afba5e30 void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<unsigned char>, std::array<char*, 1ul> >(int, at::native::FillFunctor<unsigned char>, std::array<char*, 1ul>)
[Current focus set to CUDA kernel 0, grid 9, block (17454,0,0), thread (0,0,0), device 0, sm 0, warp 1, lane 0]
#0 0x00007fc2afba5e70 in void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<unsigned char>, std::array<char*, 1ul> >(int, at::native::FillFunctor<unsigned char>, std::array<char*, 1ul>)<<<(40960,1,1),(128,1,1)>>> ()显然，这是一个 fill 函数，且 40960 的网格大小非常大。有了这些信息，我们可以轻松定位到代码行 y = from_buffer(x.data_ptr(), x.numel() * 1024 * 1024); y.fill_(1);，该行代码强制将 x 的长度扩大了一百万倍，然后将其全部填为 1，从而触发了 illegal memory access 异常。
在某些 GPU 上，这行代码可能会导致 invalid argument 错误而非 illegal memory access，因为网格大小超过了最大限制。在这种情况下，CUDA 核心转储功能无法被触发，你需要稍微调小扩大因子 1024 * 1024 以避免超过网格大小限制。
局限性与注意事项
- 理论上，CUDA 核心转储应该能够捕获由 GPU 上特定线程引起的各种异常。然而在实践中，在某些 GPU 和驱动版本上，像
operation not supported on global/shared address space这样的异常可能无法触发 CUDA 核心转储。幸运的是，illegal memory access通常可以可靠地触发核心转储，这足以满足大多数调试需求。 - 对于硬件相关错误，如
Invalid access of peer GPU memory over nvlink or a hardware error，这些不是由特定线程引起的，无法归咎于某个具体的 GPU 线程。因此，CUDA 核心转储不会为此类问题触发。 - 由不当使用驱动程序 API 引起的错误被视为非粘性错误，且与 GPU 本身无关。这些错误在驱动程序 API 层面报告，不会触发 CUDA 核心转储。一个常见的例子是在
cudaMalloc期间出现的内存不足（out-of-memory）错误，它不会导致 CUDA 核心转储。 - 对于涉及多 GPU 通信的分布式程序，经常使用内存映射将其他 GPU 的内存映射到当前 GPU。如果另一个 GPU 上的程序退出，映射的内存将变为无效，访问它会触发
illegal memory access。然而，这不属于典型的illegal memory access问题。此类问题在分布式程序关闭过程中很常见。如果 GPU 在关闭期间进行通信，关闭顺序可能导致某些 GPU 报告illegal memory access。在使用 CUDA 核心转储排查此类程序时，务必区分这些误报。 - 启用 CUDA 核心转储确实会对 CUDA 内核产生一定的性能影响（因为它需要在 GPU 线程退出时检查错误并进行归因）。因此，不建议在生产环境中启用它。建议仅在可以可靠复现
illegal memory access等错误的情况下，为了调试目的启用。 - 为了从 CUDA 核心转储中获得最大收益，建议使用调试符号重新编译 vLLM，或者至少在编译期间嵌入行信息。遗憾的是，由于二进制文件大小限制，vLLM 的默认构建版本不包含此类信息。要享受这一优势，用户必须从源码编译 vLLM，并配置环境变量
export NVCC_PREPEND_FLAGS='-lineinfo'或export NVCC_PREPEND_FLAGS='-G'。建议从-lineinfo开始，仅在-lineinfo不足以定位时才切换到-G。有了丰富的调试信息，CUDA 核心转储可以回溯到导致异常的确切代码行。
总结
本篇博客分析了 CUDA 核心转储的原理和应用场景。这种调试方法对于不当的内核启动和 CUDA 图内的内核异常等问题非常有效，使其成为调试 illegal memory access 及其他相关问题的强大工具。
举个例子，我们最近使用此技术调试了 vLLM 中的一个复杂 illegal memory access 问题，详情请见此 PR。简而言之，我们为 MRope 添加了一个 Triton 内核，但该内核隐式假设 head_size==rotary_dim（即全 Rope）。当 head_size!=rotary_dim（即部分 Rope）时，内核会触发 illegal memory access，这正是新的 GLM-4.5V 模型的情况。如果没有 CUDA 核心转储，错误会被报告为 Failed: Cuda error /workspace/csrc/custom_all_reduce.cuh:453 'an illegal memory access was encountered'，非常具有误导性。利用 CUDA 核心转储，我们可以轻松将错误定位到 MRope 内核，并加以修复。请注意，本例是由 CUDA 内核参数配置错误引起的，能够定位到引发问题的内核足以进行调试。对于更复杂的 illegal memory access 问题，我们仍需隔离内核并在最小化示例中复现问题，而不是依赖端到端示例，然后使用 Compute Sanitizer 等更专业的工具进行进一步调查。
vLLM 项目旨在为每个人提供简单、快速且廉价的 LLM 服务，便捷的调试也是其中的重要一环。未来我们将继续分享更多的调试技巧和技术，共同构建强大的 LLM 推理生态系统。要分享你与 vLLM 的故事或使用经验，请在博客仓库中提交 PR。
致谢
我们要感谢 NVIDIA 的 Ze Long、Vikram Sharma Mailthody、Jeremy Iverson 和 Sandarbh Jain 提供的有益讨论。Red Hat 的 Lucas Wilkinson 协助润色了草稿。