mirror of
https://github.com/deepseek-ai/DeepInfer.git
synced 2026-11-01 17:26:25 +08:00
docs(h20): record what a FlagTree/FlagOS image can and cannot do
Two artefacts and one finding. `scripts/flagtree_triton_probe.py` asks whether the Triton an image ships can build the features this port depends on, one feature at a time: `tl.dot`, `tl.gather` with a data-dependent index, `tl.gather` as the RoPE pairing, `fp8e4nv` with a block scale, and an int64 comparison (uint32 compares go signed in Triton, which the sampler depends on). Both vendor images pass all of them and emit `.target sm_90a`. `TODO.md` collects every open issue across both branches, with the two blockers an end-to-end run in a vendor image actually hits: `deepinfer/_C` is built against one torch ABI (`2.13.0+cu132`) so it fails to load on torch 2.11 and 2.9, and `deep_gemm`/`deep_ep` are not on PyPI in the versions `vendor/` provides. Neither is a Triton problem, and saying so plainly is the point of writing it down.
This commit is contained in:
loopyt
parent
b4f7a42652
commit
7de3bc29d6
388
TODO.md
Normal file
388
TODO.md
Normal file
@ -0,0 +1,388 @@
|
||||
# H20 / SM90 遗留问题总清单
|
||||
|
||||
截止 2026-10-22 10:00。分两条线：
|
||||
|
||||
- **TileLang PR 版** = 分支 `h20-sm90-support` → 上游 PR [#180](https://github.com/deepseek-ai/DeepInfer/pull/180)（`OPEN` / `MERGEABLE` / `CLEAN`，0 review 0 comment 0 CI）
|
||||
- **Triton 版** = 分支 `h20-sm90-triton`（基于 PR 分支，HEAD `615664e`），把推理路径上的 TileLang kernel 换成 Triton
|
||||
|
||||
两条线**共用同一个上游分支底座**，所以「底座问题」两边都有，先单列。
|
||||
|
||||
---
|
||||
|
||||
## 零、两条线共有的底座问题（不是任何一条线引入的）
|
||||
|
||||
### C1. CUDA graph capture 直接挂 —— 默认配置跑不起来 🔴 P0
|
||||
|
||||
**现象**：graph 开启（默认）时 `warm_up_model → capture_model` 8 rank 全挂：
|
||||
|
||||
```
|
||||
deepinfer/h20_ops/engram_local.py:108 torch.bincount(owner)
|
||||
RuntimeError: Cannot copy between CPU and CUDA tensors during CUDA graph capture
|
||||
unless the CPU tensor is pinned.
|
||||
```
|
||||
|
||||
**两处根因**（都在我们新写的 `LocalEngramGather`，为了替掉 SM90 上没有的 DeepEP RDMA gather）：
|
||||
|
||||
1. `torch.bincount(owner)` 里 `owner` 在 GPU，输出默认是 CPU tensor → capture 期 D2H。
|
||||
2. `_exchange_sizes` 用 `dist.all_gather_into_tensor` + `.tolist()`（**必然 host sync**），
|
||||
下游 `all_to_all_single(output_split_sizes=...)` 用**动态**切分 → 结构上不可 capture。
|
||||
|
||||
**当前绕过**：`DEEPINFER_ENFORCE_EAGER=1`（关掉整个 decode graph）。
|
||||
|
||||
**两条修法**：
|
||||
|
||||
- [ ] **C1a 快速版**：SM90 且 engram 表存在时，自动关 graph 并打一行日志说明原因。
|
||||
1 行判断 + 1 行日志，立刻让默认配置可用。代价：decode 没有 graph（慢多少没测，见 C5）。
|
||||
- [ ] **C1b 正解**：把 engram gather 改成 capture-safe —— 每行开头的 `counts` 走
|
||||
**固定 shape 的 padded 交换**（`[world, world]` 计数矩阵常驻显存，kernel 内算前缀和），
|
||||
避免 host sync 与动态 split。代价：prefill 流量 ×world，需要重新量带宽。
|
||||
|
||||
### C2. 必须手动设 `EP_DISABLE_GIN=1` 🟡 P1
|
||||
|
||||
容器的 `/dev/infiniband` **不存在**（host 上有 4 个 `uverbs*`，容器没挂），
|
||||
DeepEP 查到 `props.ginType == NCCL_GIN_TYPE_NONE` 就 assert：
|
||||
|
||||
```
|
||||
vendor/DeepEP/csrc/kernels/comm/context.cpp:125
|
||||
```
|
||||
|
||||
SM90 单机走的是 scale-up 路径，本来不需要 GIN，所以关掉无副作用（已在
|
||||
`userbuffer.py:80` 用 `allow_hybrid_mode=not PLATFORM.is_sm90` 关了一半）。
|
||||
|
||||
- [ ] **C2** 在 SM90 单节点下自动 `EP_DISABLE_GIN=1`（或检测 `/dev/infiniband` 缺失即关），
|
||||
而不是要求使用者记得加 env。否则「按 README 跑」必挂。
|
||||
|
||||
### C3. 2.7% logits 漂移 —— 硬件量化地板，**不修** ⚪ 已收敛
|
||||
|
||||
`DeepGEMM 2.8.2 csrc/apis/layout.hpp:100`：
|
||||
|
||||
```cpp
|
||||
DG_HOST_ASSERT((arch_major == 9 and gran_k == 128 and k_alignment == 128) or ...)
|
||||
```
|
||||
|
||||
SM90 强制 **per-128 + fp32 scale**，checkpoint 是 **per-32 + e8m0**。
|
||||
`gran_k == 128` 断言封死了「无损重化」这条路（`H20-DRIFT-ANALYSIS.md` §1 逐条排除）。
|
||||
|
||||
- [x] 结论：不是 bug。答案稳健性已验 —— pos14 官方 margin 只有 1.638，噪声 std 1.63，硬币。
|
||||
- [ ] **C3-doc** PR 描述里必须写清「自然提问答案正确（首 token 6127 北京）」，
|
||||
别让人误读成「移植答不对」。`H20-DRIFT-ANALYSIS.md` §5 末尾已写，PR 正文还没搬。
|
||||
|
||||
### C4. `token_logprobs` 恒为 0 🟡 P2（历史遗留，与 Triton 无关）
|
||||
|
||||
所有历史日志（含纯 TileLang 时期）都是 `[0.0] * n`。
|
||||
|
||||
**已排除**：`map_lookup` kernel（6 passed）、sampler 全部 kernel（74 passed）、
|
||||
`SampledInfo.from_reqcnt` 的 `zeros`（死代码）、sampler→engine 字段传递（直通）。
|
||||
|
||||
**剩余候选**：`process_sampling` 返回的 `logprobs` 本身为 0（`fused_sample` 的 logprob 输出路径），
|
||||
或 `build_request_outputs` 的属性回填。
|
||||
|
||||
- [ ] **C4** 跑 1 次整机（~5 min）：在 sampler 里打印 `logprobs.max()` 与 `token_logprobs_cur`。
|
||||
sampler 层就是 0 → 定位在 `fused_sample`；正常 → 定位在 output 回填。预估 1 h，最坏 2 h。
|
||||
影响面：AIME 打分、perplexity 这类评估用途。
|
||||
|
||||
### C5. 性能没有任何系统数据 🔴 未开工
|
||||
|
||||
`H20-PORTING-JOURNAL.md` §6.4：只测过两个点 ——
|
||||
稀疏 MLA prefill **104 TFLOP/s**、decode **0.23 ms/step**（batch=8, topk=1024）。
|
||||
|
||||
- [ ] **C5** 端到端吞吐（prefill / decode / 不同 seq len）、首 token 延迟、
|
||||
以及 vs vLLM / SGLang 的对比。两条线都要，Triton 版还要 Triton vs TileLang 逐 kernel 对比。
|
||||
|
||||
### C6. `index KV` 的 ue8m0 布局单测 12 failed ⚪ 已知，不影响 SM90
|
||||
|
||||
`store_kv_cache` / `load_paged_kv_cache` 全布局跑 48 passed / 12 failed，
|
||||
失败**全部**落在 `FP8_INDEX_KV_FORMAT`、`MXFP8_INDEX_KV_FORMAT` 两种 ue8m0 index 布局，
|
||||
而 SM90 不用这两个（用 `FP8_FLOAT32_SF_INDEX_KV_FORMAT`）。
|
||||
|
||||
失败点在 **ref 自身**：它用 `PLATFORM.fp8_sf_type`（SM90 是 fp32）去重解释 ue8m0 字节，
|
||||
`index_copy_` 撞 dtype。
|
||||
|
||||
- [ ] **C6** 可选：修 ref 的 dtype 推断，让 SM90 上跑全布局也全绿。属「别留红灯」的整洁性工作。
|
||||
|
||||
---
|
||||
|
||||
## 一、TileLang PR 版（`h20-sm90-support` / PR #180）
|
||||
|
||||
### 规模
|
||||
|
||||
```
|
||||
48 files changed, 4562 insertions(+), 78 deletions(-)
|
||||
```
|
||||
|
||||
新增：`deepinfer/h20_ops/`（7 文件，SM90 后端替身）、`deepinfer/debug_dump.py`、
|
||||
14 个 `scripts/h20_*.py` / `ref_*.py` 验证脚本、4 份设计/调查文档。
|
||||
|
||||
### 已完成（可作为 PR 的验证证据）
|
||||
|
||||
- [x] prefill 因果掩码 bug（首 token 从 `EVAL -` 变回 `17607`）
|
||||
- [x] 专家权重 scale 广播 bug（`requant_expert_fp4` 把 `(n_out, k/32)` 当 `(n_out/32, k/32)` 展开）
|
||||
→ w1/w3/w2 的 cos 从 0.89 到 0.9997~1.0000
|
||||
- [x] layer0 整层 rel_err 1.27e-1 → 2.88e-2，整条误差曲线压平一个数量级
|
||||
- [x] KV cache 闭环 round-trip（rel 0.027，FP8 地板）
|
||||
- [x] `_decode_p_sum_sm90` vs 独立 float64 参考（rel 1e-7）
|
||||
- [x] decode 逐步误差曲线（Δmargin mean +0.711 / std 1.628，8 位置 7 个 top-1 正确）
|
||||
- [x] 自然文本提问答案完全正确（`北京。`，8 rank 一致）
|
||||
- [x] FlashMLA SM100-only → 自己写 SM90 稀疏 MLA 后端 + DeepGEMM sparse MQA logits 替身
|
||||
|
||||
### 遗留
|
||||
|
||||
| # | 问题 | 优先级 | 状态 |
|
||||
|---|---|---|---|
|
||||
| 1.1 | 继承 C1：CUDA graph capture 挂，必须 eager | 🔴 P0 | 未修 |
|
||||
| 1.2 | 继承 C2：必须 `EP_DISABLE_GIN=1` | 🟡 P1 | 未修 |
|
||||
| 1.3 | 继承 C4：`token_logprobs` 恒 0 | 🟡 P2 | 未修，有判定路径 |
|
||||
| 1.4 | 继承 C5：性能无系统数据 | 🔴 未开工 | — |
|
||||
| 1.5 | **绕过项需要逐个 review 后再进主干**（见下表） | 🟡 P1 | 未做 |
|
||||
| 1.6 | PR 无 reviewer / 无 CI，也没写「怎么验」的 checklist | 🟡 P1 | 未做 |
|
||||
| 1.7 | 4 份 MD（~1700 行）是否都该进仓库，还是压成 1 份 + PR 正文 | ⚪ P3 | 待定 |
|
||||
|
||||
### 1.5 绕过项明细（`H20-DEBUG-PLAN.md` 备注 + §6.4）
|
||||
|
||||
| 绕过 | 原因 | 数值风险 | 待办 |
|
||||
|---|---|---|---|
|
||||
| `DEEPINFER_ENFORCE_EAGER=1` | C1 | 无（仅性能） | 见 C1a/C1b |
|
||||
| `EP_DISABLE_GIN=1` | C2 | 无（改走 NVLink all_to_all，已单测） | 见 C2 |
|
||||
| DeepSelect topk → `torch.topk` | SM90 launch 失败 | 无（契约等价，`h20_topk_check.py` 已验） | 查能否修 kernel 而不用 torch 兜底 |
|
||||
| decode indexer → `_decode_p_sum_sm90` | SM90 paged kernel 要求 64-token page | **已验**（P4.4 rel 1e-7） | 关闭 |
|
||||
| 关闭共享层候选块优化 | `fp8_fp4_sparse_mqa_logits` SM100-only | 无（稠密等价） | 记入 PR 说明 |
|
||||
| `mega_mhc` / `mega_moe` 关闭 | SM100-only | 无（有单测） | 记入 PR 说明 |
|
||||
| ViT → SDPA | FA-4 是 SM100-only | 无（多模态路径，本轮不跑） | 记入 PR 说明 |
|
||||
|
||||
### 建议动手顺序（PR 线）
|
||||
|
||||
1. **C1a**（30 分钟）→ 默认配置终于能起来，别人的第一印象取决于这个
|
||||
2. **1.6** PR 正文 + 验证 checklist（半天，纯文档）
|
||||
3. **C2** 自动检测（1 小时）
|
||||
4. **1.5** 把 6 个绕过的「数值风险」在 PR 里逐条落字（半天）
|
||||
5. **C4** logprob（1~2 小时）
|
||||
6. **C5** 性能（1 天）
|
||||
7. **C1b** engram capture-safe 重写（1~2 天，收益是 decode graph）
|
||||
8. **1.7** 文档收敛（半天）
|
||||
|
||||
---
|
||||
|
||||
## 二、Triton 版（`h20-sm90-triton`）
|
||||
|
||||
### 规模与成果
|
||||
|
||||
- 新增 `tilelang_ops/triton/` 19 个模块，与 `cuda/` 平级，`ops.py` 按 `DEEPINFER_OPS_BACKEND` 选后端
|
||||
- 全量 A/B：**22 个 kernel、约 140 用例，全绿**（`scripts/h20_ops_ab.py`）
|
||||
- 非 MTP 路径：**20/20 kernel 全 Triton，回退提示 0 条**，token 与纯 TileLang 逐 token 相同
|
||||
- MTP 开：16/20 Triton，5 条回退提示，token 同样一致
|
||||
- 采样器：**逆推出 cuRAND Philox 的 seeding**（20/20 bit-identical），engine 配置下 token 100% 一致
|
||||
|
||||
### 遗留
|
||||
|
||||
| # | 问题 | 优先级 | 状态 |
|
||||
|---|---|---|---|
|
||||
| 2.1 | 继承 C1 / C2 / C4 / C5（同底座） | 见零章 | — |
|
||||
| 2.2 | **MTP 4 个 kernel 未移植** | 🟡 P1 | 未做 |
|
||||
| 2.3 | **采样器 rejection-sampling 变体未实现** | 🟡 P1 | 显式委托 cuda |
|
||||
| 2.4 | **20 个可选路径 kernel 缺失**（导入期报 `lacks 20 kernel(s)`） | 🟡 P1 | 未做 |
|
||||
| 2.5 | 窄 cut 采样器 token 差 9~13% | ⚪ P3 | 已定量解释，不修 |
|
||||
| 2.6 | **Triton 后端 0 个 pytest 用例** | 🔴 P1 | 未做 |
|
||||
| 2.7 | 没有 Triton vs TileLang 的性能对比 | 🔴 未开工 | — |
|
||||
| 2.8 | PR 正文没写；`H20-TRITON-PLAN.md` 307 行需收敛 | ⚪ P3 | — |
|
||||
| 2.9 | `scripts/h20_state_check.py` 未提交 | ⚪ | 已生成，待 commit |
|
||||
| 2.10 | **Triton-only 环境跑不通**（见第四章） | 🟡 P1 | 部分已修 |
|
||||
|
||||
### 2.10 Triton-only 环境（vendor 镜像 / CI 镜像）
|
||||
|
||||
「技术栈换成 Triton」的隐含目标是「只装 Triton 也能跑」，但当前做不到，因为 **TileLang
|
||||
概念漏到了接口层**。已在第四章验证并修掉第一批：
|
||||
`ops.py` 加 `DEEPINFER_OPS_STRICT`、`tile_kernels` / `tilelang` 的 import 移进参考实现与单测函数。
|
||||
剩下的问题见第四章 §4.4。
|
||||
|
||||
### 2.2 MTP 4 个 kernel
|
||||
|
||||
| kernel | 难度 | 说明 |
|
||||
|---|---|---|
|
||||
| `remap_draft_topk_tl` | **高** | baseline 用一块 25856/44032 项的 shared 池按 vocab 分块扫。Triton 无共享内存池原语，照搬要换实现（位图或 global scratch），不能机械翻译 |
|
||||
| `get_accept_result_tl` | 中 | 接受判定，带 `has_golden` 变体 |
|
||||
| `finalize_verification_tl` | 中 | 带 `has_sum_smaller_probs` / `has_confidence` / `num_steps` 三个变体维度 |
|
||||
| `get_fused_dflash_get_input_tl` | 中 | MTP draft 输入组装 |
|
||||
|
||||
### 2.3 采样器 verify 变体
|
||||
|
||||
非 MTP 路径的 `fused_sample_tl` 已移植（单 tile，vocab=4096），但 **rejection sampling 变体**
|
||||
（MTP verify 路径）没实现，现在**显式委托**给 TileLang 并打印一行说明。
|
||||
|
||||
注意它是**显式委托**而不是静默回退 —— `fused_sample` 是模块导出函数，`_FallbackNotice`
|
||||
覆盖不到，所以工厂自己打印。要做需要：accept/reject 比较、residual mass、拒绝后第二次 draw。
|
||||
|
||||
### 2.4 缺失的 20 个 kernel = 16 个可选路径 + 4 个 MTP
|
||||
|
||||
对着 `cuda/` 的 43 个 `*_tl` 与 `triton/` 的 24 个做差集得到，与导入期那行
|
||||
`lacks 20 kernel(s)` 逐个对得上（first = `apply_token_bitmask_by_step_inplace_tl`）。
|
||||
|
||||
**16 个可选路径**（当前配置不执行，一旦打开就静默回退）：
|
||||
|
||||
| 模块 | kernel | 何时会用到 |
|
||||
|---|---|---|
|
||||
| `structure_output` | `apply_token_bitmask_by_step_inplace` | **引导解码 / 工具调用 / JSON 约束** ← 最可能真被用到 |
|
||||
| `image_embed_ops` | `batch_downsample` / `replace_input_embeds` / `rgb_pad_to_patches` | 多模态 |
|
||||
| `candidate_blocks` | `candidate_block_max` / `get_candidate_compact_lens` / `restore_candidate_indices_inplace` | 共享层候选块优化（SM90 当前已关） |
|
||||
| `prepare_cp_mirror` | `prepare_cp_mirror` / `yoco_cp_boundary_fused` | CP > 1 |
|
||||
| `state_cache` | `load_state_to_cache` / `select_state_slice` | state cache 复用 |
|
||||
| `schedule_mtp_varlen` | `compute_seqlens_q_by_confidence_rates` | varlen MTP |
|
||||
| `dflash_copy` | `copy_dflash_decode_input` | dflash 变体 |
|
||||
| `outlier_detection` | `detect_outlier_and_abort` | 异常检测 |
|
||||
| `random_sampling` | `random_sampling` | 旧采样路径 |
|
||||
| `sum_smaller_probs` | `sum_smaller_probs` | 采样后处理 |
|
||||
|
||||
**4 个 MTP**：`remap_draft_topk` / `get_accept_result` / `finalize_verification` /
|
||||
`get_fused_dflash_get_input`（见 2.2）。
|
||||
|
||||
建议：**重点补 `structure_output`**（唯一在正常服务里可能被打开的功能），其余按需。
|
||||
|
||||
### 2.5 窄 cut 采样器的 9~13% 差异（结论已定，不建议修）
|
||||
|
||||
```
|
||||
r = draw % total, total = sum(floor(p_i * 2^31))
|
||||
```
|
||||
|
||||
baseline 拒采样上界 `(2^32-1) // total * total` 在 `total` 跨 **2^31** 时从 `2*total` **阶跃**到 `1*total`。
|
||||
`total` 因概率和为 1 就贴着 2^31，归约顺序差几个单位就会换 draw。
|
||||
|
||||
- engine 实际配置（top_p=0.99）：**100%（128/128）**
|
||||
- top_k=128：91.4%，top_p=0.5：86.7%，top_p=0.9+top_k=256：87.5%
|
||||
- 保留集合逐列一致，logprobs 差 1e-5
|
||||
- clamp `total` 无效（baseline 自己两个方向都会跨，clamp 只是把不一致换到别的行：11 → 30）
|
||||
|
||||
### 2.6 Triton 后端没有 pytest 用例 🔴
|
||||
|
||||
`tests/` 下 **0 个**文件提到 `DEEPINFER_OPS_BACKEND` 或 triton。140 个 A/B 用例全在
|
||||
`scripts/h20_ops_ab.py` 里，是「我在跑」，不是「CI 在跑」。
|
||||
|
||||
- [ ] **2.6** 把 `scripts/ab_cases.py` 的用例搬成 `tests/` 下的参数化 pytest（需要 GPU，
|
||||
按 `PLATFORM.is_sm90` skip），至少覆盖 20 个已移植 kernel 的 smoke。
|
||||
这是这个 PR 能不能长期活下来的关键 —— 否则下一个人改 `cuda/` 侧，Triton 侧静默漂掉。
|
||||
|
||||
### 建议动手顺序（Triton 线）
|
||||
|
||||
1. **C1a**（先解决共用 P0，否则默认配置跑不起来）
|
||||
2. **2.9** 提交 `h20_state_check.py`（10 分钟）
|
||||
3. **2.6** pytest 化（1 天）← 优先级高于继续移植
|
||||
4. **2.4** 补 `structure_output`（半天）
|
||||
5. **2.3** 采样器 verify 变体（1 天，Philox 已破，剩下是拒绝逻辑）
|
||||
6. **2.2** MTP 4 个（`remap_draft_topk` 是硬骨头，单独排 1~2 天）
|
||||
7. **2.7 + C5** 性能对比（1~2 天）
|
||||
8. **2.8** PR 正文（半天）
|
||||
|
||||
---
|
||||
|
||||
## 三、两条线的关系
|
||||
|
||||
```
|
||||
main
|
||||
└── h20-sm90-support ← PR #180（TileLang，SM90 能跑）
|
||||
└── h20-sm90-triton ← 本清单第二条线（kernel 换 Triton）
|
||||
```
|
||||
|
||||
- `h20-sm90-triton` **包含** PR #180 的全部内容，所以 PR #180 的遗留问题它全都有。
|
||||
- 建议 **先把 C1/C2 在 `h20-sm90-support` 上修掉再往下走**，否则 Triton 版的每次端到端
|
||||
都得记着加两个 env，容易把「配置问题」误判成「Triton 问题」。
|
||||
- Triton 版真正的增量风险只有 2.2 / 2.3 / 2.4 / 2.5 / 2.6。
|
||||
|
||||
---
|
||||
|
||||
## 四、在 FlagTree / FlagOS 镜像里跑 Triton（2026-10-22 实测）
|
||||
|
||||
目标：确认 Triton 版能不能直接放进 vendor 镜像（FlagOS / FlagTree），不开 FlagGems。
|
||||
|
||||
### 4.1 三套环境
|
||||
|
||||
| 容器 | 镜像 | torch | triton |
|
||||
|---|---|---|---|
|
||||
| `di-h20` | `vllm/vllm-openai:v0.32.0` | 2.13.0+cu130 | 3.7.1（标准） |
|
||||
| `zjr-flWithTree-0930` | `flagtree/flagtree-py312-torch2.11.0_cu130` | 2.11.0+cu130 | **3.6.0**（`tokenspeed-triton 3.8.10.post20260721`，vendor 构建） |
|
||||
| `flagos` | `moonlight-16b-a3b-...-tree_0.5.0_3.5-gems_5.0.2` | 2.9.0+cu128 | 3.5.1 |
|
||||
|
||||
### 4.2 结论一：FlagTree 的 Triton 能编我们的 kernel ✅
|
||||
|
||||
`scripts/flagtree_triton_probe.py` 覆盖 DeepInfer 真正用到的 7 类特性，**两个镜像都 7/7 通过**：
|
||||
|
||||
| 特性 | 为什么单独测 | 结果 |
|
||||
|---|---|---|
|
||||
| elementwise add | 基本盘 | max abs err **0.00e+00** |
|
||||
| `tl.dot` matmul | tensor core 路径 | rel 3e-4（fp16 输入、fp32 累加） |
|
||||
| `tl.gather` on tile | compressor 选行、没有对应原语时最容易挂 | 行与 `pool[idx]` 全等 |
|
||||
| `tl.gather` 做 RoPE 配对 | `lanes ^ 1` / `lanes ^ half` | 逐元素相等 |
|
||||
| `fp8e4nv` cast + block scale | 全是 FP8 权重 | rel 3.2~3.4%（e4m3 理论量级） |
|
||||
| int64 CDF 比较 | uint32 走有符号是已知坑 | 完全相等 |
|
||||
| PTX target | 确认真的出 sm_90a | `.version 8.7` / `.target sm_90a` |
|
||||
|
||||
⚠️ FlagTree 镜像（cu13）需要补一条：
|
||||
`LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/root/.triton/nvidia/cudart/cuda_cudart-linux-x86_64-12.8.57-archive/lib`，
|
||||
否则报 `OSError('libcudart.so.12: cannot open shared object file')`。Triton 自己下载了 cudart
|
||||
但没把它加进搜索路径，跟 FlagTree 无关，是 cu13 环境的通病。
|
||||
|
||||
### 4.3 结论二：19 个 kernel 的输出在三套 Triton 上逐元素一致 ✅
|
||||
|
||||
新增 `scripts/h20_ops_ab.py --dump DIR` / `--compare DIR`：在没有 TileLang 基线的机器上，
|
||||
先把选中后端的输出存盘，再到另一台机器上逐元素比对。
|
||||
|
||||
```
|
||||
flagtree (triton 3.6.0) --dump → 115 用例，2 个 kernel 跳过
|
||||
di-h20 (triton 3.7.1) --dump → 125 用例
|
||||
flagtree --compare di-h20 的 dump → AB_OK
|
||||
flagos (triton 3.5.1) --compare di-h20 的 dump → AB_OK
|
||||
```
|
||||
|
||||
跳过的是 `fused_sample_tl` 和 `map_lookup_tl`，因为它们要 `deepinfer._C.engine`
|
||||
（`sampling_inputs` 的结构体）。**两个镜像上都是同一个原因，不是 Triton 的问题。**
|
||||
|
||||
顺带确认：`di-h20` 上全量 A/B 仍是 **AB_OK**，本次改动零回归。
|
||||
|
||||
### 4.4 结论三：端到端推理在两个镜像里都跑不起来 ❌
|
||||
|
||||
不是 Triton 的问题。三层依赖，按难度排序：
|
||||
|
||||
| # | 缺什么 | 影响 | 修法 |
|
||||
|---|---|---|---|
|
||||
| 1 | **`deepinfer/_C` 的 ABI** | flagtree(torch 2.11) 报 `undefined symbol: at::TensorBase::const_data_ptr<int>()`，flagos(torch 2.9) 直接找不到模块 | 用对应 torch 版本重编 `_C`；官方 pin 的是 `torch==2.13.0+cu132`，两个镜像都不是 |
|
||||
| 2 | **`deep_gemm` / `deep_ep`** | 缺 vendor CUDA 库，MoE/GEMM 无法执行。PyPI 上的 `deep-gemm 1.0.0` 不是 2.8.2 定制版，必须从 `vendor/` 编译 | 在目标镜像里 `scripts/build_vendor.sh`（需 NCCL 2.32.3 + nvcc，flagtree 工具链齐全） |
|
||||
| 3 | **`tile_kernels` 要 `tilelang>=0.1.15`** | flagtree 的 tilelang 是 **0.1.9**，缺 `tilelang.cuda` 模块 | 升级 tilelang，或彻底去掉对它的依赖（见下） |
|
||||
|
||||
**关键判断：「Triton 化」不等于「只装 Triton 就能跑」。** 真正卡住的是 ①torch 版本绑定的
|
||||
C++ 扩展、②两个 vendor CUDA 库。这两个跟 Triton 无关，换任何 kernel 语言都躲不开。
|
||||
|
||||
### 4.5 本轮为「Triton-only 环境」修掉的东西
|
||||
|
||||
原来 `import deepinfer` 会顺带拉起整个 TileLang 栈，在只装 Triton 的机器上直接 ImportError。
|
||||
本轮把 TileLang 概念从**接口层**剥出去：
|
||||
|
||||
| 文件 | 改动 |
|
||||
|---|---|
|
||||
| `ops.py` | 新增 `DEEPINFER_OPS_STRICT=1`：不 import 默认后端、不 fallback，缺 kernel 直接抛。这是「纯 Triton 环境」的前提 |
|
||||
| `norm.py` | `tilelang.language as T` 只在取 dtype 句柄时用 → 新增 `_tl_dtype()`，没装 TileLang 就返回字符串（Triton 工厂本来就忽略这个参数）；`tile_kernels` 移进 `test_norm` |
|
||||
| `compress.py` | `_enable_fast_math` 改成「有 TileLang 才设 pass_configs」，Triton 工厂的空 dict 直接跳过 |
|
||||
| `cache_common.py` | `tile_kernels` 移进 `ref_cast_from_bf16` / `ref_cast_to_bf16`（参考实现专用） |
|
||||
| `fused_transform.py` | 同上，移进 `..._ref()` |
|
||||
| `load_paged_kv_cache.py` | 同上，移进 `test_load_paged_kv_cache` |
|
||||
| `scripts/ab_cases.py` | 所有接口模块改函数内 import，缺包时该组 case 报 skip 而不是整个脚本起不来 |
|
||||
| `scripts/h20_ops_ab.py` | `--dump` / `--compare`；无基线后端时降级为「只跑选中后端」 |
|
||||
|
||||
**这是有实际价值的重构**：`tile_kernels` 原本只在参考实现和单测里用到，却卡在生产模块的
|
||||
import 路径上。修完之后 `import deepinfer` 在纯 Triton 环境不再需要 TileLang。
|
||||
|
||||
### 4.6 剩下的 Triton-only 障碍
|
||||
|
||||
- [ ] **4.6a** `deepinfer/_C` 仍是硬依赖：`sampling_inputs` → `_C.engine`。
|
||||
`fused_sample` / `map_lookup` 两组用例因此在两个镜像上都跳过。要么按目标 torch 重编，
|
||||
要么给 `sampling_inputs` 加纯 Python 回退（对 Triton 后端够用，只是慢）。
|
||||
- [ ] **4.6b** `deepinfer/platform.py` → `deep_cluster.hardware`，`ops.py` 顶层就要它。
|
||||
在 flagos 里第一次跑是 `ModuleNotFoundError: deep_cluster`，把 `deep_cluster` 也带上才行。
|
||||
应改成可选依赖。
|
||||
- [ ] **4.6c** 三套 triton 的 `autotune` / `num_stages` 行为不同，性能未测（对应 2.7）。
|
||||
|
||||
---
|
||||
|
||||
## 五、一句话结论
|
||||
|
||||
- **TileLang PR 版**：语义正确、数值已收敛到硬件地板，**卡在「默认配置跑不起来」（C1）+ 没有性能数据（C5）+ 没有 CI/验证 checklist（1.6）**。
|
||||
- **Triton 版**：非 MTP 路径已经是纯 Triton 且 token 逐位一致，**卡在 MTP 4 个 kernel + 采样器 verify 变体 + 0 个 pytest**。
|
||||
- **FlagTree / FlagOS 镜像**：Triton 编译器没问题（7/7 特性、19 个 kernel、115 用例跨 3.5.1/3.6.0/3.7.1
|
||||
逐元素一致）；**端到端跑不起来，卡在 torch 版本绑定的 `_C` 和 vendor 的 `deep_gemm`/`deep_ep`，与 Triton 无关**。
|
||||
- 两边共同的硬骨头只有 **C1（engram gather 不可 capture）**，其余都是收尾工作。
|
||||
196
scripts/flagtree_triton_probe.py
Normal file
196
scripts/flagtree_triton_probe.py
Normal file
@ -0,0 +1,196 @@
|
||||
"""Probe whether the Triton in this image can compile and run the kernel features
|
||||
DeepInfer's H20 Triton backend actually uses.
|
||||
|
||||
Each case is a miniature of a real kernel, so a failure names the feature that would
|
||||
block the port rather than just "triton is broken".
|
||||
"""
|
||||
|
||||
import os
|
||||
import sys
|
||||
|
||||
import torch
|
||||
import triton
|
||||
import triton.language as tl
|
||||
|
||||
print("=" * 72)
|
||||
print("python ", sys.version.split()[0])
|
||||
print("torch ", torch.__version__, "cuda", torch.version.cuda)
|
||||
print("triton ", triton.__version__)
|
||||
print("triton@", os.path.dirname(triton.__file__))
|
||||
print("device ", torch.cuda.get_device_name(0), torch.cuda.get_device_capability(0))
|
||||
print("flaggems imported?", "flag_gems" in sys.modules)
|
||||
try:
|
||||
from triton.runtime import driver
|
||||
|
||||
print("driver ", type(driver.active).__name__, "target", getattr(driver.active, "target", None))
|
||||
except Exception as exc: # noqa: BLE001
|
||||
print("driver probe failed:", exc)
|
||||
print("=" * 72)
|
||||
|
||||
|
||||
@triton.jit
|
||||
def _add(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):
|
||||
off = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)
|
||||
mask = off < n
|
||||
tl.store(out_ptr + off, tl.load(x_ptr + off, mask=mask) + tl.load(y_ptr + off, mask=mask), mask=mask)
|
||||
|
||||
|
||||
@triton.jit
|
||||
def _matmul(a_ptr, b_ptr, c_ptr, M, N, K, BM: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr):
|
||||
rm = tl.program_id(0) * BM + tl.arange(0, BM)
|
||||
rn = tl.program_id(1) * BN + tl.arange(0, BN)
|
||||
rk = tl.arange(0, BK)
|
||||
acc = tl.zeros((BM, BN), dtype=tl.float32)
|
||||
for k in range(0, tl.cdiv(K, BK)):
|
||||
kk = k * BK + rk
|
||||
a = tl.load(a_ptr + rm[:, None] * K + kk[None, :], mask=(rm[:, None] < M) & (kk[None, :] < K), other=0.0)
|
||||
b = tl.load(b_ptr + kk[:, None] * N + rn[None, :], mask=(kk[:, None] < K) & (rn[None, :] < N), other=0.0)
|
||||
acc = tl.dot(a, b, acc)
|
||||
tl.store(c_ptr + rm[:, None] * N + rn[None, :], acc.to(tl.float16), mask=(rm[:, None] < M) & (rn[None, :] < N))
|
||||
|
||||
|
||||
@triton.jit
|
||||
def _gather(pool_ptr, idx_ptr, out_ptr, n, BLOCK: tl.constexpr):
|
||||
"""Data-dependent gather inside a tile: this is how the compressor picks rows."""
|
||||
lanes = tl.arange(0, BLOCK)
|
||||
pool = tl.load(pool_ptr + lanes)
|
||||
off = tl.program_id(0) * BLOCK + lanes
|
||||
mask = off < n
|
||||
idx = tl.load(idx_ptr + off, mask=mask, other=0) % BLOCK
|
||||
tl.store(out_ptr + off, tl.gather(pool, idx, 0), mask=mask)
|
||||
|
||||
|
||||
@triton.jit
|
||||
def _rope_pair(x_ptr, out_ptr, n_rows, D: tl.constexpr, BLOCK: tl.constexpr):
|
||||
"""RoPE pairs (2k, 2k+1) or (k, k+D/2); both are one xor of the lane index."""
|
||||
row = tl.program_id(0)
|
||||
lanes = tl.arange(0, BLOCK)
|
||||
base = x_ptr + row * D
|
||||
left = tl.load(base + lanes)
|
||||
right = tl.gather(left, lanes ^ 1, 0)
|
||||
tl.store(out_ptr + row * D + lanes, left + right)
|
||||
|
||||
|
||||
@triton.jit
|
||||
def _fp8_quant(x_ptr, out_ptr, sf_ptr, n, BLOCK: tl.constexpr):
|
||||
pid = tl.program_id(0)
|
||||
off = pid * BLOCK + tl.arange(0, BLOCK)
|
||||
mask = off < n
|
||||
v = tl.load(x_ptr + off, mask=mask, other=0.0).to(tl.float32)
|
||||
amax = tl.max(tl.abs(v))
|
||||
sf = tl.where(amax > 0.0, amax / 448.0, 1.0)
|
||||
tl.store(sf_ptr + pid, sf)
|
||||
tl.store(out_ptr + off, (v / sf).to(tl.float8e4nv), mask=mask)
|
||||
|
||||
|
||||
@triton.jit
|
||||
def _int64_cdf(p_ptr, draw_ptr, out_ptr, n, BLOCK: tl.constexpr):
|
||||
"""The sampler's CDF is int64 on purpose: uint32 comparisons go signed in Triton."""
|
||||
off = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)
|
||||
mask = off < n
|
||||
cdf = tl.load(p_ptr + off, mask=mask, other=0).to(tl.int64)
|
||||
draw = tl.load(draw_ptr).to(tl.int64)
|
||||
tl.store(out_ptr + off, (cdf > draw).to(tl.int32), mask=mask)
|
||||
|
||||
|
||||
def run(name, fn):
|
||||
try:
|
||||
fn()
|
||||
print(f"[ OK ] {name}")
|
||||
except Exception as exc: # noqa: BLE001
|
||||
print(f"[FAIL] {name}: {type(exc).__name__}: {str(exc)[:400]}")
|
||||
|
||||
|
||||
def case_add():
|
||||
n = 1 << 20
|
||||
x, y = torch.randn(n, device="cuda"), torch.randn(n, device="cuda")
|
||||
out = torch.empty_like(x)
|
||||
_add[(triton.cdiv(n, 1024),)](x, y, out, n, 1024)
|
||||
torch.cuda.synchronize()
|
||||
err = (out - (x + y)).abs().max().item()
|
||||
assert err < 1e-5, err
|
||||
print(f" max abs err {err:.2e}")
|
||||
|
||||
|
||||
def case_matmul():
|
||||
M = N = K = 512
|
||||
a = torch.randn(M, K, device="cuda", dtype=torch.float16)
|
||||
b = torch.randn(K, N, device="cuda", dtype=torch.float16)
|
||||
c = torch.empty(M, N, device="cuda", dtype=torch.float16)
|
||||
_matmul[(triton.cdiv(M, 64), triton.cdiv(N, 64))](a, b, c, M, N, K, 64, 64, 32)
|
||||
torch.cuda.synchronize()
|
||||
ref = a.float() @ b.float()
|
||||
rel = ((c.float() - ref).abs().max() / ref.abs().max()).item()
|
||||
assert rel < 5e-2, rel
|
||||
print(f" rel err {rel:.4f} (fp16 accumulate in fp32)")
|
||||
|
||||
|
||||
def case_gather():
|
||||
n, BLOCK = 4096, 256
|
||||
pool = torch.arange(BLOCK, device="cuda", dtype=torch.int32)
|
||||
idx = torch.randint(0, BLOCK, (n,), device="cuda", dtype=torch.int32)
|
||||
out = torch.empty(n, device="cuda", dtype=torch.int32)
|
||||
_gather[(triton.cdiv(n, BLOCK),)](pool, idx, out, n, BLOCK)
|
||||
torch.cuda.synchronize()
|
||||
assert torch.equal(out, pool[idx.long()]), "gather returned wrong rows"
|
||||
print(" rows match pool[idx]")
|
||||
|
||||
|
||||
def case_rope_pair():
|
||||
rows, D = 64, 128
|
||||
x = torch.randn(rows, D, device="cuda")
|
||||
out = torch.empty_like(x)
|
||||
_rope_pair[(rows,)](x, out, rows, D, D)
|
||||
torch.cuda.synchronize()
|
||||
# `lanes ^ 1` swaps adjacent lanes, so the reference is that permutation, not a roll.
|
||||
ref = x + x[:, torch.arange(D, device="cuda") ^ 1]
|
||||
err = (out - ref).abs().max().item()
|
||||
assert err < 1e-5, err
|
||||
print(" adjacent-lane pairing exact")
|
||||
|
||||
|
||||
def case_fp8():
|
||||
n, BLOCK = 1 << 18, 1024
|
||||
x = torch.randn(n, device="cuda")
|
||||
out = torch.empty(n, device="cuda", dtype=torch.float8_e4m3fn)
|
||||
nblk = triton.cdiv(n, BLOCK)
|
||||
sf = torch.empty(nblk, device="cuda")
|
||||
_fp8_quant[(nblk,)](x, out, sf, n, BLOCK)
|
||||
torch.cuda.synchronize()
|
||||
deq = out.float() * sf.repeat_interleave(BLOCK)[:n]
|
||||
rel = ((deq - x).abs().max() / x.abs().max()).item()
|
||||
assert rel < 1e-1, rel
|
||||
print(f" fp8 e4m3 per-1024 roundtrip rel {rel:.4f}")
|
||||
|
||||
|
||||
def case_int64_cdf():
|
||||
n, BLOCK = 4096, 256
|
||||
cdf = torch.arange(n, device="cuda", dtype=torch.int64) * 1000
|
||||
draw = torch.tensor([123456789], device="cuda", dtype=torch.int64)
|
||||
out = torch.empty(n, device="cuda", dtype=torch.int32)
|
||||
_int64_cdf[(triton.cdiv(n, BLOCK),)](cdf, draw, out, n, BLOCK)
|
||||
torch.cuda.synchronize()
|
||||
expect = (cdf > draw[0]).to(torch.int32)
|
||||
assert torch.equal(out, expect), "int64 comparison went wrong"
|
||||
print(" int64 compare exact")
|
||||
|
||||
|
||||
def report_ptx():
|
||||
ck = _add.warmup(torch.empty(1, device="cuda"), torch.empty(1, device="cuda"), torch.empty(1, device="cuda"), 1, BLOCK=1024, grid=(1,))
|
||||
head = [ln.strip() for ln in ck.asm["ptx"].splitlines() if ln.strip().startswith((".version", ".target"))]
|
||||
print(" ", head)
|
||||
print(" num_warps", ck.metadata.num_warps, "shared", ck.metadata.shared)
|
||||
|
||||
|
||||
for label, fn in [
|
||||
("elementwise add", case_add),
|
||||
("tl.dot matmul (tensor core)", case_matmul),
|
||||
("tl.gather on tile (data-dependent)", case_gather),
|
||||
("tl.gather for RoPE pairing", case_rope_pair),
|
||||
("fp8e4nv cast + block scale", case_fp8),
|
||||
("int64 CDF comparison", case_int64_cdf),
|
||||
("PTX target", report_ptx),
|
||||
]:
|
||||
run(label, fn)
|
||||
|
||||
print("\nDONE")
|
||||
Loading…
x
Reference in New Issue
Block a user