AttributeError: ‘FlashAttentionForwardSm120’ object has no attribute ‘is_split_kv’`

快速结论：该报错发生在 vLLM 0.26.0 在 SM120 架构 GPU（如 RTX PRO 6000 Blackwell）上首次调用 CuTe 版 FlashAttention 前向函数时，属于 Python 层代码缺陷。优先排查这两处一行级代码问题：`mDynamicCausal` 变量作用域缺失，以及 `self.is_split_kv` 属性从未赋值。

适用环境：NVIDIA RTX PRO 6000 Blackwell Server Edition（sm_120），CUDA 13.2，Ubuntu 25.10，vLLM 0.26.0（Docker 镜像 vllm/vllm-openai:v0.26.0），亦在 0.26.1.dev1 源码构建版复现。

最快修复方案：暂无一键修复方案。需在 vendored 仓库 vllm-project/flash-attention（pin 在 caaa4eb59845388a20b1f435ecaafb4bd9517ad8）中修改 flash_attn/cute/flash_fwd.py 的两处一行代码，等待官方 bump pin 后更新 vLLM。

注意事项：此修复涉及 vLLM 构建时从外部仓库 vendored 的文件，直接修改安装目录下的文件可能在下一次重新安装时被覆盖。该缺陷与 paged KV 相关 issue #51405 相互独立，修复本问题不代表 SM120 的 paged KV 功能可用。

问题场景

在单 GPU、不加载任何模型、仅通过 ~25 行 Python 脚本直接调用 vllm.vllm_flash_attn.cute.flash_attn_varlen_func 时触发。SM80 级 GPU 走 C++ FA2/FA3 内核，不受影响；sm_120 架构强制走 CuTe DSL 路径，因此首次调用即失败，且发生在任何 kernel 实际运行之前。

报错原文

AttributeError: 'FlashAttentionForwardSm120' object has no attribute 'is_split_kv'

在修复第一处 NameError 后出现此错误，完整原始报错如下：

error: NameError in `__call__`: name 'mDynamicCausal' is not defined
   --> /usr/local/lib/python3.12/dist-packages/vllm/vllm_flash_attn/cute/flash_fwd.py:840
  = note: Caused exception: name 'mDynamicCausal' is not defined

原因分析

可能原因有两处独立缺陷：

Bug 1（NameError）：mDynamicCausal 在 __call__ 方法中作为参数定义（第 681 行），但在 @cute.kernel def kernel(...) 的 840 行被引用。由于 kernel 是类中独立方法而非 __call__ 的嵌套闭包，该参数不在作用域内；同时 launch 处（761 行）按位置传参在 output_scale 处结束，未传递该参数。

Bug 2（AttributeError）：self.is_split_kv 在整个文件中被读取 4 次、但从未赋值。在 SM120 上这类属性缺失直接抛出 AttributeError。

注意：vllm/vllm_flash_attn/cute/flash_fwd.py 并不在 vLLM 主仓库中，而是构建时从 vllm-project/flash-attention vendored 进来的，修复需在源仓库完成。

环境排查

GPU 架构：确认是否为 sm_120（如 RTX PRO 6000 Blackwell）
驱动版本：Driver 595.84，CUDA 13.2（Issue 中已确认环境）
vLLM 版本：0.26.0 或 0.26.1.dev1（均可复现）
检查 vendored 文件路径：/usr/local/lib/python3.12/dist-packages/vllm/vllm_flash_attn/cute/flash_fwd.py
确认 flash-attention 子仓库 pin：caaa4eb59845388a20b1f435ecaafb4bd9517ad8（在 cmake/external_projects/vllm_flash_attn.cmake 中）

解决步骤

确认你正处于 sm_120 架构环境且有上述报错，而非其他架构问题。
定位源文件：flash_attn/cute/flash_fwd.py（在 pin 的 flash-attention 仓库中，与 vendored 文件内容一致，行号相同）。
修复 Bug 1：在 __call__ 方法的 kernel 签名末尾追加参数 mDynamicCausal: Optional[cute.Tensor] = None（第 681 行参数列表后），并在 launch 位置参数列表 output_scale 后追加 mDynamicCausal。
修复 Bug 2：在 FlashAttentionForwardSm120 类（或其基类）中为 self.is_split_kv 添加初始化赋值。
若修改的是安装目录下的 vendored 文件，请保留修改或跟踪变更，重新安装 vLLM 时不会覆盖（实际上重新安装会覆盖，见注意事项）。
等待 vLLM bump flash-attention pin 以正式合入修复，或向 vllm-project/flash-attention 提交 PR（该仓库 issues 已关闭，可在 vLLM 仓库讨论）。

验证方法

运行 Issue 中的复现脚本（batch size 1，单 GPU，无需模型），预期打印 OK torch.Size([1024, 16, 128]) 而非任何 Python 异常。若两处修复均生效但仍有其他错误，请确认是否涉及 paged KV 或相关架构门控问题（参见 #51560 的独立修复）。