feat(ops): admit the Q6 fused gate_up shape and evaluate it unfused

docs/maintainer/examples/q6-linear.md

@@ -0,0 +1,209 @@
+# Q6 Linear 性能报告：N=34816，K=5120
+
+2026-09-19 首次测量，2026-09-20 更新 T=8～32 的分派与 SIMT 权重流策略（见文末两轮调优）。本文记录 `text/layers/*/mlp/gate_up` 在 Q6 权重下的完整 Linear Op 性能，
+供评估该 shape 和后续调优参考。
+
+## 测量对象与条件
+
+| 项目 | 本次测量 |
+|---|---|
+| 权重 | `Q6_G64_FP16`，row-split layout；每 64 个权重共享一个 FP16 scale，另有 16 字节高位平面 |
+| 数学形状 | `W[34816,5120] × X[5120,T] → Y[34816,T]` |
+| 输入、输出与 policy | BF16 输入、BF16 输出，`A16Only` |
+| GPU / 工具链 | NVIDIA GeForce RTX 5090；CUDA 13.4，Release，`sm_120a` |
+| 计时入口 | `ninfer::ops::linear`；每个 CUDA Graph 包含一次完整 Op 调用 |
+| cache / 采样 | 每个样本前清除 256 MiB L2；5 次 warmup、50 次测量，报告 median |
+| 输入 fixture | 使用公开 bench 的 Q6 packed-weight 与 BF16 activation fixture |
+| 覆盖 | T=1～128 每个整数；512、1024 两个大 T 锚点 |
+| 并发条件 | 测量时同机有一个空闲的常驻 NInfer 服务占用显存；bench 每样本清 256 MiB L2 |
+
+当前被测调用均只有一个 Graph kernel 节点，外部 workspace 为零。测量范围是纯 Linear Op。
+
+## 最终耗时曲线
+
+![Q6 Linear 最终耗时曲线](q6-linear.svg)
+
+左图包含全部 128 个实测点，圆点标出重点 T，橙色强调 T=1、4、8。
+右图分别展示 512、1024。图中连线仅连接相邻实测点。
+
+## 逻辑带宽与 Tensor Core 利用率
+
+权重包含 `34816 × 5120 / 64 × 50 = 139,264,000` 字节（32 字节低位码 + 16 字节高位平面 +
+2 字节 FP16 scale，每 64 个权重一组）。沿用
+[Linear bench 指标定义](../linear-benchmark.md#5-数学工作量与-route-neutral-指标)：
+
+```text
+model_bytes = 139,264,000 + 2 × 5120 × T + 2 × 34816 × T
+逻辑带宽 GB/s = model_bytes / seconds / 1e9
+带宽利用率 = 逻辑带宽 / 1792 GB/s × 100%
+
+有效算力 TFLOP/s = 2 × 34816 × 5120 × T / seconds / 1e12
+TC 利用率 = 有效算力 / 209.5 TFLOP/s × 100%
+```
+
+分母取自 bench 的 RTX 5090 固定规格常量。T=1 使用 GEMV，TC 利用率填 `—`；T≥2 的最终路径
+均使用 BF16 输入、FP32 累加的 dense MMA，使用对应的 **209.5 TFLOP/s** 峰值。
+当前 bench 的 Q6 行未填充 TC 字段，本文按上述公式补算；不加入 padding 或额外 tile 工作量。
+
+“READ %”列是相对 bench 另列的 1674.5 GB/s 实测持续读上限的口径。
+
+| T | 延迟 µs | 逻辑带宽 GB/s | 带宽利用率 | READ % | 有效算力 TFLOP/s | TC 利用率 |
+|---:|---:|---:|---:|---:|---:|---:|
+| 1 | 105.536 | 1320.3 | 73.68% | 78.85% | 3.38 | — |
+| 2 | 107.264 | 1299.8 | 72.53% | 77.62% | 6.65 | 3.17% |
+| 3 | 107.424 | 1298.6 | 72.47% | 77.55% | 9.96 | 4.75% |
+| 4 | 111.392 | 1253.1 | 69.93% | 74.83% | 12.80 | 6.11% |
+| 5 | 126.112 | 1107.5 | 61.80% | 66.14% | 14.13 | 6.75% |
+| 6 | 140.032 | 997.9 | 55.69% | 59.60% | 15.28 | 7.29% |
+| 7 | 150.592 | 928.5 | 51.81% | 55.45% | 16.57 | 7.91% |
+| 8 | 166.752 | 839.0 | 46.82% | 50.10% | 17.10 | 8.16% |
+| 12 | 170.560 | 822.1 | 45.88% | 49.10% | 25.08 | 11.97% |
+| 16 | 168.928 | 832.0 | 46.43% | 49.68% | 33.77 | 16.12% |
+| 20 | 169.728 | 829.9 | 46.31% | 49.56% | 42.01 | 20.05% |
+| 24 | 168.704 | 836.9 | 46.70% | 49.98% | 50.72 | 24.21% |
+| 28 | 158.400 | 893.3 | 49.85% | 53.35% | 63.02 | 30.08% |
+| 32 | 158.528 | 894.6 | 49.92% | 53.43% | 71.97 | 34.35% |
+| 40 | 218.112 | 653.1 | 36.45% | 39.01% | 65.38 | 31.21% |
+| 48 | 215.808 | 663.1 | 37.00% | 39.60% | 79.30 | 37.85% |
+| 56 | 227.008 | 633.2 | 35.33% | 37.81% | 87.95 | 41.98% |
+| 64 | 232.704 | 620.4 | 34.62% | 37.05% | 98.05 | 46.80% |
+| 80 | 252.896 | 575.9 | 32.14% | 34.39% | 112.78 | 53.83% |
+| 96 | 273.248 | 537.7 | 30.01% | 32.11% | 125.25 | 59.79% |
+| 112 | 336.672 | 440.2 | 24.57% | 26.29% | 118.60 | 56.61% |
+| 128 | 328.320 | 455.3 | 25.41% | 27.19% | 138.99 | 66.35% |
+| 512 | 997.984 | 180.5 | 10.07% | 10.78% | 182.91 | 87.31% |
+| 1024 | 1868.510 | 118.3 | 6.60% | 7.07% | 195.38 | 93.26% |
+
+## 曲线解读与验证
+
+[shape 实现](../../../src/ops/linear/q6/shapes/n34816_k5120.cpp) 按 T 选择 SIMT GEMV、
+`k128` 小块 MMA 和固定列宽 MMA 容量：
+
+- **T=1 为 105.184µs、1324.8 GB/s、
+  标称带宽利用率 73.93%**（持续读口径
+  79.11%）。权重 139.3 MB 构成该点的全部流量。
+- **T=1～7 使用 SIMT 路径**（容量 4、5、6、7），**T=8 起进入 MMA**。T=7→8 是本曲线最大的
+  相邻台阶，`delta_pct` 为 +11.2%，从 150.528µs 到
+  167.296µs。SIMT 段每增加一列约 +10µs，MMA 段则有约 165µs 的固定
+  权重量，两者组织不同，当前保留更低延迟的单 token 专用路径。
+- **T=8～32 使用 `BM=32`、`BK=256` 的 MMA tile**，容量 16、24、32；**T=33～64 仍使用
+  `r64` 系列的 `k128` 小块 MMA**。这一段是**带宽受限**的：T=8～32 的逻辑带宽
+  840.8..895.1 GB/s、标称利用率 46.9..50.0%（第二轮调优后），
+  T=33～64 为 669.4..617.0 GB/s、37.4..34.4%。
+  后一段的耗时仍然几乎与 T 无关（T=40 为 216.192µs、T=64 为
+  233.984µs），说明固定成本还是权重读取；受静态 48 KiB 共享内存预算限制，
+  `BK=256` 的 tile 在 T=48 以上放不下（`BM=16` 的变体又因行块并行度不足更慢），
+  这一段是当前实现的已知限制。
+- **T=49～128 使用固定列宽容量**，是本形状相对注册词表形状最主要的差别。词表形状
+  （N=248320）在 49～128 整段只保留一个 128 列 tile，而本形状的 N 小 7 倍，
+  `r64` tile 的行块数为 544 而非 3880。按调优规程的候选流程划分后：
+
+| T | 词表形状的阶梯 | 本形状最终阶梯 | 变化 |
+|---:|---:|---:|---:|
+| 48 | 217.408 | 213.760 | -1.7% |
+| 56 | 330.528 | 226.528 | -31.5% |
+| 64 | 330.880 | 233.984 | -29.3% |
+| 80 | 332.544 | 254.592 | -23.4% |
+| 96 | 334.496 | 275.136 | -17.7% |
+| 112 | 335.104 | 338.560 | +1.0% |
+| 128 | 325.792 | 327.712 | +0.6% |
+
+  全部 128 个实测点相对该基线的变化区间为 **-31.5% .. +1.0%**：最大降幅 -31.5%（T=56），
+  最大涨幅 +1.0%（T=112，落在逐点重复测量的波动范围内）。
+- **T=112 到 128 不再细分**：T=112 为 334.976µs，T=128 为
+  325.600µs。候选在 97～128 段测得的收益不足 1%，
+  不值得再加一条路由。
+- **大 T 锚点**：T=512 为 **997.984µs、
+  182.91 TFLOP/s、TC 利用率
+  87.31%**；T=1024 为
+  **1868.510µs、195.38 TFLOP/s、
+  TC 利用率 93.26%**。
+  两者单独测量，避免平均值掩盖其中一个的回归。
+
+`ninfer_linear_q6_a16_test` 已通过。新增的 `N=34816, K=5120` case 覆盖 ladder 区分出的
+全部边界（1、4、5、6、7、8、9、16、17、24、25、32、33、48、49、50、128、129），并以现有
+A16 误差标准对独立解码的 packed code 与 FP16 scale 作 FP64 oracle 比较。
+调优过程中每次改动 ladder 都重跑该测试，均通过。
+
+## 复现
+
+```bash
+cmake --build build -j --target ninfer_linear_bench ninfer_linear_q6_a16_test
+./build/tests/ninfer_linear_q6_a16_test
+
+./build/bench/ninfer_linear_bench \
+  --qtype q6 --policy a16 --n 34816 --k 5120 \
+  --sweep 1:128 --execution graph --warmup 5 --repeat 50 --flush-mib 256 \
+  --csv-out profiles/bench/q6_n34816_k5120/final_t1_128.csv
+./build/bench/ninfer_linear_bench \
+  --qtype q6 --policy a16 --n 34816 --k 5120 \
+  --sweep 512:1024:512 --execution graph --warmup 5 --repeat 50 --flush-mib 256 \
+  --csv-out profiles/bench/q6_n34816_k5120/final_bulk.csv
+```
+
+继续调优时，按 [Linear 调优与报告规范](../linear-tuning.md) 选择候选、验证数值并收敛分派；
+本报告提供最终实现的测量结果与曲线，不承诺这些容量适用于其他 shape。
+
+## 第二轮调优：T=8～32 换用宽 K tile（2026-09-20）
+
+第一轮之后，T=8～48 整段的耗时几乎与 T 无关（T=8 199.712µs、T=48 217.216µs），
+`k128` 小块 MMA 的固定成本是权重读取。根因是**每个 stage 每行的字节数太少**：
+`BK=128` 时一行一个 stage 只有 64 字节码（32 字节低位 + 16 字节高位按行距 2560 字节分布），
+warp 的 16 字节 `cp.async` 请求跨行跳、线利用率低。把 tile 改成 **`BM=32`、`BK=256`**
+后一行一个 stage 达到 128 字节码，同时行块数从 544 翻到 1088，读效率明显上升：
+
+| T | 第一轮阶梯 µs | 第二轮阶梯 µs | 变化 |
+|---:|---:|---:|---:|
+| 8 | 197.440 | 166.400 | -15.7% |
+| 12 | 197.376 | 169.184 | -14.3% |
+| 16 | 197.696 | 168.544 | -14.7% |
+| 20 | 203.136 | 168.992 | -16.8% |
+| 24 | 202.080 | 168.640 | -16.5% |
+| 28 | 205.568 | 158.464 | -22.9% |
+| 32 | 205.280 | 158.432 | -22.8% |
+
+T=8～32 段均值 **-17.84%**（区间
+-23.81% .. -13.80%），T=8 的逻辑带宽从
+700.5 GB/s 升到 840.8 GB/s（标称利用率 39.09% → 46.92%）。
+全部 128 点相对第一轮阶梯的均值变化为 -3.59%，
+最大降幅 -23.81%，最大涨幅 +0.72%（落在重复测量波动内），
+没有超过 1% 的回归。
+
+第一轮的 `T=7→8` 台阶随之从 +31.2% 收到 +9.9%（151.360µs → 166.400µs）。
+
+被否掉的候选（同一计时合同）：
+
+| 假设 | 实测 | 结论 |
+|---|---|---|
+| `k128` tile 的流水线深度 2 → 3 → 4 | T=8/16/32 三点在 ±1µs 内 | 不是延迟隐藏不足，加深无收益 |
+| `BM=16`、`BK=256` 覆盖 T=48～64 | T=48 从 213.8µs 涨到 251.2µs | 行块并行度不足，该段保留 `r64_c48_k128` |
+
+数值验证：`ninfer_linear_q6_a16_test` 覆盖 ladder 区分出的全部边界
+（1、4、5、6、7、8、9、16、17、24、25、32、33、48、49、50、128、129），
+以 FP64 oracle 比较独立解码的 packed code 与 FP16 scale，通过。
+
+## 第三轮调优：SIMT 权重流走 L2（2026-09-20）
+
+`T=1～7` 的 SIMT 路径把权重平面用 `cp.async` 直接送进共享内存，走的是 `Cache::ca`（缓存到 L1）。
+权重是一次性流式读，L1 里不会复用，反而挤掉激活的常驻行。改成 `Cache::cg`（只过 L2）后
+`c4` 档明显变快，`c5～c7` 档中性：
+
+| T | `ca`（µs） | `cg`（µs） | 变化 |
+|---:|---:|---:|---:|
+| 2 | 109.248 | 107.136 | -1.9% |
+| 3 | 115.680 | 107.616 | -7.0% |
+| 4 | 120.000 | 111.424 | -7.1% |
+| 5 | 131.680 | 126.272 | -4.1% |
+| 6 | 142.144 | 140.352 | -1.3% |
+| 7 | 152.384 | 150.528 | -1.2% |
+
+全部 128 点相对上一轮阶梯的均值变化为 -0.17%，最大降幅 -7.15%，最大涨幅 +0.69%，
+没有超过 1% 的回归。四处 `Cache::ca` 同时改为 `Cache::cg`，与 `r64` MMA 系列的策略一致。
+
+被否掉的候选（同一计时合同，`c4` 档）：
+
+| 假设 | 实测 | 结论 |
+|---|---|---|
+| 增大行块（`RowsPerCta` 8 → 16 → 32） | T=4 从 120.0µs 涨到 120.1 / 126.0µs | 行块并行度已够，加大只增固定开销 |
+| 加深 `cp.async` 流水线（2 → 3 → 4） | T=4 从 120.0µs 涨到 130.0 / 136.8µs | 不是延迟隐藏不足 |
+| 缩小每 stage 组数（16 → 8） | T=4 从 120.0µs 涨到 128.2µs | 每 stage 字节数太少反而更慢 |

docs/maintainer/linear-tuning.md

@@ -83,6 +83,9 @@ requirements for the large-extent region, not a mandatory position in the develo
 A retained Linear performance report describes the final implementation's absolute performance.
 Use the [Q4 6144×5120 report](examples/q4-linear.md) as a worked example, with this structure:
 
+Reports for shapes registered later live beside it; the
+[Q6 34816×5120 report](examples/q6-linear.md) records the fused gate/up projection at Q6.
+
 1. State the format/layout, N/K, input/output types, activation policy, GPU/toolchain, timing
    boundary, cache conditions, warmup/repetitions, and latency statistic.
 2. Plot the final latency at every valid T through 128 on linear axes, marking priority points