执行摘要

• 一句话：修复 FA4 后端 Gemma4 长上下文乱码
• 推荐动作：建议合入该 PR，它修复了 Gemma4 在 FA4 后端上的严重输出错误。推荐注意力后端开发者详细阅读 _make_mm_prefix_mask_mod 的修改，学习如何正确地将滑动窗口与自定义 mask_mod 结合，以及绝对位置计算在分块预填充中的关键作用。对于使用 mm_prefix 特性的模型维护者，可参考其滑动窗口 clamp 的设计。

功能与动机

根据关联 Issue #47300 和 PR body，Gemma4 视觉模型在 SM90（FlashAttention-4）后端上，当上下文长度超过滑动窗口（1024 tokens）且包含至少一张图像时，输出完全乱码（如重复的 '768368...' 或 'la la la...'）。Triton 后端在同一输入下输出正确，表明是 FA4 后端特有的 bug。根因在于 _make_mm_prefix_mask_mod 缺失滑动窗口约束和使用局部 q_idx，导致非图像 Token 的注意力分布被破坏。本 PR 旨在修复这两个根因，恢复长上下文多模态推理的正确性。

实现拆解

变更仅涉及一个文件 vllm/v1/attention/backends/flash_attn.py，按以下步骤实现：

1. 修改函数签名：在 _make_mm_prefix_mask_mod 中新增 sliding_window_left: int | None 参数，用于传递滑动窗口左边界大小（Triton 约定：1 + window_size[0]）；全局注意力层 ((-1,-1)) 传递 None。

2. 计算绝对 Q 位置：在 mask_mod 内部通过 seqlen_info 计算出 q_abs = q_idx + (seqlen_k - seqlen_q)，代替原来的局部 q_idx，使得因果比较、滑动窗口限制和 mm_prefix 范围判断都使用一致的绝对位置，与 Triton 参考路径一致。

3. 添加滑动窗口约束：当 sliding_window_left 不为 None 时，在 causal 条件上追加 (q_abs - kv_idx) < sliding_window_left，实现 (causal AND sliding_window) OR mm_prefix 语义。使用两条 @cute.jit 分支（有窗口和无窗口）避免 CuTe DSL 的运行时分支。

4. 调整调用端：在 FlashAttentionBackend.forward 中，根据 sliding_window_size 计算 sw_val = 1 + sliding_window_size[0]（当 window_size[0] >= 0 时），否则 None。该值同时用于 mm_clamp_sw（控制 mm_prefix 双向块的 clamp）和 sliding_window_left。mm_prefix_clamp_sliding_window 层属性（来自 PR #47217）决定是否启用 clamp。

5. 清理冗余：移除旧有的直接赋值和注释，统一使用 sw_val，提高可读性。

测试配套：本次改动未新增自动化测试文件，但作者在 PR body 中提供了详细的 NIAH（needle-in-a-haystack）手动复现步骤，并在 H100（SM90）上验证了修复结果。cjackal 在 Issue 评论中也确认了 dense 和 moe 模型均通过测试。

关键文件：

• vllm/v1/attention/backends/flash_attn.py（模块 注意力层；类别 source；类型 core-logic；符号 _make_mm_prefix_mask_mod, mm_prefix_mask_mod）: 唯一修改文件，实现所有修复逻辑，包括函数签名变更、绝对位置计算、滑动窗口约束和调用端调整。

关键符号：_make_mm_prefix_mask_mod, mm_prefix_mask_mod

评论区精华

审核由 Isotr0py 完成并批准，未提出修改意见。cjackal 在关联 Issue #47300 中手动验证了修复在 dense 和 moe 模型上均有效（'Thanks for a quick fix! I have verified that this PR branch passes the test script without trouble for both dense and moe model.'）。Mergify 自动化提醒分支存在冲突，但在合并前成功解决。整体没有出现设计分歧或未解决疑虑。

• 审核与批准 (other): PR 获得批准，可以合并。
• 手动验证修复有效性 (testing): 验证通过，修复有效。

风险与影响

• 风险：变更严格局限在 _make_mm_prefix_mask_mod 一个函数内，仅影响 is_mm_prefix_lm=True 且使用滑动窗口的模型（目前主要是 Gemma4），风险可控：
  • 回归风险：全局注意力层（window=(-1,-1)）传递 sw_left=None，走无窗口分支，行为与修复前基本一致（仅 q_abs 修正可能带来微小变化，理论上修复错误，不引入退化）。
  • 手动验证覆盖有限：虽然 NIAH 测试涵盖了单块预填充（32k tokens）和分块预填充（144k tokens），但未覆盖所有序列长度和图像组合组合。
  • 硬件依赖：仅 SM90+ GPU（FA4）受影响，其他硬件（如 Triton 后端）不受影响。
  • 测试配套：无新增自动化测试，回归风险需由现有 CI 覆盖（可能不足）。
  • 影响：对用户的影响：之前使用 Gemma4 模型（H100 等 SM90+ GPU）进行长上下文视觉推理时输出完全乱码，修复后恢复正常，是严重 Bug 的关键修复。对于其他模型（如 bagel、molmo2、paligemma）或非 FA4 后端，行为不变。对系统影响：无性能退化，两条 @cute.jit 分支避免了运行时分支，应保持零开销。对团队影响：为后续 mm_prefix 模型正确集成滑动窗口提供了参考实现。
  • 风险标记：仅手动验证, 无新增自动化测试, 依赖特定 GPU 架构 (SM90+), 影响范围限于 FA4 + mm_prefix + sliding_window

关联脉络

• PR #47217 PR #47217: 引入 mm_prefix_clamp_sliding_window 层属性: 本 PR 依赖该 PR 引入的 mm_prefix_clamp_sliding_window 标志来启用对 mm_prefix 双向块的滑动窗口 clamp，PR body 中明确引用。