diff --git a/TODO.md b/TODO.md
new file mode 100644
index 0000000..2239e85
--- /dev/null
+++ b/TODO.md
@@ -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）**，其余都是收尾工作。
diff --git a/scripts/flagtree_triton_probe.py b/scripts/flagtree_triton_probe.py
new file mode 100644
index 0000000..c8a362a
--- /dev/null
+++ b/scripts/flagtree_triton_probe.py
@@ -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")