# 海光 Z100（gfx906）vLLM 适配：优化与算子开发全记录（精简版） > 汇总：2026-10-11 · 覆盖 `docs/z100-todolist.md` 的 S1–S15 + 本轮工作（复核修复、block-scaled 算子、FP8 模型上线、P2P 实验） > **引用简写**：`S11f §9.3` 即 `docs/result/z100-11f-fp8-bf16-emu-kernels-result.md §9.3`（`S` 对应 `docs/result/z100--*.md`，完整对照见第六章） > **完整版**（919 行、含 103 条问题的完整论述与逐条出处）：`git show ec19a76d9:docs/海光Z100-优化与算子开发全记录.md` ## 一页速览 **平台**：4 × 海光 Z100（`gfx906`，各 32 GB HBM）· **无 MFMA、无 fp8 指令、无 bf16 运算指令** · 宿主 31 GB RAM（关键约束）· DTK 26.04 · torch 2.11 + 自建 triton 3.6 · 容器 `z100-vllm`（bridge 网络、`RestartPolicy=no`）· 宿主内核 `iommu=pt` + `pci=disable_acs_redir=`（12 桥） 目前开源在https://github.com/sibadaaa/z100-vllm **关键成果** | 项 | 结果 | 口径 / 条件 | 出处 | |---|---|---|---| | 27B W8A8 INT8 端到端 | **TP=4 graph 49.66 ms/token = 2.44×** 生产基线（121.43）；TP=2 graph 84.22（1.44×）；eager TP=2/4 为 113.5 / 126.8 | 编译模式必须固定，否则不可比 | S14 §4 | | INT8 `sdot4` 内核 | decode **3.87–3.91×**；kernel 级 GPU busy 320.5→82.5 ms/step、304 次/step 的 int8 GEMV 290.85→52.68 ms；峰值 22.2 TFLOPS | 同命令 A/B；**默认 OFF** | S11k §3.5/§3.6 | | FP8 仿真内核（per-tensor/per-channel） | decode **1.6–5.5×** vs fp16 rocBLAS（17408×5120：0.220 ms @405 GB/s） | gfx906 上 FP8 原本全体拒绝 | S11f §3.4 | | BF16/FP16 dense GEMV | **1.3–5.7×** vs rocBLAS；换成引擎真正在用的 `triton_matmul` 基线为 **2.28–8.54×** | 误差 ≤0.39 ulp | S11f §8.1 | | block-scaled FP8（本轮） | 端到端 **16.1×**（TTFT 4969→600 ms，decode 1984.8→123.0 ms/token）；**已上线 :8000** | 官方 Qwen3.8-27B-FP8、TP=4 eager | S11f §9.4 | | prefill 快路径 | unified attention 调度 kernel −22~−28%、服务 +5%/+19%；split-KV 7–9%；FLA chunked GDN 2.1–2.3×；causal-conv1d 696→36.9 ms（**S15 因质量回归默认关，prefill −6.6%**） | 逐笔 opt-in | 提交 `9bc2855da`…`c49f91925`、S15 | | 编译 / 构建 | AOT 重载缺陷修复 54.15→19.5 s；`ninja` 二次构建 986 s → **0.096 s** | — | S11m §3、S04d §3.3 | | 卡间通信 | 设备级 P2P **28.3 GB/s 可用**；RCCL 集合通信**不走 P2P**（SHM 的 0.70 GB/s 是实测最优） | 10 组开关 + `iommu=pt` 全部无效 | S11f §10 | **一句话边界**：能依赖的只有 `v_dot4_i32_i8`（INT8）、`v_dot2_f32_f16` + LUT/f32 展开（FP8/BF16）、`v_fma_mix_f32`（需 inline asm）。 **不要再投入**：MFMA/张量核、fp8/bf16 硬件转换、skinny GEMM/LLMM1/wvSplitK、ROCm paged attention、aiter/flash-attn/hipBLASLt、custom all-reduce、RCCL 的 P2P 传输。 **质量口径（重要）**：质量/精度评测一律用 **eager**；graph 更快但**改变 token 流**且**运行间不可复现**（S14b/S14d/S14e）。INT8 动态 per-token 量化 × 64 层 GDN 递推会把 1 ULP 放大 2–5× ⇒ `sdot4` 默认 OFF、生产显式开；`--enforce-eager` 禁令的历史依据已被 S7b/S7c 推翻（根因是 fp16 packed silu 分支，已修 csrc）。 **怎么读**：只看结论 → 本页；要复现/接手 → 第二、三、五章；要排查"能不能做" → 第四章问题清单（⛔ 结构性 / 🔶 有绕法 / ✅ 已修或已裁定）。 --- ## 一、性能优化历程（S1–S15） ### 1. 平台适配前置：先让 gfx906 能编译、能起服务 | 环节 | 关键实测数字 | 决策 / 结论 | |---|---|---| | 架构归一化（S01 §3.2–§3.3） | `normalize_gcn_arch()` + `_HYGON_DCU_ARCH_ALIAS` 接入 KFD/torch/env/amdsmi 四来源；KFD 表 `90106 → gfx906`，拒绝 `gfx18446744073709551615`；真实 `rocm.py` 探针 `_GCN_ARCH=gfx906`、capability `(9,0)` | stub harness **49 passed / 0 failed / 3 skipped**（基线 29 项零回归）⇒ 前置成立 | | bf16 门控（S03 §3.3–§3.7） | `check_bf16.py` **exit=0**：GEMM rel **0.229%** / linear **0.173%** / softmax **2.5e-05**；但 4096³ bf16 比 fp16 慢（**0.36×**）；同轮确证 `tools/dcu/env.sh` 的 `ROCBLAS_TENSILE_LIBPATH` 导出即任意 GEMM 崩（A/B/A） | 保留 bf16、默认 `auto`；该导出已删除 | | gfx906 构建与兼容（S04 §3.2–§3.9 · S05 §3.2–§3.9） | `pip install -e .` **不可用**（setuptools 随机临时构建目录）⇒ 统一 `build_ext --inplace`；修 **4 类** DTK 工具链缺陷（rocPRIM placement-new / `__syncwarp` 重定义 / `atomicAdd(half)` 二义 / hipify 注入 `dtk_macros.h` 触发 `#error`）；必须先 `import torch` 再 `import vllm`，否则 `std::bad_alloc` → SIGABRT(134) | 5 个 gfx906 测试 **66 passed / 16 failed**（14 个 Triton `unsupported target`、2 个替身过期）；⛔ `wvSplitK*` 与 `_rocm_C::paged_attention` 是 **MFMA-only 桩**（真机 device assert）⇒ 短路，attention 只能走 Triton | | 目标栈升级（S04b §3.2–§3.7 · S04b-2 §3.4–§3.10） | torch `2.11.0+das.opt1.dtk2604` / triton `3.6.0+das185.dtk2604` 下四扩展全部编成且只含 gfx906；新缺陷：hipify 把 `CUDA_VERSION` 改写成 `D` | 版本判定必须用 `importlib.metadata`（DTK 轮子的 `torch.__version__` 只报基础版号） | | 构建可增量性（S04d §3.3–§3.8） | `hipify_all` 加 `add_custom_command(OUTPUT …/hipify_all.stamp)` + staging 树 + 内容同步 + `.hip` 的 `OBJECT_DEPENDS`：无改动第二遍 `ninja` **986 s → 0.096 s（`ninja: no work to do.`）**，单 TU 40–43 s、全量 963 s；四原始 `.so` 与基线逐字节相同 | ⚠️ ninja 仍按 mtime 判 clean（内容变而 mtime 倒退会产出陈旧物） | | Triton gfx906 打通 + SFU hazard（S06 §2.1–§3.6） | 自编 triton v3.6.0（补丁 5 文件 48 行：VEGA ISAFamily ×3 + `driver.py` 的 DTK `hipDeviceProp_t` + `compiler.py`）：五文件 **74 passed / 0 failed**，S5 遗留 **14/14** 个 `unsupported target` 消除；容器内探针 `TRITON_OK`/`REDUCE_OK`/`INT8DOT_OK` | ✅ 打通；⛔ 但 gfx906 的 SFU 与紧邻 VALU 有 hazard（`tl.exp2`/`tl.sqrt` 第 8–11 道读回旧值，修法 = SFU 前 `s_nop 2`），**生产验证过的 triton 3.4.0 轮子同样复现 ⇒ 既有产线上是"静默算错"**；改后端 `.py`/include 后必须清 `TRITON_CACHE_DIR` | ### 2. 内存 / KV cache / CUDA graph：从"能不能起"到 TP=4 graph 的 2.44× | 优化项 | 收益 / 结论 | 口径与条件 | |---|---|---| | 单卡 KV 自动算出 + graph vs eager（S07 §3.2–§3.4） | 0.6B fp16、`maxlen 2048`、util 0.9：`Available KV cache memory` **27.18 GiB（eager）/ 26.75 GiB（graph）**、KV size 254,464 / 250,384 tokens，**未用** `--kv-cache-memory-bytes`、未触发 0.98 兜底；decode **11.605 vs 34.636 ms/token（2.98×）** | 单卡 Z100L，`mem_utils.py`/`gpu_worker.py` 零改动；⚠️ 该 eager 臂当时数值已坏（`'lock'`×16，见 §7） | | FULL 图：崩溃 → 修复 → 归因纠正（S09 §3.10.1 · S09b §2.2–§3.4） | 27B TP=2 默认 `FULL_AND_PIECEWISE` 在 FULL 捕获期抛 `AttributeError: 'VllmConfig' object has no attribute 'pad_for_cudagraph'`；`gdn_attn.py` 两处改 `m.num_actual_tokens` 后 FULL 恢复（PIECEWISE 51/51 + FULL 35/35，`--cudagraph-metrics` 证明真在 replay），回滚演练换回旧文件即重现 `AttributeError` | ✅ **纠正"decode 3.11× 是缺 FULL 图"**：FULL **378.78 / 378.43** vs PIECEWISE **380.62** ms/token ⇒ **只快 0.5%**，conc8 14.92 vs 14.94 tok/s、TTFT in=8000 3.508 vs 3.524 全在噪声内 ⇒ 欠账在 int8/GDN 内核与 TP=4，不在图模式（27B W8A8 INT8、TP=2、`maxlen 8192`、T=0） | | TP=4 graph 为何 OOM（S14 §3.3.1–§3.3.2） | 三次 OOM（S9b/S12/S14）**没有一次是 cgroup 上限触发**：`dmesg` 原文全是 `constraint=CONSTRAINT_NONE` + **`global_oom`**，无上限容器里 `memory.current` 峰值 **29.76 GB**，连 32 GiB 上限都没撞到；真因是 **graph TP=4 引擎匿名需求 ≈34–35 GB**（4 worker 各 ~7 GB，被杀者 `anon-rss` 6.76 GiB）≥ `MemTotal 31.26 GiB` + `SwapTotal 4.0 GiB`，page cache 被榨到 ~600 MB | OOM 落在 `Available KV cache memory: 16.99 GiB` 之后 **5 秒**、**第 2 遍图捕获起步时**（不是 KV 分配失败）；§3.1 据 `memory.events oom_kill` 判 cgroup OOM 已被 §3.3.1 的内核原文更正，硬约束 26 记为"已作废" | | 显存/内存构成 + 排除"图池 / sdot4"（S14 §3.7.1–§3.8.5） | 权重在**显存**（`Model loading took 7.27 GiB`×4）；每 worker `RssAnon` **4.30–4.52 GiB（TP=4 eager）/ 6.60–6.93 GiB（TP=2 graph）**、`RssFile` 0.19 / 1.27 GiB；图池 **1.67 → 0.22 GiB**、捕获 685 → 7 s **仍 OOM 且落点不变**；`VLLM_DCU_USE_INT8_SDOT4=0` 改走 `TritonInt8ScaledMMLinearKernel` **同档 OOM**（`anon-rss` 7.28 GB，与前几次 6.76/7.09/7.21/7.28 同档） | 宿主侧 `smaps_rollup` 采样；`RssAnon` 明显小于权重分片（4.4 vs 7.19 GiB）⇒ 权重不是匿名常驻；图池与 sdot4 两条为单因子 A/B | | S14c 归因 + 两个真杠杆（S14c §3.2.1–§3.11.0） | `peak − steady` = **4 ~ 1469 MiB**（C1 1469、C2 仅 +4）⇒ 削尖峰手段天生无效；TP 对齐后 graph 比 eager 每 worker 多 **≈1.4 GiB**（+1437~+1471 MiB ×4），被 EngineCore 反向 ≈1.4 GiB 抵消 ⇒ **净缺口 ≈4 GiB**（S14 首版 8.8 GB 是高估）。真杠杆：① **AOT 磁盘重载** `init engine` **183 → 87 s**、ready 351 → 230 s（峰值 `ΣRssAnon` **+1.8%，无益**）；② **`vm.swappiness` 60 → 200** ⇒ 同样 10 GiB swap 下余量 **0.29 → 2.71 GiB**、decode 零代价 | 17 个 arm 同容器同会话对照；⛔ **七个省内存因子全否证**：`MALLOC_ARENA_MAX=1`（只把 2.4–3.1 GiB 从 unnamed 搬到 `[heap]`）、`MBT=512`、`max-num-seqs 64`、`combo_kernels=false`、`expandable_segments:False`、图池压缩、sdot4 路径 | | 达标配方 C14 vs 默认图配置（S14c §3.10–§3.12.1） | ✅ C14 = 附加 **4 GiB swapfile（总 8 GiB）** + **`swappiness=200`** + AOT 重载 + **`-cc.cudagraph_mode=PIECEWISE --cudagraph-capture-sizes "1 2 4 8"`** ⇒ ready **230 s**、decode 中位 **49.65 ms/token**、`min MemAvailable` **2.35 GiB**、swap 峰值 5.40 GiB；⛔ 默认 `FULL_AND_PIECEWISE` 51+2 桶：ready **1645 s**、`min MemAvailable` **66 MiB**、KV 16.99 vs 18.47 GiB、decode **47.97 ms（只快 4.3%）**、图池 1.64 GiB/worker | 生产 flags（`maxlen 262144` / util 0.85 / MBT 2048 / `--disable-custom-all-reduce`）；"不挂 swap"未达成（最好 1.78 GiB 且 swap 打满）；默认图配置需挂 16 GiB swap 才跑得起来 ⇒ 不可用于生产 | | **2.44× 的来历** + 最大保留意见（S14 §3.8.4 · S14c §5） | 2×2 矩阵（decode 中位 TPOT，ms/token；graph-TP4 n=10、其余 n=5）：eager TP=2 **113.51** / TP=4 **126.83**（TP=4 慢 **1.117×**）；graph TP=2 **84.22** / TP=4 **49.66**（TP=4 快 **1.70×**）；对生产基线 **121.43** ⇒ **0.409× = 快 2.44×** | `maxlen 262144 / util 0.85 / MBT 2048 / VLLM_DCU_USE_INT8_SDOT4=1`；graph-TP4 额外需 4 桶 + 16 GiB swap；⚠️ 最大保留意见：TP=4 graph 与 TP=2 的图配置不同（默认配置本机起不来），"TP=4 更强"混入的图配置贡献无法在本机剥离，且推荐配方只测了 decode，conc8 / 8K / 长稳态未复测 | ### 3. W8A8 INT8 `sdot4`：decode 3.9×，本项目最大的一笔单项收益 | 优化项 | 收益 | 口径 / 条件 | |---|---|---| | P0 剖析：380 ms 花在哪 + 真瓶颈（S11b §3.5–§3.8） | decode sync **383.24 ms/step**、GPU busy **380.50（rank0）**、1448 kernels/step：**gemm/int8 370.37 ms = 97.3%**（`scaled_mm_kernel` **346.98 ms = 91.2%**、304 call/step）、nccl 5.67、gdn+mamba **0.4%**、attention **0.1%**、**CPU 3.3%**（371.19 ms 是 `hipEventSynchronize` 纯等 GPU）。真瓶颈 = **M=1 的 64× tile padding**：`BLOCK_SIZE_M=64` ⇒ **98.4% padding 冗余**，有效算力仅 **≈70 GFLOP/s**、有效权重带宽 **≈38.8 GB/s（≈3.8% HBM）**，访存下界 **13.13 ms/token** ⇒ 实测慢 **26.4×** | 27B W8A8 INT8、TP=2、fp16、`maxlen 4096`；门禁：GEMM ≫45% ⇒ **只做 P2**，GDN 0.4% ⇒ 不做 P3，CPU 3.3% ⇒ 不做 P5；纠正立项理由：本树 triton **已发 2048–6144 条 `v_dot4_i32_i8`**（0 条 fp16/MFMA），"没吃到 dot4 指令"不成立 | | 算力审计（S11d · 《性能优化方案》§4A.1） | 可用：`v_dot4_i32_i8` / `v_dot4_u32_u8` / `v_dot2_i32_i16` / `v_dot8_i32_i4`、`v_dot2_f32_f16`、`v_pk_fma_f16`、**`v_fma_mix_f32`**；不可用：`v_dot4c_i32_i8` / `v_dot2c_f32_f16`、全部 fp8/bf8/bf16 转换、`v_mfma` / `v_wmma`；机器码复用点 `v_dot4_i32_i8` **71,442**、`v_fma_mix_f32` **12,017**、`v_dot2_f32_f16` **4,608**、`v_pk_fma_f16` **4,426**、`v_mfma` **0** | 两套工具链（clang 15/18）结论逐字节一致；⛔ **汇编器接受 ≠ 硬件支持**，决策只依赖有真机证据的指令 | | S11e：逐 shape ABBA | **63/63 组 ≥1（no shape regresses）**：M=1 **×4.93** / M=2 ×4.04 / M=4 ×2.94 / M=8 ×1.62 / M=256 ×1.02，最大 kernel **×22.61**；峰值算力 **22.2 vs Triton 18.0 TFLOPS**；**63/63 组与 `triton_scaled_mm` 逐位相同**，rel_err ≤**2.988e-04**（门禁 ≤1e-2，余量 33×） | 7 个 27B 真实 (N,K) × 9 个 M；ABBA 每点 5 轮 × 50 iter 取 min，两次独立运行偏差中位 0.2%、最大 2.6% | | S11k：修编译集成 + 端到端 3.9× | `torch.library.custom_op("vllm::w8a8_sdot4_mm")` + `register_fake`、M 分支放 opaque body ⇒ **0 graph break**（原崩因：Dynamo 给符号 token 维加 `Eq(dim,hint)` → `ConstraintViolationError`）；同命令 A/B：decode **84.07 / 82.74 vs 329.05 / 320.30 ms/token = 3.91× / 3.87×**；GPU busy **3.88×**；`gemm/int8` **4.29×**；304 次/step 的 int8 GEMV **290.85 → 52.68 ms（5.52×）**；kernels/step 1448 不变、KV 11.53 GiB 不变 | 27B W8A8 INT8、TP=2、卡 0/1、`maxlen 4096`、无 `--enforce-eager`；与 S11b 卡 2/3 的绝对差约 16%，两表不可混用；`torch.compile` 88.24 s（热缓存 52.7 s） | | S11l：收益可上线的三条证据 | ① **隔离实验洗清内核**：`opaque(triton)` vs `bare triton` 照样翻转（**8/8 prompt**，首分叉位与方向与 `opaque(sdot4)` 完全重合）⇒ 扰动来自 custom_op 引起的图/编译变化，与 sdot4 数学无关，内核累计 **294/294 逐位相同**；② 16 个分叉点全部 **≤0.53 nat 贴边**（13 个 <0.5）；③ 同口径对拍 decode **84.00 vs 321.09 ms/token（3.82×）**、**比生产基线 122.47 快 1.458×**、TTFT in3800 **1.32×**、conc8 **1.94×** | 7 文件 **129 passed / 0 failed**（`=0` 逐文件计数一致）；机制补充：动态 int8 量化把 1 ULP 扰动放大 2–5× | | S11e2 / S11f：仿真内核族 | 族总开关 `VLLM_DCU_USE_GFX906_EMU`（一次打开 INT8/FP8/dense GEMV，**默认值未改**，生产仍显式 `VLLM_DCU_USE_INT8_SDOT4=1`）；FP8 W8A8 GEMV+GEMM decode **1.59–5.46×**、dense bf16/fp16 decode **1.29–5.71×**（误差 ≤0.39 ulp）、数值 sweep **384/384**；回归 `--mode check` WORST **2.988e-04** 与 S11e 逐位一致 | 内核级已落地，**端到端未做**、block-scaled 未做、M≥8 仍落后 rocBLAS；§8.3 更正 `v_fmac_f32 = 0` 是 grep 口径假象（实为 `v_fmac_f32_e32`） | ### 4. 注意力与 Mamba / GDN：收益全在 prefill，decode 侧 attention 只占 0.1% | 优化项 | 收益 | 口径 / 条件 | |---|---|---| | unified attention prefill 调度调优 | 4×3×3×3 sweep 后统一取 `BLOCK_M 16→32`、`num_warps 2→4`、TILE 32、`num_stages 1`：kernel **−22% @4k / −25% @8k / −27% @8k+8k 第二块 / −28% @16k**；服务 prefill **197.5 → 208 tok/s（4k，+5%）**、**105.9 → 126 tok/s（16k，+19%）**；数值**逐位相同**，4k prefill 后 32 token 贪心续写逐字节相同 | Qwen3.8-27B GGUF（24Q/4KV、head_dim 256）、单序列 GQA6、fp16、800-token KV page；`num_stages>1` 在 gfx906 上回退 2–3×（64 KB LDS 上限、K/V tile 无法异步 staging）；`VLLM_GFX906_TRITON_ATTN_TUNING=0` 为紧急开关；commit `9bc2855da` | | split-KV（3D）prefill 路径 | 交错 A/B：**2–4 段比 2D 快 7–9%**（每个形状），8+ 段回退、`BLOCK_M 16` 回退约 20%；段数 2 与 4 相差 ~1%，**选 2 以减半 fp32 段工作区**；服务 **226.8 → 229.2 tok/s（4k，+1%）**、**153.9 → 159.2 tok/s（16k，+3.4%）**；数值与 2D 差 **≤1 fp16 ulp（3.1e-5）**，decode 输出逐位不变 | **opt-in，默认关**（`VLLM_GFX906_TRITON_ATTN_PREFILL_3D=1`）：端到端只有低个位数收益，而工作区在 8192-token chunk 下占 **0.4 GiB**（`util 0.95` 时峰值 30.3 → 30.7 GiB）；仅限 tuned head-256/GQA6 普通因果 attention，从不进入图捕获；commit `fbbb184d3` | | FLA chunked GDN prefill 走 Triton（+ 队列纪律） | 替换 PyTorch-eager 回退（fp32 上转 + 逐 chunk 主机循环 + 63 步 UTU 替换）：单层单次调用 **42 → 18.6 ms（T=3999）**、**175 → 84 ms（T=15998）= 2.1–2.3×**；误差：输出 max abs err **6e-5 ~ 1e-3**、state **4e-4 ~ 6e-4**，GDN 套件 20 passed；同路径 launch 后 drain stream（图捕获期跳过）端到端代价 **0（1075.7 vs 1076.2 s，64K prefill）**，换来逐层计时确定 | 六个 building block 入口共用一个 kill switch `VLLM_GFX906_GDN_TRITON_PREFILL=0`（drain 用 `VLLM_GFX906_TRITON_CAUSAL_CONV1D_SYNC=0` 恢复旧行为）；commit `13452fb7a` · `c49f91925` | | causal-conv1d prefill 走 Triton | 单层单次调用：8192-token chunk **696 → 36.9 ms**、varlen 4.1k **17.9 → 15.8 ms**、15998 token **92 → 77 ms**；**0 次 `.item()`**；max abs err **2.4e-4**（conv state 更新逐位相同） | fallback 每序列每层 4 次 `.item()`，并把 8192 token / dim=16384 / fp32 激活整块物化（**≈0.5 GiB/层**，`util 0.95` 下 OOM）；commit `575eb6a97` | | **S15：conv Triton prefill 的质量回归** | 根因**不是内核算错**（服务内 **12/12 call 逐位一致**，`replay=0.00e+00`、state 0.0）——而是该快路径"确定性地把模型推向灾难性退化吸引子"：开启时 **4/4 次启动**进入"思考循环、永不输出正文"，关闭后 **0/5 次** ⇒ 修复 = 默认 **1 → 0**；代价 prefill **424.7 → 395.8 t/s（−6.6%）**、decode 噪声内（6.71 vs 7.44 t/s） | 质量电池：工具 6/6、推理 4/6、中文短文正常（修复后）；影子容器 `z100-s14`、TP=4、`=1`、`--enable-prefix-caching --mamba-cache-mode align`；⚠️ 附带发现（未修）：eager 也存在"启动变体"、**bf16 在 27B 上输出乱码**（与 packed-decode 按 fp16 元素尺寸读 `A_log`/`dt_bias` 的隐患一致）、graph 仍不可复现 | ### 5. 通信与并行度：ACS、custom all-reduce、TP 缩放 | 优化项 | 收益 / 结论 | 口径与条件 | |---|---|---| | PCIe ACS 关闭 | 4 张 DCU 父链上 **12 个桥的 ACS `0x001d` 清到 `0x0000`**；跨卡 `copy_` **6.95 → 28.3–28.4 GB/s（4.1×）**（本卡 D2D 对照 145–282 GB/s）；`can_device_access_peer` 前后都是 12/12 YES（**不反映真实带宽**） | `tools/dcu/dcu_p2p_acs.py` + systemd `dcu-p2p-acs`（`ExecStop=revert`，幂等）；已修正 S11m 占卡期间测得的 3.12 GB/s / 8–9× 污染值 | | 为什么 decode 吃不到 | `NCCL_DEBUG=INFO` 证明 **RCCL 只用 `SHM/direct/direct`：56/56 channel 全 SHM、0 条 P2P**，连 `NCCL_P2P_DISABLE=1` 都测不出差别；**7 变体 × 5 尺寸 × 3 次取 min 两态逐档相同**；对 27B：ACS 关 vs 开 decode **84.68 vs 84.59**、prefill in3800 **8686.90 vs 8684.11**、conc8 **33.45 vs 33.23** ⇒ **全部 ≤0.7%** | 两态 fabric 状态切换后各测一次；⇒ **历史 TP≥2 基线无需因 ACS 重测**（27B W8A8 INT8、TP=2、`=1`） | | 机理量化：ACS 只在大消息上有效 | 10 KB 消息的 peer 写是**延迟受限**：**59.8 µs（ACS 关）vs 62.2 µs（ACS 开），仅差 4%**；ACS 的 4.1× 要到 **≥1 MiB** 才显现（2 MiB 3.2×、128 MiB 4.07×），而 TP=2 的 AR 消息只有 KB 量级 ⇒ 原"5.7 ms/step"**高估一个数量级**；flag 不影响正确性：同 vintage 双加载臂 **A4 vs A5 = 0/505 逐位一致、Δlogprob 0.000000** ⇒ **AR flag 与编译 flag 都不进 AOT key** | `s13_pw_emul` 模拟 ⇒ **生产配方不动**（保留 `--disable-custom-all-reduce`，既无害也无用） | | custom all-reduce：双重阻断 + 去掉 flag 收益 = 0 | ① 平台门禁 `rocm.py:1239-1241` 只认 gfx94/gfx95 ⇒ `parallel.py:1104-1109` **强制** `disable_custom_all_reduce=True`（CLI 带不带 flag 是同一套配置）；② 强闯也死：`csrc/custom_all_reduce.cu:156-159` 的 `hipDeviceMallocUncached` 在 gfx906 返回 `hipErrorOutOfMemory`（同尺寸普通 `hipMalloc` 成功 ⇒ 非内存不足）；两臂 AR 实际都选 **`PYNCCL`**；不带 flag **84.12 / 8667.97（in=3854）/ conc8 33.39** vs 带 flag **83.98 / 8650.34 / 33.51** ⇒ 全部 **≤0.4%（噪声）** | 探针 + 2-rank 真实构造异常链，`file:line` 级；⇒ 设想的 ~5.7 ms/step（6–7%）**未达成且不具备条件** | | 生产栈公平基线 + 同 fabric 真实倍率 | 无上限容器内忠实复现生产栈（vendor `0.18.1+das` + phaseA + `int8_sdot4` + `gdn_decode` + 自建 triton 3.4.0），vendor 每态 n=4：ACS 关 **121.43 [117.57, 124.44]** vs 开 **122.19 [118.99, 124.64]** ⇒ **区间完全重叠、均值差 0.6%** ⇒ S12 的"关 ACS 让生产慢 5.5%"**判为噪声并撤销**；同 fabric：本 fork TP=2 **84.68** vs 生产栈 TP=4 **132.28 ms/token** ⇒ **快 1.56×**（跨口径对 122.47 是 1.446×，与 S11l 的 1.458× 吻合） | 同一 fabric 状态、S9 协议；权威生产基线 122.47（ACS 开）落在本步 ACS 关区间内；ACS 关态 | | TP 缩放与其它口径 | Qwen3.5-4B 代理：TP=1 **46.80** / TP=2 **32.09** / TP=4 **29.93 ms/token** ⇒ TP=2→4 **仅 1.07×**（远低于 ≥1.5× 判据），prefill 7897.46 → 4659.25 → 3885.29 ms；8K prefill **2.578 ms/prefill-token**（in=8154、Median TTFT 21019.66 ms）vs 生产 2.789 ⇒ **优 1.082×（−7.6%）**；ttft8k TP=2 **20838** → TP=4 **12962 ms（1.61×）**；conc8 TP=2 33.34 → TP=4 eager 35.55 → TP=4 graph **41.13 tok/s** | 4B **不是** W8A8 INT8、不走 sdot4，只衡量框架层 TP 缩放 ⇒ **不能外推到 27B**（27B 实测见 §2）；8K 口径为引擎 `maxlen 8192`、TP=2、`=1`、MBT 2048 | ### 6. 编译与 AOT / 可复现性：AOT 修复有效，但 graph 的 token 流救不回来 | 优化项 | 收益 / 结论 | 口径与条件 | |---|---|---| | AOT 重载缺陷（S11m §3.1–§3.4，已修）+ 启动三相位 | 真因：`_verify_source_unchanged` 拿产物里 `inspect.getsource` 的快照（`linecache` 给**缺尾换行**的文件补 `\n`）比磁盘**裸读**，而 `gfx906_sdot4.py` 恰好缺尾换行 ⇒ `=1` **每次启动都判 "Source code has changed"、永远冷编译**；修法 `caching.py:608 _read_traced_file()` + `decorators.py:266-289`（缓存 key 路径不动）⇒ `torch.compile` **54.15 → 19.5 s**；`=1` uncached **1663 s** → cache-hit **1553–1573 s**（编译 95.05 → 19.6 s）、`=0` 1804 → **1633 s** | ✅ 已修且 `FORCE_AOT_LOAD=1` 通过、无静默回落；**瓶颈是两轮 graph capture（1383–1424 s ≈ 23–24 min，占 `init engine` 92–95%）⇒ AOT 的价值是"能不能活 + 启动快 + 确定性"，不是省内存**；每个开关值一个独立 `VLLM_CACHE_ROOT` 子目录 | | vintage 口径修正（S11m §3.5） | token 变体由**编译产物/cache 的 vintage** 决定，不由开关决定：同开关不同 vintage **422/512** 不同；同会话内 `=0` vs `=1` **仅 2/512**（max Δ 0.118 nat）；`=1` 热启动 3/3 逐位一致（**0/512**） | 10 次独立启动至少 6 个 token 变体，与 `=0/=1` 不对应；同一份产物"编译进程自己跑 vs 后续加载跑"也有 127/512、318/490 | | S14d：graph 可复现性**硬失败** | 生产候选配方下，预热那次（现场编译，ready 352 s、`aot_load=0`）与服务那次（同 cache 重载，ready 230 s、`aot_load=4`）对 12 prompt **12/12 不逐位一致**；同上下文 `prompt_logprobs` **mean 1.04–1.76 nat、最坏 12.67 nat**、flips 9–28、覆盖 **91.6%** 位置、首个分歧中位 ≈9 token；**7 臂 / 21 对 / 252 prompt-pair 仅 2.4% 一致**，21 对 `mean` 全落在 **1.24–1.45 nat** 窄带且与配置正交；关 `combo_kernels` 无收益（1/36）；**与 TP 正交**（TP=2 graph 同样危险）；根因不是 H-AOT | 口径：TP=4 / `maxlen 262144` / `PIECEWISE` + `1 2 4 8` / AOT 重载 / 附加 4 GiB swap / `swappiness=200`；⛔ **graph 只作延迟/吞吐口径，绝不进入任何依赖 token 流的门禁/SLA** | | S14d 纠正的另两处 | ① **"eager 需要内存余量"不成立**：挂 16 GiB swapfile **133.83** vs 不挂 **131.78 ms（1.6%）**，那 16 GiB **一页未用**、两边 available 都 6.1 GiB；② **新洞：eager 批组合效应**——conc=2 与串行在第 **6** 个 token 分叉（mean **1.82** / max **15.02** nat），而无并发的串行重发**逐位一致（\|Δ\| = 0.00000）**，该效应在两个独立引擎上五个数字完全相同 | ⇒ **评测必须固定并记录 `--max-num-seqs`/批组合**；此前所有"conc=1 与 conc=8 混用"的数值对比都可能混入批组合效应 | | S14e：边界二分 + S14g：变体钉不住（否定结论） | **G1**（默认 `FULL_AND_PIECEWISE` 51+2 桶 / TP=2 / maxlen 4096，真 AOT 重载，artifact md5 `911769244fccaac536a3cf0ff699d736` 与 S11m 逐字符相同）：启动**完全同构**（ready **1542/1542/1553 s**，`oom_kill` Δ=0）却落 **2 个离散变体**，6 跑逐位一致率 **7/15 ≈ 47%**，同变体内仍有 1e-2~6e-2 nat 抖动；**G2**（省内存图配置）3 跑落 **3 个变体**、一致率 **8/36（22%）** ⇒ 省内存配置只是**抬高分叉频率**（默认 50% → 省内存 22% → TP=4/262144 **2.4%**）。S14g ① **ASLR 否证**：`seccomp=unconfined` + `setarch -R` 把 PIE 基址钉死为 `555555554000`（两 worker 逐字节相同）仍 **0/8** 一致（`max\|Δlp\|` **0.401 nat**）、落到新变体；② **autotune 排除**（7 个 GDN `@triton.autotune` kernel 两次 run config 逐字相同、全命中磁盘缓存）；③ 14 次 run 共 6 个签名，强吸引子 Va 占 **7/14 = 50%** | 生产只剩两条路：**要 token 流可复现 ⇒ eager**；**要 2.44× ⇒ 省内存 graph 配方 + SLA 写明"重启后输出可变"**；⛔ **不能选默认图配置**（慢 9×：ready 1542 s vs 170–180 s）；残余最可能 = **allocator 残余地址布局 → Triton specialization/cache key**（需编译期取证）；工程注意：`PYTHONPATH`/sitecustomize **到不了 `VLLM::Worker_TP*`**，唯一有效通道是 site-packages `.pth` | | 其它编译开关 | `-cc.inductor_compile_config.benchmark_combo_kernel=False` 生效但**换 AOT key**，只能把冷/热差异从 8/8 降到 2/8（318/490 → 58/506）⇒ **不推荐**；`VLLM_LOGGING_LEVEL=DEBUG` 会打开 inductor asserts 从而**换 AOT key** ⇒ 带缓存的机器上不要用 DEBUG 起引擎 | `=1` 下 S9 七文件 **146 passed / 0 failed**（并集 8 文件 149） | ### 7. 质量与数值结论：两次"静默算错"的根因与一次口径裁定 | 项 | 结论 | 关键实测数字 | |---|---|---| | S7 现象 + S7b NaN 根因（S07 §3.2–§3.7 · S07b §3.2–§3.8） | gfx906 上 `--enforce-eager` 时 logits 全 NaN、输出退化成 `"lock"`（eager KV 27.18 GiB / 254,464 tok、decode 34.636 ms/token；graph 26.75 GiB、11.605）；首个非有限值在 `model.layers.0.mlp.act_fn`（SiluAndMul）**6/15360** 非有限、absmax **5320**（输入 absmax 5.922）；根因是 `_C.silu_and_mul` / `silu_and_mul_with_clamp` 的 **fp16 `__half2` packed 分支**（`csrc/libtorch_stable/activation_kernels.cu:110-159`）在 DTK 26.04 + gfx906 上被编坏，且**非确定**（同输入 20 次坏集不一致，最多 **3708/196608**）；标量分支 / `mul_and_silu` / gelu / bf16 / fp32 均正常 | Python 侧兜底后 eager == `VLLM_COMPILE` == HF 逐 token 一致，逐层 385/396 → **0**；⚠️ **覆盖缺口**：MoE/GGUF 直调点仍 **1917/98304** 非有限 ⇒ 立 S7c | | S7c：csrc 根修（S07c §3.1–§3.6） | `activation_kernels.cu` 新增 host 侧守卫 `dtk_gfx906_fp16_no_packed()`（`defined(__DTK_ARCH__)` + 运行时 `gcnArchName` 前缀 `gfx906` + `dtype==Half`，**host pass 里 `__gfx906__` 未定义**），3 处 `use_vec` 全部加守卫；非有限计数：`_C.silu_and_mul`（20 次，**424–455/次**）→ **0**、`silu_and_mul_with_clamp` **107/24576 → 0**、`apply_moe_activation(SILU)` **1907/98304 → 0**；**关掉 Python 兜底**后 eager == graph == HF，逐层 **396 模块 0 非有限** | 修的是 `_C_stable_libtorch`（md5 **`93bf7b4a391ac97e2f67cc7141633cb1`**），`vllm/_C.abi3.so` 逐字节未变；负向控制：关掉 `__DTK_ARCH__` 时 host 汇编与改动前**逐指令相同**，新增回归测试 **3 passed（旧 `.so` 上 3 failed）**；⚠️ 性质是**等价规避**（走标量分支），ISA 机理仍未定性 | | S14b：graph↔eager 口径裁定（S14b §3.10–§4.1） | 判定**以 B（可解释数值）为主**：**真值锚**（未量化 4B、TP=1、CPU HF 真值，首 token 4/4 一致）显示 **eager 差真值 0.0025 nat（max 0.0080）、graph 0.0023、compile-无图 0.0027**；**graph↔eager 同上下文 max 仅 0.054 nat（mean 0.0052）** ⇒ 编译模式本身只值 0.05 nat；27B 上的 1–4 nat 来自**动态 per-token INT8 激活量化（离散取整）× 64 层 GDN 递推**的放大 | 五个 graph 臂 vs eager：**mean 1.35–1.38 nat、max 10.3–12.0 nat、每 354 位置 ~100–108 个 argmax 翻转**，但 argmax token 自身只差 **0.20–0.33 nat**，且随上下文递增（**2.94 → 4.17 nat**）；只改 TP（eager 内 2→4）也能差 **3.29 nat** | | S14b：负向控制 + A 形状缺陷（S14b §3.9–§3.11） | `VLLM_QWEN35_PACKED_DECODE` 开/关在 eager 上**逐位相同**；**cudagraph 与数值无关**（`cudagraph_mode=NONE` ≡ PIECEWISE，首 token 判定逐个相同）；前缀缓存全程 **0.0% 命中**；无 NaN/Inf；后端/kernel 选择逐项相同 ⇒ **不是"某个算子被编坏"，也不支持改 `csrc`**；graph **运行间不可复现**：P2 是**一次坏 draw**（同配置/同 cache/同 AOT hash 重跑得 `'用户'`/gap 1.047，首 token 分布差 **5.36 nat**；重载后 3/4 prompt 逐位一致，第 4 个从第 59 位起差到 **9.5 nat**） | 唯一相关因子 = "AOT 现场编译 vs 磁盘重载" ⇒ 立 S14d | | **评测纪律（已落地）** | **质量/精度参照 = eager**（逐位可复现：1024 token、8/8 flip、跨引擎与同引擎重复全部 Δ=0.000；4B 上与真值 0.0025 nat）；**graph 的 token 流不得用于任何质量/精度评测**（性能结论不受影响）；**eager 可安全用于生产**（生产配方默认 eager）；另一条来自 S14d E3：**必须固定批组合**（`--max-num-seqs`），否则 conc=1 与 conc=2 的对比不可比（第 6 token 分叉、mean 1.82 nat） | ⬜ **未决项 S14f（未做）**：eager"批组合效应"的函数形状 + 评测口径 ⇒ 此前所有 conc=1/conc=8 混用的数值对比都可能混入该效应 | ### 8. 生产配方与部署：三套配置、两条路 | 项 | 关键数字 / 内容 | 结论 / 说明 | |---|---|---| | 配置 A：**TP=2 + graph** | decode 中位 **84.22 ms**（n=5），相对生产栈 121.43 快 **1.44×** | ⚠️ 仅可用于**延迟/吞吐**口径；无需宿主改动；**S14d 判定 graph 的"运行间不可复现"与 TP 正交 ⇒ TP=2 graph 在 token 流上同样危险** | | 配置 B：**TP=4 + eager**（= `serve_27b_tp4.sh` 默认） | decode 中位 **126.83–133.83 ms**，相对生产栈 121.43 慢 4.4% | ✅ 唯一**逐位可复现**口径；必须固定批组合（S14d E3） | | 配置 C：**TP=4 + graph** | decode 中位 **49.66 ms**（n=10），快 **2.44×** | ❌ 第二关硬失败：不得用于任何依赖 token 流一致的用途；内存配方已定（4 GiB swapfile + `swappiness=200` + AOT 重载 + 4 桶）；可复现性未通过且 S14e 已排除"换默认图配置"这条路 | | **必须显式导出的环境变量**（少一个都会出事） | `. /opt/dtk/env.sh` 与 `. /src/tools/dcu/env.sh`；`PYTHONPATH=/src:/mnt/game/triton-gfx906/install`（缺了编不出 gfx906）；`TRITON_CACHE_DIR=/mnt/game/triton-gfx906/cache`；`VLLM_WORKER_MULTIPROC_METHOD=spawn`；`TOKENIZERS_PARALLELISM=false`；`PYTORCH_HIP_ALLOC_CONF=expandable_segments:True`；`PYTHONUNBUFFERED=1` | **固定 `VLLM_CACHE_ROOT`**（否则每次都是新 vintage、输出抽签）；**`HIP_VISIBLE_DEVICES` 永远显式给（绝不要 `unset`，等于占满 4 卡）**；`VLLM_DCU_USE_INT8_SDOT4=1`（仓库默认 OFF，生产必须显式开）；`VLLM_LOGGING_LEVEL=INFO`（**不要 DEBUG**，会打开 inductor asserts 并换 AOT key） | | **命令行核心参数** | torch-first 包装 `python -c 'import torch; from vllm.entrypoints.cli.main import main; main()' serve <模型>`；`--dtype float16`（**永远不要 `auto`**）；`--tensor-parallel-size <2\|4>`；`--gpu-memory-utilization 0.85`；`--max-model-len 262144`；`--max-num-batched-tokens 2048`；`--enable-prefix-caching --mamba-cache-mode align`；`--disable-custom-all-reduce`（无害也无用，保留）；`--skip-mm-profiling`；`--enable-prompt-tokens-details`；agent 三件套 `--enable-auto-tool-choice --tool-call-parser qwen3_xml --reasoning-parser qwen3`（缺了 `tool_choice:"auto"` 请求被 400 拒）；`--override-generation-config '{"temperature":0.7,"repetition_penalty":1.05}'` 与 `--default-chat-template-kwargs '{"reasoning_effort":"low"}'`（模板默认 `xhigh`，易只出 thinking/重复循环） | **配置 C 额外加** `-cc.cudagraph_mode=PIECEWISE --cudagraph-capture-sizes "1 2 4 8"`；服务参考实现 `tools/dcu/serve_z100_tp4.sh`、基准口径参考 `tools/dcu/s14_arm_env.sh`（S14 全部实测出自它，`S14_CACHE` 必填、每次 run 独立缓存根；`S14_EAGER=1` 属硬约束 13 违例、必须过 NaN 检查） | | **内存与 vintage 纪律** | C 配置达标配方 = 附加 **4 GiB swapfile（与宿主原有 4 GiB `/swap.img` 合计 8 GiB）** + **`vm.swappiness=200`** + **AOT 重载预热** + 4 桶 `PIECEWISE` ⇒ ready **230 s**、decode 中位 **49.65 ms**、`min MemAvailable` **2.35 GiB**；**不要用默认 51+2 桶**（ready 1645 s、余量 66 MiB，只换 4.3% decode） | **"先空缓存预热再重启服务"保留，但它只保证启动快、不保证输出一致**（S14e 已否证 S11m 的"预热即可复现"）；容器**不得设 cgroup 内存上限**；监控 `Available KV cache memory`、`/proc/swaps`、`dmesg \| grep global_oom` | | **回滚路径** | 内核级 `VLLM_DCU_USE_INT8_SDOT4=0` 即回 Triton 路径（同一份树内，差距 3.82–3.91×）；扩展级 `_rocm_C.abi3.so` 可回到 `5d1e50e1…`（加 sdot4 前）；**`_C_stable_libtorch` 不要回退**（当前 `93bf7b4a…` 含 S7c 的 packed-silu 修复，回退会重现 NaN）；ACS 保持部署态关闭（12×`0x0000`，systemd `dcu-p2p-acs`） | **生产若要求 token 流可复现只能用 eager（慢 4.4%）；要用 graph 就得在 SLA 写明"同一请求跨重启/跨批组合输出不保证一致"。** | ### 9. 收益总表 | 优化项 | 收益 | 口径 / 条件 | 出处 | |---|---|---|---| | W8A8 INT8 `sdot4` 内核 | decode **3.91× / 3.87×**（84.07 / 82.74 vs 329.05 / 320.30 ms/token）；GPU busy **3.88×**；304 次/step 的 int8 GEMV **5.52×** | 27B W8A8 INT8、TP=2、`maxlen 4096`、无 `--enforce-eager` | S11k §3.5.2 · §3.6.2 | | 同上，S9 协议对拍 | decode **84.00 vs 321.09（3.82×）**，比生产基线 122.47 快 **1.458×**；TTFT in3800 **1.32×**；conc8 **1.94×**；ttft8k 2.571 ms/prefill-token | 同 vintage、同客户端 | S11l §3.4 | | 并行度 TP=4（graph 口径） | 相对 TP=2 graph **1.70×**（84.22 → 49.66）；对生产基线 **2.44×**；ttft8k 1.61×；conc8 33.34 → 41.13 tok/s | 4 桶 `PIECEWISE` + 16 GiB swap；eager 口径下 TP=4 反而慢 1.12× | S14 §3.8.4 | | CUDA graph（相对完全不编译） | 0.6B 上 graph **11.605 vs eager 34.636 ms/token（2.98×）** | ⚠️ 该 eager 臂当时数值已坏；27B 上 FULL 相对 PIECEWISE 只快 0.5% | S07 §3.4；S09b §3.4 | | AOT 磁盘重载 | `torch.compile` **54.15 → 19.5 s**；`init engine` **183 → 87 s**、ready **351 → 230 s** | 对内存无益（+1.8%） | S11m §3.3；S14c §3.6 | | gfx906 prefill 快路径（3 项） | attention 调度调优 kernel **−22% ~ −28%**、服务 **+5% / +19%**；FLA chunked GDN **2.1–2.3×**；causal-conv1d Triton **8192-token 696 → 36.9 ms**（后被质量门禁关默认） | Qwen3.8-27B GGUF / prefill 口径 | commit `9bc2855da` · `13452fb7a` · `575eb6a97`；S15 §2 | | 构建可增量性 | 无改动 `ninja` **986 s → 0.096 s**；单 TU 40–43 s | 开发效率项 | S04d §3.3 · §3.5 | | ACS 关闭 | 跨卡 `copy_` **4.1×**（6.95 → 28.3–28.4 GB/s） | **对本项目 decode 零收益**（NCCL 全走 SHM、AR 消息只有 KB 量级） | S12 §3.2 · §3.3.2；S13 §3.1.6 | --- ## 二、自研 HIP 算子：方法学与四族算子 ### 1. 方法学：从 ISA 探针到归档的五步闭环 Z100（`gfx906`）上"能不能写算子"不靠文档与经验，靠**探针**：先把 ISA 能力钉死，再按钉死的能力设计内核，最后用**机器码**证明内核真的走了那条指令。三族算子（INT8 `sdot4`、FP8 仿真、BF16/FP16 dense GEMV）与 block-scaled 扩展全部走同一条闭环。 | 步 | 做法 | 为什么需要 | 产物 / 出处 | |---|---|---|---| | ① ISA 探针 | `tools/dcu/isa_probe.sh` 输出**三列表**（汇编器接受性 / `__has_builtin` / codegen 实际发射），外加 gfx906 **基线断言** | 文档会过时、内建存在性 ≠ 硬件支持、"我觉得 clang 会发这条"是假象的唯一来源；三列表把"能写"降级成"能落码"，再降级成"真落码" | `tools/dcu/isa_probe.sh`；S11d §3.1–§3.3/§3.7；S11f §3.1 | | ② 内核设计 | 按①结论选**唯一可依赖**的 MAC 指令与数据通路，形态照抄已在本硬件验证过的骨架 | 指令没得选，能动的只有阻塞方式、LDS 布局、scale 放哪里；选错指令 = 白写（`v_mad_mix_f32`、`v_dot4c_i32_i8`、`v_dot2_f32_bf16` 在 gfx906 上都不存在） | `docs/design/z100-11f-fp8-bf16-emu-kernels.md` §3；`csrc/rocm/{w8a8_sdot4,fp8_emu_gemm,bf16_emu_gemv}.cu` | | ③ 宿主集成 | `csrc/rocm/**` → CMake `VLLM_ROCM_EXT_SRC` → `torch.ops._rocm_C` → Python 内核类 → 选核表第 0 位 + env 开关 + 回退 | 内核再快，进不了选核表就等于没有；"默认打开一个新数值路径"不可接受（会无感改输出），每个算子必须能一键回到原行为 | `CMakeLists.txt:1548`；`csrc/rocm/{ops.h,torch_bindings.cpp}`；`vllm/model_executor/kernels/linear/scaled_mm/gfx906_*.py`；`vllm/envs.py` | | ④ 门禁 | 单测 → 数值 sweep → 机器码直方图 → 开关回归 → 端到端 A/B，逐层递进 | 四类失败模式各不相同：单测抓契约、sweep 抓数值、直方图抓"是否退化成标量模拟"、开关回归抓"关掉是否真回到原行为"，只做一两层必漏 | 五类门禁见下（S11f 的 `v_fmac_f32` 假象正是机器码层单独暴露的） | | ⑤ 归档 | 每个数字带**原始命令 + 原始输出**，记录产物 `md5`、回滚件、未做项 | 算子的结论依赖具体二进制，不留 md5 与回滚件则"复现失败"无法归因；不留未做项，未验证的边界会被后来者当成已验证 | `docs/result/z100-{11e,11e2,11f,11k,11l}-*.md`，每档登记 md5 与"未做" | 三列表的结论只有三条判读（逐列细节不再展开）： - **汇编器接受 ≠ 硬件支持**：`v_qsad_pk_u16_u8`（gfx10+ 指令）在 gfx900/gfx906 上也被接受 ⇒ LLVM 谓词宽松；用 gfx900/gfx1010 交叉验证才有区分度（`v_dot4/v_dot8/fma_mix` 有，QSAD 没有）。 - **内建存在 ≠ 本目标可用**：`sdot8/udot8` 报 1；真正解释全部观测的是 feature 门槛——**gfx906 有 dot1 + dot2 + dot7，没有 dot8**。 - **codegen 只证明发射，不证明语义**：S11d 20 项 `EMITTED=20 MISSING=0`、两套工具链（宿主 DTK 24.04.3/clang 15 与容器 DTK 26.04/clang 18）三张表逐字节一致，但位运算语义正确性只能靠端到端数值兜住（见坑 1、坑 3）。 门禁五层各抓不同失败：编译纪律（`strings \| grep -oE "amdgcn-amd-amdhsa--gfx[0-9a-f]+"` 只含 `gfx906`）、单测（契约矩阵：scale 形式 / bias / 奇数形状 / 错误路径 / 门控与 `can_implement`）、数值 sweep（对 fp32 反量化参考逐 shape 而非抽样）、机器码（目标指令计数 + 禁止项 `v_mfma`/`v_cvt_f32_fp8`/`v_cvt_f32_bf16`/`v_dot2_f32_bf16` 必须为 0，**按 kernel 归因**）、开关回归（默认 OFF，关闭后回到 `TritonInt8ScaledMMLinearKernel`/`TritonFp8BlockScaledMMKernel`/rocBLAS）。归档口径本身也要复核：统一口径为 **53 条基线 MATCH**，但那是脚本硬编码标签，探针表实际条目为 32 `asm` + 7 `blt` + 13 `cgo` = 52 条；基线还会随工具链漂移（clang 18 把 `u16x2 - {0xc000,0xc000}` 规范化成 `v_pk_add_u16`，`v_pk_sub_u16` 不再发射，该行已重钉）。机器码直方图口径：`llvm-objcopy --dump-section .hip_fatbin` → 按 `__CLANG_OFFLOAD_BUNDLE__` 拆包 → `clang-offload-bundler -unbundle -targets=hipv4-amdgcn-amd-amdhsa--gfx906` → `llvm-objdump -d`（S11d §3.12；S11f §3.5）。 ### 2. 三族算子与 block-scaled 扩展 | 算子 | ISA 核心 / 契约要点 | 门禁 | 收益（实测） | |---|---|---|---| | INT8 W8A8 `sdot4` | `v_dot4_i32_i8`（`__builtin_amdgcn_sdot4`，clamp=false）int32 精确累加；权重保持 `[N,K]` int8 原样，不转置不反量化；`x_s`=`[M]`/`[1]`、`w_s`=`[N]`/`[1]`（1 元素用 stride 0）；epilogue `(acc*x_s)*w_s` → cast → `+bias`，与 `triton_scaled_mm` 同序 | int32 35/35 逐位相同；`check` 63/63（worst 2.988e-04，门禁 1e-2）；63/63 `no shape regresses`；机器码 9,255 条；选核 × 开关组合单测 | decode **3.82–3.91×**（端到端，同命令 A/B）；峰值 22.2 TFLOPS vs Triton 18.0；vs 生产基线 1.458× | | FP8 W8A8 仿真（per-tensor / per-channel） | `v_dot2_f32_f16`（2 fp16 MAC/指令 + fp32 累加）+ `ds_read_u16` 256 项 LDS LUT（512 B，对 e4m3fn/e5m2 全部有限码无损）；`w` 为 `[N,K]` K 连续、`[K,N]` 视图零拷贝；**边读边转**，不物化 fp16 副本 | 28 passed；sweep 384/384（必须显式 `--rel-tol 1e-5`，worst `max\|err\|/max\|ref\|`=4.173e-04）；LUT 全 256 码无损断言；机器码 dot2 18,666 / LUT 1,092 / `v_mfma`·`v_cvt_f32_fp8` 0 | decode **1.6–5.5×** vs fp16 rocBLAS（M=1 N=17408：0.220 ms @404.8 GB/s vs 0.826 ms = 3.75×） | | BF16/FP16 dense decode GEMV | f32 路线 `v_lshlrev_b32 16` + `v_fma/v_fmac_f32`；f16 重定标路线 `(s-112<<7)<<3` + `v_dot2_f32_f16`；一 block 一输出列 + LDS 树归约；`route` 0=auto/1=fp32/2=dot2；`K` 为 2 的倍数、`M ∈ [1,8]`、`nthreads` 为 2 的幂且 ≤ 512 | 两条路线 vs fp32 参考（≤0.39 ulp）+ 窄契约反例 + 派发回归；76 kernel 机器码归因（dot2 18,666 / `ds_read_u16` 1,092 / `v_fma_f32` 24 / `v_mfma` 0） | **1.3–5.7× vs rocBLAS**（590–625 GB/s vs 215–235）；换引擎真正在用的 `triton_matmul` 为 **2.28–8.54×** | | FP8 block-scaled 128×128 | 同上 + **scale 折叠 FMA**（scale 乘在线程局部部分和上再归约）；GEMV 每线程独占整个 K 块；GEMM 双 tile 在 128-K 边界折叠；支持 padding 行 stride | 41 passed（新增 11 项 block）；`ISA_PROBE_STRICT=1` 53/53 MATCH、退出 0；选核 ON/OFF；106 kernel：dot2 37,254 / LUT 2,184 / fmac 1,258（全在 block 内核且是 scale 折叠 FMA；per-tensor 内核 fmac 仍为 0） | 端到端 **16.1×**：TTFT 4969→600 ms、decode 1984.8→123.0 ms/token（0.50→8.1 tok/s）；同请求口径 85.0→5.38 s；与生产 INT8 126.83 ms/token 同档 | #### 2.1 INT8 W8A8 `sdot4` **背景**：ROCM 的 INT8 选核表 `_POSSIBLE_INT8_KERNELS[ROCM] = [Aiter, Triton]` 在 gfx906 上全不可用（Aiter 要求 `compute_capability >= 90`；Triton 退化成 FMA 展开，实测 2.2 TFLOPS），表里没有任何 `sdot4`（S11d §3.5），而 `v_dot4_i32_i8` 是 gfx906 上唯一有生产硬件先例（前序工程 13–14 ms/token）的 INT8 指令。 - **decode（M ≤ 8）一 block 一输出列**，64 线程切 `K/4` 个 int32，`sdot4` int32 精确累加 + LDS 树形归约——不能用 `tl.dot`：最小 tile M=16，M=1 时 15/16 是掩码行（16× 算力浪费，S11d §3.6(2)）。 - **prefill（M > 8）** 用 `BM×BN` 寄存器 tile、`BK=32`（8 int32）双缓冲 global→寄存器→LDS；`cfg` 0 自动（M≤32→32×64，否则 64×64）、1..4 显式；K 非 4 的倍数则上层回落 Triton。 - **不要** `v_dot4c_i32_i8`（gfx906 缺 dot8 特性，汇编器直接拒绝）、**不要** `v_dot4_u32_u8`（W8A8 需有符号，且树内 0 次使用无先例）；`can_implement` 只要求 `input_symmetric`。 **3.9× 由三条独立证据链支撑（别只记一个数字）**：逐 shape ABBA kernel-only 63/63 组最差 ×1.022、中位 ×2.16、最好 ×22.61（N=1024 K=5120 M=1，decode 的 k/v_proj），`full` 口径（含 `scaled_int8_quant`）最差 ×1.148；27B W8A8 TP=2 端到端 **84.07/82.74 vs 329.05/320.30 ms/token = 3.91×/3.87×**；kernel 级占比 GPU busy **82.5 vs 320.5 ms/step**、304 次/step 的 int8 GEMV **52.68 vs 290.85 ms = 3.88×/5.52×**；S11l 重验收 decode **84.00 ms/token**（3.82×）、TTFT 2.234 ms/prefill-tok（1.32×）、conc8 33.12 tok/s（1.94×）（S11e §3.5；S11k §3.5.2/§3.6；S11l §3.4）。 **`custom_op` + Dynamo 的教训（本族最贵的一课）**：内核/接入/数值/逐 shape 性能全过，但 `VLLM_DCU_USE_INT8_SDOT4=1` 时 27B 在 `torch.compile` 阶段抛 `ConstraintViolationError`（`L['positions'].size()[1]` 被标成 dynamic 却特化成 2048）、150 s 内退出，`=0` 正常。三次对照后定位：去掉 Python 的 M 分支**仍崩**，同 kernel 类换成 `_triton_fallback` 则 173 s 编译通过 ⇒ 触发者是**裸 `TORCH_LIBRARY` op 调用点**——Dynamo 在那里给符号 token 维加 `Eq(dim, hint)` 守卫并 `set_replacement(..., range_refined_to_singleton)` 成常量，而 vLLM 标记该维 dynamic；**不是 Meta 缺失、不是分支选错**。修法：`torch.library.custom_op` + `register_fake` 声明输出 `(x_q.size(0), w.size(0))`，并把 GEMV/GEMM 的 M 分支搬进 opaque body（`@torch.compiler.disable` 可用但引入 1 次 graph break；`splitting_ops` 结构性无效，崩在 Dynamo 出图后的 `build_guards`；给已注册 C++ op 补 `register_fake` 被 torch 拒绝）。验证：vs 裸内核 18/18、vs `triton_scaled_mm` 126/126 逐位相同，`dynamic=True` 下 M=2048→M=1 复用同一张图、无 graph break，副作用是 `torch.compile` 更快（88.24 s vs 159 s）、KV 显存相同（11.53 GiB）。 **遗留与默认值**：逐 token 文本一致未达成——把 opaque op 的 body 换成与基线**完全相同**的 `triton_scaled_mm` 后 **8/8 prompt 照样翻转**、5 个早期分叉点与 `opaque(sdot4)` 完全重合且方向一致 ⇒ **H_graph 成立、H_kernel 否定**，16 个分叉点全部落在 ≤0.53 nat 的贴边决策上；1 ULP 输入扰动 → int8 code 变 20–120 个（5120 里）→ 输出变 2.2–3.3e-3，比不量化的 fp16 路径大 2–5×。**默认保持 OFF**（S11k 有意推翻 S11e 的"默认开"），生产显式 `VLLM_DCU_USE_INT8_SDOT4=1`；S11e2 并入族总开关 `VLLM_DCU_USE_GFX906_EMU`，显式单项值优先。 #### 2.2 FP8 W8A8 仿真（`ds_read_u16` LUT 加宽 + `v_dot2_f32_f16`） **背景**：gfx906 **没有任何 fp8 指令**（`v_cvt_f32_fp8`/`v_cvt_pk_fp8_f32`/`v_cvt_pk_f32_fp8`/`v_cvt_f32_bf8` 全部 `not supported on this GPU`）、没有 fp8 MFMA、`torch._scaled_mm` 的 FP8 路径不支持，vLLM 里**所有** FP8 选核在 gfx906 上拒绝；唯一能跑的是纯 Triton 标量 `v_fmac_f32` 模拟，实测 **≈36× 慢**（4096³：304.2 ms vs fp16 8.4 ms，机器码 `v_fmac_f32_e32` 8160 条、`v_mfma`/`v_dot2_f32_f16`/`v_dot4_i32`/`v_pk_fma_f16`/`v_cvt_f32_fp8` 全 0，S08b）。硬约束：27B FP8 ≈ 27 GiB，fp16 副本 54 GiB 装不下 ⇒ 必须**边读边转**。 - **路线选 LUT 而非位运算**：e4m3 的次正规（`\|x\| < 2^-6 ≈ 0.0156`，per-channel scaled 权重里很常见）与 fp16 次正规指数不同，无法"移位 + 加偏置"一步映射，位运算路线必须再补逐 lane 修正；而 **256 项 fp16 LUT 对 e4m3fn 与 e5m2 的全部有限码无损**，一次 `ds_read_u16` 即得精确值（注意 `ds_read_b16` 在 gfx906 上**不存在**）。默认 `VLLM_DCU_FP8_CVT=lut`，位运算只留 A/B；LUT 由宿主侧按 dtype 构建（`functools.lru_cache` + 每设备缓存，512 B），`test_lut_is_lossless_for_every_code` 对全部 256 码逐条断言。 - **GEMV（M ≤ 8）**：一 block 一输出列、`nt=64` 线程跨 `K/4` dword 连续读 + LDS 树归约；**M ≥ 2 时把 A 的加宽提到循环外**（多一个 widen kernel + `M×K×2` 字节 scratch），因为每个 block 都要重读同一份 A，不提出去的话 M=8 时 LUT 访存是 dot2 的 **9 倍**。 - **GEMM（M > 8）**：`BM×BN` 寄存器 tile + `BK=32` 双缓冲 LDS，结构照搬 `w8a8_sdot4.cu`，只是 LDS tile 单元从 int32 字换成 fp16 对（1 dword = 2 lane = 1 条 dot2），staging 时用 LUT 把 fp8 加宽成 fp16 对。 **收益与边界**：`fp8_emu_gemm.cu` 明确"只乘一次 scale"、权重常驻 fp8 不展开；Python 侧只启用实测获胜形状（单 token 全 N；`bf16 + M=4 + 128 ≤ N ≤ 1024`），M ≥ 8 与 fp16 小 N 一律回落 rocBLAS（M=512 N=5120 为 0.84×、M=64 N=48 为 0.07×，均未启用）⇒ **没有引入任何已知回退形状**。 #### 2.3 BF16/FP16 dense decode GEMV：f32 路线 vs f16 重定标路线 **背景**：gfx906 没有任何 bf16 运算指令（转换/点积都不存在），且 decode 形状的**未量化**线性层两条快路都不可用——`ops.LLMM1`（skinny GEMM）因 MFMA 被排除出 gfx906 构建，`n <= 16` 的 Triton matmul 带 `not on_gfx9()`，只能落到 rocBLAS：`[1,5120] @ [17408,5120]` bf16 GEMV 仅 **235 GB/s**（HBM 的 24%），N=512 更差（**24 GB/s**）。 - **f32 路线（bf16 默认）**：每 lane `v_lshlrev_b32 16` 精确展宽 + `v_fma_f32`/`v_fmac_f32`，对**全部** bf16 取值精确（含次正规/零/inf），就是"bf16 权重 + fp32 累加"的硬件 MMA 数值口径。 - **f16 重定标（仅 A/B，`range_safe=False` 默认不开）**：`(s - 112<<7) << 3` 得到 fp16 位型（前提 `e8 ∈ [113,143]`）再走 dot2；fp16 指数窗口比 bf16 **窄约 30 个 binade**——`\|x\| < 2^-14` 冲零、`≥ 65536` 饱和，**零值的 `e8=0` 会静默回绕（实测零 lane → 2.0）**，σ~1e-3 的权重分布被冲掉约 **6%** 的信号（阈值是绝对的）；契约要求**两个操作数的每个 lane** 都在窗口内。fp16 输入则直接走 dot2（route 0 默认，无转换、精确）。 - **S11f §8.1 修掉一个真 bug**：`on_gfx9()` 在 gfx906 上恒 `False`，`n <= 16 and not on_gfx9()` 的 Triton 分支把本内核**全部**可用形状截走（派发不可达）；修好后换用引擎真正在用的 `triton_matmul` 作基线。 **收益**：`bf16 M=8 N=17408` 为 0.51×（内核落后 rocBLAS，**未启用**），唯一启用的 M>1 形状是 `bf16 M=4 N=512`（2.14×）；`N=512` 上 bf16 达 **5.71×**、`N=48` 为 **2.46×**（rocBLAS 小 N 极差）。修 bug 后对 `triton_matmul` 收益为 **emu/triton = 2.28–8.54×**（bf16 N=17408：0.271 ms @657 GB/s vs triton 0.908 ms @196 GB/s），全部误差 ≤ **0.39 ulp**。 #### 2.4 block-scaled（128×128）FP8 扩展（2026-10-11） **动机**：端到端选用的 ModelScope `Qwen/Qwen3.8-27B-FP8` 就是 **block-scaled** 的（`activation_scheme: dynamic`、`fmt: e4m3`、`weight_block_size: [128,128]`，**407 个 scale 全部是 `weight_scale_inv`**，没有任何 per-tensor `weight_scale`）。per-tensor 内核按设计拒绝它 ⇒ 回落到 `TritonFp8BlockScaledMMKernel`（标量 `v_fmac_f32` 模拟）：模型能起、输出连贯，但 **TTFT 4969 ms / decode 1984.8 ms/token（0.50 tok/s）**，服务不可用。 ``` out[m, n] = Σ_kb ( Σ_{k ∈ kb} a[m,k] · w[n,k] ) · a_s[m, kb] · w_s[n // block_n, kb] ``` **关键洞见**：标量乘法可以分配到线程局部部分和上再归约—— ``` Σ_t Σ_kb p[t,kb]·s[kb] == Σ_kb s[kb]·Σ_t p[t,kb] ``` ⇒ 只需在**当前 K 块的累加器**上乘 scale，就**保住了原有的"一次归约树 / 一个寄存器 tile"结构**，不必先归约再乘 scale，也不必为每个 K 块开一遍归约。 - **GEMV `fp8_emu_gemv_block_kernel`**：**每线程独占整个 K 块**（`kb = t, t+nt, …`）而不是跨线程切 dword——一个 block 只碰一个 cache line（block_k 字节、每元素 1 字节），**scale 每 128 元素只取一次**而非每 dword 一次；`acc[m] += loc[m] * (a_s[m·nkb+kb]·wsv)`。 - **GEMM `fp8_emu_gemm_block_kernel`**：两个 tile——`blk[TM][TN]`（当前 K 窗口）+ `acc[TM][TN]`（已完成块），**在 128-K 边界折叠**（`fp8_gemm_fold_block_scale` 做 `acc[i][j] += blk[i][j]*sa*w_s[(col/block_n)*nkb+kb]; blk[i][j]=0`，折叠点 `((c+1) % cpb) == 0`，`cpb = block_k / 32`），epilogue 只做 cast + bias。 - **W 行 stride（padding）**：block 路径走 `process_fp8_weight_block_strategy` → `_maybe_pad_fp8_weight`（`VLLM_ROCM_FP8_PADDING` 默认开，K=5120 满足 `stride*1 % 512 == 0`），权重变成 `stride(0) = K + 256` 的**非连续**视图；内核按 `w_stride = w.stride(0)` 索引行并要求 `stride(1) == 1` 且 `w_stride % 4 == 0`。权重方向自适应：`apply_block_scaled_mm` 按 `B.shape[1] == K` 判定 `[N,K]`/`[K,N]`，两种都**零拷贝**，`stride(1) != 1` 时**明确报错**而不是偷偷 copy。 **结果与上线**：质量 4 prompt × 32 greedy token 下 2/4 逐位一致，另 2 条只在**最后 1–3 个 token**、`Δlogprob ≤ 0.32 nat` 处分叉（与该机型噪声底 0.1–0.6 nat 同量级），4 条输出全部连贯；生产槽位起在 **:8000**（`api-gateway` upstream，服务名沿用 `qwen3.8-27b`，脚本 `/mnt/game/port/serve_fp8_8000.sh`），实测 307 tok / 39.07 s（~127 ms/token），经 gateway :8080 鉴权通过，KV 18.55 GiB / 1,192,990 tokens（S11f §9.3/§9.4/§10）。**未做（明确登记）**：graph 模式端到端（只跑 `--enforce-eager`）、UE8M0/MXFP8/非 128 的 block 形状（内核通用但只测 128×128）、MTP/视觉塔/长上下文/TP=2、block GEMM 的 tile 重调优（沿用 per-tensor 的四组 cfg）。 ### 3. 踩过的真实坑（现象 → 根因 → 处置） | # | 现象 | 根因 | 处置 | |---|---|---|---| | 1 | DTK clang packed u16 移位 lane 选择错：bf16→fp16 重定标 `(s-c)<<3` 数值错，`2^-14 → 4.0` | 用 `u16x2_t`（`ext_vector_type(2)`）写位运算时，clang 15/18 都给 packed 移位加 `op_sel_hi:[0,1]` 的 lane 选择，**低 lane 实际拿到没减偏置的值** | 改成 **32 位标量逐 lane 计算**，由 `test_bf16_fast_rebase_is_accurate_for_range_checked_weights`（含 `fast == guarded` 逐位断言）锁住；FP8 侧走 LUT 不受影响 | | 2 | `V_PK_LSHLREV_B16` 按 lane 的移位量：inline asm 给标量立即数 `3` 时只有一个 lane 被移动，`2^-14 → 64.0` | 该指令移位量**按 lane 从 src0 的两个半字**取，标量立即数只覆盖一个 lane | 同上：弃用 packed/inline-asm 路线，改 32 位标量逐 lane | | 3 | `v_fmac_f32` 的 grep 词边界假象：归档记成 **0**，实际有 **28** 条 | 反汇编助记符是 `v_fmac_f32_e32`，`\bv_fmac_f32\b` 匹配不到（`_` 是 word 字符）；用归档自留的 `emu_gfx906_final2.dis` 复算：plain grep = 28、word-boundary = 0 | 机器码门禁改为**按 kernel 归因**：28 条**全部**在 `bf16_emu_gemv_f32_kernel`（设计 §3.3 写明该走 f32 路线，设计上就该有），`fp8_emu_*` 为 0 ⇒ 结论不变但口径必须正确 | | 4 | 共享头文件不改触发重编：改 `csrc/rocm/gfx906_emu_common.cuh` 后 ninja 报 `no work to do`，新逻辑没进产物 | 构建系统没有把 `csrc/rocm/*.cuh` 当作 `.cu` 的依赖，且 `hipify_stage` 里的头文件副本不刷新 | 规避：改公共头后 **`touch` 相关 `.cu`** 再构建；属构建系统遗留，未修（登记为未做项） | | 5 | 硬链接导致构建用旧源码：`cmake --build` 报 rc=0，但产物无新符号（`nm -D \| grep fp8_emu_gemv_block` = 0） | 工作区与构建树 `/mnt/game/port/gfx906-vllm-211` 是**硬链接**关系，"写新文件 + rename"的编辑器会**断开硬链接** | 改完 `cp -f` 把被改的 `csrc/**` 同步到 `${REPO_DIR}`，再 `build-211.sh _rocm_C` + `install-211.sh _rocm_C`；注意 install 做 RPATH 后处理（md5 `54e75087…` → `9818d8a2…`） | | 6 | `_maybe_pad_fp8_weight` 造成非连续权重：内核按连续 `[N,K]` 假设读会读错行 | block 路径默认开 `VLLM_ROCM_FP8_PADDING`，权重被 pad 成 `stride(0) = K + 256` 的**非连续视图**（K=5120 满足 `stride*1 % 512 == 0`） | 内核改用 `w.stride(0)` 索引行并校验 `stride(1) == 1`、`stride(0) % 4 == 0`；`test_block_accepts_padded_weight_row_stride` 断言填充视图与连续权重**逐位一致**；per-tensor 路径遇非连续明确报错拒绝 copy | | 7 | 归档写的 `=1` 起不来，但换掉"Python 的 M 分支"仍然崩 | 崩因不在 Python 逻辑，而在**裸自定义 op 调用点**给符号维加 `Eq(dim, hint)` 守卫（见 §2.1） | 见 §2.1：`custom_op` + `register_fake` + M 分支进 opaque body，并用 30 s 最小复现器（`tools/dcu/s11k_dyn_repro.py`）把各修法逐条回归 | | 8 | `ISA_PROBE_STRICT=1` 报 2 行基线过期，且第一次"修复"直接把该行变成 `ERROR` | ① clang 18 把 `u16x2 - {0xc000}` 规范化成 `v_pk_add_u16`；② **regex 字段不能含 `\|`**（那是表格列分隔符） | 重钉到实际发射形式（`pk_u16_sub → v_pk_add_u16`、`bf16_to_f16_bits → v_pk_lshlrev_b16`）并在表头注释原因；修后 53/53 MATCH、`codegen EMITTED=29 MISSING=0 ERROR=0` | | 9 | 数值 sweep 默认门限比文档松 **100 倍**：归档称"fp32 输出 < 1e-5"，日志里 `fp32_tol=0.001` | 探针 `--rel-tol` 默认 `1e-3`，归档命令没有覆盖它 | 显式 `--rel-tol 1e-5 --abs-tol 1e-3` 重跑 ⇒ 384/384 仍通过（worst 4.173e-04）⇒ 结论成立，但**必须显式传参才算测过** | | 10 | 同一开关（`=1`）两次启动的 greedy 文本不同 | `=1` 的 AOT 产物**每次启动都加载失败→重编**（`_verify_source_unchanged` 把 `gfx906_sdot4.py` 当 traced file），叠加本树 `benchmark_combo_kernel=True` 的 timing-based 选核 ⇒ 每次编译产物可不同 | 登记为未定位缺陷；建议固定 `VLLM_CACHE_ROOT` + `-cc.inductor_compile_config.benchmark_combo_kernel=False`，并明确"**逐 token 一致不能承诺**" | 运维补充（同批记录）：FP8 生产槽位必须带 `--host 0.0.0.0`（bridge 网络靠 docker 端口映射，只绑 `127.0.0.1` 时容器内自测正常、宿主 `connection reset by peer`）；`StagedWriteTensor(uva_instead_of_gpu=True)` 按 `max_num_reqs × max_model_len × 4 B` 分配 **pinned 主机内存**，宿主 31 GB RAM + 容器 `ulimit -l` 64 MB 时 maxlen 262144 会把启动顶到临界点（`UvaBuffer.__init__` 抛 `hipErrorOutOfMemory`，重试即成功）——**排障口诀：日志里 OOM 落在 `buffer_utils.py` / `pin_memory=True` 就不是显存不足**（S11f §10.1）。 ### 4. 汇总表 | 算子 | 文件 | 指令核心 | 门禁 | 实测收益 | 状态 | |---|---|---|---|---|---| | **INT8 W8A8 `sdot4`** | `csrc/rocm/w8a8_sdot4.{cu,h}`；`scaled_mm/gfx906_sdot4.py` | `v_dot4_i32_i8`（`__builtin_amdgcn_sdot4`，clamp=false）；int32 精确累加 | int32 35/35 逐位相同；`check` 63/63（worst 2.988e-04，门禁 1e-2）；63/63 `no shape regresses`；机器码 9,255 条；选核 × 开关组合单测 | decode **3.82–3.91×**（端到端，同命令 A/B）；vs 生产基线 1.458×；峰值 22.2 TFLOPS vs Triton 18.0 | ✅ 内核/集成/重验收完成；**默认 OFF**，生产显式 `=1` | | **FP8 W8A8 仿真**（per-tensor / per-channel） | `csrc/rocm/fp8_emu_gemm.{cu,h}`；`scaled_mm/gfx906_fp8.py` | `v_dot2_f32_f16` + `ds_read_u16` 256 项 LDS LUT | 28 passed；sweep 384/384（显式 1e-5）；LUT 全 256 码无损；机器码 dot2 18,666 / LUT 1,092 / `v_mfma`·`v_cvt_f32_fp8` 0 | decode **1.6–5.5×** vs fp16 rocBLAS（17408×5120：0.220 ms @405 GB/s ≈ 等效 fp16 流量 810 GB/s） | ✅ 可选内核（默认 OFF） | | **BF16/FP16 dense decode GEMV** | `csrc/rocm/bf16_emu_gemv.{cu,h}`；`layers/utils.py` 形状表 | f32 路线 `v_lshlrev_b32 16` + `v_fma/fmac_f32`；f16 路线 `v_dot2_f32_f16` | 两条路线 vs fp32 参考（≤0.39 ulp）+ 窄契约反例 + 派发回归；机器码归因 | **1.3–5.7× vs rocBLAS**（590–625 GB/s vs 215–235）；换真实基线 `triton_matmul` 为 **2.28–8.54×** | ✅ 默认 OFF，只启用实测获胜形状 | | **FP8 block-scaled 128×128** | `fp8_emu_gemm.cu`（`fp8_emu_*_block_kernel`、`fp8_gemm_fold_block_scale`、`fp8_gemm_load_block`）；`gfx906_fp8.py` | 同上 + **scale 折叠 FMA**（在线程局部部分和上乘 scale） | 41 passed（新增 11 项 block）；ISA 53/53；选核 ON/OFF；机器码 dot2 37,254 / LUT 2,184 / fmac 1,258（全在 block 内核、scale 折叠） | 端到端 **16.1×**（TTFT 4969→600 ms；decode 1984.8→123.0 ms/token）；与生产 INT8 同档 | ✅ **已上线 :8000** | 三族算子共享同一套"探针 → 设计 → 集成 → 门禁 → 归档"闭环，也共享同一套数据通路结论：**gfx906 上 `v_mfma` 恒 0、bf16 无原生转换、fp8 无硬件转换——这三条已反复确认，不要再投入**；可依赖的只有 `v_dot4_i32_i8`（INT8）、`v_dot2_f32_f16` + LUT/f32 展开（FP8/BF16），以及 `v_fma_mix_f32`（存在、可用、树内已发射 12,017 次，但**无 clang 内建，必须走 inline asm**；`v_mad_mix_f32` 在 gfx906 上**不存在**）。 --- ## 三、2026-10-11 本轮工作：复核修复 · block-scaled · 上线 · P2P ### 1. 复核修掉的 4 个问题 | # | 问题 | 根因 | 处置与验证 | |---|---|---|---| | 1 | **bf16/fp16 dense GEMV 派发不可达**（真 bug） | `_ON_GFX9 = gfx90a/gfx942/gfx950` 是上游 CDNA2/3 语义，gfx906 不在其中 ⇒ `on_gfx9()` 恒 `False`，`n<=16` 的 Triton 分支排在 GEMV 之前，把该内核全部可用形状（n∈{1,4}）截走 | 分支提到 Triton 之前 + `on_gfx906()` 守卫；新增 2 条派发回归测试。收益：比引擎真正在用的 `triton_matmul` 快 **2.28–8.54×**（N=17408 bf16：0.271 ms @657 GB/s vs 0.908 ms @196 GB/s） | | 2 | ISA 探针 2 行基线过期 | clang 18 把 `u16x2 - 0xc000` 规范化成 `v_pk_add_u16 x, x, 2.0`（0x4000 是补码增量），`v_pk_sub_u16` 不再发射 | 重钉到实际发射形式 ⇒ `ISA_PROBE_STRICT=1` **53/53 MATCH、退出 0**（注意 regex 字段不能含 `\|`，是表格列分隔符） | | 3 | 归档 `v_fmac_f32 = 0` 是假象 | 助记符是 `v_fmac_f32_e32`，`\bv_fmac_f32\b` 匹配不到；归档自己的 `.dis` 复算 = **28**，全部在 `bf16_emu_gemv_f32_kernel`（f32 路线，设计如此） | 机器码门禁改为**按 kernel 归因**；更正归档 | | 4 | sweep 门限比文档松 100× | 探针 `--rel-tol` 默认 `1e-3`，归档命令没覆盖 | `--rel-tol 1e-5` 重跑 ⇒ **384/384**（worst 4.173e-04，全在 fp16 输出） | ### 2. block-scaled（128×128）FP8 内核 **动机**：官方 `Qwen/Qwen3.8-27B-FP8` 是 block-scaled（`weight_block_size:[128,128]`、407 个 `weight_scale_inv`），原内核按设计拒绝 ⇒ 回落 Triton 标量模拟（decode 1984.8 ms/token）。 **契约**：`out[m,n] = Σ_kb ( Σ_{k∈kb} a[m,k]·w[n,k] ) · a_s[m,kb] · w_s[n//128, kb]`。 **关键洞见**：标量乘法可分配到线程局部部分和再归约（`Σ_t Σ_kb p·s = Σ_kb s·Σ_t p`）⇒ 只在**当前 K 块**的累加器上乘 scale，原有"一次归约树 / 一个寄存器 tile"结构不变。GEMV 每线程独占整个 K 块（一个 cache line/块，scale 每 128 元素取一次）；GEMM 在 128-K 边界把 `blk` 折叠进 `acc`；**W 行 stride 必须支持**（`_maybe_pad_fp8_weight` 会 pad K，权重成非连续视图）。 **门禁**（产物 md5 `fbfd45edb00713594793151a3e9ca76a`）：单测 **41 passed**（新增 11 项 block）｜ISA **53/53**｜选核 ON→新内核、OFF→Triton｜机器码 106 kernel、dot2 **37254**、LUT **2184**、`v_mfma`/`v_cvt_f32_fp8` 全 0；`v_fmac_f32` 1258 条**全是 block 内核的 scale 折叠 FMA**，per-tensor 内核仍 0。 ### 3. 端到端上线（官方 Qwen3.8-27B-FP8，TP=4） 模型 82 文件 / **30.89 GB 与 API 清单逐项一致**。两臂同配方只差开关： | 臂 | 内核 | TTFT | decode | |---|---|---|---| | 对照 | `TritonFp8BlockScaledMMKernel` | 4969 ms | **1984.8 ms/token** | | 本步 | `Gfx906Fp8BlockScaledMMLinearKernel` | 600 ms | **123.0 ms/token（8.1 tok/s）** | ⇒ **16.1×**；与生产 **INT8** eager/TP=4 的 126.83 ms/token 同档 ⇒ 剩余成本在 eager 与 GDN/attention，不在 FP8 GEMM。**质量**：4 prompt × 32 greedy token 中 2/4 逐位一致，另 2 条只在第 29/31 个 token 分叉、`max\|Δlogprob\| ≤ 0.32 nat`（本机噪声底 0.1–0.6 nat），输出全部连贯。**上线形态**：生产槽位 drop-in 起在 `:8000`（网关 upstream），服务名 `qwen3.8-27b`，KV 18.55 GiB / 1,192,990 tokens。 ### 4. 卡间通信 P2P：四层结论 | 层 | 状态 | 关键证据 | |---|---|---| | PCIe ACS | ✅ 关键位已清 | 12 桥 `ACS_Ctl` 由 `0x001d` → **`0x0011`**（ReqRedir/CmpltRedir 清零，靠内核参数 `pci=disable_acs_redir=`，开机即生效）；`dcu-p2p-acs.service` 已禁用 | | HIP peer access | ✅ 12/12 | `torch.cuda.can_device_access_peer` 全 YES | | 跨卡 `copy_` | ✅ **28.27–28.40 GB/s** | 本卡基线 396 GB/s；非 P2P 时约 6.95 GB/s | | `hipExtMallocWithFlags(Uncached, 0x3)` | ⛔ 不可用 | **`out of memory`**（4 KB 也失败、同卡 `Default` 成功 ⇒ 缺 uncached 内存池；`Finegrained`/`UncachedXdp` 同；`hipIpcGetMemHandle`/`hipHostMalloc` 正常） | | RCCL 集合通信 | ⛔ 不用 P2P | **10 组开关 + `iommu=pt` 全部 `via P2P=0`**；`NCCL_SHM_DISABLE=1` 掉到 `via NET/Socket`（1 MiB 2247→16371 µs，慢 7.3×）⇒ SHM 是最优 | **根因**：RCCL 的 `transport/p2p.cc → ncclP2pAllocateShareableBuffer` 与 vLLM custom-AR 都依赖该 flag（`nm -D librccl.so` 实证动态链接 `hipExtMallocWithFlags@HIP`）；与 ACS/IOMMU 无关（`iommu=pt` 重启实验已证伪该假设，见 S11f §10）。**命令**：`sudo python3 tools/dcu/dcu_p2p_acs.py show`、`docker exec z100-vllm bash -lc 'PYTHONPATH=/src HIP_VISIBLE_DEVICES=0,1,2,3 python tools/dcu/p2p_verify.py'`、`bash /mnt/game/port/verify_p2p_after_reboot.sh`。 **补充（PCIe 物理口径，用于判定"到底是不是 P2P"）**：4 张卡链路均为 **Gen4 x16（16 GT/s ×16 ≈ 31.5 GB/s 可用）**。据此可判定：torch 跨卡 `copy_` 的 28.3 GB/s = 链路带宽的 **90%**，而**任何"绕主机中转"的拷贝上限只有 ~15.75 GB/s**（同一份数据要在链路上走两趟）⇒ 28.3 GB/s **必然是直接 P2P**；它随 ACS redirect 位开关（S12：6.95 → 28.3 GB/s，4.1×），而这两个位只影响对端事务，正是佐证。**反例警告**：`hipMemcpyPeer`（即便 `hipDeviceEnablePeerAccess` 返回 rc=0）在本平台只有 **6.2–6.9 GB/s**，即走中转载而非 P2P —— 所以判"P2P 是否生效"要用 torch 跨卡 `copy_`（或内核里直接访问对端指针），**不要用 `hipMemcpyPeer`**（实测见 `tools/dcu/p2p_capability_probe.py`）。 ### 5. 本轮运维坑（都已定位） | 坑 | 现象 | 根因 / 处置 | |---|---|---| | pinned 内存是启动瓶颈 | 同参数启动偶发 `load_model()` 内 `hipErrorOutOfMemory` | `all_token_ids = StagedWriteTensor((max_num_reqs, max_model_len), int32, uva=True)` 是**每 worker** `max_num_reqs × max_model_len × 4 B` 的 pinned 内存；宿主 31 GB RAM + 容器 `ulimit -l`=64 MB（硬限同值、容器内无法提升），maxlen 262144 处临界（`MemAvailable` 峰值仅 ~12.8 GB），**重试即成功**。口诀：OOM 落在 `buffer_utils.py`/`pin_memory=True` 就不是显存不足 | | 宿主访问不到服务 | 容器内 curl 正常、宿主 `connection reset by peer` | bridge + docker 端口映射 ⇒ 必须 `--host 0.0.0.0` | | 构建"成功"但没有新算子 | `ninja rc=0`，`nm -D` 找不到新符号 | `/src` 与构建树是硬链接，编辑器 rename 写入会断开链接 ⇒ 构建树仍是旧源码。处置：`cp -f` 同步 `csrc/**` → `build-211.sh` → `install-211.sh`（**install 会改 RPATH，md5 与 temp 件不同**） | | 新测试不入版本控制 | `git status` 看不到 `tests/**` 新文件 | `.gitignore:301` 的 `test_*.py` 规则；已用 `git add -f` 纳入 3 个新测试，建议后续补 `!tests/...` 白名单 | | 宿主重启会换内核 | 5.15.0-194 无 DCU 驱动模块 | `GRUB_DEFAULT="1>2"` 钉在 191（有驱动）；换内核前必须确认新内核有 `hydcu` 模块 | --- ## 四、当前问题清单（103 条，按状态分档） ### 1. ⛔ 结构性无法解决（需厂商/上游） > 取材：`docs/result/` 下 35 份执行归档 + 索引 README，逐份提取「未完成项 / 阻塞与原因」「已知限制」「未做（明确登记）」「结论中的反例」「硬约束」类条目，去重后按主题归类（不按文件堆砌）。 > 分档与口径：⛔ 结构性（需厂商/上游改代码或硬件）｜🔶 有绕法或降级（含"未做 + 原因""可用但代价明确"）｜✅ 已修 / 已裁定口径；数值全部取自归档实测，归档间互相更正的按最终口径归类；出处一律简写（`S11f §x` = S11f 归档 §8–§10），每条最多 2 个出处。 | 状态 | 条目 | 限制与影响 | 处置 | 出处 | |---|---|---|---|---| | ⛔ | 无 MFMA ⇒ 无张量核 | `v_mfma_f32_16x16x16f16` 汇编器直接回 `instruction not supported on this GPU`、`__builtin_amdgcn_mfma_*` 的 HAS=0、自编 triton 出的码里 `v_mfma`/`v_wmma` 恒为 0 ⇒ 全部矩阵乘只能走 `v_dot4_i32_i8`（int8）/ `v_dot2_f32_f16`（fp16）/ 标量 `v_fmac_f32` | 接受为既定前提：P2 改用 `v_dot4_i32_i8`、P6 用 dot2+LUT 模拟；任何"上张量核"方案在此平台不成立 | S11d §3.1；S11b §3.8(c) | | ⛔ | custom all-reduce 第一重阻断：平台门禁 | `vllm/platforms/rocm.py:1241` 的 `use_custom_allreduce()` 只认 gfx94/gfx95（MI300 系），`disable_custom_all_reduce` 被强制 True（与命令行给不给 flag 无关，日志是 debug 级）；`QuickAllReduce` 依赖同一谓词 | 生产配方固定带 `--disable-custom-all-reduce`；A1 臂（不带 flag）的配置回显逐字证明两臂 AR 实现完全一致 | S13 §3.1.4 | | ⛔ | custom all-reduce 第二重阻断：分配层不支持 | `csrc/custom_all_reduce.cu:156-159` 走 `hipExtMallocWithFlags(..., hipDeviceMallocUncached)`，该 flag 在 gfx906 返回 `hipErrorOutOfMemory`（空闲卡、1 MiB 同样失败 ⇒ flag 不支持而非 OOM）；`create_shared_buffer` 收下 `uncached` 却从不传底层，**无 Python 级绕法** | 不实施；`custom_all_reduce.py:249-273` 无 try/except 兜底 ⇒ 一旦真被选中，worker 初始化在第一个 `meta_ptrs` 分配即崩，不存在"静默回落 NCCL" | S13 §3.1.1 | | ⛔ | RCCL 根本不走 P2P | 7 变体 × 5 尺寸 × 两态扫描全部 `SHM/direct/direct`（56/56 channel，0 条 P2P）；ACS 从 `0x001d` 清到 `0x0000` 让跨卡 `copy_` 快 4.1×（6.95→28.4 GB/s），但对 all-reduce 与 TP=2 同会话 A/B ≤0.7% | ✅ 已裁定"历史基线不必重测"；只剩生产栈 TP=4 单流 decode 的 ~5.5% 差异待钉死（n=2、态内抖动 6.6%，S13 复测判为噪声） | S12 §3.2.1；S13 §3.5.2 | | ⛔ | skinny GEMM / LLMM1 / wvSplitK* 不可用 | `CMakeLists.txt` 已在 gfx906 上 `REMOVE_ITEM` `skinny_gemms*.cu`（注释原文："gfx906 无 MFMA，注册只会暴露没有可用路径的 op"），`_rocm_C` 注册面只剩 `paged_attention`；旧件上 `wvSplitK` 的 `switch(N_in)` 只实现 1..5，超出即 device assert + VM fault | 构建期摘除 + 运行期 `ROCM_SKINNY_GEMM_ENABLED=hasattr(...)` 探针 + `tools/dcu/env.sh` 兜底导出 `VLLM_ROCM_USE_SKINNY_GEMM=0` | S5 §3.6；S8 §3.4(c) | | ⛔ | `ROCM_ATTN` / `_rocm_C::paged_attention` 是 MFMA-only 桩 | `csrc/rocm/attention.cu:3293-3307` 在 gfx906 上退化成 `UNREACHABLE_CODE`，真机调用得到 device-side assertion + VM fault | `use_rocm_custom_paged_attention()` 在 gfx906 恒 False（两个 head_size 实测均 False），自动选后端落到 `TRITON_ATTN`；禁用清单第 1 行：不要显式指定 `--attention-backend` | S8 §3.3；S11a §4（F2） | | ⛔ | FP8 per-tensor / per-channel 被 arch 否决 | `ROCmFP8ScaledMMLinearKernel.is_supported()` 明确回 `False`（"wvSplitKQ (skinny GEMM) is not built on gfx906 (no MFMA)"）；而 `rocm_per_tensor_float_w8a8_scaled_mm_impl` 自身不设防（M≤4 无条件调 `wvSplitKQ`），只靠调用方被否决兜底 | 列为残余风险：将来若被直接注册成自定义 op 绕过选核，会踩 S5 已证实的 device assert；当前是"响亮失败（AttributeError）而非静默算错" | S8 §3.6.3；S8b §5.3-2 | | ⛔ | packed SiLU 的 ISA 级机理仍未定性 | 未取得 live kernel 反汇编（`hipcc -S` 用错了 TU）；现有证据只到"VALU→SFU 紧邻且缺等待槽"的可疑级别，静态代码在 `packed_compute` 里是对称的 | 当前是等价规避（关分支走已验证正确的标量实现），不是"修好 `__half2` 紧凑路径"；要闭合需从 `.so` 抽 gfx906 code object 反汇编定位调度缺陷 | S7b §5.2；S7c §5 第 3 条 | | ⛔ | graph ↔ eager 换一条 token 流（是常态不是偶发） | 19 个同上下文 prompt 里有 10 个（53%）首 token 不同，同上下文 Δlogprob 0.60–4.28 nat；TP=2 上 graph↔eager 同上下文 mean 1.38 / max 10.3 nat，主导 logit 只差 0.20–0.33 nat，且随上下文长度递增（prefix 电池 2.94→4.17 nat） | 已裁定口径：质量 / 精度 / 回归 / 审计一律用 eager，graph 只用于延迟、吞吐、内存测量；判据是同架构族未量化 4B + CPU HF 真值上编译模式本身只值 0.054 nat | S14b §3.8；S14b §7 | | ⛔ | graph 模式运行间不可复现 | 同配置、同 cache、同 AOT hash、同冷首请求的两次 TP=2 graph 运行给出 `'The'`/gap 4.750 vs `'用户'`/gap 1.047（同上下文差 5.36 nat）；S14d 在 252 个 prompt-pair 上量到 246 个（97.6%）token 流不一致，P(两次相同) ≈2.4%/prompt，前 10–20 个 token 内必定分叉 | 已排除 cudagraph、packed decode、prefix caching（hit rate 全程 0.0%）、`combo_kernels`、内存形态/swap，**根因未定位**；降级方案：eager 做质量、graph 做性能，SLA 写明"同一请求跨重启/跨批组合的输出不保证一致" | S14b §3.7；S14d §4.1 | | ⛔ | 默认图配置只有两个离散变体，每次启动随机抽一个 | 3 次逐字相同启动落在 2 个变体上（`{G1a,G1c}` 与 `{G1b,S2,S3,S4}`）；变体间 403/512 位置不同（≈1 nat），变体内只剩 1e-2–6e-2 nat 抖动；与 S11m 合并后 7/15 = 47% 的启动对逐位相同 ⇒ 近似"一次 50% 的抛硬币" | 与"省内存图配置"同一处置：不用于任何要求"同请求同输出"的场景；性能结论不受影响（graph TP=2 84.22 ms、TP=4 49.66 ms 仍成立） | S14e §3.3.1 | | ⛔ | "重启变体"钉不住（否定结论） | 14 次 run 共 6 个签名、一个强吸引子占 7/14（50%）；同配置两次 run 的 token 流不一致率 Pa/Pb 9/12、A1/A2 8/12；H2/H3/H4（recompiles 日志、profiler 直方图 diff、`PYTHONHASHSEED`/线程数/`taskset`）因预算未做 | 复盘口径：**2.44× 的 graph 收益与"输出可复现"在当前 gfx906/vLLM 组合下不可兼得**；不要承诺"重启后输出可复现"，生产保持实例常驻、避免频繁重启 | S14g §7.1；S14g §8 | | ⛔ | "编译进程执行的代码" ≠ "它存下来给后续加载的代码"（未定位） | 空缓存冷编译跑出变体 A；同一进程落盘的产物被后续 3 次启动加载（`Directly load AOT compilation from path …`、`ind_hit=1`、`ind_save=0`）后给出变体 B（三次彼此完全一致，但 B ≠ A，127/512、318/490 位置不同）；`=0` 的冷跑与其产物加载同样差 127/512 ⇒ 是这套栈的通用属性，与 sdot4 开关无关 | 已排除"缓存没命中""内核数学""采样随机性/前缀缓存"；可疑点已列（`InductorAdaptor.load` 打补丁 `AlwaysHitShapeEnv`、piecewise 子图重切分映射、`benchmark_combo_kernel` 冷编译按计时挑核）但**未定位**，建议与 S11l §5.2 合并单独立项 | S11m §3.5.3；S11m §5.2 | | ⛔ | FULL 图对 decode 几乎零收益 + 与 PIECEWISE 一致性未达成 | FULL 相对 PIECEWISE 只有 ~0.5% 收益；27B **不满足**"与 PIECEWISE 逐 token 一致"（两模式必然落到不同 torch.compile 产物，且 PIECEWISE 自身对 `--max-model-len` 已翻转 token），0.6B 满足 | **建议暂不上生产**：PIECEWISE 已是既有的安全区；要下结论必须补独立参照（HF transformers fp16 + 4 卡 device_map 逐 token 对照），未做 | S9b §3.4；S9b §5 | | ⛔ | 27B W8A8 INT8 的 TP=4 在 31 GB 宿主上不可行 | 29 GB 权重（`du -sh` 实测）+ 4 进程匿名内存过不了 31 GB；**多次复现**：S9b `oom_kill 1`、S12 两次、S14 第 1 次、S14c 第 5/6 次；失败点固定在 PIECEWISE 捕获中途 / 第 2 遍捕获起步、KV cache 分配后 6 秒 | 本机结论：TP=4 需 `--memory=0` 的容器或 **≥48 GiB RAM**（生产容器本身就是 `MemLimit=0`）；真正绕法只剩加真实 swap / 扩内存 / 换 W4A16 版 27B（本机无权重），**需用户拍板** | S14 §3.3.1；S9b §3.7 | | ⛔ | 每 worker 的 unnamed mmap（2.9–4.0 GiB）仍未定位到分配者 | 稳态每 worker `RssAnon ≈5.7 GiB`（`[heap]` 2.2–3.2 GiB + unnamed 2.9–4.0 GiB）+ EngineCore/APIServer 1.5–3.1 GiB；`[heap]` 在权重加载期就已 2.0 GiB/worker；VMA 形状只到"1 GiB + 256 MiB + ~10×42–46 MiB@64 MiB 步长 + ~265×~4 MiB/worker" | 未定位；`MALLOC_ARENA_MAX=1` 已否证"arena 浪费"；要钉死需 `LD_PRELOAD` 钩子或 gdb 抓 `malloc` 调用栈——建议只用 `malloc_info()` / `gdb -p` 只读抓栈；这是"不挂 swap 也留 2 GB"的唯一剩余方向（需削 ≈1.3 GiB/worker） | S14c §3.2.1；S14c §3.12.1 | | ⛔ | 27B 的 HF/CPU 真值参照不可得 | 权重 28.76 GiB，本机 31 GiB RAM + 27B 引擎同时在跑时放不下；装不进单卡 32 GiB（`du -sh` 29 GB）；vendor 栈无法逐 token 对照（版本/tokenizer 归一化不同） | 真值参照只能用 **Qwen3.5-4B**（同架构族、未量化、通用路径）⇒ 只能回答"通用编译模式差有多大"，**不能**回答"27B INT8 的打包路径是否算错"；27B 的独立数值参照是"FULL 数值是否可信"的唯一缺口 | S14b §5-2；S9b §5 | | ⛔ | bf16 在 27B 上输出乱码（未修） | S15 的 Q1 臂：27B 用 bf16 输出乱码；与"gfx906 packed-decode 内核按 fp16 元素尺寸读 `A_log`/`dt_bias`"的隐患一致（`gdn_linear_attn.py:367-374` 注释） | 属独立 bug，未修；与硬约束 18（一律 `--dtype float16`）相互印证 | S15 §4.4-2 | | ⛔ | 变体"吸引子"选择的底层机理未定性 | S14g 的 14 次 run 落在 6 个签名上、强吸引子 50%，是"少数几个稳定解之间的选择"而不是随机噪声；H2/H3/H4 未做，S15 又发现 eager 侧也存在同类现象 | 遗留方向：allocator → Triton specialization；若要可复现需另立专项（优先：allocator 布局固定手段 / kernel 级非确定性审计） | S14g §7.2；S15 §5 | | ⛔ | `v_dot8_i32_i4` / `v_mqsad_pk_u16_u8` 只有汇编器 + codegen 证据 | 硬约束禁止运行自编译 HIP 可执行程序 ⇒ 无真机数值验证；已有反例警示：`v_qsad_pk_u16_u8` 在 gfx900/gfx906 上都被汇编器接受（实为 gfx10+ 指令）⇒ **接受性不足为凭** | 结论里**只把 `v_dot4_i32_i8` 当作可依赖项**（前序工程已在生产硬件上验证 13–14 ms/token），`v_dot8`（INT4）留作"待真机验证"候选，不影响 P2/P6 决策 | S11d §5-2 | | ⛔ | gfx906 SFU hazard 的硬件/驱动根因未定性 | 修复是"每条 SFU 指令前多一个 `s_nop 2`"（4 字节、2 周期），本步只给出"插等待槽即消除"的经验证据，未定性是 Z100 的 SIMD16 流水特性还是 DTK 运行时行为 | 未做吞吐对比（SFU 在 attention softmax 里不是瓶颈）；若后续 DTK 升级改变 SFU 时序，该补丁需重新评估（判定脚本 `repro4.py`、`repro9.py` 一眼可判） | S6 §5.4 | | ⛔ | 未先 `import torch` 直接 `import vllm` → `std::bad_alloc` / SIGABRT | fork 的 `vllm/__init__.py` + 上游 0.21 import 链与 DTK triton 3.4.0 的交互问题；与 S4 构建系统无关 | 未修、归属待 Lead 决策；统一口径：**所有脚本/文档先 `import torch`**（生产容器实际就是这个顺序）；`gate-211.sh:2b` 照抄设计文档字面写法的 rc=134 是同一现象 | S4 §3.11；S4b-2 §5-5 | | ⛔ | RCCL 的 P2P 传输：所有开关穷尽后仍不可用 | 4 rank all-reduce 上试了 **10 组**环境组合（`NCCL_P2P_DISABLE` / `NCCL_P2P_USE_CUDA_MEMCPY` / `NCCL_P2P_LEVEL=SYS` / `NCCL_P2P_DIRECT_DISABLE` / `NCCL_IGNORE_DISABLED_P2P` / `NCCL_SHM_DISABLE` 及两两组合）：`via P2P` 恒为 0；`NCCL_SHM_DISABLE=1` 未换来 P2P 而是掉到 `via NET/Socket`（1 MiB 全归约 2247 → 16371 µs，慢 7.3×） | 接受：**不要**设 `NCCL_SHM_DISABLE=1`，TP 通信维持默认（集合通信只有 SHM 可用且已是实测最优：1 MiB 0.70–0.72 GB/s）；设备级 P2P 仍可用（跨卡 `copy_` 28.30–28.39 GB/s） | S11f §10 | | ⛔ | `hipDeviceMallocUncached` 在本平台不可用（`out of memory`，与显存余量无关） | ctypes 直探 `hipExtMallocWithFlags`：`Default(0x0) rc=0`，而 **`Uncached(0x3) rc=2 out of memory`**（4 KB 也失败、同卡 `Default` 成功 ⇒ 缺 uncached 内存池），`Finegrained(0x1)`/`UncachedXdp(0x1000)` 同为 rc=2；RCCL 的 `ncclP2pAllocateShareableBuffer` 与 vLLM custom-AR 都依赖这个 flag | 这是 **RCCL P2P 与 custom-AR 的共同根因**，且与 IOMMU 模式无关；要做只能向厂商提该内存池支持（`hipIpcGetMemHandle` 可用说明理论上存在 IPC 路径，但 RCCL 不走、自研代价过大） | S11f §10 | ### 2. 🔶 有绕法或降级方案（含"未做 + 原因"） | 状态 | 条目 | 限制与影响 | 处置 | 出处 | |---|---|---|---|---| | 🔶 | 用 ctypes 假造分配层去跑真 kernel | 用 cached 内存替代 uncached 可能让 kernel 的信号自旋永远看不到更新 ⇒ hang / 驱动风险；且收益上限已被 emulation 证明很小（ACS 的 4.1× 在 decode 消息尺寸上吃不到） | **主动放弃**；若将来要做，必须先在隔离环境 + 离线探针验证信号可见性 | S13 §5 第 2 条 | | 🔶 | bf16 设备探针不稳定 | 同容器、同卡、同脚本两次运行结论相反 ⇒ `--dtype auto` 可能在 fp16/bf16 之间摇摆 | 硬约束 18：一律显式 `--dtype float16`；S3 结论是"bf16 保留可选（`auto`），但不作默认、也不作性能手段"（bf16 4096³ 只有 fp16 的 0.36×） | S8 §3.12；S3 §4.3 | | 🔶 | `VLLM_FORCE_AOT_LOAD` / DEBUG 会换 AOT key | `VLLM_FORCE_AOT_LOAD=1` 暴露了"DEBUG 日志级别会换 AOT key"这第二个坑；env 变化同样可能换 key | 定位手段可用，但**不要**在同一套服务里换 `VLLM_CACHE_ROOT` / 删缓存 / 改 `compilation_config`（会换 key ⇒ 换变体）；不要用 `PYTHONPATH`/sitecustomize 往 worker 注入 | S11m §3.7.2；S14g §8-3 | | 🔶 | `cudagraph_mode=NONE` 对照口径被 `ConstraintViolationError` 阻塞 | `profile_run → _dummy_run` 抛 `ConstraintViolationError`（`inputs_embeds`/`positions` 被特化成常量 2048 vs `mark_dynamic`）；两任执行者、两台卡位各复现一次（`RUN_EXIT=1`，无 trace） | 判为**确定性树缺陷、不阻塞主结论**（生产形态是图开启、CPU 仅 3.3%）；修复需改 `vllm/**`，应单独立项（同族问题在新 sdot4 op 上已由 S11k 修好） | S11b §5 第 1 条；S11e §5-1 | | 🔶 | eager 也有"请求上下文敏感性"（可复现但依赖批组合） | eager 跨引擎逐位可复现（12/12、Δ=0.00000 nat），**但 conc=2 与 conc=1 不同**：第 6 个 token 分叉、同上下文 mean 1.82 / max 15.02 nat；同 prompt 串行重发则逐位一致 | 评测配方：eager + **固定批组合**，必须记录并固定 `--max-num-seqs`，否则跨评测不可比；"eager 也会变"与"graph 不可复现"要分清 | S14d §3.3.3；S14d §5-3 | | 🔶 | swap 只能桥接启动瞬时峰值，不能替代内存 | 启用 `/mnt/game/swapfile` 16 GiB 后同一 arm 可跑通（启动期换出 7.4–8.8 GB，引擎一关 swap 立刻回落到 ~0.4 GB ⇒ 被换出的确实不是每 token 都碰的热页）；worker 各被换出 223–266 MiB，换出主力是 EngineCore/APIServer（1523 MiB） | 达标配方：`TP=4 + graph(PIECEWISE + 1 2 4 8) + AOT 重载 + 附加 4 GiB swapfile（总 8 GiB）+ swappiness=200` ⇒ ready 230 s、decode 49.65 ms（基线 1.00×）、最低 `MemAvailable` 2.35 GiB；不挂附加 swap 的最好成绩只有 1.78 GiB 且 swap 100% 用满 | S14 §3.7.2；S14c §3.12.1 | | 🔶 | KV 池与 graph 池的取舍（显存侧） | graph 模式不捕获图省下约 **1.7 GiB/rank** 显存：`Available KV cache memory` 18.69 GiB（eager）vs 16.99 GiB（graph）；降 `max_num_seqs` 反而让 KV 池从 18.47 涨到 18.83 GiB（说明它主要影响显存侧 KV 池，不是宿主匿名） | 生产按口径选：要 KV 容量/可复现 ⇒ eager（代价约 2.4× 慢）；要延迟 ⇒ graph + 省内存图桶，并在 SLA 写明不可复现 | S14 §3.4.1；S14c §3.11.2 | | 🔶 | 生产容器 `MemLimit=0`，本机影子容器有上限 | 生产 `qwen38-vllm` 无 cgroup 内存上限（baseline 的 TP=4 27B 就是这么跑的），而 `z100-s9` 是 `--memory=32g`/`MemSwap=32g` | 复现生产口径必须用 `--memory=0`（或 ≥48 GiB 的机器）；放开 `z100-s9` 的上限会让**宿主整体** OOM、连累同机其它容器，被判定为不可接受的风险 | S9b §3.7；S14 §3.1 | | 🔶 | UVA pinned 内存与宿主 RAM 的关系（启动期 OOM 的隐藏来源） | `vllm/v1/worker/gpu/states.py` 的 `StagedWriteTensor((max_num_reqs, max_model_len), int32, uva_instead_of_gpu=True)` 在**每个 worker** 上分配 `max_num_reqs × max_model_len × 4 B` 的 pinned 主机内存；本机宿主只有 31 GB RAM、容器 `ulimit -l` 只有 64 MB，maxlen 262144 时该缓冲是 GB 级 ⇒ 启动处在临界点（实测一次 `UvaBuffer.__init__` 抛 `hipErrorOutOfMemory`，同参数重试即成功） | 排障口诀：日志里 OOM 落在 `buffer_utils.py` / `pin_memory=True` 就**不是显存不足**，重试或降 `--max-model-len` | S11f §10.1-2 | | 🔶 | 峰值相位在 warmup，不是"启动尖峰" | TP=4 graph 峰值 `ΣRssAnon = 26395 MiB` 只比稳态高 1469 MiB（5.9%），峰值相位是 KV cache 分配后的 `P4_warmup`；稳态本身就是 31.26 GiB RAM 的 78% | 调参方向应从"削启动尖峰"改为"削稳态常驻"；权重本体在**显存**（`Model loading took 7.27 GiB` ×4），宿主 `RssAnon` 明显小于权重分片且与分片大小不正比 | S14c §3.2.1；S14 §3.7.2 | | 🔶 | "智能下降"的根因是 conv Triton 的间接副作用，不是 conv 数学 | 独立探针 8 组：Triton vs fallback `out max ≤ 9.8e-4`（fp16 1 ulp）、`state max = 0.0`；服务内 dump 回放 **12/12 call 逐位一致（replay=0.00e+00）** ⇒ 退化来自该路径的间接副作用（布局/内核调度 → 后续 Triton 特化变体） | 处置：关掉该 conv Triton 路径后主症状消除（4/4 → 0/5，prefill −6.6%、decode 噪声内）；`causal_conv1d` Triton 的 GPU 监控环 wedge 历史未复现，禁用后风险一并消失 | S15 §4.3；S15 §5 | | 🔶 | 创作类 prompt 在采样下仍有"草稿重写"循环 | temp 0.7（rp=1.0 亦然）下约 2/4 的创作类 prompt 仍循环；旧栈同 prompt 无数据（用户要求不恢复旧服务） | 未定论：需后续用 DSH 真实 prompt 对拍判定是否模型本身属性 | S15 §5 | | 🔶 | GEMV 在 M≥8 上落后 rocBLAS | M=8 时每个 block 对同一 A 行要取 2 dword/次，A 的 LDS 访存与 dot2 条数接近 **1:1** ⇒ 内核反而落后（bf16 M=8 N=17408：0.51×；fp8 M=8 N=7168：1.39×；fp8 M=512：0.84×）⇒ 瓶颈是 **A 的 LDS 取数**，不是带宽 | Python 侧只启用实测获胜的形状（`_gfx906_gemv_shape_ok`：单 token 全部 N；`bf16 + M=4 + 128 ≤ N ≤ 1024`），M ≥ 8 与 fp16 小 N 一律回落 rocBLAS ⇒ 没有引入任何已知回退形状；下一步方向是 `uint4` 宽 LDS 读 + 提高 TM/TN | S11f §3.4；S11f §5.3 | | 🔶 | 小 N 大 M 的 GEMM tile 利用率（未做 split-K） | `M=64 N=48 K=5120` 只有 **0.07×**（因 CTA 数 = 1–2）；`M=512 N=5120` 0.84× | **未做 split-K**，明确登记为未做项；M=1 主战场另有独立优化（LDS 归约 + A 加宽外提，1.6–5.5× 于 fp16 rocBLAS） | S11f §3.4；S11f §6-3 | | 🔶 | K 循环疑似未双缓冲（推断，未证实） | decode 用的 `(64,64,256)` 变体 `shared = 32768 B`，恰好等于**单份** A+B 缓冲（32 KB）；双缓冲需 64 KB，正好顶到 gfx906 的 `sharedMemPerBlock = 65536` ⇒ 流水被 LDS 容量压掉、权重流延迟暴露 | 明确标注为**推断**，需 P2 用微基准证伪/证实；当前不作为任何结论的依据 | S11b §3.8(d)-2 | | 🔶 | block-scaled 之外的 FP8 变体未覆盖 | **UE8M0 / MXFP8 / 非 128 的 block 形状未测**（内核通用：`block_n`/`block_k` 运行期参数、要求 `block_k % 32 == 0`，但只测了 128×128）；block GEMM 的 tile 沿用 per-tensor 的四组 cfg，未针对 block 路径重新 tune | 已实现并验证的是 **128×128 block-scaled**（端到端 Qwen3.8-27B-FP8 / TP=4：TTFT 4969→600 ms、decode **1984.8 → 123.0 ms/token = 16.1×**）；其余变体登记为未做 | S11f §9.3；S11f §9.5-2 | | 🔶 | `supportsDirectToLdsLoadBitWidth` 未给 VEGA 开 32 位 | nlzy 的 gfx906 分支有开，本步保守未采用 | 明确登记：若 S8 出现 direct-to-lds 相关失败可从此处入手；当前无实测需求 | S6 §5.4 | | 🔶 | MTP / GDN 等 fp16-only 派发未扩展 bf16 | `VLLM_DISPATCH_HALF_TYPES` 未扩展 bf16；bf16 在 27B 上另有乱码 bug（见 ⛔ 组） | 登记为未做；与硬约束 18（生产一律 fp16）一致，暂不阻塞 | S11f §6-5 | | 🔶 | INT8 `sdot4` 族的覆盖缺口（三项） | ① **非对称激活量化未实现**（kernel 无 azp 修正项，`can_implement` 仍要求 `input_symmetric=True`）；② **静态输入 scale 的实际覆盖率未端到端验证**（本机只有对称动态量化的 27B W8A8）；③ **INT8 的 MoE/GGUF 变体未做**（GGUF 的 `__dp4a` 路径在 `csrc/quantization/gguf` 中独立存在） | 命中不支持配置时回落 `TritonInt8ScaledMMLinearKernel`（安全但慢）；若后续拿到静态输入量化权重，按设计文档 §失败处置 检查 | S11e2 §5-2；S11e2 §5-4 | | 🔶 | 没有 Z100 的 triton FP8 调优配置 | 配置目录下没有 `*device_name=Z100,dtype=fp8_w8a8,block_shape=[128,128].json` ⇒ 每次打 `Using default W8A8 Block FP8 kernel config. Performance might be sub-optimal!` | 对已被否决性能价值的路径不值得补配置，记录备查（S11f 上线自研内核后该回退路的定位进一步下降） | S8b §5.3-3 | | 🔶 | ninja 的 mtime 语义边界未在构建系统内闭合 | touch 源码（内容不变）会让 stamp 语义与"内容相同即 no-op"冲突；要闭合就得每次构建跑内容哈希守卫，这会让 `ninja` 不再打印 `ninja: no work to do.`（与验收直接冲突） | 选择"保留 stamp 语义 + 文档化 + 规避命令"（本步即用它恢复自洽，复验 `cmp IDENTICAL`）；未做 CUDA 构建验证（本机无 CUDA 工具链） | S4d §3.8；S4d §5-5 | | 🔶 | 共享头文件的构建依赖缺失 | 改共享头文件 `csrc/rocm/gfx906_emu_common.cuh` **不会**触发重编（ninja 报 `no work to do`，且 `hipify_stage` 里的头文件副本不刷新）⇒ 会出现"构建 rc=0 但产物没有新行为"的假成功 | **未修**（需要 `touch` 相关 `.cu` 再构建）；判为 S4d 增量性修复留下的依赖缺口，已实测记录 | S11f §3.2；S11f §6-6 | | 🔶 | 镜像里的 `dccobjdump` 是断链 | `/opt/dtk/bin/dccobjdump` 与 `/opt/dtk-26.04-DCC2602-0317/bin/dccobjdump` 都指向不存在的 target ⇒ 设计文档的门禁命令 `dccobjdump --offloading …` **不可执行** | 已用"枚举内嵌 target id `amdgcn-amd-amdhsa--gfx`"替代（并用 vendor 多架构轮子验证判据有区分力）；建议门禁表改成 `strings` 回退写法或修镜像软链 | S4 §5.2 | | 🔶 | `requirements/dcu*.txt` 固定的轮子版本 ≠ 镜像实际安装版本 | 文件按设计文档固定 torch 2.11.0+das / triton 3.6.0+das（实测可下载、PEP 508 解析通过），镜像自带却是 torch 2.10.0+das / triton 3.4.0 | 有意不装（DTK 与社区轮子不可混用，擅自升级风险更大）；是否升级由 Lead 决定；后续实际路线是"只装两枚目标栈轮子 `--no-deps`" | S4 §5.5；S4b §3.1 | | 🔶 | 镜像依赖整链比 fork 的 requirements 旧（多处缺口） | `tests/conftest.py` 无法导入：`tblib` 缺失、`mistral_common 1.10.0 < 1.11.5`、`openai 2.21.0`（代码要 3.x 的 `NamespaceTool`）、`xgrammar 0.1.33 < 0.2.1`；`compressed-tensors` 需 ≥0.15.0.1（镜像自带 0.14.0.1 不够，缺它则**任何量化模型都加载不了**） | 全部只在**影子容器内**补装（纯 Python，未触碰 DTK/torch），生产容器未受影响；`compressed-tensors` 只进了 requirements 文本、镜像未重建 ⇒ 每次起容器仍需补装 | S5 §5.4；S8b §5.3-5 | | 🔶 | `torchvision` ABI 与 torch 2.11 不匹配 | `torchvision 0.25.0+das.opt1.dtk2604.torch2100`（为 torch 2.10 构建）⇒ `import torchvision` 抛 `RuntimeError: operator torchvision::nms does not exist`；而 `import vllm.model_executor.kernels.linear` 链会经 `_aiter_ops → transformers → torchvision` 触发它 | 凡会 import 到该链的探针/门禁必须换到含 `torchvision 0.26.0` 的容器（`z100-s9`/`z100-eager`）跑；`tests/model_executor/test_qwen3_5_gdn.py` 在 `z100-build-211` 因此**采集期即失败**，与 S6/triton 无关 | S6 §5.3；S11e §5-3 | | 🔶 | `.gitignore` 把 gfx906 交付物吃掉 | `.gitignore:301` 的 `test_*.py` 让新增回归测试在 `git status` 里**不可见**（文件确实落盘、pytest 直接指定路径能跑到）；`*.sh` 规则吃掉 `tools/dcu/*.sh`（含 `s8b_gate.sh`），整个 `tools/dcu/` 一度未跟踪 | S9b 已加白名单（`!tools/dcu/*.sh`、`!tests/kernels/test_gfx906_packed_silu_fp16.py` 等，并保留 `__pycache__` 忽略）；未加白名单前需 `git add -f` | S7c §5-7；S9 §5.1 | | 🔶 | `VLLM_CACHE_ROOT` 不显式设置会把编译缓存写进容器层 | 本步实测容器层被写入 **2.1 GB**（同盘还会随容器层膨胀） | 硬约束 17：显式 `VLLM_CACHE_ROOT=/mnt/game/...`；同理 `TRITON_CACHE_DIR` 指向 `/mnt/game/triton-gfx906/cache` | S8 §5.4-5 | | 🔶 | 两轮 CUDA graph capture ⇒ 单次启动 26–39 min | 27B 在 `=1` 下启动是**两轮** capture（第一轮 ~683–703 s 捕 51 段 PIECEWISE + 一轮 FULL、打印 `Graph capturing finished`、随后分配 KV cache；**然后再来一轮** 51 段 ~17–18 min）；实测 ready 时间 1583 / 1542 / 1542 / 1553 s | 既有行为、非某步引入；编排脚本等待上限必须 ≥2700 s（1500 s 会把第 2 轮 38/51 处的引擎误杀）；S14e 记录默认图配置单格 ≈25.7 min | S11l §5.1；S14c §3.11.0 | | 🔶 | 若干测试替身/期望值过期（非本步引入） | ① `test_w4a16_kernel_selection.py:32-40` 不知道 `Gfx906GPTQWNA16LinearKernel` 的优先级、在 gfx906 上**结构性不可过**；② `test_gdn_metadata_builder.py` 5 个 CPU 用例用旧 API 与旧 stub；③ S5 两个 `test_qwen3_5_gdn.py` 替身缺 `seq_lens` / `ignore_unexpected_prefixes` | ①已在 S8b 修好（2 failed → 3 passed）；②③仍挂在 S5/S9b 名下未修（写作用域外）；判为"过期测试替身"类 | S8b §3.3；S9b §5 | | 🔶 | 单卡跑 27B W8A8 不可行（环境性阻塞） | 单卡权重 28.51 GiB / 卡 31.98 GiB，`torch.compile` 阶段停摆（`--memory=10g` 时 `Model loading took 28.51 GiB` 后挂死，只能 `kill -9`） | 口径：27B W8A8 类模型用 **TP≥2**（生产口径）；S8 的门禁改在 0.6B 上做、27B 交 S9（S9 用 TP=2 通过） | S8 §3.9；S9 §3.10 | | 🔶 | `hy-smi` 在本宿主不可用于定量 | 读 `mkfd` 失败、只报 `VRAM% = 0%`（无 MiB 明细）；`HCU%` 瞬时采样不可靠（同窗口 100% 与 0% 交替）；0/1 与 2/3 卡 `PwrCap` 不同（280 vs 250 W） | 设计文档门禁 4（4 卡显存偏差 <5%）因此**未执行**，改用 `torch.cuda.mem_get_info` 的空闲对照（4 卡均 31.73 GiB）；hy-smi 仅用于"引擎在哪张卡"的定性判定 | S7 §5.3；S11b §5-3 | | 🔶 | "非 DTK 行为不变"存在两处按 arch 生效的门控（F1/F2） | **F1**：`_rocm_supported_dtypes()` 的 bf16 分支不按 `is_hygon_dtk()` 限定 ⇒ 社区 gfx906 上 `supported_dtypes` 会从 `[f16, f32]` 变成 `[f16, bf16, f32]`；**F2**：`use_rocm_custom_paged_attention()` 的 `on_gfx906() → False` 同样不按 DTK 限定（但 HEAD 的 True 在社区 gfx906 上同样是错的，会撞 MFMA-only 桩） | F1 建议修（一行 `and is_hygon_dtk()`，本步未落地）；F2 建议保留并只在文档/禁用清单注明"按 arch 而非按 DTK 生效"；两者都需 Lead 表态 | S11a §4；S11a §5-2 | | 🔶 | "真社区 ROCm 机上 0 差异"仍是推断 | 本步是在 DTK 镜像内隐藏判据的**模拟**：amdsmi 由 DTK 镜像提供并报 Hygon 的 `gfx916`、`/opt/dtk` 仍在磁盘上、`rocminfo_candidates()` 仍会列出 `/opt/dtk/bin/rocminfo` ⇒ 模拟 ≠ 真社区机 | 已按"触发条件"逐条给出差异归属（D1–D5 全部有解释、无未解释差异）；`tests/kernels/` 的 A7 抽样因 torch 版本错配（镜像 2.10 vs 扩展编于 2.11 的 `undefined symbol`）**未跑**，是欠账 | S11a §5-1；S11a §5-3 | | 🔶 | `/mnt/game` 在某些会话不可写 ⇒ 证据落 `/tmp` | 子 agent 的文件沙箱为 workspace-write，`touch /mnt/game/...` 被拒（`Permission denied` / `[sandbox: file access denied]`），canonical 树里 15 个文件为 `root:600` | 已有绕法：经容器内 root（`docker exec`）写 `/mnt/game`，或把证据原文内嵌进归档；对结论无影响，但要求"产物落 `/mnt/game`"时必须由主 agent 或容器代落 | S11d §5-1；S11k §5.1 | | 🔶 | 外部参考（1Cat-vLLM）的数字未在本机复核 | "首次启动 271 s → 重启 95 s"的原始出处**未找到**（已全量 grep `271`/`95`，命中均为无关数值）；其文档里多数原始 artifact 在私有路径、不在仓库 | 明确"不编造、标记为未找到"；所有引用数字均为**引用他们文档的原文**，非本机实测；未把任何 1Cat 代码拷入本仓库 | S11g §5 | | 🔶 | 宿主重启会换到无驱动的内核（既存陷阱） | `GRUB_DEFAULT="1>2"` 钉在 5.15.0-191，而**已安装的 5.15.0-194 没有任何 `hydcu` 驱动模块**（`/lib/modules/5.15.0-194-generic` 下 0 个）⇒ 若默认项被改动/升序规则变化，重启即**整机无 GPU** | 保持 `GRUB_DEFAULT="1>2"`（本轮未改）；换内核前必须先确认新内核有 DCU 驱动模块 | 本轮实测（`ls /lib/modules/*/`、`dkms status` 为空） | | 🔶 | pinned 主机内存才是 FP8 服务启动瓶颈（公式化，与 🔶 组 UVA pinned 条同源） | 每 worker `max_num_reqs × max_model_len × 4 B` 的 pinned 内存；本机宿主仅 31 GB RAM、容器 `ulimit -l` 64 MB 且**硬限制同为 64、容器内无法提升**，maxlen 262144 时处临界（实测一次 `UvaBuffer.__init__` 抛 `hipErrorOutOfMemory`，同参数重试即成功） | 排障口诀：OOM 落在 `buffer_utils.py` / `pin_memory=True` 就不是显存不足 ⇒ 重试或降 `--max-model-len`；要根治需在**创建容器时**给 `--ulimit memlock=-1` 或宿主加内存 | S11f §10.1 | | 🔶 | block-scaled FP8 的端到端只覆盖 eager | 上线臂用 `--enforce-eager`（生产配方），**graph/compile 模式未跑**；其余未测：UE8M0/MXFP8/非 128 的 block 形状、MTP/视觉塔/长上下文/TP=2、block GEMM 的 tile 重调优 | 登记为未做；graph 模式在本平台另有"改变 token 流"的既有裁定，因此质量口径仍建议 eager | S11f §9.5 | ### 3. ✅ 已修 / 已裁定口径 | 状态 | 条目 | 限制与影响 | 处置 | 出处 | |---|---|---|---|---| | ✅ | hipBLASLt 在 DTK 里没有 gfx906 目标码 | `/opt/dtk/lib/hipblaslt/library/` 只有 `*gfx928*`；运行期 `rocblaslt warning: No paths matched ...*gfx906*co` 后自动退回 rocBLAS Tensile ⇒ **不需要任何额外开关** | 唯一会主动选 hipBLASLt 的 aiter hipBMM 路径已被 aiter 关闭挡住；注意**不要**导出 `ROCBLAS_TENSILE_LIBPATH`（导出即任意 GEMM 崩） | S8 §3.4(b) | | ✅ | aiter 全族不可用 | `is_aiter_found_and_supported()` 要求 `get_cdna_version()>2`；真机 6 个 `is_*_enabled()` 全为 `None`，`import aiter` 失败 | 保持全部子开关默认关闭，不为提速打开；禁用清单第 3 行 | S8 §3.2 | | ✅ | flash-attn 不可用 | `fa_utils.py:22-27` 对 gfx906 早退并打 warning；`flash_attn_triton_available()` / `is_flash_attn_varlen_func_available()` 均 `False` | 不用；不要设 `FLASH_ATTENTION_TRITON_AMD_ENABLE=TRUE`（禁用清单第 4 行） | S8 §3.4(d) | | ✅ | FP8 block-scaled 无 arch 门控（静默走 36× 慢路） | `TritonFp8BlockScaledMMKernel.is_supported()` 在 gfx906 返回 `(True, None)`，用户拿 block-FP8 权重能静默进入纯 `v_fmac_f32` 标量模拟（4096³ **304 ms vs fp16 8.4 ms ≈36×**） | **已裁定**：数值干净（17 组 `mean_rel_err ≤1.465e-07`）⇒ **不加门控**，改为 `is_supported()` 打一次性告警 + 禁用清单写明"性能上禁用（Z100 上不要用 FP8 做加速）"；最小门控补丁已写好但未应用 | S8b §3.2.5；S9 §2.1 | | ✅ | `--enforce-eager` 禁令的立论与被推翻 | 立论（S7）：`CompilationMode.NONE` 下 logits 全 NaN、输出退化成 `"lock"×N`，6 个 prompt 全崩 ⇒ 写成硬约束 13；推翻链：S7b 定位根因 = `_C.silu_and_mul` 的 fp16 `__half2` packed 分支被 DTK 编译坏 → S7c 在 csrc 根修 + 回归单测 → S14 的 P4e（TP=4 + eager）**实测非 NaN、非 `"lock"`、连贯中文** | **禁令的事实基础已被移除**（违反登记在案、经 Lead 明确授权）；现口径：eager 可用，且是唯一可作质量/精度参照的臂 | S7b §3.2；S14 §3.4.1 | | ✅ | S7b 的修复只覆盖 `CustomOp` 路径 | `fused_moe` 等直调点仍会踩雷（`apply_moe_activation(SILU)` 仍有 1917/98304 非有限），影响面含 fused_moe/gguf/minimax 多处 | 由 S7c 在 csrc 侧关掉 gfx906 的 fp16 向量化分支（`use_vec` 守卫）覆盖**全部调用者**：MoE 直调点 1907→0 非有限，裸算子 20 次 nonfinite=0；新单测在旧 `.so` 上必失败作负向控制 | S7b §3.8；S7c §3.3 | | ✅ | "冷启动预热一次 ⇒ 此后逐 token 可复现"这条配方被更正 | S11m 给出的部署流程断言"预热写入 AOT 产物后每次重启逐 token 可复现"；S14e 用**同一份产物、同一 md5** 的两次重启实测 **403/512 不同** ⇒ 该断言不成立 | 配方降级为"**只保证启动更快**（ready 352→230 s、compilation 117→25 s、对数值无影响），**不保证重启不换输出**"；生产选型时按此写入部署说明 | S14e §4.5；S14d §4.3③ | | ✅ | AOT 重载每次判"源码变了"的真因与修复 | 产物里的源码快照**带尾换行**、磁盘文件没有 ⇒ 逐字节不一致 ⇒ `_verify_source_unchanged` 每次报 `Source code has changed … Recompiling`；`=1` 因此每次启动重编（白花 ~32 s，叠加 timing autotune 变成"同开关两次启动文本都不同"） | **已逐字节定位并修复**（54.15 → 19.5 s；`=1` 热启动 3/3 逐位一致）；`=0` 的加载路径同样可复现（Z2 == Z3，0/512） | S11m §3.1；S11m §3.5.6 | | ✅ | token 变体由**产物/cache vintage** 决定，不由开关决定 | 同开关、不同 vintage 产物：**422/512** 不同（Z2 vs arm2，都是 `=0`）；不同开关、同 vintage：**2/512** 不同 ⇒ S11l 那句"`=0` vs `=1` 的 428/512 翻转"无法与"产物 vintage 不同"区分开 | **口径修正**：S11l 真正站得住的是**同进程内**的结论（内核逐位相同 294/294）；对生产的意义是"只要缓存固定，输出就稳定" | S11m §3.5.6；S11m §5.2 | | ✅ | `VllmConfig.pad_for_cudagraph` 缺失 | FULL cudagraph + GDN/mamba 混合模型必崩：`AttributeError: 'VllmConfig' object has no attribute 'pad_for_cudagraph'`（`gdn_attn.py:351,400`）；S9 只能绕道 PIECEWISE | S9b 已修并**确证语义**：旧方法 = `bs_to_padded_graph_size[bs]`，该映射对捕获桶是恒等映射且 padding 已由 model runner 提前完成 ⇒ 新旧返回值**逐位相同**，因此**不**在 gfx906/GDN 路径硬编码降级；"GDN 生产继续用 PIECEWISE"只作部署建议 | S9 §5.2-1；S9b §2.2 | | ✅ | 新自定义 op 的 `torch.compile` 集成崩 | 裸 `_rocm_C.w8a8_sdot4_gemm` 调用点给符号维加 `Eq(dim, hint)` ⇒ 该维被 `set_replacement(..., range_refined_to_singleton)` 变常量 ⇒ `produce_guards_verbose()` 抛 `ConstraintViolationError`（150 s 内退出，端到端被否决） | S11k 用 `torch.library.custom_op` + `register_fake`（opaque body）修好：18/18 与裸内核逐位相同、`dynamic=True` 无 graph break、63/63 M 档逐位一致；端到端 **decode 3.9×**（84.07/82.74 vs 329.05/320.30 ms/token） | S11e §5-1；S11k §3.5.2 | | ✅ | `combo_kernels` / `benchmark_combo_kernel` 关闭：无收益且更慢 | 关掉后峰值 `ΣRssAnon` 27430 vs 对照 26395（+1035 MiB，略超噪声底）、decode **+1.5%**、尾延迟明显变差（P90 56.81 / P99 57.14 vs 50.32 / 50.41），且它**会换 AOT key**（必须新 cache） | **不建议关闭**；但它作为"冷/热差异的降级手段"能减小不能消除差异（8/8 → 2/8 prompt；318/490 → 58/506），**不推荐进生产** | S14c §3.11.1；S11m §5.3 | | ✅ | "已排除 cgroup 上限"是误判（口径更正） | S12 用 `docker update --memory-swap=-1` 后断言"32 GiB ≈ 宿主 31 GiB + 3 GiB swap，本来就不是瓶颈"；真因是 **cgroup v2 把页缓存计入 `memory.current`**——29 GB 的 mmap 权重页 + 4 个 worker 匿名内存必然撞上 32 GiB，`memory.events` 的 `oom_kill 2` 是直接证据；`--memory=0` 还被 docker 当成"不修改" | 已在 S14 §3.1 逐条更正并以 `docker inspect` / `memory.max` / `memory.events` 原文坐实；判据：`--memory` 与 `--memory-swap` 同值时 **swap 也被限死** | S12 §3.4.1；S14 §3.1 | | ✅ | 图捕获池不是 graph-TP=4 OOM 的主因（假设被否证） | 把捕获桶从 51 桶（1.67 GiB / 685 s）砍到 `[1,2,4,8]`（0.22 GiB / 7 s，**7.6×**），**死亡点、死亡原因、`anon-rss` 全部不变**，仍在第 2 遍捕获起步 | ⇒ 调捕获桶救不了，`cudagraph-capture-sizes` 之外继续降 footprint 的调参方向**已失效**（Lead 要求立即停、不再调参）；graph 模式 TP=4 的宿主匿名峰值是编译/切分执行路径本身的常驻内存（graph 每 worker 6.6 GiB vs eager 4.4 GiB） | S14 §3.8.1；S14 §5-4 | | ✅ | 七个"压缩宿主内存"的因子全部失败 | 对照 C1 逐条实测：`MALLOC_ARENA_MAX=1`（仍 OOM，桶间搬家 2.4–3.1 GiB 但总量不减 ⇒ 否证"arena 浪费"）、AOT 重载（时间 −121 s，内存无收益）、MBT 2048→512（零/负收益）、`max-num-seqs 64`（+639 MiB、decode 差 4.2%）、`combo_kernels=false`（+1035 MiB、decode 差 1.5%）、`PYTORCH_HIP_ALLOC_CONF`（中性）、默认图配置 51+2 桶（可行但 ready +1294 s、余量仅 66 MiB ⇒ 不值） | **真正的两个杠杆**：AOT 重载（把换出需求从 8.8 GiB 降到 ~3 GiB）与 `vm.swappiness=200`（余量 +1.06 ~ +2.65 GiB、decode 零代价）；**两者都不削需求，只改变"需求如何在 RAM 与 swap 之间分配"** | S14c §3.11.0；S14c §3.12.1 | | ✅ | INT8 动态 per-token 量化 × GDN 递推的放大效应 | 一个 fp16 ULP（4.9e-04 相对）的输入扰动会让 int8 code 变 20–120 个（5120 个里），使该层输出变化 ~2.2–3.3e-3 相对，**比"同样扰动、不量化"的 fp16 路径大 2–5×**，且是离散跳变；96 层残差流累积足以在 <1 nat 的近似并列点翻转贪心 | **已裁定为机制 B（可解释数值差，不是"某算子被编坏"）**：判据是主导 logit 只差 0.2–0.3 nat、差集中在竞争结构、无 NaN/Inf、后端与 kernel 逐项相同、packed kernel 与 cudagraph 都做了负向控制 | S14b §4.1；S11l §3.2.2 | | ✅ | 编译模式本身在通用路径上只值 0.05 nat | 未量化 4B（TP=1）+ CPU HF 真值：graph vs eager 同上下文 max 0.054 nat（mean 0.0052），两臂各自与真值差 0.0023–0.0025 nat、首 token 4/4 一致；而 27B 上是 1–4 nat（尾部 10 nat） | ⇒ "换编译模式 ⇒ 换 token 流"的**幅度在 27B 上来自 INT8 动态量化 + 64 层 GDN 递推的放大**，不是编译模式把通用算子算错；评测口径据此固定为 eager | S14b §3.10；S14b §7 | | ✅ | greedy token 会被 sdot4 改变 ⇒ 仓库默认关、生产显式开 | 8/8 prompt 在 ≤0.53 nat 的贴边决策上翻转；两条续写都通顺、语义等价，但**输出确实不同**；且 `=1` 的 AOT 产物修复前每次加载失败重编 + timing autotune ⇒ 同开关两次启动文本都不一样 | **默认值保持 `False`（OFF）**：默认路径应当最保守、可复现（`=0` 跨引擎逐位可复现）；生产**显式** `VLLM_DCU_USE_INT8_SDOT4=1`（收益：decode 3.82×、TTFT 1.32×、conc8 1.94×）；发布说明必须提示"开启后 greedy 输出与 Triton 路径不同" | S11l §6.1；S11e2 §5-1 | | ✅ | S14 的"噪声底 0.5 nat"被严重低估（口径更正） | 同配置跨引擎重启本身就有 0.1–0.6 nat 噪声（P4g2 vs P4g3 首个分歧在 token 2、间距 0.078/0.562 nat）；S14 用它给 G2 判"NOT A NEAR-TIE (≥0.5 nat)"，其阈值低于真实噪声底 | 已更正：S14 §3.6.3 的结论"TP=2 下 graph↔eager 是系统性模式差、graph 模式改变输出质量"**必须改成**"graph 路径运行间不可复现，它给出的 token 流是一次抽签；eager 逐位可复现" | S14b §3.1；S14 §3.8.3 | | ✅ | eager 的逐位可复现（及它的边界） | eager 跨引擎 12/12 逐位一致、Δ=0.00000 nat；同 prompt 串行重发 0.00000 nat；两个独立引擎的 conc2 五个数字完全相同 | ✅ 已作为"唯一可作质量/精度参照的臂"；但**边界**是"同请求序列 + 同批组合"（见 🔶 组"eager 请求上下文敏感性"条） | S14b §3.4；S14d §3.3.3 | | ✅ | M=1 decode 的一级浪费：64× tile padding 冗余 | 启发式对 M=1 一律取 `BLOCK_SIZE_M = 64`，`masks_am` 把第 1..63 行全 mask 掉 ⇒ **63/64 = 98.4% 的 tile 计算是 padding 冗余**；`scaled_mm_kernel` 独占 346.98 ms/step（304 次/step），实测只吃到约 3.8% 的 HBM 带宽、比访存下界慢 **26.4×** | 已由 P2（sdot4 GEMV）取代，端到端 decode 从 329/320 降到 84/82 ms/token；顺带**更正方案假设**：本树自编 triton 3.6 已把 int8 `tl.dot` 正确降到 `v_dot4_i32_i8`（0 条 fp16/fp32 dot、0 条 MFMA/WMMA）⇒ "缺指令"不是慢的原因 | S11b §3.8(a)；S11k §3.5.2 | | ✅ | (per-tensor) FP8 从"完全不可用"变为"有可选 HIP 内核" | 支持 e4m3fn/e5m2、per-tensor/per-channel/per-token、fp16/bf16/fp32 输出；decode（M=1）比 fp16 rocBLAS 快 1.6–5.5×（17408×5120：0.220 ms @ 405 GB/s，等效 fp16 流量 ~810 GB/s，接近 HBM 上限） | 数值门禁：384/384 用例通过（`fp32_tol=1e-5` 收紧后 worst 相对误差 4.173e-04）；**全部开关默认 OFF**，关闭即回到原行为；S8b 那条"36× 慢"的告警仍适用于未开此内核时的 Triton 回退路 | S11f §4；S11f §8.4 | | ✅ | bf16 快路径默认关闭（窄契约） | `bf16x2_to_f16x2_fast`（f16 重定标）要求逐 lane 落在窗口内，**零值会回绕**；DTK clang 的 packed 16 位路径还有两个真实陷阱（`op_sel_hi` lane 选择让 2^-14 → 4.0；inline asm 立即数只移动一个 lane，2^-14 → 64.0） | 最终用 **32 位标量逐 lane** 实现快路径，并用"`fast == guarded` 逐位断言"的单测锁住；该路径**不在 Python 侧开放**（登记为未做） | S11f §5.1；S11f §6-4 | | ✅ | bf16/fp16 dense GEMV 派发不可达（真 bug，已修） | `_ON_GFX9` 只含 `gfx90a/gfx942/gfx950`（上游 CDNA 语义），gfx906 上 `on_gfx9()` 恒 False，而 `n <= 16 and not on_gfx9()` 的 Triton 分支排在 GEMV 分支**之前** ⇒ 把该内核支持的全部形状都截走（原注释"这个目标是 gfx9 所以 Triton 快路被跳过"是反的） | 已修：GEMV 分支提到 Triton 分支之前 + 加 `on_gfx906()` 守卫；修复后 bf16/fp16 decode 比引擎真正在用的 `triton_matmul` 快 **2.3–8.5×**（rocBLAS 不是生产口径，只作参照） | S11f §8.1 | | ✅ | 机器码门禁 / 门限 / 探针基线的三处口径假象已更正 | ① `\bv_fmac_f32\b` 匹配不到 `v_fmac_f32_e32`（`_` 是 word 字符）⇒ 归档原文记 0、实为 28 条（**全部**在 `bf16_emu_gemv_f32_kernel`，设计如此）⇒ 机器码门禁**必须按 kernel 归因**；② 数值 sweep 的 `--rel-tol` 默认 `1e-3` 比文档写的 `1e-5` 松 100 倍，原"fp32 输出用 1e-5 门限"当时**并未被验证** | 已用归档留下的反汇编复算并显式收紧门限重跑（384/384 通过，worst 4.173e-04）；ISA 探针 2 行过期基线也一并修正（钉 `v_pk_add_u16`/`v_pk_lshlrev_b16` 并注明原因） | S11f §8.2；S11f §8.4 | | ✅ | DTK 26.04 rocPRIM placement-new 缺陷（阻塞 `_C` 编译） | clang 18 + rocPRIM：`selective_scan_fwd.hip`（经 `selective_scan_hip.h:233` 的 hipcub `BlockLoad`）报 `operator new` 错误，`_C` 连第一步都过不去 | S5 用 shim（`-include csrc/dtk_rocprim_compat.h`）+ 逐文件修复（`__syncwarp`/`atomicAdd`/hipify 注入的 `dtk_macros.h`）解决；configure 输出含 `DTK rocPRIM placement-new workaround enabled` | S4 §3.8；S5 §3.1 | | ✅ | `_C_stable_libtorch` 被 torch 版本卡住 | `csrc/libtorch_stable/cuda_view.cu` 用 `torch::stable::Tensor::layout()`，镜像的 torch **2.10.0** 里不存在，而 `CMakeLists.txt:1282` 按 torch 2.11 的 ABI 编译（`TORCH_TARGET_VERSION=0x020B…`） | 由 S4b 升到 torch 2.11.0+das 后编成（两个 `.so` 都只含 gfx906、都能 import）；属"上游代码 + torch 版本"问题，非 csrc/DTK 工具链缺陷 | S5 §5.1；S4b §3.7 | | ✅ | torch 2.11 的 `CUDA_MACROS_MAP` 裸字符串值（DTK 工具链缺陷 #5） | `cmake/hipify.py` 消费该映射时崩；影响面是 `csrc/libtorch_stable/**` 里所有用 `CUDA_VERSION` 的 `.cu` 全部无法编译，`_C`/`_rocm_C` 不受影响 | S4b 在构建树内临时打了 `isinstance(_v, str)` 守卫补丁，**S4c 已正式落进仓库**（`cmake/hipify.py`）并撤掉 wrapper；全量 ninja rc=0（17 分 32 秒），四个 `.so` 与 S4b 基线逐字节同源 | S4b §5.1；S4b-2 §3.4 | | ✅ | S5 漏修的孪生文件 `fused_qknorm_rope_kernel.cu` | `csrc/libtorch_stable/fused_qknorm_rope_kernel.cu:43` 缺少与 `csrc/fused_qknorm_rope_kernel.cu:45-46` 一致的 `HIP_VERSION` 守卫（与 torch 版本、与 gfx906 都无关） | S4c 已落进仓库；此后再无 `_C_stable_libtorch` 编译失败 | S4b §5.1-② | | ✅ | triton 3.6 轮子的 Python 与 `libtriton.so` ABI 不一致 | `amd.passes.ttgpuir.add_convert_to_buffer_ops` 是 Python 5 参 / 原生 4 参 ⇒ 直接崩 | 绕法：`AMDGCN_USE_BUFFER_OPS=0`（S6 处理 triton 前必须先知道的前置）；S6 最终改为自编带 gfx906 补丁的 triton 并前置到 `PYTHONPATH`（不 pip install 到系统、不碰 DTK 自带 triton） | S4b §5.2；S6 §3.3 | | ✅ | 构建增量性：`hipify_all` stamp 从不创建 ⇒ 每次 ninja 全量重编 | 同一棵树连跑两次 `ninja`：改动前 16 min 26 s（986 s）、改动后 `ninja: no work to do.`（<0.2 s）；根因是 hipify 步骤缺 stamp、且"复制源"语义导致 `.hip` 每次都更新 mtime | S4d 在 `cmake/utils.cmake` 加 stamp + 修正依赖；**产物四个原始 `.so` 与 S4c 逐字节相同**；单 TU 改动只重编该 TU（~40 s）；⚠️ 写作用域偏差（规则实际在 `cmake/utils.cmake` 而非任务书写的 `CMakeLists.txt`）已上报 Lead | S4d §3.1；S4d §5-1 | | ✅ | 工作区与构建树是硬链接 ⇒ 编辑器"写新文件 + rename"会静默断开 | 构建树仍是旧内容 ⇒ `cmake --build` 报 rc=0 但产物里没有新符号（实测 `nm -D` 里 `fp8_emu_gemv_block` 计数 = 0） | 规避：改完 `cp -f` 同步被改的 `csrc/**` 到 `${REPO_DIR}`，再 `build-211.sh` + `install-211.sh`；注意 **install 会做 RPATH 后处理，install 件 md5 与中间件不同** | S11f §8.5 | | ✅ | 容器 bridge 网络必须 `--host 0.0.0.0` | 只绑 `127.0.0.1` 时容器内自测正常、宿主却得到 `connection reset by peer`（docker-proxy 连不上回环） | 生产原命令行正是 `--host 127.0.0.1 --port 8000 --host 0.0.0.0`（后者生效）；从宿主访问用 `http://172.17.0.2:` 而不是 `127.0.0.1` | S11f §10.1；S7 §6.1 | | ✅ | 生产容器纪律：全程只读、`RestartPolicy=no` | 生产 `qwen38-vllm` 的 `RestartPolicy=no`、`RestartCount=0`、`StartedAt`/`Pid` 前后完全一致 ⇒ 从未被重启/重建；探针期 HTTP 200 正常 | 硬性口径：对生产容器只允许 `docker ps`/`inspect`/`stats`/`cp` 只读观察，**不许 exec / start / kill / restart / commit / 改配置**；其启停由用户执行 | S3 §3.10；S4 §3.12 | | ✅ | `iommu=pt` 假设已实验证伪 | RCCL 启动警告 `Missing "iommu=pt"`；改 GRUB 加 `iommu=pt` 并重启（内核仍钉 5.15.0-191）：`iommu: Default domain type: **Passthrough**` 生效，但 `Uncached`(0x3) 仍 `out of memory`、RCCL 仍 `via P2P=0`、延迟在噪声内 ⇒ 排除 IOMMU 假设 | `iommu=pt` **保留**（RCCL 自述缺它可能不稳定，且本轮全链路已验证）；GRUB 菜单超时已恢复 `hidden`/`0`；回滚：`sudo cp /etc/default/grub.bak-20261011-0312 /etc/default/grub && sudo update-grub` | S11f §10 | | ✅ | bridge 网络下必须 `--host 0.0.0.0`（S11f 复述，与 ✅ 组同名条同源） | 容器是 bridge + docker 端口映射（`8000/tcp → 0.0.0.0:8000`），vLLM 只绑 `127.0.0.1` 时容器内自测正常、宿主却 `connection reset by peer` ⇒ 服务对外不可达 | 启动脚本固定带 `--host 0.0.0.0`（生产原命令行也是靠后一个 `--host` 生效） | S11f §10.1 | | ✅ | 硬链接导致"构建成功但没有新算子"（S11f 复述，与 ✅ 组同名条同源） | 工作区 `/src` 与构建树 `/mnt/game/port/gfx906-vllm-211` 是硬链接；用"写新文件 + rename"的编辑器改源码会断开链接，构建树仍是旧内容 ⇒ `ninja rc=0` 但新算子不在产物里 | 改完 `cp -f` 同步被改的 `csrc/**` 到 `${REPO_DIR}`，再 `build-211.sh _rocm_C` + `install-211.sh _rocm_C`；注意 install 会做 RPATH 后处理，install 件 md5 与中间件不同 | S11f §8.5 | | ✅ | 新测试不入版本控制（S11f 复述，与 ✅ 组 `.gitignore` 条同源） | 仓库 `.gitignore:301` 有 `test_*.py` 规则（该 fork 为"测试/验证脚本不入库"添加），会连 `tests/**` 下 vLLM 官方测试一起忽略 ⇒ `git status` 里看不到本轮新增的 `tests/kernels/**` 文件，提交时容易漏 | 提交前加 `!tests/...` 白名单或 `git add -f`；已在问题清单登记（未改 `.gitignore`） | 本轮实测（`git check-ignore -v`） | | ✅ | ISA 探针自报条数与实际不符 | 归档统一口径是"53 条基线 MATCH"，但 `tools/dcu/isa_probe.sh` 的 `BASELINE` 表**实际 52 条**（32 `asm` + 7 `blt` + 13 `cgo`），`baseline:` 行的数字来自硬编码 `${ISA_PROBE_BASELINE_N:-53}` ⇒ 门禁自报数字本身也要复核 | 未改脚本（避免与已归档的输出字符串不一致）；建议后续把它改成由表长动态计算，本次仅在文档登记 | S11f §8.2 | | ✅ | bf16/fp16 dense GEMV 派发不可达（已修，S11f 复述，与 ✅ 组同名条同源） | `_ON_GFX9 = gfx90a/gfx942/gfx950` 是上游 CDNA2/3 语义 ⇒ `on_gfx9()` 在 gfx906 恒 `False`，`n <= 16 and not on_gfx9()` 的 Triton 分支排在 GEMV 之前，把该内核**全部**可用形状截走 ⇒ `VLLM_DCU_USE_BF16_GEMV=1` 也从未生效（死代码） | 已修（分支提前 + `on_gfx906()` 守卫）+ 2 条派发回归测试；换用引擎真正在用的 `triton_matmul` 作基线后收益 **2.28–8.54×** | S11f §8.1 | ### 4. 条目与状态统计 - **条目总数：103 条** = §1–§6 的 92 条（§1 平台/生态 13 · §2 编译与图模式 18 · §3 内存 10 · §4 数值与质量 10 · §5 内核能力边界 15 · §6 运维/构建/依赖 26）+ §7 本轮（2026-10-11）新增 11 条。 - **分档计数：⛔ 24 条（原 22 + 本轮 2）、🔶 39 条（原 36 + 本轮 3）、✅ 40 条（原 34 + 本轮 6）**；原文 §8 的"22 / 36 / 34"只统计 §1–§6 的 92 条、未含 §7，故合表后为 24 / 39 / 40（总计 103）。 - **跨档说明**：同一条目可能在执行过程中跨档迁移（最典型 `--enforce-eager`：⛔ 禁令 → 根因定位 → csrc 根修 → ✅ 禁令事实基础移除），本表按**最终状态**归一次、不重复计数。 完整 103 条的原始版见 git `ec19a76d9` 与本文件同目录的归档（`part3_archived_issues.md` 的 §1–§6 92 条 + `part6_new_issues.md` 的 §7 11 条）。 --- ## 五、速查：开关、工具、复现命令 ### 1. 仿真内核开关（默认全 OFF，关闭即回到原行为） | 变量 | 默认 | 作用 / 备注 | |---|---|---| | `VLLM_DCU_USE_GFX906_EMU` | `False` | **总开关**，一次打开 INT8/FP8/dense GEMV 三族；显式单项设置优先（未设置→跟总开关，显式 0→关） | | `VLLM_DCU_USE_INT8_SDOT4` | `False` | INT8 `sdot4`；默认关是**数值理由**（改 greedy token） | | `VLLM_DCU_USE_FP8_EMU` | `False` | FP8 per-tensor/per-channel **+ block-scaled 128×128** | | `VLLM_DCU_FP8_CVT` | `"lut"` | fp8→fp16 加宽路线，固定 LUT（位运算仅作 A/B） | | `VLLM_DCU_USE_BF16_GEMV` | `False` | dense fp16/bf16 decode GEMV，只启用实测获胜形状（单 token 全 N；`bf16 && M=4 && 128≤N≤1024`） | | `VLLM_DCU_BF16_ROUTE` | `"auto"` | `auto`=fp16→dot2 / bf16→fp32；`f32` 强制精确；`f16` 重定标快路径（默认不开放） | | `VLLM_ROCM_USE_SKINNY_GEMM` | `0`（env.sh 兜底） | gfx906 无 MFMA，skinny GEMM 无用 | | `VLLM_GFX906_TRITON_CAUSAL_CONV1D` | `0` | 置 1 才启用 Triton conv（质量回归，见 S15） | ### 2. 工具 | 工具 | 用途 | |---|---| | `tools/dcu/isa_probe.sh` | ISA 三列表 + 53 条基线；`ISA_PROBE_STRICT=1` 不符即 exit 1 | | `tools/dcu/p2p_verify.py` | P2P 能力矩阵 + 跨卡 `copy_` 带宽（容器内跑，显式指定 `HIP_VISIBLE_DEVICES`） | | `tools/dcu/dcu_p2p_acs.py` | PCIe ACS `show` / `apply` / `revert`（需 root，配 `dcu-p2p-acs.service`） | | `tools/dcu/s13_car_probe1.py` | custom AR 管线探测（`meta_size` / uncached 分配） | | `tools/dcu/s11f_fp8_emu_probe.py` / `s11f_bf16_gemv_probe.py` | FP8 数值 sweep + 性能对照 / dense GEMV 加速比 | | `tools/dcu/s11f_e2e_fp8.sh`、`/mnt/game/port/serve_fp8_8000.sh` | 端到端起服务 / 生产槽位 drop-in（:8000，服务名 `qwen3.8-27b`） | | `/mnt/game/port/{hip_p2p_alloc.py,verify_p2p_after_reboot.sh,P2P_REBOOT_NOTE.md}` | P2P 探针与复验、实验记录 | ### 3. 可复制命令 ```bash # 容器环境（所有 GPU 命令的前置） docker exec -it z100-vllm bash . /opt/dtk/env.sh && . /src/tools/dcu/env.sh export PYTHONPATH=/src:/mnt/game/triton-gfx906/install HIP_VISIBLE_DEVICES=0 PYTHONDONTWRITEBYTECODE=1 cd /src # 门禁 ISA_PROBE_STRICT=1 bash tools/dcu/isa_probe.sh /tmp/isa_probe python -m pytest -q --noconftest -p no:cacheprovider \ tests/kernels/quantization/test_gfx906_fp8_emu.py tests/kernels/test_gfx906_bf16_gemv.py \ tests/kernels/quantization/test_gfx906_int8_sdot4.py python tools/dcu/s11f_fp8_emu_probe.py --mode check --max-m 512 --rel-tol 1e-5 --abs-tol 1e-3 # 端到端 / P2P docker exec -d z100-vllm bash /mnt/game/port/serve_fp8_8000.sh # :8000 sudo python3 tools/dcu/dcu_p2p_acs.py show docker exec z100-vllm bash -lc 'PYTHONPATH=/src HIP_VISIBLE_DEVICES=0,1,2,3 python tools/dcu/p2p_verify.py' # 改 csrc 后构建（先同步硬链接，再 build + install） cp -f csrc/rocm/*.cu csrc/rocm/*.h /mnt/game/port/gfx906-vllm-211/csrc/rocm/ docker exec -d z100-vllm bash /mnt/game/port/build-211.sh _rocm_C docker exec z100-vllm bash /mnt/game/port/install-211.sh _rocm_C ``` 机器码门禁（拆 `.hip_fatbin` + 按 kernel 归因）：`llvm-objcopy --dump-section .hip_fatbin=fat.bin $SO null.so` → 按 `__CLANG_OFFLOAD_BUNDLE__` 拆包 → `clang-offload-bundler -unbundle -targets=hipv4-amdgcn-amd-amdhsa--gfx906` → `llvm-objdump -d`，判据见 S11f §3.5/§9.3。 ### 4. 启动参数（生产/测试） `--tensor-parallel-size 4`（graph 下才有 2.44×）· `--gpu-memory-utilization 0.85` · `--max-model-len 262144`（受 pinned 内存约束）· `--max-num-batched-tokens 2048` · `--disable-custom-all-reduce`（gfx906 双重阻断，显式带上免日志噪音）· `--skip-mm-profiling --enable-prefix-caching --mamba-cache-mode align` · `--host 0.0.0.0`（容器内必须）· `--enforce-eager` 视口径（质量参照用 eager） --- ## 六、归档索引 | 步骤 | 归档（`docs/result/…`） | |---|---| | S01 | [`S01`](result/z100-01-arch-normalization-result.md) | | S03 | [`S03`](result/z100-03-bf16-verification-result.md) | | S04 | [`S04`](result/z100-04-build-system-dtk-result.md) | | S04b | [`S04b`](result/z100-04b-stack-upgrade-result-2.md) · [`S04b`](result/z100-04b-stack-upgrade-result.md) | | S04d | [`S04d`](result/z100-04d-build-incrementality-result.md) | | S05 | [`S05`](result/z100-05-csrc-kernels-result.md) | | S06 | [`S06`](result/z100-06-triton-gfx906-result.md) | | S07 | [`S07`](result/z100-07-memory-kvcache-graph-result.md) | | S07b | [`S07b`](result/z100-07b-eager-nan-result.md) | | S07c | [`S07c`](result/z100-07c-packed-silu-csrc-fix-result.md) | | S08 | [`S08`](result/z100-08-attention-and-quant-paths-result.md) | | S08b | [`S08b`](result/z100-08b-gating-and-requirements-result.md) | | S09 | [`S09`](result/z100-09-acceptance-and-bench-result.md) | | S09b | [`S09b`](result/z100-09b-full-cudagraph-result.md) | | S11a | [`S11a`](result/z100-11a-a7-non-dtk-result.md) | | S11b | [`S11b`](result/z100-11b-perf-profiling-result.md) | | S11d | [`S11d`](result/z100-11d-z100-isa-capability-audit-result.md) | | S11e | [`S11e`](result/z100-11e-w8a8-sdot4-result.md) | | S11e2 | [`S11e2`](result/z100-11e2-int8-family-integration-result.md) | | S11f | [`S11f`](result/z100-11f-fp8-bf16-emu-kernels-result.md) | | S11g | [`S11g`](result/z100-11g-1cat-vllm-reference-analysis-result.md) | | S11k | [`S11k`](result/z100-11k-sdot4-compile-integration-result.md) | | S11l | [`S11l`](result/z100-11l-sdot4-reacceptance-result.md) | | S11m | [`S11m`](result/z100-11m-aot-restart-reproducibility-result.md) | | S12 | [`S12`](result/z100-12-p2p-acs-impact-result.md) | | S13 | [`S13`](result/z100-13-custom-allreduce-result.md) | | S14 | [`S14`](result/z100-14-27b-tp4-result.md) | | S14b | [`S14b`](result/z100-14b-graph-vs-eager-result.md) | | S14c | [`S14c`](result/z100-14c-lead-crosscheck.md) · [`S14c`](result/z100-14c-tp4-graph-memory-result.md) | | S14d | [`S14d`](result/z100-14d-determinism-result.md) | | S14e | [`S14e`](result/z100-14e-reproducibility-boundary-result.md) | | S14g | [`S14g`](result/z100-14g-variant-pinning-result.md) | | S15 | [`S15`](result/z100-15-conv-prefill-quality-result.md) | > 完整索引（含状态与归档时间）见 [`docs/result/README.md`](result/README.md)；**完整版文档**（919 行）见 `git show ec19a76d9:docs/海光Z100-优化与算子开发全记录.md`。