git:20260926.8b6690c to git:20260926.6c7f4cf

---

name: ninfer-perf-tuning

description: 在本仓库（Ternary Bonsai 2 27B on NINFER / Ada sm_89）上做性能调优与 kernel 改动时使用。包含测量纪律、档位探针、A/B 设计、数值门，以及一系列会静默骗人的坑。改动任何 dispatch 阈值、kernel schedule 或量化档位之前先读这个。

---

# NINFER 三元引擎性能调优

这个仓库是在别人的基础上做的 Ada 移植 + 本机调优。调优的难点不在改代码，在**别把噪声当结论**。

下面每一条都是踩过之后写的，不是泛泛的原则。

## 一、动手之前：四项必须对齐

比较任何两个配置之前，先确认这四项一致。不对齐的话量到的不是引擎。

| 变量 | 默认值 | 影响 |

|---|---|---|

| 思考模式 | **开**（`--no-thinking` 关闭） | decode **+31%** |

| 采样 | temp 1.0（`--greedy` 关闭） | decode **+12%** |

| 任务 | — | 接受率 **44% ↔ 95.7%** |

| 生成长度 | — | decode 速度随 `--max-new` 变 |

另外两条：

- **显示器必须在主板核显上**（不在独显）。这条值约 10%，且是环境状态不是参数。

- **GPU 温度会漂**：同一夹具 p2048 在 1.12k~1.22k 之间晃。所以**必须真交替**——

每一轮把 A、B 都跑一遍取中位数，而不是先跑完 A 再跑 B。后者的漂移会整段记到 B 头上。

## 二、先问"走了哪一档"，再问"快不快"

**这是本仓库最重要的一条。** 档位由一串 env 可覆盖的阈值决定，阈值配错时速度差异看起来像"内核不行"。

改动 dispatch 之后，第一件事是开探针确认分支：

```bash

NINFER_TERNARY_S8_DEBUG=1 NINFER_TERNARY_S8_DEBUG_MIN_T=5 \

NINFER_TERNARY_S8_DEBUG_BUDGET=100000 \

<bin> <model> --messages <fixture> --max-new 4 --no-thinking 2>&1 \

| grep -o 'rung=[a-z0-9_]*' | sort | uniq -c

```

**每一条可走的分支都必须有探针。** 缺一条的代价是真金白银的：分块 GEMV 兜底路径曾经没有探针，

于是一个走那条路的夹具打不出任何 prefill 档位，我把它误判成"夹具异常"查了很久。

现在 `small_t` / `gemv_tile` / `s8` / `small_t_tiled` / `short_mma` / `wide_mma` /

`block_gemv` / `reference` 八条都会报出自己。

### 档位图（当前默认）

| T | 档位 | 阈值 / 开关 |

|---|---|---|

| 1 | `small_t`（T=1 入口） | `NINFER_TERNARY_DECODE=gemv` 回退 |

| 2..8 | `small_t` | `verify_token_cap()`（钳在 8） |

| 9..16 | `small_t_tiled` | `NINFER_TERNARY_GAP_SMALL_T_MAX`（16） |

| 17..∞ | `s8` | `NINFER_TERNARY_S8_MIN_TOKENS`（17）；`NINFER_TERNARY_S8=0` 关闭 |

| （s8 关时）≥32 | `short_mma` / `wide_mma` | `NINFER_TERNARY_MMA_MIN_TOKENS`（32）、`NINFER_TERNARY_SHORT_TILE_TOKENS`（64） |

| 兜底 | `block_gemv` / `reference` | `NINFER_TERNARY_PREFILL`、`NINFER_TERNARY_ROWS`、`NINFER_TERNARY_TILE` |

**注意 T 不是 prompt 长度。** 实测规律（七个夹具一致）：

```

线性层看到的 T = prompt_tokens − 4

```

所以 `mma_min_tokens = 32` 实际覆盖的是 **prompt ≥ 36**。把阈值和 prompt 长度混起来会算错覆盖面。

## 三、A/B 怎么做

- **同一个二进制、只用环境变量切**。两个二进制比会混进编译差异。

- 新加的档位必须**从第一天就做成 env 可切**，否则每个变体都要重编。

- **交替跑**（见上）。

- **同时记绝对值**。短提示（T<64）的 prefill t/s 被固定开销淹没——T=35 时 154 t/s 意味着整趟

0.23 s，而 6.65 GiB 权重在 637 GB/s 下只要 10.4 ms。**比值会骗人，用 `model elapsed` 的差值看。**

- 短提示场景要**专门造夹具**。现成夹具的 T 与预期经常差很远，造完先用探针量一遍真实 T。

- **跨版本比较时，"各自默认"和"同精度"必须分成两行。** 本仓库的 int8 档默认开、

而更早的版本根本没有这一档——把两者混在一张表里，量到的是**精度差异**而不是内核差异。

同一张表要同时给出「各自默认」（现实部署的样子）与「全部关掉 int8」（纯内核对比）。

- **长提示要单独设 `--max-context`**（KV 池按它分配），沿用短场景的 8192 会放不下；

长跑一次十几秒，也建议单独一个脚本，别让短场景的轮次干等。

- **跨版本比较前先确认对方的 CLI 认不认你要用的档位。** 正式版的 CLI 不认 `rk4v4`

（只有 serve 认），所以 >64k 的三方对比用 CLI 做不了——这是接口限制，不是没测。

## 四、数值门

改内核 / 改量化档**必须**过数值门。标准口径：

```bash

ninfer-perplexity <model> --text fixtures/wiki-slice.txt --context 512 --stride 256

# → 110 窗口 / 28,160 token

```

| 判据 | 期望 |

|---|---|

| 旧二进制与 `NINFER_TERNARY_S8=0` | **逐位相同**（9.69192）——移植惰性的硬证据 |

| 改数值的档位（s8 等） | 偏移小且**可解释**（s8 是 +0.0015%，量级 = int8 激活量化误差） |

| 任何情况 | **不是几百**。"几百" = 数值坏了，不是"慢" |

**两条容易搞错的地方**：

1. **标准口径用的是 `--context 512`，线性层看到 T=508。** 要验小 T 段的数值（比如 T=9..31 的空档），

必须用 `--context 32 --stride 16` 单独跑一次，否则那个带根本没被覆盖。

2. **正式版文档写的金判据是 PPL ≈ 6.445**，那是**它的语料和协议**下的值。本机同口径是 9.69。

逐位对账要跟本机自己的历史基线比，不要跟它的数字比。

## 五、会静默骗人的坑

按"骗过我的次数"排序。

1. **两个本该不同的臂给出同一个数 → 先怀疑它没跑，而不是它没用。**

我加 `GAP=s8` 这条入口时忘了 scratch 的分配条件是**策略阈值**（33）而不是"可被请求的下界"，

于是 T<33 上永远拿不到 scratch、静默落回 GEMV，读数正好等于基线臂。

**不报错、不崩、PPL 也不动。** 唯一的线索就是那个"一模一样"。

2. **缺探针的分支 = 隐形的分支。** 见 §二。

3. **基准跑着的时候不要重编二进制、不要起第二个基准实例。** 两个实例争 GPU 会让两边都废；

中途换二进制会让同一轮里前后用的不是同一个东西。**改任何东西之前先确认基准已停**。

**占用检测必须锚定到可执行文件路径**，用裸子串会误报。本机实测：`grep -E '[n]infer'`

被一个打开了名为 `极速档（ninfer）` 目录的 dolphin 文件管理器窗口挡下来过。可用的形式：

```bash

ps -eo pid,args | grep -E '(^|/)(ninfer|ninfer-perplexity|ninfer-serve)( |$)' | grep -v <本脚本名>

```

误报的方向是安全的（拒绝启动，不是静默继续），但会白卡一轮——而如果反过来写成

"没匹配到就继续"，一个写坏的 pattern 就会让第二个实例静默起来，那才是真的毁数据。

**任何占用闸门都要用"宁可误报"的极性。**

4. **zsh 的 `noclobber` 会让 `> file` 静默失败**（报「文件已存在」）。脚本里一律用 `>|`。

不显眼是因为它看起来像"这一格没跑"。

5. **管道与重定向是块缓冲的**，后台任务的进度看起来会卡住。写独立日志文件，或等完成通知。

6. **`pkill -f "ninfer"` 会杀掉自己刚起的命令。** 用 `pkill -x ninfer`。

7. **`pgrep` 匹配不上长进程名**（comm 截断到 15 字符，`ninfer-perplexity` 变 `ninfer-perplexi`）。

用 `ps -eo pid,args | grep` 更可靠。

8. **arena 容量必须同步。** 任何从 workspace 分配的新 scratch 都要计入

`ternary_rotation_workspace_bytes()`（它是三元线性的容量声明）。漏了就在 CUDA graph 里越界。

溢出会抛 `std::bad_alloc`（响亮失败），但别依赖这个兜底。

9. **漏写参数会让所有臂落回同一个默认值，给出 N 个一模一样的数。** 这条是 §五.1 的变体，

但更阴——它**不报错、不崩、数也不动**。实例：`ninfer-perplexity` 的 `--kv-dtype` 默认是

`fp8`，五路对照脚本漏写它就会测出五个 fp8。`--kv-dtype` 与 `--max-context` 必须**逐臂显式写**，

并在脚本末尾**断言 N 条臂给出 N 个不同的数**。同样地，`ninfer` CLI 的 KV 默认是 `bf16`、

而 `ninfer-perplexity` 是 `fp8`——**两个入口的默认不一样**。

10. **单位（GB/GiB、MB/MiB）在跨文档引用时会静默漂移。** 本仓库踩过：`637` 这个带宽天花板

一处当 GB/s 用、一处当 GiB/s 用，差 7%；`6.70 GiB` 其实是**含 MTP 的整个制品**，

而每 token 只读 6.33 GiB。**引用任何速率/占比之前，先把分子分母的单位对齐。**

11. **`model elapsed` 的单位随量级变：快的是 `NNN ms`，慢的是 `N.Ns`。**

只抓数字不抓单位，`1.0`（秒）就会和 `111`（毫秒）混进同一列，而且**中位数会被多数值主导**，

看起来像"所有夹具一样快"。抓取必须带上单位并归一：

`grep -m1 'model elapsed' | sed -E 's/.*elapsed +//'` 然后

**先判 `" ms"` 再判 `s`**（`*s)` 会把 `"111 ms"` 也匹配上，因为 `ms` 以 `s` 结尾）。

12. **bash 的 `declare -A` 在循环里不会清空已存在的数组。** 写法 `for case in ...; do declare -A v; ...`

会让**上一个夹具的值留在 `v` 里**，于是几个夹具被池化成一个中位数。

**每个夹具要么用 `unset v`，要么把循环体放进子 shell。** 本仓库的两个 A/B 脚本都栽在这上面，

症状是"四个夹具读数一模一样"——而那正是 §五.1 那个"先怀疑它没跑"的信号。

13. **重编之后、或 GPU 空闲几分钟之后的头一两次运行偏慢，要丢掉。**

实测：同一夹具在改前重编后立刻跑出 **123 / 123 ms**，再跑六次是 **109~112 ms**（中位数 110，

与改前记录一致）。机制是 GPU 时钟从空闲态爬升，不是代码变了。**A/B 的真交替 + 取中位数

能自动吸收一次冷样本**，但如果你只看"第一次跑几秒"就会得出完全错误的结论。

**★ 这条我踩了两次，第二次更贵**：我用"重编后第一次跑"的 128.6 t/s 当基线，

比一个热的 140.2，得出"`--lm-head-draft` 值 +9%"。规范 A/B 测出来是 **−3.6%**（关 145.2）。

**单次粗测只能用来决定"值不值得做规范 A/B"，不能用来下结论。**

一个开关值多少，要三夹具 × 真交替 × 取中位数，脚本见 `~/ninfer-work/lmd-ab.py`。

14. **仓库 → 构建树是手工拷贝的，有两个都会骗人的坑**：

- **漏拷一个文件**时，唯一信号是编译报错，而报的常常是"模板参数不匹配"这类看着像代码错的信息。

改完先 `diff -q` 每个改过的文件，再 `make`。

- **★ `cp -a` 保留 mtime，而 make 只看 mtime。** 拷贝过去之后 `make` 会打印

`Built target ninfer` **却一行 `Building ...` 都没有**——那就是什么都没重编，

改动根本没进二进制。**我把这个当成"改好了、测了、没用"，白扔了一轮结论。**

做法：**拷完 `touch` 一下**（或用不带 `-a` 的 `cp`），并且**每次 `make` 都确认

真的出现了 `Building CXX/CUDA` 行**——"编译成功"和"编译了"是两件事。

15. **动手前先把"这个改动能不能改变量"算清楚。** 本仓库最贵的一次近失：ncu 显示小 n 层只有

10~12% 活跃 warp，我打算把 `kRowsPerCta` 从 64 砍到 32 让 grid 翻倍。**那是算术上无效的**——

每个 warp 独占 16 行、独自走完整个 K，所以 `warp 总数 = n/16` 是个**不变量**，改 CTA 形状只是

重新分组。n=5120 恒为 320 个 warp / 66 SM = 4.85 warp/SM，砍到 32 行也还是它。

**要加 warp 只能沿 K 切。** 这一条省下的不是一次编译，是一整轮"改了、测了、没用"。

判据：**先写不变量，再看改动碰不碰它。**

16. **文档里"不要做"的理由可能是错的，而且错得便宜可测。** `优化空间` 把"把 T 放进 grid 维"列为

不要做，理由是"地板 ×ceil(T/8)"。那个地板**这个内核从来没有过**：`stage_weights` 在 chunk 循环

里、chunk 循环在 token 循环里 ⇒ 权重在 token tile 之间**不复用**，现实现就已经付了

`ceil(T/kTokens)` 遍。所以把 T 铺进 `blockIdx.y` 流量一模一样，纯粹是把 CTA 数乘以 tile 数。

实测 **+26.2%，逐位相同**。**一个"不要做"条目如果理由是算术，就把它当可证伪的预测来读。**

17. **"两臂给出同一个数"要先分清是"臂没跑"还是"臂在这个区间本来就该是恒等"。** 第 1 条讲的是前者。

后者的实例：token-in-grid 这一臂在 `T <= 64` 时只有一个 token tile、`gridDim.y` 恒为 1，

**两臂必然逐位相同**——短夹具（T≈28）上测不出差别是**门控正确**，不是臂坏了。

**判据不是"数有没有动"，而是"去 trace 里看 grid 有没有动"**：`gridY` 从 1 变成 14 才是臂活着的证据。

同理指标要选在臂的作用域内——这个臂必须用长 T 的 prefill 吞吐测，短夹具测的是一段恒等区。

18. **`--cuda-graph-trace=node` 会把 `cudaGraphLaunch` 的 host 时间放大两个数量级，

于是"轮首 GPU 空转"是假的。** trace 里 `cuda_graph.launch` 74 次、均值 **792 µs**，

首个 kernel 起于 +0.741 ms，看着像每轮白等 0.8 ms（曾经的结论是"比 fold 还大 1.6 倍，值得立项"）。

**clean 微基准实测：1307 节点的一次 `cudaGraphLaunch` = 2.12 µs**

（`~/ninfer-work/nsys/graphlaunch-bench.cu`，节点数 100→2000 线性 1.30→2.37 µs）。

同一次 trace 里 `cudaLaunchKernel` 均值 21.66 µs、clean 1.24 µs ⇒ **放大 17~370 倍，图越大越狠**

（CUPTI 要为逐节点关联插桩，放大的正是这个调用本身）。

**规矩：凡"轮首 / 提交 / launch 窗口"的结论，一律不能从这个开关下的 trace 里读。**

kernel 时间轴仍然可用（那是 GPU 侧测量）。**先量后做，省下的是一整条立项。**

19. **ncu 的利用率差要当"有空间"读，不能当"有多少空间"读。** ncu 说 mid 桶 DRAM 只到 63.9%

而大桶 87.6%（本卡实际天花板），我按比例算出"能省 27% 的时间"。**实测只省了 5.2%**

（那个桶 4360 → 4135 µs/轮）。利用率是个比例，不是一条可外推的直线。

**做法：拿利用率定"值不值得动"，拿 A/B 定"动了多少"。**

20. **全局扫描出来的"中性"，可能是逐形状的响应互相抵消。** `SMALL_T_ROWS` 全局扫 16/32/48

得"16 慢 3%、48 慢 2%，默认已最优"，这条一度被写进"不要做"。**按形状分桶之后它是非单调的**：

n ≤ 5120 要 16、n > 5120 要 32，合起来 337 µs/轮（引擎 A/B **+1.8%**）。

**"一个全局旋钮没用"和"这个旋钮没用"是两件事**——先按尺寸分桶再下结论。

21. **一个结构性改动亏了，先找它赢的那一小撮，别急着回退。** small_t 的 split-K 第一版对所有

"行网格填不满"的形状都切 ⇒ **−1%**。拆开看：每片还剩 4~5 个 K 步的形状 **−5.2%**，

只剩 1~2 步的两个形状 +3~6%（prologue 摊不掉）。加上"每片至少 4 步"这道门 ⇒ **+1.3%**。

**判据不是"这个改动好不好"，是"它在什么条件下好"。**

22. **同一夹具换生成长度 = 换内容段落 ⇒ 接受率不同**（`en-code` 上 `--max-new 200` 是 2.73、

`300` 是 2.98，同一二进制同一参数）。**A/B 的结论可以跨长度复用，绝对值不能**——

而且"用哪一段内容"本身是口径的一部分，记录时要连长度一起写。

## 六、本机（4070 Ti SUPER 16 GB / 66 SM）实测的常量

换卡要全部重测。这些值是**这台机器**的。

| 量 | 值 |

|---|---|

| 只读带宽天花板 | **637**（实测）。**单位是 GB/s 还是 GiB/s 尚未钉死**——两种用法差 7%，**钉死前别引用百分比** |

| 权重每趟 | **6.80 GB = 6.33 GiB**（64 层 6461.8 MB + `output_head` 337.7 MB，逐张量累加）。**不是 6.70 GiB——那个数是含 MTP 的整个制品**，而 MTP 只在 draft 段读 |

| 三元 prefill 各档代价 | `s8`：T=15 → 44.3 ms、T=28 → 51.3 ms。**不是"近平"**——mma 条数与 T 无关（tile 固定 64 宽），**T=15 时 76% 的 mma 算在零上**。`small_t`：**T=9..16** 约 **20.5 ms × ceil(T/8)**，**T=1..8 完全平**（13.75 / 13.3 ms，同 kernel 同权重）——**不要把 20.5 外推到 T≤8**。`wide_mma`：T=1024 的正确估计 **468~862 ms**，**不是 149**（那需要 352 TFLOP/s，超过本卡峰值；149 出自 T=39/128 的整模型值） |

| small_t / s8 交叉点 | **T=17** |

| 128 宽 / 64 宽 tile 交叉点 | **T=64**（T=65 时 64 宽要第二趟） |

| decode 带宽 | **6.80 GB / 13.75 ms = 494 GB/s**（约占 637 的 78%） |

| 每个 decode token 的 kernel 数 | **1224 个**（**全在一张 CUDA graph 里**，帧间间隙 **64 ns**，真空闲仅 3.7%）。**不要假设 launch 间隙是毫秒级** |

| 那 22% 带宽缺口的主因 | **三元解码的 ALU，不是访存。** 内核作者自己的对照实验（`small_t.cuh:63-65`）：把解码算术编译掉、其余全不动 → `34816×5120` 跑到 **644 GB/s**。解码占 kernel **24%** ⇒ 0.24 × 13.75 = **3.30 ms** ≈ 缺口。**这条路已攻过三次**（整字切片 −6%、指针上提 0%、拆 LUT −12.6%），不要重开 |

| 每个 decode token 的非矩阵乘开销 | 约 2.4 ms，其中**旋转 1.1 ms**。但旋转的**流量地板只有 0.0136 ms——高 81 倍**，且耗时**对 k 平、对 block 数也平** ⇒ 纯发射/延迟开销，不是带宽。并行度卡在算法映射上（一个 warp 负责一个 (1024-block, token) 对，T=4 时只有 20 个 warp） |

| **decode 轮内预算**（MTP draft 3，clean 20.65 ms/轮，73 轮） | **`ternary_small_t_mma` 420.6 次/轮 × 37.2 µs = 15.6 ms = 79%**；`ternary_rotate_bf16` 274 次 × 3.3 = 905 µs（4.6%）；`recurrent_record` 577 µs；`pq2_mma_s8` 386 µs；`residual_add` 141 次 = 141 µs；`gdn_norm_gating_27` 195 µs。**decode 的杠杆全在这 79% 里** |

| **small_t 逐形状带宽**（decode T=4，`gridX = div_up(n, rows)`） | n=34816 → **580 GB/s**（ncu DRAM **87.6%**，本卡实际天花板）、n=248320 → 593、n=6144 → 487（68.7%）、**n=5120 → 419（63.9%）**、n=4096 → 416（61.4%）、n=1024 → 180（25.8%）。**L2 全程只有 22~27% ⇒ 不是 L2-bound**；mid 桶 36% occupancy 而 DRAM 只到 64% ⇒ **是并发不够**（见坑 19：别按比例外推） |

| small_t vs s8 的行/片旋钮 | 行块：**n ≤ 5120 用 16、n > 5120 用 32**（全局值怎么选都错，见坑 20）。K 片：只在 **T ≤ 8 且 n ≤ 6144 且每片还剩 ≥4 个 K 步**时切（本模型只有 5120×17408 满足）；门开到"所有填不满的形状"会从 +1.3% 变成 −1%（见坑 21）。**每个 reduce 约 1.1 µs，切片数 × n ≤ 25280 是 arena 声明的上限** |

| **跨时段漂移** | **±1~2%，同一机器同一夹具同一二进制。** 45 分钟两轮全表：PRE 长提示 1.25k → 1.22k（−2.4%）、decode 146.9 → 145.4（−1.0%），三版同幅下移。⇒ **比值稳定、绝对值不稳定；A/B 必须当场两臂都跑。** 桌面进程挂在独显上时更明显（kwin + VS Code + Obsidian 实测占 1144 MiB） |

## 七、改 kernel 时的硬约束

- **静态 shared 上限 48 KiB**（sm_89）。超了要用动态 shared + `cudaFuncSetAttribute`，

而**那个调用在 CUDA graph 捕获里是非法的**。所以新 schedule 先算 `kSharedBytes`。

- **不要用 `__launch_bounds__` 的 min-blocks 换 occupancy。** 三处独立实测都指向"单调反向"

（正式版 wide_t +34%、本机 GDN record 单调变慢、本机分块 GEMV 4 CTA 时崩到 68.2 t/s）。

先减 shared，再看寄存器。

- **不要加宽 small_t 的 tile**——正式版实测是回归（45.2 → 19.0 t/s，累加器吃掉了驻留 CTA）。

要覆盖面就用 token 循环。

- **`TernaryLaunch` 是一个函数指针类型**，两个入口点（t1/t8）签名必须一致。加参数要同时改。

- **prefill 是 eager 执行的**（只有三个 decode-batch 捕获入口），所以选档可以按当次 T 走，

不需要在捕获时冻结；但 decode 路径上的选择**必须进程级冻结**（env 只读一次）。