diff --git a/H20-BRANCHES.md b/H20-BRANCHES.md new file mode 100644 index 0000000..9de9f99 --- /dev/null +++ b/H20-BRANCHES.md @@ -0,0 +1,69 @@ +# 分支记录(SM90 / H20 这条线) + +截止 2026-10-25 14:30。所有分支都已推到 `tonyh168/DeepInfer`。 + +## 一句话地图 + +``` +main (8b16149) + └── h20-sm90-support 5004c77 TileLang PR 版:SM90 attention 用 TileLang 实现 + └── h20-sm90-triton 73b8090 把引擎 kernel 换成 Triton 后端 + └── h20-sm90-triton-flagtree c6a65fe 同上,在 flagtree 镜像里重编 vendor 库 + └── h20-sm90-sparse-mla-triton f25e5d8 ← 当前 HEAD + 再换掉 attention:无 tilelang 端到端 +``` + +另一条独立的线:`h20-sm90-lossless-requant`(74ba468),做无损 requant 与 debug_dump 挪位,不在这条链上。 + +## 逐个说明 + +| 分支 | HEAD | 这条分支解决了什么 | 关键限制 | +|---|---|---|---| +| `h20-sm90-support` | `5004c77` | 上游 PR #180 的基础:FlashMLA 是 SM100-only,SM90 上按 `patch_imports` 钩子换成自己实现的 attention | attention 仍是 TileLang | +| `h20-sm90-triton` | `73b8090` | 引擎侧 29 个 kernel 加 Triton 实现(17 个可用),`DEEPINFER_OPS_BACKEND=triton` 切换 | 融合算子库 `tile_kernels` 仍要 TileLang | +| `h20-sm90-triton-flagtree` | `c6a65fe` | 在 flagtree 镜像(torch 2.11 + cu13)里重编 `deepinfer._C` / `deep_gemm` / `deep_ep` / `deep_select`;`tile_kernels` 加 compat 层走 torch 参考,从而不 import TileLang;修掉参考实现在 sync 守卫下抛错的问题 | attention 仍要 TileLang | +| **`h20-sm90-sparse-mla-triton`** | **`f25e5d8`** | **attention 也换成 FlagGems 的 Triton kernel,于是「屏蔽 tilelang 也能端到端」成立** | 采样器的 rejection-sampling 变体仍要 TileLang(关 MTP 可绕);不与 TileLang 版逐 token 一致 | + +## `f25e5d8`(当前版本)到底做了什么 + +新增三个文件、改一处 import: + +| 文件 | 作用 | +|---|---| +| `deepinfer/h20_ops/vendor/flashmla_sparse.py` | 内联 PyPI `flag_gems` 5.0.2 的 sparse MLA Triton kernel(329 行,Apache-2.0,文件头记了 md5 和两处本地改动) | +| `deepinfer/h20_ops/sparse_mla_triton.py` | 适配层:`(out, max_logits, lse)` → `(out, lse)`;lse 自然对数 → log2;两种调用约定(共享 KV / 每 query 独立 KV) | +| `deepinfer/h20_ops/sparse_mla.py` | 分发层:`DEEPINFER_SPARSE_MLA_BACKEND=tilelang\|triton`,不设则能 import tilelang 就用它 | +| `deepinfer/h20_ops/attn_sm90.py` | 只改一行 import,两个调用点未动 | + +验证(每步先过再往下): + +| 步 | 内容 | 结果 | +|---|---|---| +| 1 | 内联 kernel vs TileLang,7 形状 | `ALL_MATCH`,≤2 bf16 ULP | +| 2 | decode 形态(fp8 cache,含 SWA+full 双缓存),5 用例 | `ALL_MATCH` | +| 3 | 适配层契约,5 组 | `ALL_PASSED` | +| 4 | 分发层 + `attn_sm90` 两个入口,4 组 | `ALL_PASSED` | +| 5 | 屏蔽 tilelang 端到端 | `exit=0`,0 回退,8/8 DP 一致 | +| 6 | 有 tilelang 对照 | `exit=0`,0 回退 | + +复现: + +```bash +DEEPINFER_SIMULATE_NO_TILELANG=1 DEEPINFER_KERNELS=ref DEEPINFER_OPS_STRICT=1 \ +DEEPINFER_USE_NEXTN_DECODE=0 bash scripts/h20_e2e_report.sh +``` + +## 怎么切分支 + +```bash +git fetch tonyh168 +git checkout h20-sm90-sparse-mla-triton # 最新最全 +git checkout h20-sm90-triton-flagtree # 只要 engine kernel 是 Triton、attention 留 TileLang +git checkout h20-sm90-support # 只要 SM90 能跑(attention=TileLang,工程最省) +``` + +## 已知未做 + +- **性能一次没测**(prefill/decode 吞吐、TTFT)。所有结论都是「能跑、数值对」,不是「跑得快」。 +- 与 TileLang 版**不逐 token 一致**(prompt A 前 15 个 token 相同、第 16 分叉)。控制实验已证明非采样噪声。 +- `fused_sample` 的 rejection-sampling 变体仍依赖 TileLang,故纯 Triton 需关投机解码。