追踪挂起和复杂的 GPU 内核，直达源代码

2025年12月3日13 分钟阅读

Kaichao You (vLLM)

几个月前，我们发表了一篇关于 CUDA 核心转储：调试内存访问问题及其他问题的有效工具的博客文章，介绍了一种调试 CUDA 内核非法内存访问问题的强大技术。这是 GPU 内核调试领域的一个重要里程碑，因为它使开发者能够精确定位导致故障的内核。在此之前，由于 GPU 执行的异步特性，识别问题内核几乎是不可能的，且错误消息往往具有误导性。

随着 CUDA 核心转储技术的广泛采用，开发者们表示需要更细粒度的信息——特别是触发问题的具体代码行。在本篇博客中，我们将填补这一空白，首先介绍如何识别挂起的内核，然后演示如何将出现问题的内核追踪回其源代码。

如何寻找挂起的内核

GPU 计算能力呈指数级增长，但内存带宽的发展却没能跟上。这种失衡导致了日益复杂的内存访问模式。近年来，旗舰级数据中心 GPU 引入了异步内存访问模式，这在实现高性能内核时需要复杂的同步操作。这些同步机制特别容易在复杂的代码库中产生竞态条件和死锁。

当 GPU 内核挂起时，程序通常会冻结或变得无响应——即使按下 Ctrl-C 也无法停止。最直接的解决办法是强制杀掉进程，但这种方法无法提供关于根本原因的任何信息。开发者只能盲目猜测，通过二分法排查代码变更并重复进行测试，直到找出问题所在。

注意： 为什么当 CUDA 内核挂起时按下 Ctrl-C 无法停止进程？按下 Ctrl-C 会向进程发送 SIGINT 信号。如果进程正在运行 Python 代码，SIGINT 信号会被 Python 解释器捕获，转变为 KeyboardInterrupt 异常，并排队等待进程返回运行 Python 代码时处理。然而，如果进程正在运行 CUDA 内核并等待 GPU 完成，此时它正在等待底层的 CUDA API 返回，并没有运行任何 Python 代码，因此无法引发 KeyboardInterrupt 异常。在接下来的 conditional_hang.py 示例中，如果您想通过 Ctrl-C 终止进程，需要在脚本开头添加 import signal; signal.signal(signal.SIGINT, signal.SIG_DFL)，这样 Python 解释器就不会捕获 SIGINT 信号，Ctrl-C 就能成功终止进程。缺点是当被 Ctrl-C 停止时，Python 解释器将无法显示错误堆栈。

幸运的是，有一个更好的方法。CUDA 驱动程序包含一项名为 用户引发的 GPU 核心转储生成 的功能：驱动程序会在操作系统中打开管道，允许用户通过写入管道来触发核心转储。一旦触发，CUDA 驱动程序会将 GPU 状态转储到核心转储文件中，从而能够检查 GPU 内部发生的情况，更重要的是，识别出哪个 GPU 内核发生了挂起。

考虑一个条件挂起内核的简单示例

# 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()

执行此代码将无限期挂起。为了调试该问题，我们可以启用用户引发的 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

当代码处于无限期运行状态时，我们可以通过向管道写入数据来触发 CUDA 核心转储

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

我们向管道写入 1MB 的零数据以触发 CUDA 核心转储。注意，由于管道缓冲机制，简单的 echo 命令可能无法生效。

触发核心转储后，运行 python conditional_hang.py 的原始终端将显示核心转储进度

[01:39:15.256278] coredump: Writing ELF file to /tmp/cuda_coredump_hostname.3000837.1764236276
[01:39:15.256350] coredump: Writing out global memory (0 bytes)
[01:39:15.256354] coredump: Writing out device table
[01:39:15.292027] coredump: Writing out metadata
[01:39:15.292039] coredump: Finalizing
[01:39:15.292124] coredump: Writing done
[01:39:15.292128] coredump: All done (took 00s)

然后，我们可以使用 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），还能精确定位到其挂起的具体代码行。与之前甚至无法识别问题内核（更不用说导致挂起的特定行）的情况相比，这有显著的改进。

一个小的不便之处是核心转储管道的路径是由 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

如何追踪复杂内核的源代码

在上一篇博客中，我们提到编译时添加 export NVCC_PREPEND_FLAGS='-lineinfo' 环境变量可以将行信息嵌入到编译后的二进制文件中，使我们能够追踪导致问题的确切代码行。经过讨论和对多个实际问题的调试，我们发现 cuda-gdb 显示行信息的默认方式并不完美

对于某些复杂的内核，即使编译后的二进制文件中嵌入了行信息，cuda-gdb 也无法找到导致问题的正确代码行。
即使 cuda-gdb 能够找到正确的代码行，它也只会显示编译器内联后的最后一行，这可能并非导致问题的实际行。由于 C++ 代码严重依赖内联来消除运行时函数调用的开销，我们需要完整的内联堆栈才能理解问题所在。

让我们用一个具体的例子来说明。下面的 Python 脚本演示了一个非法内存访问问题

# 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])

使用 PyTorch >= 2.9.0 运行此代码（请特别确保包含此提交；否则您会看到类似 RuntimeError: The specified pointer resides on host memory and is not registered with any CUDA device. 的错误）。这将触发非法内存访问错误。

首先，让我们在开启 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" \
python illegal_memory_access.py

核心转储进度将明确识别出导致问题的内核

_ZN2at6native24index_elementwise_kernelILi128ELi4EZNS0_16gpu_index_kernelIZNS0_17index_kernel_implINS0_10OpaqueTypeILi1EEEEEvRNS_18TensorIteratorBaseEN3c108ArrayRefIlEESA_EUlPcPKclE_EEvS7_SA_SA_RKT_bEUliE_EEvlT1_

从内核名称中我们可以看出，该问题是由 PyTorch 的 index_elementwise_kernel 引起的。为了定位导致问题的具体代码行，我们需要在设置 export NVCC_PREPEND_FLAGS='-lineinfo' 环境变量的情况下从源码构建 PyTorch，然后再次运行代码。

当编译后的 GPU 内核嵌入了行信息时，我们可以使用 cuda-gdb 打开核心转储文件，确切地查看是哪一行代码导致了问题

(cuda-gdb) target cudacore /tmp/cuda_coredump_flow-matic.3756036.1764250282
Opening GPU coredump: /tmp/cuda_coredump_flow-matic.3756036.1764250282
[Current focus set to CUDA kernel 0, grid 4, block (0,0,0), thread (0,0,0), device 0, sm 124, warp 3, lane 0]
 
CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x7ff533bb91d0  ...
#0  void at::native::index_elementwise_kernel<128, 4, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<1> >(at
::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<
long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<1> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayR
ef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel<at::native::index_kernel_imp
l<at::native::OpaqueType<1> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::Ten
sorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<1> >(at::TensorIteratorBase&,
c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1})<<<(1,1,1),(128,1,1)>>> ()
    at /data/youkaichao/pytorch/aten/src/ATen/native/cuda/IndexKernel.cu:203 in _ZZN2at6native17index_kernel_implINS0_10OpaqueTypeILi1EEEEEvRNS
_18TensorIteratorBaseEN3c108ArrayRefIlEES8_ENKUlPcPKclE_clES9_SB_l inlined from IndexKernel.cu:118
203         *reinterpret_cast<scalar_t*>(out_data) = *reinterpret_cast<const scalar_t*>(in_data + offset);

接下来，在 cuda-gdb 中，我们可以使用 info symbol $errorpc 获取关于错误位置的更多信息

(cuda-gdb) info symbol $errorpc
void at::native::index_elementwise_kernel<128, 4, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<1> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<1> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<1> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<1> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}) + 11472 in section .text._ZN2at6native24index_elementwise_kernelILi128ELi4EZNS0_16gpu_index_kernelIZNS0_17index_kernel_implINS0_10OpaqueTypeILi1EEEEEvRNS_18TensorIteratorBaseEN3c108ArrayRefIlEESA_EUlPcPKclE_EEvS7_SA_SA_RKT_bEUliE_EEvlT1_ of /tmp/cuda-dbg/2123124/session1/elf.21407f80.24fe2940.o.4gyLzn

这提供了关于错误位置的更多信息。cuda-gdb 解包了编译后的二进制文件，/tmp/cuda-dbg/2123124/session1/elf.21407f80.24fe2940.o.4gyLzn 是一个包含 index_elementwise_kernel 的 cubin 文件。错误发生在 cubin 文件的 0x7ff533bb91d0 位置。我们可以使用 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] ;
...

现在我们可以看到导致问题的完整内联堆栈。默认情况下，cuda-gdb 只显示最后一次内联展开。

命令简要说明

-ndf：反汇编后禁用数据流分析器。
-c：仅打印代码段。
-gi：使用从 .debug_line 段获取的源代码行信息以及函数内联信息（如果存在）来标注反汇编结果。
-C20：一个 grep 参数，显示程序计数器地址 7ff533bb91d0 周围 20 行的上下文。

如果 cubin 文件包含多个具有相同程序计数器地址的内核（即 grep 显示多次匹配），我们需要进一步过滤信息

$ cuobjdump -elf /tmp/cuda-dbg/2123124/session1/elf.21407f80.24fe2940.o.4gyLzn > elf.txt
$ cat elf.txt | grep ".text._ZN2at6native24index_elementwise_kernelILi128ELi4EZNS0_16gpu_index_kernelIZNS0_17index_kernel_implINS0_10OpaqueTypeILi1EEEEEvRNS_18TensorIteratorBaseEN3c108ArrayRefIlEESA_EUlPcPKclE_EEvS7_SA_SA_RKT_bEUliE_EEvlT1_" | grep PROGBITS
 
  1ac 1b83f80   b200  0 80                     PROGBITS        6    3      26a .text._ZN2at6native24index_elementwise_kernelILi128ELi4EZNS0_16gpu_index_kernelIZNS0_17index_kernel_implINS0_10OpaqueTypeILi1EEEEEvRNS_18TensorIteratorBaseEN3c108ArrayRefIlEESA_EUlPcPKclE_EEvS7_SA_SA_RKT_bEUliE_EEvlT1_
 
$ nvdisasm -ndf -c -gi -fun 0x26a /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] ;
...

主要区别在于通过在 ELF 段中搜索函数（本例中为 26a）从 cuobjdump 获取 CUDA 函数索引（-fun 参数）。

注意，这是一个演示该技术的简化示例。现实世界的内核可能要复杂得多。例如，这里有一个复杂的内联案例

	//## 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/cute/arch/util.hpp", line 185 inlined at "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/atom/copy_traits.hpp", line 133
	//## File "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/atom/copy_traits.hpp", line 133 inlined at "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/atom/copy_atom.hpp", line 103
	//## File "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/atom/copy_atom.hpp", line 103 inlined at "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/atom/copy_atom.hpp", line 124
	//## File "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/atom/copy_atom.hpp", line 124 inlined at "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/algorithm/copy.hpp", line 211
	//## File "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/algorithm/copy.hpp", line 211 inlined at "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/algorithm/copy.hpp", line 412
	//## File "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cute/algorithm/copy.hpp", line 412 inlined at "/data/youkaichao/data/vllm_flash_attn/hopper/epilogue_fwd.hpp", line 265
	//## File "/data/youkaichao/data/vllm_flash_attn/hopper/epilogue_fwd.hpp", line 265 inlined at "/data/youkaichao/data/vllm_flash_attn/hopper/flash_fwd_kernel_sm90.h", line 454
	//## File "/data/youkaichao/data/vllm_flash_attn/hopper/flash_fwd_kernel_sm90.h", line 454 inlined at "/data/youkaichao/data/vllm_flash_attn/hopper/utils.h", line 41
	//## File "/data/youkaichao/data/vllm_flash_attn/hopper/utils.h", line 41 inlined at "/data/youkaichao/data/vllm_flash_attn/csrc/cutlass/include/cutlass/device_kernel.h", line 122
	//## 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 无法正确关联该行。实际上，它在错误位置周围没有显示任何行信息。即使它显示了正确的行，也只显示最后一个内联帧，即 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——这是 CUTLASS 函数内部的内联展开，对于调试根本问题仍然无济于事。

通过上述方法，我们可以揭示源代码的完整内联链，并仔细检查每一帧，以确定到底是哪一行代码导致了错误。

警告： 为了最大限度地利用 CUDA 核心转储，行信息至关重要。建议使用 export NVCC_PREPEND_FLAGS='-lineinfo' 环境变量进行编译，因为它透明地应用于所有已编译的内核，无需修改编译脚本。然而，这种透明度意味着如果您使用诸如 ccache 之类的编译缓存机制，它可能会忽略该标志并直接重用先前编译的结果而不进行实际编译。从源码编译时，请确保禁用了编译缓存机制。如果您使用即时（Just-In-Time）编译，请查阅您的即时编译工具文档，了解如何添加行信息。

总结

本篇博客介绍了两种用于 CUDA 内核的高级调试技术。第一种技术使用用户触发的核心转储来识别挂起的内核，第二种技术通过利用嵌入在已编译二进制文件中的行信息将复杂内核追踪回其源代码。这些技术是调试 CUDA 内核中复杂问题的强大工具，尤其是非法内存访问问题。通过结合使用这两者，我们最近得以调试出一个难以复现且棘手的 CUTLASS MLA 注意力后端挂起问题，该问题实际上源于上游 CUTLASS 代码示例，并已在 v4.3.0 版本中修复。

vLLM 项目旨在为每个人提供简单、快速且经济的 LLM 服务，而无障碍调试是这一使命的重要方面。未来我们将继续分享更多的调试技巧和技术，共同构建强大的 LLM 推理生态系统。要分享您与 vLLM 的故事或使用经验，请在博客仓库中提交 PR。

致谢

感谢 NVIDIA 的 Ze Long 和 Sandarbh Jain 提供的有益讨论。Moonshot AI 的 Chao Hong 帮助提供了激励性示例。Red Hat 的 Lucas Wilkinson 协助润色了草稿。