CUDA 核心转储：调试内存访问问题及其他问题的有效工具
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），却不知道如何调试？我们在开发 vLLM（一个用于 LLM 模型的高性能推理引擎）时，一次又一次地感受到了这种痛苦。
如果您是遇到此问题的开发人员之一，那么这篇博客文章就是为您准备的！我们将揭示我们使用的一些高级调试技术，这些技术可以帮助用户调试 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 状态，找出哪个内核导致了问题以及非法内存访问是什么。
什么是核心转储？
GPU 本质上是一个大规模并行处理器，它的许多概念都可以在 CPU 中找到对应物。
核心转储是 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 线程同时出错时是必需的。
文档还提到，将 skip_abort 添加到 CUDA_COREDUMP_GENERATION_FLAGS 可以防止 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 图中：
# 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通常可以可靠地触发 CUDA 核心转储，这满足了大多数调试需求。 - 对于硬件相关错误，例如
Invalid access of peer GPU memory over nvlink or a hardware error，这些不是由特定线程引起的，也不能归因于特定的 GPU 线程。因此，CUDA 核心转储不会为此类问题触发。 - 由于不当使用驱动程序 API 导致的错误被认为是非粘滞错误，与 GPU 本身无关。这些错误在驱动程序 API 级别报告，不会触发 CUDA 核心转储。一个常见的例子是
cudaMalloc期间的内存不足错误，这不会导致 CUDA 核心转储。 - 对于涉及多 GPU 通信的分布式程序，通常使用内存映射将其他 GPU 的内存映射到当前 GPU。如果另一个 GPU 上的程序退出，映射的内存将变为无效，访问它将触发
illegal memory access。然而，这不属于典型的illegal memory access问题。此类问题在分布式程序的关机过程中很常见。如果在关机期间 GPU 正在通信，关机顺序可能导致某些 GPU 报告illegal memory access。当为此类程序使用 CUDA 核心转储时，区分这些误报很重要。 - 启用 CUDA 核心转储确实会对 CUDA 内核产生一些性能影响（因为它需要在 GPU 线程退出时检查错误并对其进行归因）。因此，不建议在生产环境中启用 CUDA 核心转储。建议仅在可以可靠地重现
illegal memory access等错误时才启用 CUDA 核心转储进行调试。 - 为了最大程度地利用 CUDA 核心转储，建议使用调试符号重新编译 vLLM，或者至少在编译期间嵌入行信息。不幸的是，由于二进制文件大小限制，vLLM 的默认构建不包含此类信息。要享受此优势，用户必须使用环境变量
export NVCC_PREPEND_FLAGS='-lineinfo'或export NVCC_PREPEND_FLAGS='-G'从源代码编译 vLLM。建议从-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 帮助润色了草稿。