将挂起且复杂的 GPU Kernel 追溯至源代码

几个月前，我们发布了一篇关于 CUDA Core Dump：调试内存访问问题及其他问题的有效工具 的博客文章，介绍了一种调试 CUDA kernel 中非法内存访问问题的强大技术。这是 GPU kernel 调试的一个重要里程碑，因为它使开发人员能够精确定位导致失败的 kernel。此前，由于 GPU 执行的异步特性，识别有问题的 kernel 几乎是不可能的，而且错误消息通常具有误导性。

随着 CUDA core dump 技术的普及，开发人员表达了对更细粒度信息的需求——具体来说，就是触发问题的确切源代码行。在这篇博文中，我们将填补这一空白，首先介绍如何识别挂起的 kernel，然后演示如何将有问题的 kernel 追溯到其源代码。

如何寻找挂起的 kernel

GPU 的算力呈指数级增长，但内存带宽却未能跟上步伐。这种不平衡导致了日益复杂的内存访问模式。近年来，旗舰级数据中心 GPU 引入了异步内存访问模式，在实现高性能 kernel 时需要复杂的同步。这些同步机制容易产生竞争条件（race conditions）和死锁，特别是在复杂的代码库中。

当 GPU kernel 挂起时，程序通常会冻结或无响应——即使按下 Ctrl-C 也无法停止。最直接的解决方案是杀掉进程，但这种方法无法提供有关根本原因的任何信息。开发人员只能盲目猜测，反复二分代码更改并运行测试，直到确定问题所在。

说明

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

幸运的是，有一种更好的方法。CUDA 驱动程序包含一个名为 用户诱导型 GPU core dump 生成 (user induced GPU core dump generation) 的功能：驱动程序在操作系统中打开管道（pipes），允许用户通过向其写入数据来触发 core dump。触发后，CUDA 驱动程序会将 GPU 状态转储到 core dump 文件中，从而可以检查 GPU 内部发生的情况，最重要的是，可以识别哪个 GPU kernel 发生了挂起。

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

# 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 core dump 生成

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 core dump

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

我们向管道写入 1MB 的零值以触发 CUDA core dump。请注意，简单的 echo 命令可能由于管道缓冲而不起作用。

触发 core dump 后，运行 python conditional_hang.py 的原始终端将显示 core dump 进度

[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 打开 core dump 文件，查看 kernel 具体在何处挂起

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)

这种方法不仅能让我们识别挂起的 kernel（conditional_hang_kernel），还能精确定位挂起的代码行。这相对于以前无法识别问题 kernel、更不用说定位挂起具体行的情况来说，是一个巨大的进步。

一个小遗憾是，core dump 管道的路径是由 CUDA 驱动程序动态生成的，难以定位。我们可以通过使用 CUDA_COREDUMP_PIPE 环境变量来指定 core dump 管道的模板路径，从而通过检查进程的文件描述符轻松找到它

$ 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

如何追溯复杂 kernel 的源代码

在之前的博客文章中，我们提到使用 export NVCC_PREPEND_FLAGS='-lineinfo' 环境变量进行编译会将行信息嵌入到编译后的二进制文件中，使我们能够追溯到导致问题的确切代码行。在讨论和调试了几个实际问题后，我们发现 cuda-gdb 默认显示行信息的方式并不完美

1. 对于一些复杂的 kernel，即使编译后的二进制文件中嵌入了行信息，cuda-gdb 也无法找到导致问题的正确代码行。

2. 即使 cuda-gdb 能找到正确的代码行，它也只显示编译器内联（inlining）后的最后一行，这可能不是导致问题的实际行。由于 C++ 代码严重依赖内联来消除运行时函数调用开销，我们需要完整的内联堆栈（inline stack）来理解问题。

让我们用一个具体的例子来说明。以下 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 core dump 的情况下运行代码

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

core dump 进度将明确识别出导致问题的 kernel

_ZN2at6native24index_elementwise_kernelILi128ELi4EZNS0_16gpu_index_kernelIZNS0_17index_kernel_implINS0_10OpaqueTypeILi1EEEEEvRNS_18TensorIteratorBaseEN3c108ArrayRefIlEESA_EUlPcPKclE_EEvS7_SA_SA_RKT_bEUliE_EEvlT1_

从 kernel 名称可以看出，问题是由 PyTorch 的 index_elementwise_kernel 引起的。要定位导致问题的确切代码行，我们需要使用 export NVCC_PREPEND_FLAGS='-lineinfo' 环境变量从源码构建 PyTorch，然后再次运行代码。

当编译后的 GPU kernel 嵌入了行信息后，我们可以使用 cuda-gdb 打开 core dump 文件，查看具体是哪行代码导致了问题

(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 参数，显示找到的程序计数器（PC）地址 7ff533bb91d0 前后的 20 行上下文。

如果 cubin 文件包含多个具有相同程序计数器地址的 kernel（即 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 段，从 cuobjdump 获取 CUDA 函数索引（-fun 参数），在本例中为 26a。

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

	//## 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 ;

在这种情况下，有问题的代码是

attention kernel 中一行被污染的代码。

出错的源代码调用了一些 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 core dump 的效益，行信息至关重要。建议使用 export NVCC_PREPEND_FLAGS='-lineinfo' 环境变量进行编译，因为这可以透明地应用于所有编译的 kernel，而无需修改编译脚本。然而，这种透明性意味着如果你使用类似 ccache 的编译缓存机制，它可能会忽略该标志并重用之前编译的结果，而不进行实际编译。从源代码编译时，请确保禁用编译缓存机制。如果你使用即时编译（Just-In-Time compilation），请查阅即时编译工具的文档，了解如何添加行信息。

结论

这篇博文介绍了 CUDA kernel 的两种高级调试技术。第一种技术使用用户触发的 core dump 来识别挂起的 kernel，而第二种技术利用嵌入在编译二进制文件中的行信息，将复杂的 kernel 追溯到其源代码。这些技术是调试 CUDA kernel 复杂问题的强大工具，尤其是非法内存访问问题。通过结合使用这两种技术，我们最近成功调试了 CUTLASS MLA attention 后端中一个难以重现且棘手的挂起问题，该问题实际上源于上游 CUTLASS 代码示例，此后已在 v4.3.0 中修复。

vLLM 项目旨在为每个人提供简单、快速且负担得起的 LLM 服务，而便捷的调试是这一使命的重要组成部分。未来我们将继续分享更多调试技巧和技术，共同构建强大的 LLM 推理生态系统。要分享您的故事或 vLLM 使用经验，请在 blogpost 仓库提交 PR。

致谢

我们要感谢来自 NVIDIA 的 Ze Long 和 Sandarbh Jain 提供的有益讨论。来自月之暗面（Moonshot AI）的 Chao Hong 协助提供了启发性示例。来自 Red Hat 的 Lucas Wilkinson 协助润色了草案。