A Anaxagore
← 所有项目

Contribution Tasks

Dao-AILab/flash-attention

候选人擅长 CUDA/Triton kernel、算子融合与性能建模，但当前无 GPU。仓库活跃开发集中在 FA4 CuTe DSL（flash_attn/cute/），主战场是 SM100/SM120 kernel 与调度层。最佳切入点是：(1) 数值精度契约文档与回归测试，(2) 调度策略的 CPU 端微基准与行为一致性矩阵，(3) Python dispatch 开销的纯逻辑优化，(4) 依赖兼容性矩阵与 CI 补全。这些方向都能在 Mac 上用 CPU/PyTorch reference 实现，直接服务 FA4 beta 每周发版的痛点。

当前方向：FA4（flash_attn/cute/）是当前活跃开发重心，围绕 SM100 Blackwell（B200/SM103）推进 hd256 前向/反向统一、S/P ping-pong 重叠、MLA 1CTA/稀疏 top-k、2CTA 调度命名化；FA3（hopper/）标记 beta 仅用于 H100/H800 测试；FA2（csrc/）仍是 pip install 默认目标但新功能少。发布节奏为每周一 beta（fa4-v4.0.0.beta32–34），API 仍在收敛。

★ 25119Fork 31494 个候选任务Gemini：LongCat-2.0

更新于 2026-11-05T07:57:30+00:00 · 本次生成失败，展示上次结果 · 打开仓库 ↗

一、项目定位

FlashAttention 是 Dao-AILab 维护的高效精确注意力（exact attention）实现，围绕 Tri Dao 等人的 FlashAttention 系列论文演进：FA1（IO-aware exact attention）、FA2（更优并行与工作切分）、FA3（Hopper 优化，beta 版在 hopper/ 目录）、FA4（用 NVIDIA CuTeDSL 重写，面向 Hopper SM90 与 Blackwell SM100/SM110，位于 flash_attn/cute/）。仓库同时保留旧代实现（FA2 在 csrc/、FA3 在 hopper/），但活跃开发已明确切到 FA4（CLAUDE.md 直接写“active development is on FA4 in flash_attn/cute/”）。 从最近提交热点看，主战场是 CuTe DSL 上的 SM100（Blackwell B200/SM103）适配：MLA 前向/反向、hd256 反向统一、S/P ping-pong 重叠、decode 路径修复、2CTA 调度命名化、稀疏 MLA top-k 等。FA3 仍标记为 beta release，用于 H100/H800 的测试与 benchmark。 仓库规模可观：csrc/ 197 个文件（含 composable_kernel、cutlass 子树）、hopper/ 507 个文件、flash_attn/ 110 个文件、tests/ 55 个、training/ 145 个，整体是“多代并存 + 训练/推理/基准测试/CI 工具链齐全”的中大型项目。发布节奏上，FA4 已进入 beta 序列（fa4-v4.0.0.beta32–beta34，2026-09-23 至 2026-10-07 连续三周发版），说明 API 仍在收敛、但已进入可集成测试阶段。 对 GPU kernel / 训练性能工程师来说，这是一个“把 IO-aware 算法 + CuTe/TMA/warp 级调度做到生产级”的标杆仓库；切入点主要在 CuTe DSL 内核、MLA/稀疏注意力、分布式训练集成，以及端侧/推理侧的 paged-kv、量化等扩展。

解决什么问题、给谁用

解决 Transformer 中标准 O(N²) 注意力在 GPU HBM–SRAM 之间 IO 开销过大的问题：在不牺牲数值精确性（exact attention）的前提下，通过分块（tiling）、重计算（recompute softmax 中间态）、warp/CTA 级并行与 Hopper 的 TMA/GMMA 原语，把 attention 的内存占用和延迟显著降低，同时支持 causal mask、ALiBi、sliding window、MQA/GQA、varlen、dropout、rotary、paged-KV、block-sparse、MLA、FP8 等训练/推理常见变体。 目标用户：大模型训练/推理框架开发者、分布式训练工程师、推理栈（vLLM/SGLang/TensorRT-LLM 等）集成方、以及需要自己魔改 attention 的算法研究员。典型场景：替换 nn.MultiheadAttention 或 F.scaled_dot_product_attention 以降低显存、提升吞吐；在 Hopper/Blackwell 上追求极致 MFU；为扩散模型/DiT/视频生成做长序列 attention 加速；在 ROCm 或端侧通过 Triton/CK 后端获得可移植实现。

同类项目与差别

核心能力

能力在哪成熟度
FA4 CuTeDSL 精确注意力内核（Hopper SM90 + Blackwell SM100/SM110）flash_attn/cute/、flash_attn/__init__.py实验
FA3 Hopper 优化内核（FP16/BF16、FP8 前向、前向+反向）hopper/（flash_api.cpp、flash_fwd_kernel_sm90.h、mainloop_fwd_sm90_tma_gmma_ws.hpp 等）实验
FA2 Ampere/Ada/Hopper 通用 CUDA 实现与 Python 接口csrc/flash_attn/、flash_attn/flash_attn_interface.py成熟
AMD ROCm 双后端（composable_kernel 默认 + Triton/aiter）csrc/flash_attn_ck/、flash_attn/flash_attn_triton_amd.py、third_party/aiter成熟
MLA（Multi-head Latent Attention）密集/稀疏/Paged，1CTA/2CTAflash_attn/cute/、tests/cute/test_flash_attn_mla.py、hopper/flash_attn_3/实验
Varlen、mask_mod、score_mod、block-sparse、ALiBi、sliding window、dropout、rotaryflash_attn/flash_attn_interface.py、flash_attn/cute/、tests/cute/成熟
Paged KV / decode 优化（推理服务场景）hopper/paged_kv.h、benchmarks/benchmark_mla_decode.py、benchmarks/benchmark_split_kv.py成熟
训练集成（GPT 训练曲线、分布式、FSDP）training/（145 文件、configs/、src/、run.py）成熟
CI / 两阶段测试（FakeTensor 编译缓存 + xdist 并行）tools/two_pass_tests.sh、.github/workflows/、AGENTS.md/CLAUDE.md 中的 two-pass 说明成熟
Benchmark 与调优工具（SM90/SM100、GEMM、varlen、CLC）benchmarks/（bench_sm90.py、benchmark_flash_attention.py、clc_bench.py 等 12 个文件）成熟
Windows 编译（CUDA 13 + PyTorch 2.15 修复已合）setup.py、最近提交 Fix Windows builds with CUDA 13 and PyTorch 2.15 (#2872)实验

阶段：FA2/FA3 训练与推理接口已相对成熟并广泛被下游采用；FA4 处于活跃迭代阶段（README 与 CLAUDE.md 均标注为 beta/开发中），API 仍在收敛（beta32–beta34 三周连发），是当前主开发线。

技术栈：Python（接口、测试、benchmark）；C++ / CUDA（FA2 在 csrc/、FA3 在 hopper/）；CuTeDSL / PTX（FA4 在 flash_attn/cute/，依赖 nvidia-cutlass-dsl>=4.6.2）；PyTorch（torch.utils.cpp_extension、FakeTensorMode）；NVIDIA CUTLASS / composable_kernel（csrc/cutlass、csrc/composable_kernel）；Triton + aiter（ROCm Triton 后端，third_party/aiter）；einops、apache-tvm-ffi、quack-kernels>=0.5.3；ninja 构建、pytest + pytest-xdist、GitHub Actions CI

规模：仓库结构庞大：csrc/ 197 文件、hopper/ 507 文件、flash_attn/ 110 文件、tests/ 55 文件、training/ 145 文件、benchmarks/ 12 文件。FA4 发布已进入 beta 序列，最近三周连续发布 fa4-v4.0.0.beta32（2026-09-23）、beta33（2026-09-30）、beta34（2026-10-07），显示高频迭代。最近提交集中在 CuTe/SM100 适配（MLA、hd256 bwd、S/P ping-pong、decode 修复）和测试/格式化/Windows 编译修复。具体 star 数与月提交频率需验证。

二、架构与代码地图

仓库按“代际 × 作用”自然切成五层。 1) 接入层（Python API / 模型集成）：以 flash_attn/__init__.py、flash_attn/cute/__init__.py、flash_attn/flash_attn_interface.py、flash_attn/cute/interface.py 以及 flash_attn/models、flash_attn/modules、flash_attn/layers 为主，给用户暴露 flash_attn_func / flash_attn_varlen_func / MHA / rotary / GPT/LLaMA/BERT 等模型级封装，同时把训练侧（training/）和推理侧（examples/inference）也接入进来。 2) 调度层（tile/CTA/work 调度）：集中在 flash_attn/cute/tile_scheduler.py、flash_attn/cute/prepare_scheduler.py，以及 hopper/flash_prepare_scheduler.cu / hopper/instantiations 等。负责把 (batch, seqlen, nheads, headdim, causal, varlen, paged-kv) 切成 tile、选 persistent / CLC / dynamic / 2CTA 调度、产生 grid/CTA 坐标与流水线 producer/consumer 状态。 3) 执行层（FA4 CuTe DSL 内核）：flash_attn/cute/flash_fwd*.py、flash_bwd*.py、mask.py、softmax.py、pack_gqa.py、paged_kv.py、topk_gather_kv.py、pipeline.py、named_barrier.py 等，用 CuTeDSL 写 SM90/SM100/SM120 的前向、反向、MLA、稀疏 MLA、hd256 2CTA 等 kernel；是提交热点（flash_attn/cute 37 次、tests/cute 17 次）。 4) 执行层（FA2/FA3 旧代 C++ 内核）：csrc/flash_attn/src/*（FA2、SM80、CUTLASS/CuTe 老路径）、hopper/flash_fwd_kernel_sm90.h、hopper/flash_bwd_kernel_sm90.h、hopper/mainloop_*_tma_gmma_ws.hpp（FA3 Hopper TMA/GMMA 路径）。仍是 FA2/FA3 用户的生产路径，但 CLAUDE.md 明确“active development is on FA4”。 5) 工具/生态层：setup.py、benchmarks/、tests/、tools/、AI/（调试笔记与脚本）、third_party/aiter、training/。负责构建（多 arch、多代编译）、回归、性能对标、CI、以及与分布式训练 / 推理框架的集成验证。

接入层调度层FA4 执行层FA2/FA3 执工具/生态层flash_attn_interface → csrc_fa2：torch.opstorch.opscute_interface → flash_fwd：QKV 参数QKV 参数cute_interface → flash_bwd：QKV dO 参数QKV dO 参数flash_fwd → tile_scheduler：tile 坐标tile 坐标flash_bwd → tile_scheduler：tile 坐标tile 坐标flash_fwd → prepare_scheduler：元数据元数据flash_bwd → prepare_scheduler：元数据元数据tile_scheduler → utils_and_helpers：pipeline/bpipeline/bflash_fwd → utils_and_helpers：gemm/mask/flash_bwd → utils_and_helpers：gemm/mask/gemm/mask/model_layers → flash_attn_interface：MHA 调用MHA 调用model_layers → cute_interface：MHA 调用hopper_fa3 → csrc_fa2：共用 C++ 工具aiter_rocm → flash_attn_ck：ROCm 后端并列tests → flash_fwd：回归验证回归验证benchmarks → flash_fwd：性能对标性能对标flash_attn_interfaceflash_attn_interfacecute_interfacecute_interfacemodel_layersmodel_layerstrainingtrainingtile_schedulertile_schedulerprepare_schedulerprepare_schedulerutils_and_helpersutils_and_helpersflash_fwdflash_fwdflash_bwdflash_bwdcsrc_fa2csrc_fa2hopper_fa3hopper_fa3flash_attn_ckflash_attn_ckaiter_rocmaiter_rocmbenchmarksbenchmarksteststests
FlashAttention 多代并存架构：左侧 FA4（flash_attn/cute）为当前活跃开发路径，右侧 FA2/FA3 仍服务生产；顶层接入层与底层工具链跨代共享。
模块 / 路径职责 · 入口 · 依赖
flash_attn_interface
flash_attn/flash_attn_interface.py
1 文件 + csrc/flash_attn 系列 .cu/.h		FA2 时代的 Python → C++ 包装层，暴露 flash_attn_func/flash_attn_varlen_func/flash_attn_qkvpacked_func 等，内部通过 torch.ops 加载 csrc/flash_attn 编译出的 CUDA op。
入口：flash_attn_func, flash_attn_varlen_func, flash_attn_qkvpacked_func, flash_attn_with_kvcache
依赖：csrc/flash_attn (FA2 C++ 扩展)
README 与 usage.md 主要示例入口；PyTorch 训练代码（HuggingFace/Megatron）大多走这里。
cute_interface
flash_attn/cute/interface.py
1 文件（模块总 110 文件）		FA4 的 Python 入口，把 Q/K/V、causal、window、alibi、paged-kv、MLA 等参数转成 CuTeDSL kernel 调用。
入口：flash_attn_func, flash_attn_varlen_func, flash_attn_with_kvcache (FA4 路径)
依赖：flash_attn/cute/flash_fwd.py, flash_bwd.py, mla_*, tile_scheduler.py, prepare_scheduler.py
CLAUDE.md 与 AGENTS.md 指定为当前活跃入口；与 flash_attn_interface 并存，分别对应 FA4 / FA2。
flash_fwd
flash_attn/cute/flash_fwd.py
~1 主文件 + 多个变体 flash_fwd_sm90/sm100/sm120.py、flash_fwd_mla_sm100.py 等		FA4 前向主 kernel（SM90/SM100/SM120、pack-gqa、paged-kv、varlen），组装 TMA/UMMA/MMA、mask、softmax、scheduler。
入口：flash_fwd, FlashAttentionForwardSm90, FlashAttentionForwardSm100, FlashAttentionForwardSm120 等（需验证具体类名）
依赖：tile_scheduler.py, prepare_scheduler.py, mask.py, softmax.py, pack_gqa.py, paged_kv.py, pipeline.py, named_barrier.py, ampere_helpers.py, blackwell_helpers.py
提交热点区（hd256 fwd、S/P ping-pong、decode 修复、MLA fwd）集中在此。
flash_bwd
flash_attn/cute/flash_bwd.py
~1 主文件 + flash_bwd_sm90/sm100/sm120.py、flash_bwd_mla_sm100.py、sm100_hd256_2cta_fmha_backward_*.py		FA4 反向主 kernel（标准 + MLA + hd256 2CTA），把 dQ/dK/dV 拆成 DK/DV 与 DQ 两阶段或融合。
入口：flash_bwd, FlashAttentionBackwardSm90/SM100/SM120, BlackwellFusedMultiHeadAttentionBackwardDKDVKernel, BlackwellFusedMultiHeadAttentionBackwardDQKernel
依赖：tile_scheduler.py, mask.py, softmax.py, pipeline.py, copy_utils.py, fast_math.py
最近提交集中在 hd256 backward 统一、deterministic=True、2CTA dQaccum load、S/P ping-pong。
tile_scheduler
flash_attn/cute/tile_scheduler.py
1 大文件		FA4 的 tile 调度核心：static/dynamic/CLC/persistent、varlen、2CTA 命名化调度、与 CLC 硬件调度器（ClcDynamicPersistentTileScheduler）交互。
入口：SchedulerState, ClcSchedulerState, DynamicPersistentSchedulerState, SingleTileVarlenScheduler, Sm100FmhaClcDynamicTileScheduler, Sm100FmhaStaticTileScheduler, compute_sm100_fmha_varlen_grid, SchedulingMode
依赖：cutlass.utils.ClcDynamicPersistentTileScheduler, quack.cute_dsl_utils.ParamsBase, flash_attn.cute.utils
提交热点（named 2CTA schedules、SM100_TMEM_CAPACITY_COLUMNS、varlen grid）直接改这里。
prepare_scheduler
flash_attn/cute/prepare_scheduler.py
1 文件		FA4 的 host 侧调度预处理：把 (batch, seqlen, nheads, tile) 变成 cu_total_m_blocks / cu_total_splits_m_blocks / blocks_to_batch_idx 等元数据 tensor，供 kernel 启动。
入口：FlashPrepareScheduler, SchedulerMetadataTensorsTorch
依赖：flash_attn.cute.utils, cutlass.cute.FastDivmodDivisorV2
注释明确说明是 hopper/flash_prepare_scheduler.cu 的 CuTeDSL 重写版；varlen/paged-kv 路径必过。
flash_attn_ck
csrc/flash_attn_ck/
~7 文件		FA2 的 ROCm Composable Kernel 后端（AMD GPU），提供 mha_fwd/mha_bwd/mha_varlen_fwd/mha_varlen_bwd/mha_fwd_kvcache。
入口：flash_api.cpp 中注册的 mha_fwd/mha_bwd/mha_varlen_fwd/mha_bwd/kvcache
依赖：composable_kernel 子树
README 指定 ROCm CK 为默认后端；与 csrc/flash_attn（CUDA）并列。
hopper_fa3
hopper/
507 文件		FA3（Hopper beta）完整实现：C++ 版的 SM90 TMA/GMMA 前向/反向、softcap、pack-gqa、paged-kv、split-kv、MLA decode、benchmark。
入口：flash_api.cpp, flash_fwd_kernel_sm90.h, flash_bwd_kernel_sm90.h, mainloop_fwd_sm90_tma_gmma_ws.hpp, mainloop_bwd_sm90_tma_gmma_ws.hpp, generate_kernels.py
依赖：csrc/cutlass（老 C++ 子树）, 部分 csrc/flash_attn 共用工具
README 标记为 beta；仍被 H100/H800 用户和 benchmark 使用，但活跃开发已切到 FA4。
csrc_fa2
csrc/flash_attn/
~197 文件（含子树）		FA2（Ampere/Ada）CUDA 实现，仍是 PyTorch 默认 pip install flash-attn 的编译目标。
入口：flash_api.cpp, flash_fwd_kernel.h, flash_bwd_kernel.h, generate_kernels.py, src/flash_fwd_*.cu, src/flash_bwd_*.cu
依赖：csrc/cutlass, csrc/composable_kernel（子树）
与 flash_attn_interface 配套；支持 SM80/SM90（非 TMA 路径）。
model_layers
flash_attn/modules/, flash_attn/layers/, flash_attn/models/
~30 文件		把 flash attention 包装成 nn.Module（MHA、embedding、rotary、block），并提供 GPT/LLaMA/Baichuan/Falcon/BERT/ViT 等模型级替换。
入口：modules/mha.py:MultiheadAttention, layers/rotary.py, models/gpt.py, models/llama.py, models/bert.py, models/vit.py
依赖：flash_attn_interface 或 cute_interface（取决于 import 路径）
训练/推理集成的主要扩展点；新模型通常在这里加一个 models/xxx.py。
benchmarks
benchmarks/
12 文件		性能对标：FA2/FA3/FA4、SM90、MLA paged-kv、sparse MLA fwd、varlen scheduling、ALiBi、GEMM、CLC bench。
入口：benchmark_flash_attention.py, benchmark_mla_paged_kv.py, benchmark_sparse_mla_fwd.py, benchmark_varlen_sched.py, bench_sm90.py, clc_bench.py
依赖：flash_attn_interface, cute_interface, hopper_fa3
README 与发布说明的性能数字来源；新 kernel 提 PR 时通常要补对应 benchmark。
tests
tests/
55 文件		回归测试：FA2/FA3/FA4、Triton-AMD、MLA、varlen、mask/score mod、block sparsity。
入口：tests/cute/test_flash_attn.py, test_flash_attn_mla.py, test_flash_attn_varlen.py, test_mask_mod.py, test_score_mod.py, test_block_sparsity.py, test_flash_attn.py, test_flash_attn_ck.py
依赖：各代接口
AGENTS.md 给出 FakeTensor 两-pass 编译/执行工作流；新扩展必须补 test。
training
training/
145 文件		分布式训练集成：GPT/GPT3 训练脚本、FSDP/TP/PP/EP 配置、Dockerfile。
入口：run.py, configs/, src/
依赖：flash_attn_interface 或 cute_interface
与 Megatron/DeepSpeed 生态衔接；新模型/新分布式策略可在此加配置验证。
utils_and_helpers
flash_attn/utils/, flash_attn/cute/utils.py, flash_attn/cute/copy_utils.py, flash_attn/cute/fast_math.py, flash_attn/cute/barrier.py, flash_attn/cute/named_barrier.py, flash_attn/cute/pipeline.py, flash_attn/cute/mma_sm100_desc.py
~15 文件		公共工具：分布式 wrapper、generation、pretrained、benchmark、CuTeDSL 的 barrier/pipeline/divmod/TMEM/PTX 辅助。
入口：utils/distributed.py, utils/generation.py, cute/utils.py:constexpr 工具, cute/pipeline.py:PipelineStateSimple, cute/barrier.py:ld_acquire/wait_eq/arrive_inc
依赖：cutlass.cute, quack-kernels, apache-tvm-ffi
跨代共享；FA4 的 producer/consumer pipeline 与 named barrier 都在此实现。
aiter_rocm
third_party/aiter
子模块		ROCm Triton 后端子模块，提供 FA2 级别的 Triton kernel（CDNA + RDNA、FP8、sliding-window WIP）。
入口：aiter 包内 Triton kernel（具体符号需验证）
依赖：Triton, PyTorch ROCm
README 与 FLASH_ATTENTION_TRITON_AMD_ENABLE 流程说明其用途；默认不参与 CUDA 构建。
目录树（按文件数）
  • hopper/ 507 个文件
    __init__.py, benchmark_attn.py, benchmark_flash_attention_fp8.py, benchmark_mla_decode.py, benchmark_split_kv.py, block.h, copy_sm90_bulk_reduce.hpp, cuda_check.h, epilogue_bwd.hpp, epilogue_fwd.hpp, flash.h, flash_api.cpp, flash_api_stable.cpp, flash_attn_3
  • csrc/ 197 个文件
    composable_kernel, cutlass, flash_attn, flash_attn_ck, fused_dense_lib, layer_norm
  • training/ 145 个文件
    Dockerfile, README.md, configs, run.py, src, tests
  • flash_attn/ 110 个文件
    __init__.py, bert_padding.py, cute, flash_attn_interface.py, flash_attn_triton.py, flash_attn_triton_og.py, flash_blocksparse_attention.py, flash_blocksparse_attn_interface.py, layers, losses, models, modules, ops, pyproject.toml
  • tests/ 55 个文件
    cute, layers, losses, models, modules, ops, pyproject.toml, test_flash_attn.py, test_flash_attn_ck.py, test_flash_attn_triton_amd.py, test_rotary.py, test_util.py
  • assets/ 17 个文件
    fa4_paper.pdf, flash2_a100_fwd_bwd_benchmark.png, flash2_h100_fwd_bwd_benchmark.png, flash3_fp16_fwd.png, flashattention_logo.png, flashattn_banner.jpg, flashattn_banner.pdf, flashattn_memory.jpg, flashattn_speedup.jpg, flashattn_speedup_3090.jpg, flashattn_speedup_a100_d128.jpg, flashattn_speedup_t4.jpg, flashattn_speedup_t4_fwd.jpg, gpt2_training_curve.jpg
  • AI/ 16 个文件
    CLC_TRACE_DEBUG.md, DEBUG_2CTA.md, DEBUG_METHODOLOGY.md, RACECHECK_TMA_HAZARD.md, SASS_MMA_ANALYSIS.md, SM90_BLOCK_SIZE_TUNING.md, SM90_R2P_MASKING_SASS.md, SPARSE_MLA_1CTA.md, SPARSE_MLA_64H.md, SPARSE_MLA_DPSUM_PRECISION.md, SPARSE_MLA_EXACT_SOFTMAX_MAX.md, SPARSE_MLA_RECOMPUTE_P.md, VARLEN_PREPROCESS_TILE_BUG.md, parse_clc_log.py
  • benchmarks/ 12 个文件
    bench_sm90.py, benchmark_alibi.py, benchmark_attn.py, benchmark_causal.py, benchmark_flash_attention.py, benchmark_gemm.py, benchmark_mla_paged_kv.py, benchmark_sparse_mla_fwd.py, benchmark_varlen_sched.py, clc_bench.py, configs, tune_ex2_emu.py
  • tools/ 11 个文件
    ci, sass_diff.py, two_pass_tests.sh
  • .github/ 10 个文件
    actions, scripts, workflows
  • .gitignore/ 1 个文件
  • .gitmodules/ 1 个文件
  • .pre-commit-config.yaml/ 1 个文件
  • AGENTS.md/ 1 个文件
  • AUTHORS/ 1 个文件
  • CLAUDE.md/ 1 个文件

一次调用怎么流过这些模块

一条典型请求（以 FA4 推理为例）从 Python 入口 flash_attn.cute.interface.flash_attn_func 进入： 1) 接口层（flash_attn/cute/interface.py）把 Q/K/V（shape: batch, seqlen, nheads, headdim）、causal、window_size、alibi_slopes、softmax_scale、dropout_p、deterministic、paged-kv（block_table/page_cache_dtype）、MLA 标志等打包成 kernel 参数对象，并决定走 flash_fwd / flash_fwd_mla / flash_fwd_sm{90,100,120} 中的哪条路径。 2) host 侧调度预处理（flash_attn/cute/prepare_scheduler.py:FlashPrepareScheduler）把 batch/seqlen/nheads/tile 转成 SchedulerMetadataTensorsTorch：cu_total_m_blocks、cu_total_splits_m_blocks、blocks_to_batch_idx、num_splits_dynamic、tile_count_semaphore 等元数据 tensor，并计算 grid shape（CTA 数）。varlen 路径会先把 cu_seqlen 转成 block 索引（flash_attn/cute/cu_blocks_kernel.py:CuSeqlensToBlocksKernel）。 3) kernel 启动（flash_fwd.py）时，tile_scheduler.py 根据 (headdim, causal, varlen, paged_kv, is_mla, 2cta) 选择 SchedulingMode：STATIC / DYNAMIC / CLC / NONE，构造 SchedulerState（ClcSchedulerState 或 DynamicPersistentSchedulerState），建立 producer/consumer PipelineAsync / PipelineClcFetchAsync 与 named barrier。 4) 进入 kernel 主体（flash_fwd_sm100.py 等）：producer warp group 通过 cpasync / TMA 把 Q/K/V tile 从 HBM 搬到 smem，consumer warp group 用 blackwell_helpers.gemm / tcgen05.mma 做 QK^T → mask.py 应用 causal/window/ALiBi → softmax.py 在线 softmax（Flash 的 recompute 技巧）→ 与 V 做 MMA 得到 O partial sum，同时 fast_math.py 与 mma_sm100_desc.py 处理 TMEM 累加与 UMMA 描述符。pack_gqa.py 在 GQA 下把 Q head 折叠进 tile，paged_kv.py 按 block_table 间接取 KV，topk_gather_kv.py 做稀疏 MLA top-k。 5) 反向路径（flash_bwd.py）把 dO 输入，拆成 DK/DV 阶段（BlackwellFusedMultiHeadAttentionBackwardDKDVKernel）与 DQ 阶段（BlackwellFusedMultiHeadAttentionBackwardDQKernel），hd256 时用 2CTA 协作，通过 named_barrier.py / pipeline.py 同步；dQaccum 先写 HBM 再经 postprocess 加载（最近提交优化了“load 2-CTA dQaccum directly into registers”）。 6) 输出 O（以及可选 softmax_lse、rng_state）写回 HBM，Python 层返回给调用方；训练侧则把 O 送入后续 layer_norm/mlp，反向把 dQ/dK/dV 传回 optimizer。 关键数据结构：Q/K/V/O/dO/dQ/dK/dV Tensor；SchedulerMetadataTensorsTorch（cu_total_m_blocks、blocks_to_batch_idx、num_splits_dynamic、tile_count_semaphore）；SchedulerState / ClcSchedulerState（_pipeline, _consumer_state, _producer_state, hw_scheduler）；PipelineStateSimple / PipelineClcFetchAsync；FusedMask / MaskType；paged-kv 的 block_table；MLA 的 top-k 索引；TMEM 累加器与 mbarrier。

1Python 入口flash_attn_func (FA4) · flash_attn/cute/interface.py2调度预处理FlashPrepareScheduler · flash_attn/cute/prepare_scheduler.py3varlen 转 blockCuSeqlensToBlocksKernel · flash_attn/cute/cu_blocks_kernel.py4选调度模式SchedulerState / ClcSchedulerState · flash_attn/cute/tile_scheduler.py5前向 kernelFlashAttentionForwardSm100 · flash_attn/cute/flash_fwd_sm100.py6反向 DK/DVBlackwellFusedMultiHeadAttentionBackward · flash_attn/cute/sm100_hd256_2cta_fmha_backw7反向 DQBlackwellFusedMultiHeadAttentionBackward · flash_attn/cute/sm100_hd256_2cta_fmha_backw8回归验证pytest tests/cute · tests/cute/test_flash_attn.py9性能对标benchmark_mla_paged_kv · benchmarks/benchmark_mla_paged_kv.py

关键类型与函数

名称路径用途
flash_attn_func (FA4)flash_attn/cute/interface.pyFA4 的 Python 入口，把 Q/K/V 等参数派发给 CuTeDSL 前向 kernel。
FlashPrepareSchedulerflash_attn/cute/prepare_scheduler.pyhost 侧调度预处理，生成 cu_total_m_blocks、blocks_to_batch_idx 等元数据。
SchedulerState / ClcSchedulerStateflash_attn/cute/tile_scheduler.pyFA4 的 tile 调度运行时状态，封装 async pipeline 与 CLC 硬件调度器。
FlashAttentionForwardSm100flash_attn/cute/flash_fwd_sm100.py（需验证具体类名）SM100 Blackwell 前向主 kernel，使用 tcgen05 UMMA/TMEM。
BlackwellFusedMultiHeadAttentionBackwardDKDVKernelflash_attn/cute/sm100_hd256_2cta_fmha_backward_dkdvkernel.pyBlackwell hd256 反向 DK/DV 阶段 kernel（2CTA 协作）。
BlackwellFusedMultiHeadAttentionBackwardDQKernelflash_attn/cute/sm100_hd256_2cta_fmha_backward_dqkernel.pyBlackwell hd256 反向 DQ 阶段 kernel。
PipelineStateSimpleflash_attn/cute/pipeline.pyFA4 自定义的轻量 pipeline 状态（index + phase），用于 producer/consumer 同步。
FusedMask / Sm100FusedMaskflash_attn/cute/mask.py封装 causal / window / ALiBi / block-sparsity mask，在 MMA 循环中应用。
CuSeqlensToBlocksKernelflash_attn/cute/cu_blocks_kernel.pyvarlen 下把 cu_seqlen 转成 block 索引的 CuTe kernel。
flash_fwd_kernel_sm90.h (FA3)hopper/flash_fwd_kernel_sm90.hFA3 Hopper 前向 C++ 模板（TMA/GMMA），仍用于 H100/H800 生产。
flash_bwd_kernel_sm90.h (FA3)hopper/flash_bwd_kernel_sm90.hFA3 Hopper 反向 C++ 模板。
mha_fwd / mha_bwd (CK)csrc/flash_attn_ck/mha_fwd.cpp, mha_bwd.cppROCm Composable Kernel 后端的前向/反向实现。
MultiheadAttentionflash_attn/modules/mha.pynn.Module 级 MHA 封装，供训练/推理模型直接替换 torch.nn.MultiheadAttention。
flash_attn_func (FA2)flash_attn/flash_attn_interface.pyFA2 的 Python → C++ 入口，通过 torch.ops 调用 csrc/flash_attn 编译的 CUDA op。

扩展点

最近在动的地方

建议阅读顺序

  1. CLAUDE.md 或 AGENTS.md（项目总览、构建、测试、active development 在 flash_attn/cute/）
  2. flash_attn/cute/__init__.py + flash_attn/cute/interface.py（FA4 入口）
  3. flash_attn/cute/flash_fwd.py + flash_attn/cute/flash_fwd_sm100.py（前向主 kernel）
  4. flash_attn/cute/tile_scheduler.py + prepare_scheduler.py（调度层）
  5. flash_attn/cute/mask.py + softmax.py + pipeline.py + named_barrier.py（执行层原语）
  6. flash_attn/cute/flash_bwd.py + sm100_hd256_2cta_fmha_backward_*.py（反向与 2CTA）
  7. flash_attn/cute/flash_fwd_mla_sm100.py + flash_bwd_mla_sm100.py + topk_gather_kv.py（MLA / 稀疏）
  8. tests/cute/test_flash_attn.py + test_flash_attn_mla.py + test_flash_attn_varlen.py（回归测试）
  9. benchmarks/benchmark_flash_attention.py + benchmark_mla_paged_kv.py + benchmark_sparse_mla_fwd.py（性能对标）
  10. hopper/flash_fwd_kernel_sm90.h + flash_bwd_kernel_sm90.h（FA3 Hopper 路径，与 FA4 对照）

三、本地跑起来（没有 GPU 的 Mac）

安装

  1. # 1. 创建 Python 虚拟环境（推荐 Python 3.10–3.12，与 PyTorch 2.2+ 兼容）
  2. python3 -m venv fa4_env && source fa4_env/bin/activate
  4. # 2. 安装 PyTorch CPU/MPS 版本（Apple Silicon 用 MPS 后端，无 CUDA）
  5. pip install torch torchvision torchaudio
  7. # 3. 安装 FA4 依赖（跳过 CUDA 编译，仅安装 Python 包用于读代码/静态分析）
  8. # 注意：nvidia-cutlass-dsl 和 quack-kernels 需要 NVIDIA GPU 才能实际运行 kernel
  9. # 但安装它们可以导入 flash_attn.cute 模块进行代码阅读与静态检查
  10. pip install nvidia-cutlass-dsl>=4.6.2 || echo 'WARNING: nvidia-cutlass-dsl install may fail without CUDA'
  11. pip install einops apache-tvm-ffi quack-kernels>=0.5.3 || echo 'WARNING: quack-kernels may need CUDA'
  13. # 4. 开发模式安装 flash_attn/cute（不编译 CUDA kernel）
  14. pip install -e "flash_attn/cute[dev]" || pip install -e flash_attn/cute
  16. # 5. 验证安装（无 GPU 时只能验证 Python 导入，不能执行 kernel）
  17. python -c "from flash_attn.cute import flash_attn_func; print('Import OK')"
  19. # 替代方案：若 nvidia-cutlass-dsl 安装失败，仅安装纯 Python 依赖用于代码阅读
  20. pip install torch einops packaging psutil ninja
  22. # 注意：FA2 (csrc/) 和 FA3 (hopper/) 需要 CUDA toolkit + nvcc，Mac 上完全无法编译
  23. # 若只想阅读代码，直接 git clone 即可，无需安装

哪些路径能真跑

['可执行路径（CPU/MPS）：', '- flash_attn/flash_attn_triton.py — Triton 实现，理论上可在 CPU 模拟但极慢，实际需 GPU', '- flash_attn/flash_attn_triton_og.py — 原始 Triton 实现，同上', '- flash_attn/ops/triton/*.py — Triton 算子（cross_entropy, layer_norm, linear, mlp, rotary），需 GPU', '- flash_attn/models/*.py — 模型定义（GPT, LLaMA, BERT 等），可 import 但推理需 CUDA', '- flash_attn/modules/*.py — MHA, MLP, Embedding, Block 模块，可 import 但 forward 需 CUDA', '- flash_attn/layers/*.py — Rotary, PatchEmbed，可 import', '- flash_attn/losses/*.py — CrossEntropy loss，可 import', '- flash_attn/utils/*.py — 工具函数，部分可 CPU 运行', '', '仅可读代码路径（必须 NVIDIA GPU 才能执行）：', '- flash_attn/cute/*.py — FA4 CuTe DSL kernel，编译到 PTX/CUBIN，需 SM90/SM100/SM120 GPU', '- flash_attn/cute/flash_fwd*.py, flash_bwd*.py — 前向/反向 kernel', '- flash_attn/cute/sm100_hd256_2cta_fmha_backward_*.py — SM100 hd256 2CTA 反向', '- flash_attn/cute/tile_scheduler.py, prepare_scheduler.py — 调度器', '- csrc/flash_attn/src/*.cu — FA2 CUDA kernel（SM80/SM90），需 nvcc + GPU', '- csrc/flash_attn_ck/* — Composable Kernel backend，需 ROCm 或 CUDA', '- hopper/*.h, *.cu, instantiations/*.cu — FA3 Hopper kernel，需 H100/H800 + CUDA >= 12.3', '- csrc/fused_dense_lib/*, csrc/layer_norm/* — 融合算子，需 CUDA', '', 'Colab T4 验证路径（T4 为 SM75，支持 FA2 不支持 FA4）：', '- FA2 (csrc/) 可在 T4 上运行，但性能不如 A100/H100', '- FA4 (flash_attn/cute/) 需要 SM90+（H100/B200），T4 不支持', '- FA3 (hopper/) 需要 SM90，T4 不支持']

最小可运行

  1. # 冒烟测试 1：验证 Python 导入（无 GPU，仅检查模块结构）
  2. python -c "from flash_attn.cute import flash_attn_func, flash_attn_varlen_func; print('FA4 interface import OK')"
  4. # 冒烟测试 2：验证 FA2 接口导入（csrc/ 路径，需 CUDA 编译后才能执行）
  5. python -c "from flash_attn import flash_attn_func; print('FA2 interface import OK')"
  7. # 冒烟测试 3：验证模型模块可 import（不执行 forward）
  8. python -c "from flash_attn.models import llama; print('LLaMA model import OK')"
  10. # 冒烟测试 4：验证 Triton ops 可 import（不执行 kernel）
  11. python -c "from flash_attn.ops.triton import rotary; print('Triton rotary import OK')"
  13. # 冒烟测试 5：在 Colab T4 上运行 FA2 测试（需 GPU）
  14. # 仅验证 FA2，FA4 需要 SM90+ GPU
  15. # pytest tests/test_flash_attn.py -k "test_flash_attn_output" -x --timeout=300
  17. # 冒烟测试 6：验证 FA4 代码路径可读（无 GPU 时只能语法检查）
  18. # python -c "import ast; ast.parse(open('flash_attn/cute/flash_fwd.py').read()); print('FA4 syntax OK')"

测试

['测试框架：pytest + pytest-xdist（并行测试）', '', '主要测试目录：', '- tests/cute/ — FA4 CuTe DSL 测试（17 次提交热点）', ' - test_flash_attn.py — 标准 attention 测试', ' - test_flash_attn_mla.py — MLA（Multi-Latent Attention）测试', ' - test_flash_attn_varlen.py — 变长序列测试', ' - test_mask_mod.py, test_score_mod.py — mask/score 修饰符测试', ' - test_block_sparsity.py — 块稀疏测试', '- tests/test_flash_attn.py — FA2 测试', '- tests/test_flash_attn_ck.py — Composable Kernel backend 测试', '- tests/test_flash_attn_triton_amd.py — Triton AMD 测试', '- tests/layers/, tests/modules/, tests/models/, tests/ops/ — 模块级测试', '', 'CPU 子集测试（无 GPU 时可运行）：', '# 纯 Python 工具测试（不依赖 CUDA）', 'pytest tests/test_util.py -x', '', '# 模型结构测试（import 检查，不执行 forward）', 'pytest tests/models/ -x --co -q # 仅收集测试，不执行', '', '# 代码语法检查（自定义脚本）', 'python -c "import ast; [ast.parse(open(f).read()) for f in [\'flash_attn/cute/flash_fwd.py\', \'flash_attn/cute/flash_bwd.py\']]"', '', '完整测试（需 GPU）：', '# 两-pass 测试（编译 + 执行分离，见 CLAUDE.md）', 'FLASH_ATTENTION_FAKE_TENSOR=1 FLASH_ATTENTION_CUTE_DSL_CACHE_ENABLED=1 pytest -n 64 -x tests/cute/test_flash_attn.py', 'FLASH_ATTENTION_FAKE_TENSOR=0 FLASH_ATTENTION_CUTE_DSL_CACHE_ENABLED=1 pytest -x tests/cute/test_flash_attn.py', '', '# 单测试示例', 'pytest tests/cute/test_flash_attn.py -k "test_flash_attn_output" -x', '', '耗时估算：', '- 完整 FA4 测试套件：数小时（GPU 上）', '- 编译 pass（FakeTensorMode）：30 分钟–1 小时（64 workers）', '- 执行 pass（cached kernels）：10–30 分钟', '- CPU 子集（仅 import 检查）：< 1 分钟']

调试

CI

GitHub Actions（.github/workflows/），CI 矩阵覆盖：

坑

四、维护者与社区

高频迭代：最近三周（2026-09-23 → 2026-10-07）连续发布 FA4 beta32 → beta33 → beta34，对应 releases 时间戳 2026-09-23T08:21、2026-09-30T08:22、2026-10-07T08:23，几乎一周一版。提交热点集中在 flash_attn/cute（37 次）与 tests/cute（17 次），说明 FA4 内核与测试是当前开发主线；hopper/ 与 csrc/ 已不再是热点。

谁角色依据
Tri Dao项目创始人 / 方向决策者README 与 AUTHORS 中论文署名（FA1/FA2/FA3 第一作者）；setup.py 头部注释 Copyright (c) 2023, Tri Dao；仓库归属 Dao-AILab。
tzadouriFA4 CuTe 内核 / 调度器主力Issue 2456（[RFC] FA4 head_dim=256/512）assignees 中出现；多轮 hd256 前向/反向、MLA、2CTA 相关提交与 PR 作者。
JohnsonmsFA4 数值精度 / MLA / 稀疏注意力Issue 2456 assignees；Issue 2969（dsink loses precision: dpsum uses bf16-rounded O）assignee；SPARSE_MLA 系列 AI/ 文档与 sparse_mla 提交。
micmelesseROCm / CK / aiter 后端维护Issue 2933（[ROCm/CK] gfx950 asm backward）assignee；多个 ROCm/CK 相关 PR 与提交。
社区贡献者（head_dim=256、SM100/SM120、Windows、Triton 后端等）特性 / 平台扩展 / Bug 修复Issue 2581（hd=512 on SM89）、Issue 2413（RTX 6000 Pro Blackwell）、Issue 2307（SM120 support）、PR #2872（Windows builds with CUDA 13）等。

流程与 Review 风格

仓库没有独立的 CONTRIBUTING.md（contributing 字段为空）。CLAUDE.md 与 AGENTS.md 提供开发指引，强调“active development is on FA4 in flash_attn/cute/”、依赖 nvidia-cutlass-dsl>=4.6.2、torch、einops、apache-tvm-ffi、quack-kernels>=0.5.3，并给出两-pass 测试流程（FLASH_ATTENTION_FAKE_TENSOR=1 编译 + 缓存执行）。.pre-commit-config.yaml 存在，配置 ruff 等代码格式检查（见提交 'ruff format (#2940)'）。没有看到 CLA/DCO 要求。Issue 模板与 PR 流程未在证据中显式描述，但常见做法是先开 Issue 描述 bug/feature，再提 PR 关联 issue number（如 #2978、#2964、#2951）。

技术导向、快速响应。从最近 PR 看，review 关注点集中在：数值精度（Issue 2969 dsink precision、Issue 2946 finite-sink overflow）、内存正确性（Issue 2987 illegal memory access、Issue 2860 varlen illegal memory access、Issue 2957 paged KV wrong output）、性能优化（PR #2990 overlap loads with epilogue、PR #2989 LPT scheduler L2 sizing、Issue 2426 Python dispatch overhead）。维护者会直接 assign 给具体负责人（如 Johnsonms、micmelesse、tzadouri），说明责任分工明确。对 RFC 类 Issue（如 #2456 hd256/512）会持续讨论多轮。

渠道

这里的规矩

维护者现在最想要的帮助

五、切入方案

建议长期负责：FA4 调度层（flash_attn/cute/tile_scheduler.py + prepare_scheduler.py）与 MLA/稀疏 MLA 路径的“设计契约文档 + 回归测试 + 微基准”三件套
tile_scheduler.py 是提交热点（与 37 次 flash_attn/cute 提交直接相关），且正在经历 2CTA 命名化 (#2964)、S/P ping-pong (#2928)、CLC dynamic 调度三条主线演进；调度器输出是纯 Python 数据结构（SchedulerState、ClcSchedulerState、DynamicPersistentSchedulerState），可在 CPU 上独立测试与 benchmark。MLA/稀疏 MLA 路径（flash_fwd_mla_sm100.py、topk_gather_kv.py、sparse_mla 系列）是另一条提交热点，且精度问题（Issue 2969、AI/SPARSE_MLA_*.md）迫切需要“文档化设计决策 + 回归锚点”。两者结合，既能深入 CuTe DSL 核心，又能在无 GPU 条件下通过 Python 模拟、测试、文档产出持续贡献，通向 reviewer 身份的路径清晰。

它现在缺什么（你无 GPU 也能补）

缺口依据为什么是你
稀疏 MLA 数值精度文档与回归测试覆盖不足AI/SPARSE_MLA_DPSUM_PRECISION.md、AI/SPARSE_MLA_EXACT_SOFTMAX_MAX.md、AI/SPARSE_MLA_RECOMPUTE_P.md 表明稀疏 MLA 的 dpsum 精度、softmax max 计算、P recompute 策略是最近才梳理清楚的设计决策；Issue 2969（dsink loses precision: dpsum uses bf16-rounded O）说明该问题曾导致精度回退。tests/cute/ 中虽有 test_flash_attn_mla.py，但针对 sparse top-k + bf16-rounded O 的回归测试分散在多个测试文件中，缺少一条“精度契约文档 + 对应 pytest 回归锚点”的闭环。这是典型的“设计决策散落 AI/ 目录、未沉淀为可执行契约”的缺口，纯 Python/CPU 即可推动：整理 AI/*.md 为 docs/ 下的设计备忘录，补充 pytest 锚点（用 torch 在 CPU 上验证 bf16-rounded O 的数值边界），不依赖 GPU 即可验证逻辑正确性。
FA4 调度策略（CLC / 2CTA / ping-pong）缺少可独立运行的 benchmark 微基准提交热点显示 S/P ping-pong (#2928)、2CTA 调度命名化 (#2964)、hd256 前后向统一是 2026-10 的主旋律；benchmarks/ 目录仅有 benchmark_varlen_sched.py、bench_sm90.py 等少数脚本，且多数依赖真实 GPU 才能跑出有意义的数据。缺少一套“在 CPU 上用 FakeTensor / 模拟 tile 延迟来比较调度策略（static vs CLC vs 2CTA）优劣”的微基准。调度策略比较本质是 tile → CTA 的映射逻辑，可在 CPU 上用 Python 复现调度器输出（tile_scheduler.py 的输出结构），对比不同 scheduler 在 synthetic workload 下的 load balance 与 pipeline bubble。纯 Python 工作，产出却是性能团队高度依赖的决策依据。
FA2/FA3 旧代代码与 FA4 之间的行为一致性对照表缺失仓库同时保留 csrc/（FA2）、hopper/（FA3）、flash_attn/cute/（FA4）三代实现，CLAUDE.md 明确指出“active development is on FA4”，但 Issue 2456（RFC: FA4 head_dim=256/512）等讨论中频繁出现“FA4 行为需与 FA2/FA3 对齐”的诉求。目前没有一份文档列出 (causal, softcap, GQA, paged-kv, varlen) × (fp16, bf16, fp8) 组合下三代实现的行为差异与已知偏差。整理对照表只需读代码 + 现有测试用例，无需 GPU。产出可直接作为 Issue 2456 等 RFC 的决策附件，也能帮助训练侧（training/）用户迁移时评估风险。
training/ 侧分布式训练集成缺少 FA4 适配的性能回归基准training/ 目录含 145 个文件，覆盖 GPT/LLaMA 训练；但 FA4 的 Python 接口（flash_attn/cute/interface.py）与 FA2 的 C++ 接口（flash_attn/flash_attn_interface.py）在 varlen、GQA、MLA 路径上存在差异（参见 flash_attn_interface.py 与 interface.py 的签名对比）。目前缺少一套“在 CPU 上跑通端到端训练 step、对比 FA2 vs FA4 接口行为与计时”的回归脚本。可在 CPU 上用 torch 的 dummy tensor 跑通 training/ 的 run.py 数据流，验证 FA4 接口在训练路径（forward + backward + gradient accumulation）中的行为一致性。纯 PyTorch CPU 工作，产出是训练侧迁移的安全网。
Triton 后端（flash_attn_triton.py）与 CuTe DSL 前向的数值一致性测试缺失flash_attn/flash_attn_triton.py 是仓库自带的 Triton 实现，README 提到 Triton 后端支持 AMD ROCm；但 tests/ 中没有直接对比 Triton 与 CuTe DSL 前向输出数值一致性的测试。Issue 2969 涉及精度问题，若有 Triton ↔ CuTe 的交叉验证，可快速定位是算法问题还是实现问题。Triton kernel 可在 CPU 上用 Triton 的 CPU backend 或 interpreter 模式运行（或直接用 PyTorch reference 替代），与 CuTe DSL 的 CPU 模拟输出做对比。纯 Python/CPU 可完成，产出是双后端一致性的安全网。
CI 中 Windows + CUDA 13 路径的自动化覆盖薄弱PR #2872（Fix Windows builds with CUDA 13 and PyTorch 2.15）于 2026-09-30 合并，说明 Windows 支持是近期才补齐的；.github/workflows/ 中需验证是否覆盖 Windows + CUDA 13 + PyTorch 2.15 的组合。Windows 用户（Issue 2307 SM120、Issue 2413 RTX 6000 Pro Blackwell）在增长。CI 配置审查与补充是纯文本工作，无需 GPU。可通过阅读 .github/workflows/ 现有 YAML、对照 setup.py 的 Windows 分支逻辑，提出补全 Windows + CUDA 13 矩阵的 PR。
FA4 的 head_dim=256/512 支持状态与已知限制未形成单一权威文档Issue 2456（RFC: FA4 head_dim=256/512）仍在讨论，最近提交 #2964、#2962、#2951 在推进 hd256 前向/反向统一；但 README 与 docs/ 中没有一处集中列出“hd=128/192/256/512 在各代（FA2/FA3/FA4）× 各架构（SM80/SM90/SM100/SM120）的支持状态与已知限制”。用户在 Issue 2581（hd=512 on SM89）等平台相关提问中反复遇到信息分散的问题。整理支持矩阵只需读代码（csrc/flash_attn/src/ 的 .cu 文件名模式、flash_attn/cute/flash_fwd_sm*.py 的存在性）+ 现有测试参数化列表，无需 GPU。产出直接服务社区，减少重复 Issue。
FA4 的 quack-kernels / apache-tvm-ffi 依赖升级路径缺少兼容性矩阵CLAUDE.md 列出依赖 nvidia-cutlass-dsl>=4.6.2、quack-kernels>=0.5.3、apache-tvm-ffi，最近提交 #2948 允许 apache-tvm-ffi 0.1.11；但 setup.py 与各 flash_attn/cute/*.py 中对这些依赖的版本敏感（例如 tile_scheduler.py 引用 quack.cute_dsl_utils.ParamsBase）。缺少一份“依赖版本 × PyTorch 版本 × CUDA 版本”的兼容性矩阵与 CI 测试覆盖。可在 CPU 上用 pip install + import 测试不同依赖组合下的 API 兼容性（ParamsBase 签名、cutlass.cute API 变化），纯 Python 虚拟环境工作。产出是 release 工程师（一周一 beta）的决策依据。

第 1–30 天：看懂并露面

第 31–60 天：稳定产出

第 61–90 天：接管一块

第一批 PR

题目范围为什么安全
[Docs] Consolidate sparse MLA precision design notes into docs/design/sparse-mla-numerics.md将 AI/SPARSE_MLA_DPSUM_PRECISION.md、AI/SPARSE_MLA_EXACT_SOFTMAX_MAX.md、AI/SPARSE_MLA_RECOMPUTE_P.md 合并为一份权威文档，补充 Issue 2969 的结论与回归测试锚点纯文档工作，不触及 kernel 代码；引用已有的 AI/*.md 和 Issue 2969 证据，风险为零；产出直接服务精度敏感的 MLA 用户
[Tests] Add bf16-rounded O regression anchors in test_flash_attn_mla.py在 tests/cute/test_flash_attn_mla.py 中补充针对 Issue 2969（dsink loses precision: dpsum uses bf16-rounded O）的回归测试，用 torch CPU 验证边界值仅新增测试用例，不修改现有 kernel；测试用 CPU 即可运行，不依赖 GPU；失败只会暴露已有问题，不会引入新回归
[Benchmarks] Add benchmark_schedulers.py for CPU-side scheduling strategy comparison在 benchmarks/ 新增脚本，用 synthetic workload 在 CPU 上比较 STATIC / DYNAMIC / CLC 三种 tile_scheduler 的 load balance 与 pipeline bubble独立新文件，不影响现有代码；纯 Python 逻辑，无 GPU 依赖；输出是决策数据，不改变任何 kernel 行为
[Docs] Add FA2/FA3/FA4 behavior comparison matrix (causal, GQA, paged-kv, varlen, MLA, sparse)docs/matrix-attention-behavior.md 列出三代实现在 (causal, softcap, GQA, paged-kv, varlen, MLA, sparse) × (fp16, bf16, fp8) 的行为差异与 Issue 链接纯文档 + 代码阅读产出；不修改任何实现；引用已有测试与 Issue 作为证据；用户可独立验证
[CI] Extend Windows + CUDA 13 matrix in .github/workflows/基于 PR #2872 的修复，在 CI 配置中补全 Windows + CUDA 13 + PyTorch 2.15 的测试矩阵仅修改 CI YAML，不触及源码；失败只会暴露已有 Windows 构建问题；PR #2872 已证明该方向正确

怎么知道自己站住了

风险与对策
  • 无 GPU 导致无法验证 kernel 在真实硬件上的数值正确性：所有 kernel 级 PR 均附带 Colab T4 链接作为证据；在 PR 描述中明确标注“CPU 逻辑验证通过，GPU 验证由 reviewer/CI 完成”；优先做文档/测试/基准/调度器 CPU 模拟，kernel 改动留给有 GPU 的合作者
  • FA4 API 仍在收敛（beta32→beta34 一周一版），文档与测试可能因接口变更而失效：在 PR 中明确标注基于的 fa4 版本号（如 fa4-v4.0.0.beta34）；优先做“设计契约”而非“API 绑定”级别的文档；关注 setup.py 与 flash_attn/cute/interface.py 的变更，及时调整
  • 核心维护者（tzadouri、Johnsonms）活跃但 PR review 队列长，贡献可能被淹没：从“小、可独立合并、不阻塞主线”的 PR 开始（文档、测试、基准）；在 Issue 中先提供有价值的诊断信息再提 PR；用 GitHub Discussion 提前公开思路，获取非正式 feedback
  • Mac Apple Silicon 16GB 内存受限，跑大型测试或依赖安装可能 OOM：用虚拟环境隔离依赖；对大型测试用 pytest -k 过滤子集；依赖安装失败时回退到“仅阅读代码 + 静态分析”模式；Colab 仅用于最终验证，控制 T4 使用时间
  • FA2/FA3 旧代代码（csrc/、hopper/）与 FA4 的 C++/CuTe DSL 知识门槛高，阅读速度慢：优先聚焦 FA4（flash_attn/cute/），旧代代码仅做“行为对照”级别的浅读；利用 AI/*.md 目录中的设计笔记加速理解；在个人笔记中建立“FA4 概念 ↔ FA2/FA3 实现”的映射表
  • Issue 2456 等 RFC 仍在讨论中，过早整理“支持矩阵”可能与最终决策冲突：在文档中明确标注“当前状态基于 commit X / beta Y，可能随 RFC 演变而更新”；将支持矩阵设计为可增量更新的表格；在 RFC 讨论中主动提供“当前事实”而非“最终结论”
  • Triton 后端（flash_attn_triton.py）在 Mac 上无法运行，交叉验证受限：用 PyTorch reference 实现替代 Triton CPU backend；在 PR 中明确标注“Triton 端验证由 CI/ reviewer 完成”；优先验证 CuTe DSL 端的数值边界
  • CI 中 Windows 配置复杂，PR 可能因 CI 矩阵扩展而暴露已有 Windows 构建问题：在 PR 中仅添加“允许失败”（continue-on-error）的 Windows 作业，避免阻塞其他 CI；在 PR 描述中附上 Windows 构建日志供维护者参考；与 PR #2872 的作者（如有）私下沟通确认方向

六、怎么介入这个项目

建议顺序

成为长期维护者的路径
  • Dao-AILab 维护者（Tri Dao 等）在 CLAUDE.md / AGENTS.md 明确指定 FA4 为活跃开发路径
  • 每周一 beta 发版（fa4-v4.0.0.beta3x），说明对回归测试与契约文档需求迫切
  • Issue 2969（dsink 精度）已有 assignee Johnsonms 和多个 open PR，说明精度是热点但文档化不足
  • Issue 2456（RFC hd256/512）指派 Johnsonms、tzadouri，有多个 open PR，说明大 head_dim 是长期方向
  • CI 工具链（tools/two_pass_tests.sh、.github/scripts/bump_beta_tag.py）显示对自动化测试投入大

七、任务卡

任务 1优先 · medium · CPU · 2-3 个晚上

[Docs] 合并稀疏 MLA 精度设计笔记并补回归锚点

Issue #2969 · [CuTe, SM100] dsink loses precision: dpsum uses the bf16-rounded O ↗

已指派 Johnsonms已有 PR #29680 条评论更新 2026-10-05

用到的专长：利用算子数值正确性专长，将散落的精度决策整理为可执行契约

目标：将 AI/SPARSE_MLA_DPSUM_PRECISION.md、AI/SPARSE_MLA_EXACT_SOFTMAX_MAX.md、AI/SPARSE_MLA_RECOMPUTE_P.md 合并为 docs/design/sparse-mla-numerics.md，并在 tests/cute/test_flash_attn_mla.py 新增 bf16-rounded O 边界值的 pytest 锚点。

为什么值得长期做：稀疏 MLA 是当前攻坚方向（AI/SPARSE_MLA_*.md 密集更新），但精度决策散落各处。整理为权威文档 + 回归测试是 FA4 beta 发版的安全网，直接通向 ownership_target 中的'设计契约文档'职责。

怎么介入：Issue #2969 已有 assignee Johnsonms 和多个 open PR（#2968/#2967/#2965/#2966/#2971），但无文档化 PR。建议以'文档 + 回归测试'切入，不与 kernel 修复冲突。在 Issue 下评论：'我能否将 AI/ 笔记整理为正式文档并补回归锚点？'
第一个 PR 的边界：第一个 PR 仅包含 docs/design/sparse-mla-numerics.md + tests/cute/test_flash_attn_mla.py 新增 2-3 个 pytest 用例，不修改任何 kernel 代码。
第一步：读取 AI/ 目录下三篇 .md 与 Issue #2969 正文，提炼 dpsum 精度、softmax max 计算、P recompute 策略三条设计决策及其数学依据。
本机怎么复现 / 验证：在 Mac 上 git checkout main && pytest tests/cute/test_flash_attn_mla.py -k mla -x 确认现有测试框架可在 CPU 运行；新增锚点用 torch.bfloat16 CPU tensor 构造 sink=0/±极值 case。
认领留言（英文，可直接贴到 Issue）
Issue #2969's precision findings are scattered across AI/*.md. I'd like to consolidate them into docs/design/sparse-mla-numerics.md and add bf16-rounded O regression anchors in tests/cute/test_flash_attn_mla.py. This won't touch the kernel fixes (PRs #2968/#2971 etc.) — just document their conclusions. Should I keep the AI/ originals or replace them? What bf16 ulp tolerance do you prefer for the anchors? I can have a PR ready in 2-3 evenings.
大致实施方案
  • 在 docs/ 下创建 design/sparse-mla-numerics.md，按'问题→决策→数学依据→回归锚点'四段式组织
  • 引用 Issue #2969 的失败日志与 PR #2971 的修复方向作为证据
  • 在 tests/cute/test_flash_attn_mla.py 新增 test_mla_dsink_bf16_rounding_anchor，用 torch CPU 构造边界 case（sink=0/负/正、bf16-rounded O 极值）
  • 本地 pytest 验证新增测试在 CPU 下通过
  • PR 标题 [Docs] Consolidate sparse MLA precision design notes + regression anchors
可能涉及的目录或文件
  • docs/design/ 需新建目录
  • tests/cute/test_flash_attn_mla.py 已有 mla 测试框架可复用
  • AI/ 目录笔记维护者可能保留原件或删除——需在 PR 中说明
验收方式
  • pytest tests/cute/test_flash_attn_mla.py -k dsink_bf16 在 Mac CPU 上通过
  • docs/design/sparse-mla-numerics.md 被 PR review 确认覆盖 Issue #2969 所有结论
  • 新增测试在 bf16-rounded O 边界值上 fail-without-fix / pass-with-fix
开工前问题与风险

向维护者确认

  • 维护者希望 AI/ 目录原件保留还是迁移？
  • 回归锚点的容差应设为多少 bf16 ulp？
  • 是否需要同步更新 CLAUDE.md 引用新文档路径？

风险

  • 若维护者认为 AI/ 是内部笔记不愿公开，需调整文档定位
  • bf16-rounded O 边界值构造不当可能测试不稳定
任务 2可选 · medium · CPU · 1-2 个晚上

[Tests] 为 #2986 新增 pre-softmax 最大 logit CPU 参考实现与接口契约测试

Issue #2986 · Expose maximum pre-softmax attention logit from FlashAttention forward pass ↗

无人认领0 条评论更新 2026-10-10

用到的专长：利用 attention 算法与数值计算专长定义精确的 CPU 参考实现

目标：在 tests/cute/ 新增 test_premax_logit_contract.py，用 PyTorch CPU 实现 pre-softmax 最大 logit 的参考计算（含 causal/mask/GQA），定义期望的 API 契约（返回形状、dtype、可选性）。

为什么值得长期做：用户希望复用 kernel 内已有的 running row maxima 做训练监控。此 Issue 无 assignee、无 open PR，是纯需求定义阶段。先提供 CPU 参考实现 + 接口契约测试，可帮助维护者确定 API 形状（return_max_logit? max_logit_per_head?），是 kernel 实现前的必要铺垫。

怎么介入：Issue #2986 无 assignee、无 open PR、0 条评论，是空白地带。先提供 CPU 参考实现 + API 建议，在 Issue 下评论征求意见，再决定是否推进 kernel 实现。
第一个 PR 的边界：第一个 PR 仅含 tests/cute/test_premax_logit_contract.py（新增），不修改任何接口或 kernel。
第一步：分析 Issue #2986 需求：用户要的是 max(QK^T) 而非 log-sum-exp，需支持 causal/padding/GQA。读取 flash_attn/cute/softmax.py 确认 kernel 是否已有 row maxima 中间值。
本机怎么复现 / 验证：在 Mac 上 pytest tests/cute/test_premax_logit_contract.py -x 即可运行，全部用 torch.float32/bfloat16 CPU tensor，无需 GPU。
认领留言（英文，可直接贴到 Issue）
Issue #2986: I'd like to help define the API contract for exposing pre-softmax max logits. I've written a CPU reference implementation (tests/cute/test_premax_logit_contract.py) covering causal/GQA/ALiBi/window cases. My suggestion: add return_max_logit: bool = False to flash_attn_func, returning per-head maxima of shape (batch, nheads). This test file can serve as the contract for the future kernel implementation. Does this API shape work for you? I can PR the tests first.
大致实施方案
  • 在 tests/cute/test_premax_logit_contract.py 实现 reference_max_pre_softmax_logit(q, k, causal, window_size, alibi_slopes)
  • 构造 3-5 个边界 case：全 causal、GQA、ALiBi、sliding window、混合
  • 用 torch.testing.assert_allclose 定义容差
  • 在 Issue #2986 下评论：提供参考实现 + 建议 API 形状（如 return_max_logit: bool = False, per_head: bool = True）
  • 若维护者认可，PR 仅含测试文件，留待后续 kernel 实现
可能涉及的目录或文件
  • tests/cute/ 已有 test_flash_attn.py 可参考 pytest 风格
  • API 契约需与 flash_attn/cute/interface.py 的 flash_attn_func 签名对齐
验收方式
  • pytest tests/cute/test_premax_logit_contract.py 在 Mac CPU 通过
  • 覆盖 causal/GQA/ALiBi/window 四种 mask 组合
  • 测试文档明确说明期望的返回形状与 dtype
开工前问题与风险

向维护者确认

  • 维
  • 护
  • 者
  • 倾
  • 向
  • r
  • e
  • t
  • u
  • r
  • n
  • _
  • m
  • a
  • x
  • _
  • l
  • o
  • g
  • i
  • t
  • =
  • T
  • r
  • u
  • e
  • 返
  • 回
  • 全
  • 局
  • 最
  • 大
  • ，
  • 还
  • 是
  • p
  • e
  • r
  • -
  • h
  • e
  • a
  • d
  • 最
  • 大
  • ？
  • 是
  • 否
  • 要
  • 同
  • 时
  • 返
  • 回
  • t
  • o
  • k
  • e
  • n
  • -
  • p
  • a
  • i
  • r
  • 位
  • 置
  • ？
  • 是
  • 否
  • 只
  • 在
  • f
  • o
  • r
  • w
  • a
  • r
  • d
  • 暴
  • 露
  • （
  • b
  • a
  • c
  • k
  • w
  • a
  • r
  • d
  • 不
  • 需
  • 要
  • ）
  • ？

风险

  • 若
  • k
  • e
  • r
  • n
  • e
  • l
  • 内
  • r
  • o
  • w
  • m
  • a
  • x
  • i
  • m
  • a
  • 并
  • 未
  • 以
  • 可
  • 导
  • 出
  • 形
  • 式
  • 存
  • 在
  • ，
  • 后
  • 续
  • 实
  • 现
  • 需
  • 额
  • 外
  • s
  • h
  • a
  • r
  • e
  • d
  • m
  • e
  • m
  • o
  • r
  • y
  • ，
  • 但
  • 测
  • 试
  • 契
  • 约
  • 不
  • 受
  • 影
  • 响
任务 3可选 · low · CPU · 2-3 个晚上

[Docs] 新增 FA2/FA3/FA4 行为一致性对照矩阵

Issue #2456 · [RFC] FA4 — head_dim=256 & head_dim=512 ↗

已指派 Johnsonms, tzadouri已有 PR #15 条评论更新 2026-09-29

用到的专长：利用跨代 kernel 代码阅读专长快速建立功能映射

目标：创建 docs/matrix-attention-behavior.md，列出三代实现在各功能组合下的支持状态（✅/❌/⚠️）与已知偏差，附 Issue/PR 链接。

为什么值得长期做：Issue #2456 是 RFC，正文频繁出现'FA4 需与 FA2/FA3 对齐'。当前无一份文档列出 (causal, softcap, GQA, paged-kv, varlen, MLA, sparse) × (fp16, bf16, fp8) 组合下三代实现的差异。这是纯代码阅读 + 测试参数化列表的产出，无 GPU 依赖。

怎么介入：Issue #2456 有 assignee（Johnsonms、tzadouri）和多个 open PR，但无人做文档整理。以'纯文档贡献'切入，在 Issue 下评论：'我整理了行为对照矩阵草稿，能否帮你们确认准确性？'不与 kernel 实现冲突。
第一个 PR 的边界：第一个 PR 仅含 docs/matrix-attention-behavior.md（新增），不修改任何代码。
第一步：读取 csrc/flash_attn/src/ 的 .cu 文件名模式（推断支持的 head_dim/dtype）、flash_attn/cute/flash_fwd_sm*.py 的存在性、hopper/ 的 README 标注，建立支持矩阵骨架。
本机怎么复现 / 验证：在 Mac 上无需运行代码，纯阅读 csrc/flash_attn/src/*.cu、flash_attn/cute/flash_fwd*.py、hopper/ 文件名模式，用 grep 提取支持矩阵。
认领留言（英文，可直接贴到 Issue）
Issue #2456: While the hd256/512 kernel work is in progress (PRs #2843/#2976 etc.), I've compiled a behavior comparison matrix (docs/matrix-attention-behavior.md) mapping FA2/FA3/FA4 support across (causal, GQA, paged-kv, varlen, MLA, sparse) × (fp16, bf16, fp8) × (SM80/SM100/SM120). Each gap links to its tracking issue. This is pure doc + code reading — no kernel changes. Could you review the accuracy? Happy to adjust format if you prefer it in README.
大致实施方案
  • 建立功能维度：causal, softcap, ALiBi, GQA, paged-kv, varlen, MLA, sparse, block-sparse
  • 建立 dtype 维度：fp16, bf16, fp8
  • 建立架构维度：SM80(FA2), SM90(FA3/FA4), SM100(FA4), SM120(FA4)
  • 用 ✅/❌/⚠️ 标注，⚠️ 附 Issue 链接（如 #2456 hd512, #2852 hd80/SM103）
  • 在 Issue #2456 下评论提供矩阵草稿供维护者补充
  • PR 仅含 docs/ 文件
可能涉及的目录或文件
  • tests/ 下现有 pytest 参数化列表是验证支持矩阵的直接证据
  • flash_attn/cute/flash_fwd.py 的 dispatch 逻辑揭示 FA4 实际支持组合
验收方式
  • 矩阵覆盖 Issue #2456 提及的所有 head_dim 组合
  • 每个 ⚠️/❌ 都有对应 Issue 或代码引用
  • 维护者在 Issue #2456 确认矩阵准确性
开工前问题与风险

向维护者确认

  • 维
  • 护
  • 者
  • 是
  • 否
  • 希
  • 望
  • 矩
  • 阵
  • 放
  • 在
  • d
  • o
  • c
  • s
  • /
  • 还
  • 是
  • 根
  • 目
  • 录
  • R
  • E
  • A
  • D
  • M
  • E
  • 子
  • 章
  • 节
  • ？
  • 是
  • 否
  • 需
  • 要
  • 同
  • 步
  • 更
  • 新
  • u
  • s
  • a
  • g
  • e
  • .
  • m
  • d
  • ？

风险

  • F
  • A
  • 4
  • A
  • P
  • I
  • 仍
  • 在
  • 收
  • 敛
  • ，
  • 矩
  • 阵
  • 需
  • 标
  • 注
  • '
  • 截
  • 至
  • c
  • o
  • m
  • m
  • i
  • t
  • x
  • x
  • x
  • '
  • 的
  • 快
  • 照
  • 日
  • 期
任务 4候补 · low · CPU · 1 个晚上

FA4: ~20-25us Python dispatch overhead per call — optimization targets

Issue #2426 · FA4: ~20-25us Python dispatch overhead per call — optimization targets ↗

无人认领3 条评论更新 2026-10-01

用到的专长：利用性能建模专长建立 Python 开销的离线评估方法

目标：新增 benchmarks/benchmark_dispatch_overhead.py，在 Mac CPU 上测量 is_fake_mode / get_device_properties / torch.empty 的单独耗时，验证 Issue #2426 列出的优化目标。

为什么值得长期做：Issue #2426 详细列出了 Python dispatch 开销（is_fake_mode 3x/call、get_device_properties 每 call、6 tensor allocs in backward），但 PR #2516 已关闭。维护者明确标记了优化目标但无后续。提供 CPU 端微基准可离线验证各优化项的效果（如 lru_cache is_fake_mode、workspace pre-alloc）。

怎么介入：Issue #2426 无 assignee，PR #2516 已关闭但无后续。以'提供 CPU 微基准数据'切入，在 Issue 下评论：'我可以在 Mac 上测 is_fake_mode / get_device_properties / torch.empty 的相对优化比，帮你们排序 4 个优化目标的优先级。'不与任何实现冲突。
第一个 PR 的边界：第一个 PR 仅含 benchmarks/benchmark_dispatch_overhead.py（新增），不修改任何接口或 kernel。
第一步：读取 Issue #2426 的表格，确认 4 个优化方向：(1) cache is_fake_mode, (2) cache get_device_properties, (3) fast-path for repeated shapes, (4) pre-alloc backward workspace。
本机怎么复现 / 验证：在 Mac 上 python benchmarks/benchmark_dispatch_overhead.py 即可运行，全部用 torch CPU + timeit，无需 GPU。
认领留言（英文，可直接贴到 Issue）
Issue #2426: PR #2516 was closed but the optimization targets remain. I'd like to add benchmarks/benchmark_dispatch_overhead.py that measures is_fake_mode, get_device_properties, and torch.empty on CPU, comparing cached vs uncached variants. While absolute us values differ from B200, the relative speedup ratios should be representative for prioritizing the 4 optimization targets. Is there interest? I can have the script ready in 1-2 evenings.
大致实施方案
  • 构造 micro-benchmark：单独测 is_fake_mode()、get_device_properties()、torch.empty_like() 的 us/call
  • 用 timeit.repeat 统计 1000 次调用的均值/P50/P99
  • 输出与 Issue #2426 表格对齐的对比格式
  • 新开 Issue 或评论 #2426 提供数据
  • PR 仅含 benchmarks/ 文件
可能涉及的目录或文件
  • benchmarks/ 已有 benchmark_varlen_sched.py 可参考风格
  • torch.utils.benchmark 是更专业的计时工具
验收方式
  • p
  • y
  • t
  • h
  • o
  • n
  • b
  • e
  • n
  • c
  • h
  • m
  • a
  • r
  • k
  • s
  • /
  • b
  • e
  • n
  • c
  • h
  • m
  • a
  • r
  • k
  • _
  • d
  • i
  • s
  • p
  • a
  • t
  • c
  • h
  • _
  • o
  • v
  • e
  • r
  • h
  • e
  • a
  • d
  • .
  • p
  • y
  • 在
  • M
  • a
  • c
  • 输
  • 出
  • u
  • s
  • /
  • c
  • a
  • l
  • l
  • 表
  • 格
  • ，
  • 覆
  • 盖
  • I
  • s
  • s
  • u
  • e
  • #
  • 2
  • 4
  • 2
  • 6
  • 列
  • 出
  • 的
  • 4
  • 个
  • 函
  • 数
  • 。
开工前问题与风险

向维护者确认

  • I
  • s
  • s
  • u
  • e
  • #
  • 2
  • 4
  • 2
  • 6
  • 的
  • 基
  • 准
  • 数
  • 据
  • 是
  • 在
  • B
  • 2
  • 0
  • 0
  • 上
  • 测
  • 的
  • ，
  • M
  • a
  • c
  • C
  • P
  • U
  • 上
  • 的
  • 绝
  • 对
  • 值
  • 无
  • 意
  • 义
  • —
  • —
  • 维
  • 护
  • 者
  • 是
  • 否
  • 接
  • 受
  • 相
  • 对
  • 比
  • 较
  • （
  • 如
  • i
  • s
  • _
  • f
  • a
  • k
  • e
  • _
  • m
  • o
  • d
  • e
  • v
  • s
  • c
  • a
  • c
  • h
  • e
  • d
  • 的
  • 加
  • 速
  • 比
  • ）
  • ？

风险

  • M
  • a
  • c
  • C
  • P
  • U
  • 上
  • 的
  • 绝
  • 对
  • u
  • s
  • 值
  • 无
  • 法
  • 与
  • G
  • P
  • U
  • 环
  • 境
  • 对
  • 比
  • ，
  • 需
  • 明
  • 确
  • 标
  • 注
  • 为
  • '
  • 相
  • 对
  • 优
  • 化
  • 比
  • '