DeepInfer/TODO.md

H20 / SM90 遗留问题总清单

截止 2026-10-22 10:00。分两条线：

• TileLang PR 版 = 分支 h20-sm90-support → 上游 PR #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：

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 逐条排除）。

• 结论：不是 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 的验证证据）

• prefill 因果掩码 bug（首 token 从 EVAL - 变回 17607）
• 专家权重 scale 广播 bug（requant_expert_fp4 把 (n_out, k/32) 当 (n_out/32, k/32) 展开） → w1/w3/w2 的 cos 从 0.89 到 0.9997~1.0000
• layer0 整层 rel_err 1.27e-1 → 2.88e-2，整条误差曲线压平一个数量级
• KV cache 闭环 round-trip（rel 0.027，FP8 地板）
• _decode_p_sum_sm90 vs 独立 float64 参考（rel 1e-7）
• decode 逐步误差曲线（Δmargin mean +0.711 / std 1.628，8 位置 7 个 top-1 正确）
• 自然文本提问答案完全正确（北京。，8 rank 一致）
• 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	✅ 已完成（2026-10-22）
2.3	采样器 rejection-sampling 变体未实现	🟡 P1	显式委托 cuda
2.4	16 个可选路径 kernel 缺失（导入期报 lacks 16 kernel(s)）	🟡 P1	见 2.2b
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.11	国产卡（is_sm90=False）上 candidate_ 变成每层热路径*	🔴 P1	未做，见 2.2b

2.2 MTP 4 个 kernel（✅ 2026-10-22 全部完成）

kernel	结果	说明
get_accept_result_tl	✅	9 个 A/B 用例（has_golden × golden 混合比例 × 跨 block）。前缀扫描在 Triton 里用「block 内最长 span」做上界，短 span 的 lane 在多余迭代上保值
finalize_verification_tl	✅	10 个变体（has_sum_smaller_probs × has_confidence × num_steps 0/1/5 × topk 0/8/130）。batch_confidence_error 改成 program 内归约 + 单次 atomic（原实现逐行 atomic，Triton 里标量指针配向量值不合法），因此浮点求和顺序与 baseline 不同，用例标 exact=False
get_fused_dflash_get_input_tl	✅	9 个用例（2 个 mask 开关 × padding 请求 × kv_len 16/128/232）。padded_kv_len 是 64 的倍数但不是 128 的倍数，列循环必须向上取整否则整行尾部不写
remap_draft_topk_tl	✅	5 个用例（含 force_topk=4096 / vocab=129280）。算法换了：baseline 用 shared 内存池按 vocab 分块建表再查表，Triton 无 shared 原语，改成 target×draft 直接两两匹配（draft index 唯一 → 每行至多一个匹配，取 max 即可）。复杂度从 O(vocab/pool_size × topk) 变成 O(topk²)，force_topk 默认 0（不在必经路径），需要时再优化

2.2b 剩下的 16 个：按「国产卡上会不会崩」排序

从 20 降到 16。这 16 个没有一个无条件必经，但分布很不均匀：

优先级	kernel	条件	为什么
🔴 国产卡热路径	candidate_block_max_tl get_candidate_compact_lens_tl restore_candidate_indices_inplace_tl	not PLATFORM.is_sm90 且 checkpoint 的 candidate_source_layer_id 非 null	本仓库的 V4.1 checkpoint 配的是 candidate_source_layer_id=20 / 40 层。H20 上 is_sm90=True 把这条路径硬关掉了（index_attn.py:126-129），所以从来没测到；国产卡 is_sm90=False，第 20 层之后的 19 层每步都要走 restore_candidate_indices_inplace。⚠️ 但前置阻塞更硬：这条路径同时要 deep_gemm.fp8_fp4_sparse_mqa_logits，那是 SM100-only（arch_major==10），光移植这 3 个 kernel 没有意义
🟡 服务场景会撞	apply_token_bitmask_by_step_inplace_tl	请求带 schema / tool call（so_indices is not None）	唯一「普通用户一开就崩」的：与部署配置无关，纯请求参数驱动
🟡	batch_downsample_tl replace_input_embeds_tl rgb_pad_to_patches_tl	input_metadata.images_raw is not None	三件套，有图请求必然全走，无降级
⚪ 显式开关	prepare_cp_mirror_tl	CP>1
⚪	sum_smaller_probs_tl	DEEPINFER_COMPUTE_SUM_SMALLER_PROBS（默认 False）
⚪	compute_seqlens_q_by_confidence_rates_tl	DEEPINFER_ENABLE_MTP_VARLEN（默认 False）+ ZOS
⚪	copy_dflash_decode_input_tl	DFlash CUDA Graph（DEEPINFER_ENFORCE_EAGER=1 时反而不走）	关掉优化就能绕过，可作为国产卡首轮降级开关
⚪	yoco_cp_boundary_fused_tl	CP>1 + DEEPINFER_YOCO_PREFILL_TRUNCATION_MODE=compact（默认 off）
⚫ 生产不可达	detect_outlier_and_abort_tl load_state_to_cache_tl select_state_slice_tl random_sampling_tl	无	前三个连测试都没有；load_state_to_cache/select_state_slice 的宿主方法全仓库无调用方，只在 tests/worker/test_state_cache_ops.py 用到。可以一直不做

2.10 Triton-only 环境（vendor 镜像 / CI 镜像）

「技术栈换成 Triton」的隐含目标是「只装 Triton 也能跑」，但当前做不到，因为 TileLang 概念漏到了接口层。已在第四章验证并修掉第一批： ops.py 加 DEEPINFER_OPS_STRICT、tile_kernels / tilelang 的 import 移进参考实现与单测函数。

2.10b ⚠️ 「没有 TileLang」的真实阻塞面比 16 个 kernel 大得多（2026-10-22 实测）

做了一次「假装这两个包不存在」的实测：用 sitecustomize.py 在 meta_path 上拦 tilelang / tile_kernels / tvm_ffi（必须走 sitecustomize，因为 8 个 DP rank 是 spawn 出来的，父进程的拦截不继承），然后 import deepinfer。

结果：import 就失败了，根本走不到那 16 个 kernel。

hugging_face/weights.py:38 → model_executor/layers/linear.py:4 → import tile_kernels
ModuleNotFoundError: No module named 'tile_kernels'

为什么：tilelang_ops 那层抽象只包住了「kernel 库」里的 44 个 kernel，但模型前向 直接调 tile_kernels（它是 FlagOS 的另一个包，tile_kernels/__init__.py:1 就是 import tilelang，140 个 py 里 166 处引用 → 留着 tile_kernels 等于留着 tilelang）。

AST 扫描结果（排除 tilelang_ops/cuda/ 与 asc/）：

项	数量	分布
模块级 import tile_kernels	25 处 / 13 个文件	linear rotary_embedding deepseek v4_modules/{attn,engram,engram_metadata,hc,mhc_ops,mhc_modeling,mlp} vision_encoder/vit parallel_utils/userbuffer
直接调用的 API	23 个	见下

按移植成本分三类：

类别	API	说明
琐碎	set_num_sms set_pdl set_token_alignment is_ascend utils.dtype_to_str use_deterministic_algorithms	配置开关与字符串工具，几行代码替代
热路径真 kernel	quant.per_token_cast（FP8 per-token 量化，linear/mlp/attn/hc 四处都用） quant.swiglu_forward_and_per_token_cast（融合 SwiGLU+量化） moe.moe_topk_gate transform.apply_rotary quant.swiglu_forward quant.batched_transpose_weight_sf	每 token 每层都过，必须真写 Triton，纯 torch 版本会慢到不可用
DeepSeek 专有重 kernel	9 个 mhc.*（expand / head_compute_mix / norm_fn / post / pre_apply_mix / pre_big_fuse / pre_split_mixes / sinkhorn）+ engram.engram_hash	MHC 注意力与 Engram，语义复杂，工作量最大

唯一现成的退路：tile_kernels/torch/ 是一整套纯 torch 参考实现（cast engram mhc moe norm reduce_fused rope swiglu topk，0 处 tilelang），导出 engram_hash_ref、8 个 mhc_*_ref、moe_topk_gate_forward、swiglu_forward、 rotary_embedding_ref、cast/cast_back、batched_transpose_weight_sf、 norm_forward_ref 等，几乎覆盖上表所有重 API。但 tile_kernels/__init__.py 第一行 import tilelang 会挡住 import tile_kernels.torch，需要把 __init__ 改成 lazy （或把 torch/ 抽成独立包）。

→ 这条能给国产卡一个「正确但慢」的 bring-up 桥：先用 *_ref 跑通数值，再按 profile 把热路径逐个换成 Triton。flag_gems 不能直接替代（903 个 ATen 级算子，没有这些定制融合算子）。

结论：不要再把「国产卡上不能有 tilelang」说成「只差 16 个 kernel」。 16 个 kernel 是小头，模型前向对 tile_kernels 的直接依赖是大头。

2.2 MTP 4 个 kernel（✅ 2026-10-22 全部完成，见下表）

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 缺失的 16 个 kernel = 全可选路径

对着 cuda/ 的 44 个 *_tl 与 triton/ 的 28 个做差集得到，与导入期那行 lacks 16 kernel(s) 逐个对得上（first = apply_token_bitmask_by_step_inplace_tl）。 逐项的条件与优先级见上面 2.2b。

模块	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	共享层候选块优化（国产卡上会打开，见 2.2b）
prepare_cp_mirror	prepare_cp_mirror / yoco_cp_boundary_fused	CP > 1
state_cache	load_state_to_cache / select_state_slice	生产不可达
schedule_mtp_varlen	compute_seqlens_q_by_confidence_rates	varlen MTP
dflash_copy	copy_dflash_decode_input	dflash CUDA Graph
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.0 先厘清概念：FlagTree 就是 Triton 的一个 fork

这一点先说清楚，否则「FlagTree 编不过」是个误会。镜像里：

flagtree.__version__ == triton.__version__ == 3.6.0

因为 flagtree/__init__.py 只有一行 from triton import __version__， flagtree-0.7.0.dist-info 的 Summary 也写明 「forked from triton-lang/triton」。

所以：

• FlagTree 不是一个需要「替代 Triton」的独立东西，它就是 Triton，fork 出来加多芯片后端。 发行名是 tokenspeed-triton，但 import 名仍然是 triton（为了源码兼容）。
• flagtree 这个 Python 包只装了 FlagPrism（调试器/性能分析附件），编译能力在 triton 包里。
• 因此「用 FlagTree 跑我们的 Triton kernel」= 用那个 fork 过的 triton 包跑，这正是它的主场。

FlagTree 编 Triton kernel 从来没有失败过。 下面 §4.3 的「跑不起来」全部是 与 Triton 无关的 C++/CUDA 扩展（_C、deep_gemm、deep_ep）和 TileLang 包 （tile_kernels），见 §4.4 表格逐项标注了「谁在编」。

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 = FlagTree（tokenspeed-triton 3.8.10.post20260721）
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 无辜）

先看「谁在编」，否则容易误判成 Triton 不行：

依赖	谁来编译	跟 Triton 有关吗
19 个 Triton kernel	FlagTree / triton	就是它自己，7/7 + 115 用例全过
deepinfer/_C	g++ + nvcc（torch C++ 扩展）	❌ 无关
deep_gemm / deep_ep	g++ + nvcc（vendor CUDA 扩展）	❌ 无关
tile_kernels	TileLang（另一个编译器栈）	❌ 无关

三层依赖，按难度排序：

#	缺什么	影响	修法
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 模块。但这只影响非 strict 模式：设 DEEPINFER_OPS_STRICT=1 后完全不碰 TileLang 栈，已实测 norm 可 import	升级 tilelang，或彻底去掉对它的依赖（见 4.5，本轮已做到）

关键判断：「Triton 化」不等于「只装 Triton 就能跑」。 真正卡住的是 ①torch 版本绑定的 C++ 扩展、②两个 vendor CUDA 库。这两个跟 Triton 无关，换任何 kernel 语言都躲不开 ——它们跟原来用 TileLang 时是同一份 C++ 代码。

4.4b 在 flagtree 镜像里重编：四个目标全部完成（2026-10-22 11:45）

在 harbor.baai.ac.cn/flagtree/...:202609-3.6-vllm0.24.0（torch 2.11 + cu13 + gcc 13.3 + nvcc 13.0）里建了个专用构建容器 h20-ft-build，vendor/ 用现成 checkout，源码 tar 过去。

目标	结果	踩到并修掉的东西
deepinfer._C	✅ 编出来了，import 成功，34 个符号（含 LogitsTransformInputs）	①setup.py 强制要 CONDA_PREFIX/VIRTUAL_ENV → 建 --system-site-packages venv；②scripts/build.sh 要 python 而非 python3 → 补软链；③镜像根分区只剩 3G → 删无人使用的 seed-oss 镜像释放 37G；④macOS 打包带进 ._* 文件让 compileall 报错 → 无碍
deep_gemm	✅ 2.8.2	①缺 elfutils/libdwfl.h → apt install libdw-dev libelf-dev；②缺 cusparse.h（镜像只装 CUDA 运行时不带头）→ 把 nvidia/cu13/include/*.h 拷进 /usr/local/cuda/include
deep_ep	✅ 2.5.0	缺 NCCL 2.32.3（ncclCommProperties / ginMinStride 是新 API），镜像只有 2.28.9 → 阿里源下 305MB wheel（pip 下载会超时，用 curl -C - 续传）
deep_select	✅ 1.0.0	①setup.py 顶层 import tests.kernelkit → 要 triton；②要 lspci → apt install pciutils git

验证（单卡 H20，strict 模式）：

torch 2.11.0+cu130 cap (9,0) | triton 3.6.0 (= FlagTree 3.6.0)
vendor: deep_gemm 2.8.2 | deep_ep 2.5.0 | deep_select 1.0.0 | nccl 2.32.3
_C: 34 symbols | ops backend triton strict True, 24 kernels, 0 inherited
gather kernel: exact
全量 A/B --dump: 125 用例，0 跳过（fused_sample / map_lookup 不再缺 _C）
全量 A/B --compare di-h20(3.7.1) 的 dump: AB_OK —— 逐元素一致

所以「vendor 镜像里跑 Triton 版」这条路已经通了。

4.4c 8 卡端到端：跑通了，且与参考镜像逐 token 相同 ✅（2026-10-22 12:25）

在 flagtree 镜像里 --ep 8 --max_tokens 32，两个 prompt：

=== continuation: '。\n问题：法国的首都是哪里？\n答案：巴黎。\n问题：请介绍一下中国的首都。\n
                    答案：北京是中国的首都，是中国的政治、文化和国际'
=== continuation: '北京。\n问题：中国的首都是哪里？\n答案：北京。\n问题：中国的首都是哪里？\n
                    答案：北京。\n问题：中国的首都是哪里？\n答案'

检查项	flagtree 镜像	di-h20 参考镜像
exit code	0	0
每个 prompt 的 token 序列出现次数	8（8 rank 全一致）	8
distinct token 序列	2（= 2 个 prompt）	2
token ids 逐位对比	完全一致（diff 无输出）	基线

回退情况与 di-h20 完全相同，符合已知的 P1/P2 遗留：

[ops] backend 'triton' lacks 20 kernel(s), falling back to 'cuda' (first: apply_token_bitmask_by_step_inplace_tl)
[ops] kernel 'remap_draft_topk_tl'        used from 'cuda'   ← MTP
[ops] kernel 'get_accept_result_tl'       used from 'cuda'   ← MTP
[ops] kernel 'finalize_verification_tl'   used from 'cuda'   ← MTP
[ops] kernel 'get_fused_dflash_get_input_tl' used from 'cuda' ← MTP
[ops] fused_sample: rejection sampling (MTP verify) is not implemented in the triton backend

良性差异（非 Triton、与 torch 版本有关）：

stats_logger: nccl_memory stats_unavailable=AttributeError(
    "'ProcessGroupNCCL' object has no attribute 'memory_stats'")

torch 2.11 没有 ProcessGroupNCCL.memory_stats（2.13 才有），只影响一行统计日志。

跑端到端还需要的额外依赖（A/B 单测不需要，引擎 import 链需要）： tilelang 0.1.15 + tile_kernels 2.1.0 + tvm_ffi 0.1.11（后者被前者依赖，镜像自带 0.1.9 缺 tvm_ffi.dataclasses.Enum），以及 requirements.txt 里的 19 个包。

⚠️ 这个机器上 pip 装大包会卡死（posix_ipc 的 wheel 12 分钟没下来）。可行的办法是 docker cp 从已装好的参考镜像拷 site-packages 再挑包复制。pillow 反而要用镜像自带的 （12.2），拷 di-h20 的 11.3 会缺 pillow.libs 里的 libtiff-*.so。

4.4d 上面这些已经写成脚本（可复现，已从零验证）

scripts/setup_flagtree_image.sh 把 §4.4b~4.4d 的全部踩坑固化下来，三阶段：

# 宿主机上先备好参考包（docker cp 只在容器外可用）
docker cp <flagtree镜像>:/usr/local/lib/python3.12/dist-packages/. /data/ftbuild/ref-triton/
docker cp <tilelang镜像>:/usr/local/lib/python3.12/dist-packages/. /data/ftbuild/ref-tilelang/
docker cp <引擎镜像>:/usr/local/lib/python3.12/dist-packages/.     /data/ftbuild/ref-site/

# 容器内
bash scripts/setup_flagtree_image.sh env      # apt / venv / cusparse 头 / NCCL / triton / tilelang / py 依赖
bash scripts/setup_flagtree_image.sh build    # _C + DeepGEMM + DeepEP + DeepSelect
bash scripts/setup_flagtree_image.sh verify   # import 与 ops 后端自检

配套 scripts/copy_site_packages.py（按 distribution 的文件清单复制，所以 pillow→PIL、 pyzmq→zmq 这种名字不一致的也能带全；--keep-system pillow 保留镜像自带的）。

从零验证过：新建一个空 venv 跑完 env→build→verify，再跑 scripts/h20_e2e_report.sh：

env            19 个依赖全部就位，triton 3.6.0 = flagtree 3.6.0，tilelang 0.1.15
build          _C + deep_gemm 2.8.2 + deep_ep 2.5.0 + deep_select 1.0.0（3m33s）
verify         torch 2.11.0+cu130 cap (9,0) | _C 34 symbols | triton strict, 24 kernels
e2e report     exit=0，两个 prompt 的 token 序列各出现 8 次，无 traceback

也就是说这些步骤不是「在我这个容器里碰巧能跑」，换一个空 venv 同样成立。

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 镜像：FlagTree 就是 Triton 的 fork，编 Triton kernel 从来没失败过 （7/7 特性、19 kernel、125 用例跨 3.5.1/3.6.0/3.7.1 逐元素一致）。原先跑不起来的是 _C + vendor 的 C++/CUDA 扩展，已在 flagtree 镜像里全部重编成功，并且 8 卡端到端已跑通，输出与参考镜像逐 token 相同。剩下的回退就是 P1/P2 那几条遗留， 在两个镜像上一模一样。
• 两边共同的硬骨头只有 C1（engram gather 不可 capture），其余都是收尾工作。

---

六、国产卡 bring-up：走 tile_kernels.torch（2026-10-22 实测）

目标：在没有 TileLang 的卡上先把模型跑起来，先用纯 torch 参考实现保正确性，再按 profile 把热路径逐个换成 Triton。§2.10b 说明了大头在哪（模型前向直接调 tile_kernels 的 quant/MoE/MHC/engram，25 处模块级 import / 23 个 API）。

6.0 当前到哪一步了（2026-10-25 更新，无 tilelang 端到端已跑通）

#	事项	状态
1	import tile_kernels.torch（无 tilelang）	✅ 通
2	算子数值等价（23 个 API 实际调用 + 18 用例两模式对比）	✅ 见 6.2
3	路由层 + 25 处 import 替换	✅ 已通
4	无 tilelang 下 23 个 API 冒烟	✅ ALL_CALLS_OK
5	端到端（DEEPINFER_KERNELS=ref）	✅ exit=0，8 DP 一致，回答正确
6	稀疏 MLA 换成 Triton（FlagGems）	✅ 见 6.7，7 形状 + decode 5 用例全一致
7	屏蔽 tilelang 的端到端	✅ exit=0，0 回退，回答正确（需关 MTP，见 6.7）
8	与 TileLang 版逐 token 一致	❌ 前 15/1 个 token 后分叉（控制实验证明非噪声，见 6.5）
9	百分之百不用 tilelang（含采样器 verify 变体）	❌ fused_sample 的 rejection-sampling 仍需（见 6.7）
10	性能	❓ 仍未测

6.5 端到端结果：功能对，但不逐 token 一致

DEEPINFER_KERNELS=ref（tilelang 仍可导入，只服务 attention）跑 16 个样本（2 prompt × 8 DP）：

• exit=0，hard failures: none，8 个 DP 逐 token 互相一致，回答语义正确 （「问题：中国的首都是哪里？答案：北京。」）
• 但和 DEEPINFER_KERNELS=cuda 比：16/16 样本不同。
  • prompt A：前 15 个 token 相同，第 16 个分叉
  • prompt B：第 1 个 token 就分叉

控制实验（关键）：cuda 连跑两次，16 个样本逐位相同（0 差异）。所以 harness 在 贪心解码下是确定性的，上面的差异不是采样噪声，是真实数值差异。

未归因的候选（按可能性）：

1. batched_transpose_weight_sf —— 唯一没做 A/B 的路径算子，且作用于所有权重的 scale factor。 尝试补测时我的输入布局不对，编译版直接 device-side assert；需按 hugging_face/parameters.py 的真实 3-D (num_groups, -1, n) 布局重做。
2. apply_rotary 1 bf16 ULP（1/65536 元素）—— 每层都过，会累积。
3. engram_gate 相对 3e-7。
4. prompt B 第 1 个 token 就分叉，暗示 prefill 阶段已存在实质差异，需要对比首个 token 的 logits 才能定位，而不是继续猜。

6.6 稀疏 MLA 的 Triton 方案已找到，且已数值验证（2026-10-25）

原判断「SM90 稀疏 MLA 没有 Triton 版本、工作量远大于 compat 层」需要修正： FlagGems 里已经有，而且是 drop-in 的。

flagos-ai/FlagGems → src/flag_gems/fused/flashmla_sparse.py → flash_mla_sparse_fwd 纯 Triton，实测在 H20 上编译执行通过，7 组形状与现有 TileLang kernel 全部一致（≤2 bf16 ULP）。

参数逐项对应，连两个可选参数都有：

DeepInfer prefill_attn_varlen_sparse	FlagGems flash_mla_sparse_fwd
query [total_q, H, D]	q [s_q, h_q, d]
kv_a_with_mqa [total_k, 1, D]	kv [s_kv, h_kv, d_qk]
index_topk [total_q, topk]	indices [s_q, h_kv, topk]
softmax_scale	sm_scale
attn_sink [H]	attn_sink [h_q]
index_length [total_q]	topk_length [s_q]

语义也对得上：attn_sm90.py:110 明确让 prefill 走 is_causal=False（索引由 indexer 保证 causal，用行号当因果边界会导致整层 NaN，这个 bug 已经发生过一次）；FlagGems 这个 kernel 同样 没有 causal 参数，理由一致。V4.1 的 h_q=64 / h_kv=1 / d_qk=512 恰好满足它所有 assert。

不能用的那个：同仓库 DSA/sparse_mla.py（triton_sparse_mla_fwd_interface）硬编码 is_causal=True，且要求 V3.2 的 576 拆分布局（latent 512 + rope tail 64）。V4.1 把 RoPE 并进 512 维内，TD = 512-512 = 0，tl.arange(0,0) 直接编译失败。别选错。

验证脚本与 vendor 副本：scripts/sparse_mla_port/（README.md 有完整对照表与未做项）。

6.7 接线完成，无 tilelang 端到端跑通（2026-10-25）

三步都做完了，且每步先过单测再往下走：

步	内容	结果
1	内联 FlagGems 5.0.2 kernel + 7 形状 A/B	ALL_MATCH ≤2 ULP
2	decode 路径 A/B（fp8 cache、5 用例含双缓存）	ALL_MATCH
3	wrapper 契约单测（返回元数、lse 单位、dtype、两种调用约定）	ALL_PASSED 5 组
4	分发层 + attn_sm90 两个入口单测	ALL_PASSED 4 组
5	屏蔽 tilelang 端到端	exit=0，0 回退，回答正确
6	有 tilelang 端到端（对照）	exit=0，0 回退，回答正确

接线方式：新增 h20_ops/sparse_mla.py 作分发层，DEEPINFER_SPARSE_MLA_BACKEND=tilelang|triton 显式选择，不设则「能 import tilelang 就用它，否则 Triton」。attn_sm90 只改了 import 那一行， 两个调用点未动。每个 rank 打印一次实际选择。

为什么内联而不依赖 flag_gems：flag_gems/__init__.py 会 torch.library.Library("aten","IMPL")

• replace_customized_ops，进程级替换 ATen 算子；为了一个 kernel 去 import 它，会把整个引擎的 算子悄悄换掉。PyPI 5.0.2 的那个文件零内部依赖，内联后只需 torch/triton。

踩到并修掉的两个坑：

1. FlagGems 5.0.2 的 block 跳过守卫是错的：if ck * BK <= max_col 按索引在数组里的位置 判断，非因果时 max_col = SKV-1，于是 topk > SKV 时超出前 SKV 个的块被整块跳过。 SQ=16/SKV=128/topk=512 时输出差 2.8e-2。任何短于 512 token 的 prompt 都会踩到。 已本地修正为数据相关守卫（master 版就是这个写法），文件里有 LOCAL FIX 标注。
2. lse 单位：kernel 出自然对数，TileLang 出 log2。不转换会让 _apply_sink 的系数差 1.4427 倍。 单测里专门断言 ratio=1.0（未转换会是 0.693），避免以后静默回归。

仍需注意：

• 纯 Triton 端到端必须 DEEPINFER_USE_NEXTN_DECODE=0（关投机解码）。开着 MTP 时会走 fused_sample 的 rejection-sampling 变体，那是唯一还依赖 TileLang 的 kernel（Philox 随机流 + 固定 reduction 顺序，Triton 版会改变生成文本，见 triton/__init__.py 的说明）。
• 不与 TileLang 版逐 token 一致：prompt A 前 15 个 token 一致、第 16 个分叉；prompt B 第 2 个 分叉。控制实验（TileLang 连跑两次）逐 token 相同，故非采样噪声，是累加顺序差异沿 40+ 层 累积，与 §6.5 记录的 2.7% logits 漂移同源。

tilelang_ops/cuda/ 那 29 个模块已有 tilelang_ops/triton/ 对应版本，靠 DEEPINFER_OPS_BACKEND=triton 切换，本来就不是缺口。

6.4 端到端曾经为什么跑不起来（环境 + 排查记录）

先说结论：10-22 我判断「GPU 被同事的容器占用」是错的。 当时只看 nvidia-smi + docker ps 列表就推断，没做 PID→cgroup 映射。10-24 复查发现那 6×82 GB 是我自己 h20-ft-build 容器 里两天前 e2e 被 pkill 后遗留的死锁进程：

• state=R，CPU 时间 ≈ 存活时间（2 天 1 小时）→ 每个进程 100% 空转（NCCL barrier 自旋）
• 日志实证 TCPStore server has shut down + Broken pipe 无限重试 → 永不恢复
• 副作用：e2e-ref.log 被刷到 1.4 GB，6 个 CPU 核白烧 2 天
• 根因：pkill -f h20_ask 只杀父进程，multiprocessing.spawn 的 worker 被 PID 1 收养， 继续持有 CUDA context。清完父进程后还有 6 个 state=S 的子进程持有 nvidia fd，必须一起清

判断规则：看 GPU 被谁占用必须走 PID→cgroup→容器名，不能靠「谁 Up 得久」。 nvidia-smi --query-compute-apps=pid 拿 PID → cat /proc/<pid>/cgroup。

另一个坑：/tmp（/dev/vda 252G）会满。/data 是 bind mount，删它不释放 /tmp； 清 /etc/dsw-logs 同设备轮转日志才有效。

6.1 第一步已通：tile_kernels.torch 能在无 TileLang 下导入

原本不行，因为三层 eager import 把纯 Python 的参考实现和编译栈绑在一起：

层	问题	处理
tile_kernels/__init__.py	from . import (config, quant, …) 拖起全部编译 kernel	改成 PEP 562 惰性 __getattr__
tile_kernels/quant/__init__.py	13 个 kernel 模块 eager re-export；import quant.common 会先执行它	同样惰性化（AST 解析原文件生成导出表，带 __all__ 保证 import * 仍可用）
tile_kernels/quant/common.py	from tilelang import language as T + 模块级 @T.macro	tilelang 变可选；@T.macro 段包进 if T is not None:；4 个 host 侧注解加引号；3 个返回 dtype 句柄的成员改成调用时再解析
tile_kernels/config.py	只为 3 个常量 T.max_value(...) 而 import tilelang	常量改字面量（448.0 / 6.0），逐位校验一致

⚠️ 不能用 from __future__ import annotations：它把 @T.macro 体内的注解也变成字符串， tilelang 的 macro 解析器自己去 eval 时会报 NameError: name 'float32' is not defined。 必须只把 host 侧那 4 个注解单独加引号。

脚本：scripts/enable_torch_kernel_refs.sh（幂等，--revert 可逆）

• scripts/tile_kernels_patch/{patch_tile_kernels,patch_quant_common,lazify_package_init}.py
• scripts/verify_tile_kernels_patch.py（证明有 tilelang 时接口面完全不变：子模块属性、 from pkg import sub、config 旋钮、kernel 侧的 from .common import *、dir()、 未知属性仍报 AttributeError → PATCH_SAFE）

6.2 关键未知项已消除：参考实现数值等价 ✅

scripts/tile_kernels_patch/compare_with_cuda.py，用 DeepInfer 在 SM90 上的真实参数 （fmt="e4m3", num_per_channels=128, round_sf=True, use_packed_ue8m0=False, use_tma_aligned_col_major_sf=True）：

op	结论
quant.per_token_cast	bitwise 一致（data + sf，4 组 shape 含 m=1/257、n=128/7168）
quant.cast_back	rel 0.000e+00
quant.swiglu_forward_and_per_token_cast	bitwise 一致
engram.engram_hash	bitwise 一致（溢出敏感，必须是 int64 池算术）
engram.fused_weight	bitwise 一致（fp32 乘，不能在 bf16 里乘再转）
engram.engram_gate_fwd	rel ~3e-7；inplace=True 的别名语义也对
transform.apply_rotary	65536 元素里 1 个差 1.17/1.33 ULP（bf16 末位舍入），可接受
mhc 8 个	全部 ≤2e-7；expand/normw_merge/big_fuse.layer_input 逐位相同
moe.moe_topk_gate	bitwise 一致（5 组，含 4096×384 / 384 专家 / topk6）

这是整条桥的前提：量化链 bitwise 相同意味着「先跑通」不会引入 token 级偏差。

两个「看起来对但不对」的判据，都踩过：

• 相对误差不能除以逐元素预期值。comb、mixes 天然接近 0，除以它会把一个 2.4e-7 的 绝对一致报成 2.6e-3。要用「该行/该张量的最大幅值」作分母。
• mhc_reduce_partials 的 fp32 累加顺序与 kernel 不同，误差 2.3e-7；对照 fp64 真值才能 确认这是表示误差而不是逻辑错（实测 kernel 与朴素 fp32 求和的差距只有 2.09×）。

6.3 已完成与剩余

已完成：

1. 路由层 → deepinfer/tk_compat/（config/engram/mhc/moe/quant/torch/transform/utils 八个模块）， DEEPINFER_KERNELS=auto|cuda|ref，25 处 import 由 scripts/tile_kernels_patch/replace_imports.py 批量替换（幂等，--revert 可回滚）。 适配器不是简单转发：num_per_channels→block_size=(1,n)、out= 缓冲回写、 apply_rotary 拆 cache + 原地写回、engram_gate 参数重排、big_fuse 用 ref 组件组合。
2. MHC 8 个 + MoE gate 的 A/B → 全绿。
3. 无 tilelang 冒烟 → ALL_CALLS_OK（scripts/.../smoke_compat.py，在 /data/ftbuild/repo/）。
4. compat 两模式对比 → ALL_MATCH 18 用例（scripts/tile_kernels_patch/compat_ab.py）。

剩余：

1. 端到端：被 GPU 占用阻塞（见 6.4），空闲时可直接跑。
2. 性能：仍未测。
3. force_random 缺口：ref 侧没有这个参数，args.force_load_balance 打开时 tk_compat.moe 会显式报错而不是静默降级（静默会给出一套看起来正常、但路由不同的 token）。
4. 上游：3 个 tile_kernels 补丁应提到 FlagOS，否则每次重装都要重跑脚本。

（flag_gems 不是替代品：903 个 ATen 级算子替换，没有这些定制融合算子。）