Tracing Hanging and Complicated GPU Kernels Down To The Source Code

发表时间: 2025-12 · Blog post by Kaichao You (vLLM) (blog.vllm.ai)

原文: https://blog.vllm.ai/blog/2025-12-03-improved-cuda-debugging

作者/机构：Kaichao You (vLLM)

A1 主要贡献

此前，作者曾介绍过使用CUDA Core Dump作为调试CUDA内核非法内存访问问题的有效工具，它使开发者能够精确定位导致故障的具体内核，从而解决了GPU异步执行导致错误信息具有误导性的难题。随着该技术的普及，开发者提出了对更细粒度信息的需求——即希望定位到触发问题的确切源代码行。

本文的核心贡献在于填补了这一空白，具体提出并演示了两种高级调试技术：

1. 识别并定位挂起（Hanging）的GPU内核：解决了常规方法（如Ctrl-C）无法终止并提供挂起原因的问题。
2. 追踪复杂内核的源代码：克服了默认调试工具在面对现代C++代码中大量使用的编译器内联（Inlining）时，无法准确显示真实错误代码行的局限性。

A3 背景知识与关键Observation

GPU内存带宽不平衡与挂起问题
GPU的计算能力呈指数级增长，但内存带宽并未保持同步。这种不平衡导致了越来越复杂的内存访问模式。近年来，旗舰级数据中心GPU引入了异步内存访问模式，在实现高性能内核时需要复杂的同步机制。这些同步机制在复杂的代码库中极易引发竞争条件（Race conditions）和死锁（Deadlocks）。当GPU内核挂起时，程序通常会冻结或无响应。

为什么按 Ctrl-C 无法停止挂起的CUDA内核？
按下Ctrl-C会向进程发送SIGINT信号。如果进程正在运行Python代码，Python解释器会捕获该信号，将其转换为KeyboardInterrupt异常，并将其排队等候在进程返回运行Python代码后处理。然而，如果进程正在运行CUDA内核并等待GPU完成，它实际上是在等待底层的CUDA API返回，此时没有Python代码在运行，因此无法抛出KeyboardInterrupt异常。
解决此问题的权宜之计：可以在脚本开头添加 import signal; signal.signal(signal.SIGINT, signal.SIG_DFL)，使Python解释器不捕获SIGINT信号，从而让Ctrl-C成功终止进程。但缺点是Python解释器在被Ctrl-C停止时将无法显示错误堆栈。

默认调试工具在复杂内核中的局限性
即使使用-lineinfo环境变量将行信息嵌入到编译后的二进制文件中，cuda-gdb 默认显示行信息的方式仍存在缺陷：

1. 对于某些复杂的内核，cuda-gdb 无法找到引发问题的正确代码行。
2. 即使找到了代码行，它通常只显示编译器内联（Inlining）后的最后一行，这往往不是实际引发问题的原始代码行。由于C++代码严重依赖内联来消除运行时函数调用的开销，开发者必须获取完整的内联栈（Inline stack）才能真正理解问题。

A2 方法细节

定位挂起的内核

CUDA驱动程序包含一个名为用户触发的GPU核心转储生成（user induced GPU core dump generation）的功能：驱动程序在操作系统中打开管道（pipes），允许用户通过向其写入数据来触发核心转储。触发后，CUDA驱动会将GPU状态转储到核心转储文件中，从而允许开发者检查GPU内部状态并确定哪个内核处于挂起状态。

1. 构造条件挂起内核示例
作者提供了一个基于Triton的条件挂起内核，当传入的flag为0时，将陷入无限循环：

# save as conditional_hang.py

import triton
import triton.language as tl
import torch

@triton.jit
def conditional_hang_kernel(x_ptr,
                            flag,          # int32 scalar
                            n_elements,    # int32 scalar
                            BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offs < n_elements

    # Load values
    x = tl.load(x_ptr + offs, mask=mask, other=0)

    # If flag == 1: do a normal "+1" update
    if flag == 1:
        x = x + 1
        tl.store(x_ptr + offs, x, mask=mask)
    else:
        # Else: non-terminating loop, no break.
        # The loop condition depends on `flag`, which is invariant,
        # so this is effectively an infinite loop when flag == 0.
        while flag == 0:
            # do something trivial so the loop isn't optimized away
            x = x + 1
            tl.store(x_ptr + offs, x, mask=mask)

x = torch.ones(16, dtype=torch.float32, device="cuda")
n_elements = x.numel()
BLOCK_SIZE = 16

# 1) Normal behavior: increment by 1
conditional_hang_kernel[(1,)](
   x,
   flag=1,
   n_elements=n_elements,
   BLOCK_SIZE=BLOCK_SIZE,
)
print("After flag=1:", x)  # should be all 2s

# 2) Hanging behavior: this will spin forever
conditional_hang_kernel[(1,)](
   x,
   flag=0,
   n_elements=n_elements,
   BLOCK_SIZE=BLOCK_SIZE,
)

# this print will hang, because printing x will synchronize the device,
# and the kernel will never finish.
print("After flag=0:", x)

# the following line will never be reached
x = x + 2
torch.cuda.synchronize()

2. 启用环境变量并运行
为了调试挂起问题，通过配置环境变量启用用户触发的GPU核心转储机制并运行脚本：

CUDA_ENABLE_USER_TRIGGERED_COREDUMP=1 \
CUDA_COREDUMP_PIPE="/tmp/cuda_coredump_pipe_%h.%p.%t" \
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" \
python conditional_hang.py

3. 触发核心转储
在代码无限期运行时，通过向管道写入数据来触发CUDA核心转储。由于管道缓冲的原因，简单的echo命令可能不起作用，因此使用dd命令写入1MB的零数据：

dd if=/dev/zero bs=1M count=1 > /tmp/cuda_coredump_pipe_hostname.3000837.1764236276

4. 分析转储文件
触发转储后，原始终端会打印转储进度日志。随后，使用 cuda-gdb 打开生成的核心转储文件：

Opening GPU coredump: /tmp/cuda_coredump_hostname.3000837.1764236276
[Current focus set to CUDA kernel 0, grid 53, block (0,0,0), thread (0,0,0), device 0, sm 124, warp 0, lane 0]
#0  0x00007f2e6fbff300 in conditional_hang_kernel<<<(1,1,1),(128,1,1)>>> () at conditional_hang.py:31
31                  tl.store(x_ptr + offs, x, mask=mask)

此方法不仅能识别出挂起的内核是conditional_hang_kernel，还能精确定位到其挂起的具体代码行（第31行）。

5. 寻找管道路径的技巧
由于核心转储管道的路径是由CUDA驱动动态生成的，较难直接找到。可以通过 CUDA_COREDUMP_PIPE 环境变量指定管道的模板路径，然后通过检查进程的文件描述符来轻松定位：

$ ls /proc/3037675/fd/ -alth | grep /tmp/cuda_coredump_pipe_
lr-x------ 1 user user 64 Nov 27 01:50 98 -> /tmp/cuda_coredump_pipe_hostname.3037675.1764237014

追踪复杂内核的源代码

对于复杂的C++内核，仅依靠cuda-gdb往往只能看到内联后的最底层代码。为了获取完整的内联栈，需要结合使用多个底层工具。

1. 构造非法内存访问示例
以下Python脚本通过PyTorch构造了一个非法的内存访问异常：

# save as illegal_memory_access.py

from dataclasses import dataclass
import torch

@dataclass
class TensorWrapper:
    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, device: str, dtype: torch.dtype) -> torch.Tensor:
    return torch.as_tensor(TensorWrapper(data_ptr, size_in_bytes), device=device).view(dtype)

data = from_buffer(123456, 1024, device="cuda:0", dtype=torch.uint8)

index = torch.ones(10, device="cuda", dtype=torch.int32) + 100
print(data[index])

2. 获取引发异常的内核名称
启用CUDA core dump环境变量运行上述脚本，转储进度日志中会明确暴露出引发异常的内核名称（例如 _ZN2at6native24index_elementwise_kernel...），这表明问题出在PyTorch的 index_elementwise_kernel。

3. 带行信息编译并获取错误PC地址
使用 export NVCC_PREPEND_FLAGS='-lineinfo' 环境变量从源码重新编译PyTorch。再次运行并生成Core dump后，使用 cuda-gdb 打开转储文件：

(cuda-gdb) target cudacore /tmp/cuda_coredump_flow-matic.3756036.1764250282
...
CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x7ff533bb91d0  ...

此时cuda-gdb会显示异常为 Warp Illegal Address，并给出程序计数器（Program Counter, PC）地址为 0x7ff533bb91d0。

4. 提取错误位置符号与Cubin文件路径
在 cuda-gdb 中，使用 info symbol $errorpc 命令获取更多信息：

(cuda-gdb) info symbol $errorpc
void at::native::index_elementwise_kernel... + 11472 in section .text._ZN2at6native24index_elementwise_kernel... of /tmp/cuda-dbg/2123124/session1/elf.21407f80.24fe2940.o.4gyLzn

该输出表明 cuda-gdb 解包了编译后的二进制文件，并生成了包含内核的cubin文件（路径为 /tmp/cuda-dbg/.../elf...o...），错误精确发生在该cubin文件的 0x7ff533bb91d0 位置。

5. 使用 nvdisasm 反汇编并查看完整内联栈
利用 nvdisasm 工具对生成的cubin文件进行反汇编，以查看确切的代码执行栈：

$ nvdisasm -ndf -c -gi /tmp/cuda-dbg/2123124/session1/elf.21407f80.24fe2940.o.4gyLzn > output.txt
$ grep -C20 7ff533bb91d0 output.txt
...
        /*7ff533bb9190*/                   IMAD.IADD R19, R23, 0x1, R3 ;
.L_x_27840:
	//## File "/data/youkaichao/pytorch/aten/src/ATen/native/cuda/IndexKernel.cu", line 203 inlined at "/data/youkaichao/pytorch/aten/src/ATen/native/cuda/IndexKernel.cu", line 118
	//## File "/data/youkaichao/pytorch/aten/src/ATen/native/cuda/IndexKernel.cu", line 118 inlined at "/data/youkaichao/pytorch/aten/src/ATen/native/cuda/IndexKernel.cu", line 37
	//## File "/data/youkaichao/pytorch/aten/src/ATen/native/cuda/IndexKernel.cu", line 37
        /*7ff533bb91a0*/                   ULDC.64 UR4, c[0x0][0x480] ;
        /*7ff533bb91b0*/                   IADD3 R2, P0, P1, R22, UR4, R2 ;
        /*7ff533bb91c0*/                   IADD3.X R3, R19, UR5, RZ, P0, P1 ;
        /*7ff533bb91d0*/                   LDG.E.U8 R3, desc[UR36][R2.64] ;
...

命令参数解释：

• -ndf: 反汇编后禁用数据流分析器。
• -c: 仅打印代码段。
• -gi: 使用从 .debug_line 段获取的源行信息以及函数内联信息（如果存在）来注释反汇编结果。
• -C20: grep参数，显示匹配PC地址 7ff533bb91d0 前后的20行上下文。

通过这种方式，开发者可以看到引发问题的完整内联栈，而默认的 cuda-gdb 仅会显示最后一次内联展开。

6. 处理多内核PC地址冲突
如果cubin文件包含多个具有相同PC地址的内核（即 grep 结果显示多个匹配项），需要使用 cuobjdump 进一步过滤信息：

$ cuobjdump -elf /tmp/cuda-dbg/2123124/session1/elf.21407f80.24fe2940.o.4gyLzn > elf.txt
$ cat elf.txt | grep ".text._ZN2at6native24index_elementwise_kernel..." | grep PROGBITS
 
  1ac 1b83f80   b200  0 80                     PROGBITS        6    3      26a .text._ZN2at6native24index_elementwise_kernel...
 
$ nvdisasm -ndf -c -gi -fun 0x26a /tmp/cuda-dbg/2123124/session1/elf.21407f80.24fe2940.o.4gyLzn > output.txt
$ grep -C20 7ff533bb91d0 output.txt

这里的核心区别在于，通过 cuobjdump 搜索函数的ELF段，获取了CUDA函数的索引（在此例中为 26a），然后通过 -fun 0x26a 参数传递给 nvdisasm，从而精准定位目标内核。

7. 真实复杂内联场景分析
在真实的复杂内核中（如CUTLASS），内联栈可能会非常深：

	//## File "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/arch/copy_sm90.hpp", line 93 inlined at "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/arch/util.hpp", line 158
	//## File "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/arch/util.hpp", line 158 inlined at "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/arch/util.hpp", line 185
... (多层内联省略)
	//## File "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cutlass/device_kernel.h", line 122
        /*7eebf5e9eb80*/                   STSM.16.M88.4 [R13], R4 ;
        /*7eebf5e9eb90*/                   MOV R34, R26 ;

在上述案例中，有缺陷的源代码调用了CUTLASS函数，且包含该缺陷的函数又被上层调用者内联。此时 cuda-gdb 无法正确关联行信息，或者只能显示CUTLASS内部深处的展开行（如 copy_sm90.hpp 第93行），这对调试毫无帮助。通过本文揭示的方法，开发者能够获取完整的内联链，逐帧检查并最终定位到真实的错误代码行。

A4 实验环境

• 软件配置：

  • PyTorch >= 2.9.0（必须包含特定commit以修复Host memory未向CUDA设备注册的报错）。
  • Triton 框架。
  • CUDA 底层工具链：cuda-gdb, nvdisasm, cuobjdump。
  • CUTLASS 库。

• 硬件配置：提及了具有异步内存访问特性的旗舰级数据中心GPU（Flagship datacenter GPUs）。

A4 实验结果

• 针对挂起内核的实验：成功通过向管道写入1MB零数据触发了Core Dump，利用 cuda-gdb 精确到了基于Triton编写的条件挂起脚本（conditional_hang.py）的第31行（执行 tl.store 的位置）。
• 针对复杂内核的实验：在运行PyTorch非法内存访问脚本时，成功通过 nvdisasm 配合 -gi 参数，反汇编了由 cuda-gdb 提取出的cubin文件，还原了由于C++高度内联导致 cuda-gdb 无法有效显示的完整源代码内联栈，精确定位到了 IndexKernel.cu 的第203行。
• 真实场景应用：结合上述两种技术，作者团队最近成功调试了CUTLASS MLA attention后端中一个难以复现且极为棘手的挂起问题。该问题实际源自上游的CUTLASS代码示例，目前已在CUTLASS v4.3.0中得到修复。

A5 结论

本文介绍了两种针对CUDA内核的高级调试技术：第一种技术利用用户触发的核心转储来识别处于挂起状态的内核；第二种技术通过利用嵌入在编译二进制文件中的行信息，将高度复杂的内核追踪回其确切的源代码行。这两种技术协同使用，构成了调试CUDA内核复杂问题（尤其是在非法内存访问和并发死锁方面）的强大工具链。vLLM项目致力于提供简单、快速且经济的LLM服务，而降低底层调试门槛是实现这一使命的重要环节。未来，团队将继续分享更多调试技巧，共同构建强大的LLM推理生态系统。

A7 补充细节

关于编译缓存的严重警告
为了最大化CUDA核心转储的优势，行信息至关重要。强烈推荐使用环境变量 export NVCC_PREPEND_FLAGS='-lineinfo' 进行编译，因为它无需修改编译脚本即可透明地应用于所有编译内核。然而，这种透明性也带来了一个隐患：如果使用了编译缓存机制（如 ccache），它可能会忽略该环境变量标志，直接复用以前编译好（但不包含行信息）的缓存结果，而不进行实际编译。

• 源码编译：从源码编译时，必须确保禁用所有的编译缓存机制。
• JIT编译：如果使用即时编译（Just-In-Time compilation），开发者需要查阅相应的JIT工具文档，了解如何显式添加行信息。