diff --git a/TODO.md b/TODO.md
index 13bf2fe..be80a88 100644
--- a/TODO.md
+++ b/TODO.md
@@ -595,19 +595,45 @@ import 路径上。修完之后 `import deepinfer` 在纯 Triton 环境不再需
 把热路径逐个换成 Triton。§2.10b 说明了大头在哪（模型前向直接调 `tile_kernels` 的
 quant/MoE/MHC/engram，25 处模块级 import / 23 个 API）。
 
-### 6.0 ⚠️ 当前到哪一步了（先说清楚，别误读成「已能跑」）
+### 6.0 当前到哪一步了（2026-10-22 晚更新）
 
 | # | 事项 | 状态 |
 |---|---|---|
 | 1 | `import tile_kernels.torch`（无 tilelang） | ✅ 通 |
-| 2 | `import deepinfer`（无 tilelang） | ✅ 通（lazy init 的副作用，不等于能跑） |
-| 3 | 算子数值等价（5 个） | ✅ 4 个 bitwise + rope 1 ULP |
-| 4 | **模型前向实际调用 `tile_kernels.quant.*`** | ❌ **仍然抛 `ModuleNotFoundError: tilelang`** |
-| 5 | **端到端推理** | ❌ **从未跑过** |
-| 6 | **性能** | ❓ **未测**（跑不起来就无从测，退化是推测） |
+| 2 | `import deepinfer`（无 tilelang） | ✅ 通 |
+| 3 | 算子数值等价（全部 23 个 API 的实际调用） | ✅ 见 6.2 |
+| 4 | **模型前向调用 `tile_kernels.*`** | ✅ **已通**（`tk_compat` 路由层 + 25 处 import 替换） |
+| 5 | 无 tilelang 下 23 个 API 冒烟 | ✅ `ALL_CALLS_OK`（唯一 skip 是 `force_random`，见下） |
+| 6 | compat 两模式逐函数对比 | ✅ **`ALL_MATCH` 18 用例，12 项逐位相同** |
+| 7 | **端到端推理** | ⚠️ **被 GPU 占用阻塞，不是代码问题**（见 6.4） |
+| 8 | **性能** | ❓ 仍未测 |
 
-第 4 行是当前的断点：deepinfer 里 25 处 `import tile_kernels` 指向的还是**编译版**，
-`tk_compat` 路由层没写，所以调用一进去就要求 tilelang。**「导进去了」和「调得动」是两件事。**
+**「导进去了」和「调得动」现在都是 ✅**，且两模式在真实调用签名下逐函数一致。
+
+### 6.4 端到端为什么还没跑完（环境问题，不是代码）
+
+第 4 次端到端（`DEEPINFER_KERNELS=ref` + 屏蔽 tilelang）**通过了权重加载**（31s，80 GiB）、
+YOCO/EPLB 初始化，然后卡在 `allocate_userbuffer` 18 分钟无输出。日志里是：
+
+```
+torch.OutOfMemoryError: CUDA out of memory. Tried to allocate 540.00 MiB.
+GPU 0 has a total capacity of 139.80 GiB of which 423.62 MiB is free.
+```
+
+查 `nvidia-smi`：**GPU 1/2/4/5/6/7 各被别的容器占了 82 GiB**（`zjr-flWithTree` 等 Up 3 周），
+只有 GPU 0/3 空闲。8 卡 EP 每卡要约 80 GiB，装不下 → 部分 rank OOM 退出，其余 rank 卡在
+通信 barrier。**与 ref 实现无关。**
+
+两次环境干扰记录（下次遇到先查这两个）：
+- `/tmp`（`/dev/vda`，252G）被 `/tmp/docker-data`（217G）占满 → 容器 `docker exec` 全部失败
+  （`write /tmp/runc-process...: no space left`）。清理 `/etc/dsw-logs` 下轮转日志（同设备）
+  释放 ~2.6G 才恢复。**注意 `/data` 是 bind mount，删它不释放 `/tmp`。**
+- 丢弃 `deepinfer/` 目录重新同步会**丢掉编译产物 `_C*.so`**（它不在源码里）。
+  可从 `build/nvidia-x86_64/lib.linux-x86_64-cpython-312/deepinfer/_C*.so` 复制回来。
+
+端到端命令（GPU 空闲时可直接执行）：
+`PYTHONPATH=scripts/simulate_no_tilelang:... DEEPINFER_SIMULATE_NO_TILELANG=1
+DEEPINFER_KERNELS=ref DEEPINFER_OPS_STRICT=1 bash scripts/h20_e2e_report.sh`
 
 ### 6.1 第一步已通：`tile_kernels.torch` 能在无 TileLang 下导入
 
@@ -642,22 +668,37 @@ use_tma_aligned_col_major_sf=True`）：
 | `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 级偏差。
 
-### 6.3 还没做的（下一步，按顺序）
+两个「看起来对但不对」的判据，都踩过：
+- 相对误差不能除以**逐元素**预期值。`comb`、`mixes` 天然接近 0，除以它会把一个 2.4e-7 的
+  绝对一致报成 2.6e-3。要用「该行/该张量的最大幅值」作分母。
+- `mhc_reduce_partials` 的 fp32 累加顺序与 kernel 不同，误差 2.3e-7；对照 fp64 真值才能
+  确认这是表示误差而不是逻辑错（实测 kernel 与朴素 fp32 求和的差距只有 2.09×）。
 
-1. **路由层（唯一的硬阻塞）**：把 deepinfer 里 25 处 `import tile_kernels` 改成走
-   `deepinfer/tk_compat/`，按有无 tilelang 或 `DEEPINFER_KERNELS=ref` 选编译版还是 `_ref`。
-   最小改法：只换 import 那一行（`import tile_kernels` → `from deepinfer import tk_compat as
-   tile_kernels`），`tile_kernels.quant.xxx(...)` 的调用点一行都不用动。
-   **ref 侧覆盖率已核对**：11 个主 API 里 8 个有直接对应（另 3 个是 `config` 开关，本来就在
-   `tile_kernels.config` 里，不需要 ref）；8 个 `mhc_*_ref` 全在。唯一不明的是
-   `mhc_pre_big_fuse_fwd`（只有个 `reduce_fused` 名字像，未确认）。
-2. **MHC 8 个 + MoE gate 的数值 A/B**：ref 存在 ≠ 数值等价，还没验过，做之前不能假设。
-3. **端到端**：路由层 + 上面的 A/B 都过之后才能真正跑，然后再谈 token 一致性。
-4. **性能**：现在说「退化严重」只是推测，没测过。跑通后才有基准。
-5. **上游**：以上 3 个补丁应该提到 FlagOS，否则每次重装都要重跑脚本。
+### 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 级算子替换，没有这些定制融合算子。）