执行摘要

• 一句话：修复NVFP4 swizzled scale buffer零初始化，恢复Blackwell吞吐
• 推荐动作：值得所有涉及量化内核或硬件适配的开发者精读：1）学习如何通过隔离实验定位内存初始化问题；2）理解Tensor Core tile padding对buffer分配的影响；3）设计决策“只还原必要的分支”体现了精准修复的思维方式。建议将相关注释纳入团队知识库。

功能与动机

Blackwell GPU (B300/GB300) 上NVFP4量化解码遭遇严重性能回归：吞吐量下降87-93%，输出长度坍塌（请求1000 token，实际仅~65-127 token）。二分定位到vllm/_custom_ops.py中create_fp4_scale_tensor函数的swizzled scale buffer分配从torch.zeros改为torch.empty（PR #42988）。该buffer因Tensor Core tile要求被填充（round_up至128×4），但量化内核并未写入所有填充元素，torch.empty导致GEMM读取未初始化的scale因子，产生错误结果。详见关联issue #45741。

实现拆解

1. 问题定位：通过二分法和隔离实验，确定create_fp4_scale_tensor中swizzled int32 buffer的torch.empty是唯一致因。
2. 安全分析：逐一审查其他同批修改的非swizzled uint8 buffer（精确尺寸无填充）、packed-FP8 scale、TurboQuant attn_out、GDN z_out等，确认其分配尺寸准确且内核全覆盖，无需零初始化。
3. 实施修改：将vllm/_custom_ops.py第60行的torch.empty恢复为torch.zeros，并添加详细注释解释padding与内核不完全写入的关系，防止未来误改。
4. 隔离验证：通过分别恢复各分支的对照实验（swizzled only、non-swizzled only、all），证明仅还原swizzled一行即可完全复原吞吐和输出长度。

关键文件：

• vllm/_custom_ops.py（模块 核心算子；类别 source；类型 core-logic；符号 create_fp4_scale_tensor）: 唯一变更文件，核心NVFP4 scale buffer分配函数，修复了性能回归

关键符号：create_fp4_scale_tensor

关键源码片段

vllm/_custom_ops.py

唯一变更文件，核心NVFP4 scale buffer分配函数，修复了性能回归

def create_fp4_scale_tensor(
    m: int,
    n: int,
    device: torch.device,
    is_sf_swizzled_layout: bool,
) -> torch.Tensor:
    """
    Allocate the output scale tensor for scaled_fp4_quant.
​
    When is_sf_swizzled_layout=True, we use rounded values to store the
    swizzled scales. Due to the requirement of the Tensor Core, the minimum
    tile is 128x4 for the scales. So, we first pad the scales to multiples
    of 128 (rows) and 4 (cols). Then, the scales (in float8_e4m3fn) are
    packed into an int32 for every 4 values. More:
    https://docs.nvidia.com/cuda/parallel-thread-execution/
    #tcgen05-mma-scale-factor-b-layout-4x
    """
    from vllm.utils.math_utils import round_up
​
    block_size = 16
    if is_sf_swizzled_layout:
        rounded_m = round_up(m, 128)
        scale_n = n // block_size
        rounded_n = round_up(scale_n, 4)
        # Must be zero-initialized: the swizzled scale buffer is padded to
        # (round_up(m, 128), round_up(scale_n, 4) // 4) but the NVFP4 quant
        # kernel does not write every padded element that the downstream
        # NVFP4 GEMM reads. torch.empty leaves those padded scale factors
        # uninitialized, which corrupts dequantization and causes a severe
        # Blackwell NVFP4 decode throughput/output-length regression.
        return torch.zeros(
            (rounded_m, rounded_n // 4), device=device, dtype=torch.int32
        )
    else:
        return torch.empty((m, n // block_size), device=device, dtype=torch.uint8)

评论区精华

• vadiklyutiy 注意到仅还原swizzled分支后仍与完全恢复有~20%差异，但 qiching 重跑后差异消失，确认是测量噪声。
• vadiklyutiy 询问非swizzled uint8分支是否也需要改为zeros，qiching 分析该分支分配精确尺寸 (m, n//16)，内核 cvt_fp16_to_fp4_sf_major 全部覆盖，不存在未初始化读取，故保持torch.empty。
• yewentao256 要求提供完整复现命令和日志，qiching 提供了vllm serve和aiperf的详细命令及输出。

• 针对根因是否为NaN导致性能慢的疑问，qiching 深入分析发现7×吞吐下降实为aiperf测量伪像，真实问题是正确性退化（输出退化），离线decode速度测试显示empty和zeros的每token延迟相同。

• 非swizzled分支是否需要同样修正 (design): 确认非swizzled分支无需修改，维持 torch.empty。

• 请求复现基准命令和日志 (testing): 提供了完整复现方法和结果日志。
• 根因分析与7x吞吐下降的误解 (correctness): 确认根本问题是正确性 Bug，而非内核性能差。

风险与影响

• 风险：该PR仅修改一行代码，将分配从empty恢复为zeros，回归风险极低。风险在于：若未来有开发者不理解padding语义而再次将zeros改为empty，会重新引入问题；当前修改已通过注释明确说明。恢复zeros会引入微小的初始化开销（约一次对round_up(m,128)*round_up(n/64,1)大小tensor的零填充），相比原先的吞吐量回归可忽略不计。
• 影响：对使用NVFP4量化（--quantization modelopt_fp4）的Blackwell架构用户（B300/GB300）：解编码吞吐量从~19 tok/s恢复至~140 tok/s以上，输出长度从~65-127恢复至正常值~1000 token。对非Blackwell或其他量化方案用户无影响。团队应关注create_fp4_scale_tensor的分配注释，避免后续优化中误改该行为。
• 风险标记：核心量化路径变更, 未初始化内存回归, 仅单行修复, 测试未直接覆盖

关联脉络

• PR #42988 [Perf] zeros -> empty to remove additional fill: 本PR部分恢复了该PR引入的回退，该PR将swizzled scale buffer从 zeros 改为 empty 导致回归。
• PR #45741 [Perf]: Severe NVFP4 decode throughput regression on Blackwell (B300/GB300): uninitialized swizzled scale buffer in create_fp4_scale_tensor: 关联issue，报告了同样的回归现象和根因分析。