把 NInfer 推理引擎移植到 NVIDIA L20（Ada / sm_89 / 92 SM / 48 GB GDDR6 / 864 GB/s）的补丁集、构建工具与实测数据。
基线是 sergiuszm/ninfer-4090（sm_89 / RTX 4090 移植， 源自 Don-Chad/ninfer-3090 → Neroued/ninfer）。
L20 与 RTX 4090 同为 sm_89，所以这是同代移植；差别在 SM 数、带宽与显存：
| RTX 4090 | L20 | |
|---|---|---|
| SM 数 | 128 | 92（−28%） |
| 带宽 | 1008 GB/s | 864 GB/s（−14%） |
| 显存 | 24 GB | 48 GB（+100%） |
| 架构 | sm_89 | sm_89 |
这两条差异决定了移植的全部内容：网格尺寸必须按 92 SM 校准；KV/并发 profile 应当放宽而不是收紧。
单请求，L20，官方 groupwise 产物（16.96 GiB）+ INT8 KV @ 262144 + --prefill-chunk 1024：
| 用例 | 采样 | tok/s | 草稿接受率 | TTFT |
|---|---|---|---|---|
| code | 贪心 | 100.11 | 65.5% | 209 ms |
| qa | 贪心 | 99.98 | 65.1% | 191 ms |
| qa | temp0.7 | 91.39 | 56.7% | 63 ms |
| code | temp0.7 | 82.63 | 48.9% | 63 ms |
| prose | 贪心 | 79.91 | 45.2% | 203 ms |
| prose | temp0.7 | 78.34 | 43.8% | 63 ms |
全部（--spec none 对照） |
— | 38.9 – 39.3 | — | — |
MTP 加速比：code 贪心 2.57× / code temp0.7 2.12× / prose 2.05×。 显存：开 MTP 27,509 MiB，不开 26,197 MiB。
| 指标 | RTX 4090（上游 README） | L20 实测 | 占 4090 |
|---|---|---|---|
| 无投机 | 50.5 | 39.0 | 77% |
| code MTP3 | 142.9 – 148.6 | 100.11 | 68 – 70% |
L20 的带宽是 4090 的 85.7%、SM 数是 71.9%，实测落在 68–77% —— 与硬件规格吻合，移植没有引入额外损失。
| 引擎 | 产物 | 权重 | 无投机 | code MTP3 |
|---|---|---|---|---|
| llama.cpp | Q6_K_XL GGUF | 24.13 GiB | 25.45 | 59.58 |
| NInfer-L20 | container-v2 .ninfer |
16.96 GiB | 39.0 | 100.11 |
两者不是同一量化，所以 1.68× 的差距里混着"产物更小"和"引擎更强"两部分。按字节归一化才能分离引擎贡献：
| 引擎 | 每 token 字节 | ms/token | 有效带宽 | 占 864 GB/s |
|---|---|---|---|---|
| llama.cpp | 25.91 GB | 39.3 | 659 GB/s | 76% |
| NInfer-L20 | 18.21 GB | 25.6 | 710 GB/s | 82% |
→ 同为 Ada、同一张卡，NInfer 的引擎效率高约 8%（与上游自称的 10% 同向）。
补丁共 23 个文件，全部在 patches/l20-port.patch，
由 patches/apply-l20-port.ps1 幂等应用。
| 补丁 | 内容 | 性质 |
|---|---|---|
| P1 | 顶层 CMakeLists.txt 按架构定义 NINFER_SM86 / NINFER_SM89 |
架构身份 |
| P2 | 删除 src/CMakeLists.txt 里硬编码的 NINFER_SM86=1 |
架构身份 |
| P3 | 16 个文件里 28 处兼容类开关改为 SM86 || SM89（行为保持） |
架构身份 |
| P4 | core/device.h：新增 device_sm_count() 声明 |
运行时宽度 |
| P5 | core/device.cu：device_sm_count() 实现（查询一次并缓存） |
运行时宽度 |
| P6 | gdn/chunked/output.cu：波尺寸从硬编码 170 SM 改为运行时 device_sm_count() |
真缺陷 A |
| P7 | sparse_moe/prefill：同类修正（仅 35B-A3B 路径） |
同类缺陷 |
| P8 | bf16_gdn_gating_proj_kernels.cu：split4/2 的协同驻留数 2 → 1 |
真缺陷 C（L20 真实崩溃） |
ninfer_gdn_gating_proj_test ......... Subprocess aborted
bf16_gdn_gating_proj_kernels.cu:310: CUDA_CHECK(cudaLaunchKernelEx(...)) failed:
cudaErrorCooperativeLaunchTooLarge: too many blocks in cooperative launch
根因：cooperative_resident_ctas_per_sm<Bf16Gdn27Geometry, SplitK>() 对所有 SplitK 都返回 2，
但只有 split8 是 8 warps（256 线程 / 65 寄存器），split4、split2 走默认 kBf16GdnWarps = 16
（512 线程 / 74 寄存器）：
| SplitK | 线程 | 寄存器 | 每 CTA | ×2 | Ada 寄存器文件 65,536 |
|---|---|---|---|---|---|
| split8 | 256 | 65 | 16,640 | 33,280 | 放得下 → 2 CTA/SM |
| split4 / split2 | 512 | 74 | 37,888 | 75,776 | 超了 → 只能 1 CTA/SM |
该常量被发射钳制读取，用来决定单次协同发射覆盖多少 token tile。高估 → 网格超出真实驻留容量。
92 SM 上 split2 算出 max_token_tiles = 92×2 ÷ 6 = 30，T=2688 单次发射生成 126 CTA 网格，
真实容量只有 92 → 驱动拒绝。
为什么只在 L20 暴露：4090 有 128 SM，同样的高估（128×2 = 256）刚好盖住测试的网格规模。
关键旁证：bf16_gdn_gating_proj_plan.cpp 的 resident_ctas_27() 本来就是对的
（split8→2、split4/2→1）。两个文件互相矛盾，而只有内核文件那个常量被发射钳制真正读取——
静态断言读的是对的那个，所以问题一直没被发现。
修复是一行：return SplitK == 8 ? 2 : 1;
| 修复前 | 修复后 | |
|---|---|---|
ninfer_gdn_gating_proj_test |
Subprocess aborted | Passed 1.75 s |
| 全套件可运行测试 | 98 中 1 失败 | 98 全通过 / 0 失败 |
constexpr std::int64_t kRtx5090SmCount = 170;
constexpr std::int64_t kTargetCtas = kRtx5090SmCount * kCtasPerSm; // 680L20 真实单波是 92 × 4 = 368。目标值偏大 → jobs_per_block 偏大 → 网格偏窄 → 每 CTA 串行更多。
而且 jobs_per_block == 1 决定 launch_fixed<false> / <true>（MULTI_JOB 特化），
目标值错了会连内核特化一起选错。该路径在 prefill 时每个 GDN 层都跑，而 Qwen3.8-27B 的
65 层里有 49 层是 GDN。
需要精确说明危害边界：w8_config.h 那个分支是 compat 调度，而 Ada 本来就该走 compat 分支
（#else 是 Blackwell 专用）；pdl.cuh 的 SM86 分支是"无 PDL 的普通发射"，Ada 也确实没有 PDL。
所以它当前不改变生成代码，不是性能 bug。
真正的危害是架构身份失效：任何将来为 Ampere 添加的 NINFER_SM86 专用调优会静默作用于 Ada，
而 w8_config.h 恰恰管着 MTP 投影的 W8 GEMM，也就是直接卡在解码关键路径上。
需要 WSL2 + Ubuntu 24.04（22.04 不可用，见下）、CUDA ≥ 12.8、GCC 13、CMake ≥ 3.28。
# 1) 依赖（脚本会把项目的真实版本下限做成硬门禁）
bash scripts/wsl-install-deps.sh
# 2) 克隆上游 + 打补丁 + 配置 + 编译
NINFER_L20_REPO=/mnt/d/ninfer-L20 bash scripts/wsl-build.shwsl-build.sh 会 clone sergiuszm/ninfer-4090@rtx4090-port 到 /opt/ninfer-L20/src、
git apply 本仓库的补丁、用 Ninja 配置（-DCMAKE_CUDA_ARCHITECTURES=89，GCC 13 作 host compiler），
然后 cmake --build -j。
源码放 WSL 原生文件系统（
/opt），不放/mnt/d——跨 9p 做 1400 文件的 CUDA 构建会慢很多倍。
顶层 CMakeLists.txt 的下限是写死的：
pkg_check_modules(FFMPEG REQUIRED IMPORTED_TARGET
libavformat>=60 libavcodec>=60 libavutil>=58 libswscale>=7)
pkg_check_modules(LIBCURL REQUIRED IMPORTED_TARGET libcurl>=7.85)22.04 的 FFmpeg 停在 4.4（libavformat 58.76）、libcurl 7.81，一次 configure 就失败。
24.04 原生满足全部：FFmpeg 6.1（60/60/58/7）+ libcurl 8.5 + GCC 13 默认。
CUDA toolkit 装在某个 WSL 发行版里、而构建要在另一个发行版里做，可以用共享的 /mnt/wsl
（16 GB tmpfs）搬运，实测 4.9 GB 只需几秒：
# 在已有 toolkit 的发行版里
bash scripts/cuda-stage-2204.sh
# 在目标发行版里
bash scripts/setup-2404.shbash scripts/ninfer-start.sh 8090L20 的推荐 profile（脚本内已固化）：
ninfer-serve qwen3_8_27b.ninfer \
--host 0.0.0.0 --port 8090 \
--max-context 262144 --kv-capacity 262144 \
--max-concurrency 1 --max-pending-requests 16 --pending-timeout-ms 600000 \
--prefill-chunk 1024 --kv-dtype int8 \
--spec mtp --draft-tokens 3 --lm-head-draft这里体现了 48 GB 的价值：4090 在 24 GB 下 INT8 KV 只能到 172,032 （196,608 启动即被拒），被迫改用 4-bit 的 E8 模式并承担 5.7% 解码税。 L20 的 48 GB 让 INT8 KV 直接开满原生 262,144，精度更高、无额外税。
--prefill-chunk 保持在 2688 以下（超过会路由到 unsplit 调度，onset 精度差约 1e-5）。
HF 的 main 已经是 v3 容器，基线 fork 只认 v2，不钉 revision 会直接启动失败：
FATAL server failed during startup | artifact magic is not NInfer v2
| revision | 提交日期 | 说明 | 容器版本 | 体积 | 可读 |
|---|---|---|---|---|---|
51630a0c0f |
2026-09-15 | "Publish v3 artifact" ← main 现在指向它 |
3 | 19.03 GiB | ❌ |
dc370fb629 |
2026-09-06 | 加入 DFlash2 陪伴权重 | 2 | 19.03 GiB | ✅ |
3526913004 |
2026-08-14 | 初始产物 | 2 | 16.96 GiB | ✅ |
rtx4090-port 分支最后推送是 2026-09-12，比 v3 发布早三天——这就是漂移来源。
3526913004 的 18,210,531,328 字节 = 16.96 GiB，正是基线 README 引用的那个体积。
M=https://huggingface.co/neroued/Qwen3.8-27B-NInfer/resolve/3526913004
curl -fsSL -o SHA256SUMS.v2 "$M/SHA256SUMS"
curl -L -C - --retry 8 --retry-all-errors -o qwen3_8_27b.ninfer "$M/qwen3_8_27b.ninfer"
sha256sum -c SHA256SUMS.v2 # eec39564993d6e9c7d5e383382a760f093465c9d163ec9a1bd6b80199514bf3e用 curl -r 0-7 只读头部 8 字节就能验证任意 revision 的容器版本，无需整包下载。
中国大陆网络下
huggingface.co可能不可达，把域名换成hf-mirror.com即可。
# 引擎侧（NInfer）：含无投机对照组
python3 scripts/bench-ninfer-l20.py --model /opt/ninfer-L20/models/qwen3_8_27b.ninfer \
--kv-dtype int8 --max-context 262144 --spec mtp --draft-tokens 3
# llama.cpp 侧对照：解码 + 多深度扫描
$env:LLAMA_DIR = "C:\path\to\llama.cpp"
.\bench\bench-baseline.ps1
.\bench\bench-mtp-matrix.ps1bench/results/ 里保留了产生本文所有数字的原始 JSON。
docs/LLAMACPP-BAT-OPTIMIZATION.md 记录了在同一台机器上
对 llama.cpp 启动脚本的实测调优，其中几条与引擎无关、可普遍适用：
--no-mmproj-offload在 48 GB 卡上是净损失：视觉塔钉在 CPU 时 视觉 TTFT 1525 ms → 398 ms、 视觉 prefill 165 → 631 tok/s，代价只有 1106 MiB 显存--threads要按物理核数给：i9-13900K 是 24 物理核，给 16 会闲置 8 个- 每个响应都带隐藏思考：
reasoning_effort只接受low/medium/xhigh（默认xhigh）， 拒绝none；要关思考用顶层{"enable_thinking": false}—— 否则用"答案 token ÷ 总耗时"算速率会远低于引擎实际解码速率
- 未做深上下文扫描。本文所有数字都是浅上下文（prompt 5–81 tokens）。上游数据显示 MTP 接受率 随深度下降（code 81% → 72.9% @128K），L20 的深度曲线需要长 prompt 单独测。
- P6 的收益未做前后 A/B（预期集中在 prefill）。
draft-tokens最优 K 未扫（本文固定 3；上游姊妹 fork 报 MTP4/MTP7 在深上下文更优）。- 35B-A3B 的协同驻留常量未审计：
Bf16Gdn35Geometry返回SplitK == 32 ? 2 : 4， 而 4 CTA × 512 线程 = 2048 线程/SM 超过 Ada 的 1536 线程/SM 上限，同样可疑。 Qwen3.8-27B 是稠密模型不走该路径，所以未触及。 - 6 个
*_real_test未运行（需要真实产物）。 - 未经上游 review。这是个人移植，不是上游官方支持的平台。
.
├─ patches/
│ ├─ l20-port.patch 23 文件补丁（git apply 可直接使用）
│ └─ apply-l20-port.ps1 幂等补丁脚本（逐处精确匹配，失败即报）
├─ scripts/
│ ├─ wsl-probe.sh WSL 环境探测
│ ├─ wsl-install-deps.sh 依赖安装 + 项目版本下限硬门禁
│ ├─ cuda-stage-2204.sh 跨发行版搬运 CUDA toolkit（共享 /mnt/wsl）
│ ├─ setup-2404.sh Ubuntu 24.04 侧接收 + 依赖
│ ├─ wsl-build.sh 克隆 + 打补丁 + 配置 + 编译
│ ├─ ninfer-start.sh 按 L20 profile 启动引擎
│ ├─ bench-ninfer-l20.py 引擎侧基准（含无投机对照组）
│ └─ run-l20-measurement.ps1 端到端测量编排
├─ bench/
│ ├─ bench-baseline.ps1 llama-bench 基线 + MTP 测试
│ ├─ bench-mtp-matrix.ps1 MTP 有效性矩阵
│ ├─ bench-llamacpp-3.8.ps1 多深度解码/prefill 扫描
│ ├─ bench-llamacpp-vision-thinking.ps1 视觉 oracle + 思考开销
│ └─ results/ 原始实测输出
└─ docs/
├─ L20-PORT-SPEC.md 逐补丁技术规格与后续工作
└─ LLAMACPP-BAT-OPTIMIZATION.md llama.cpp 启动脚本调优实测
本项目是 Apache-2.0 许可，且是以下 Apache-2.0 项目的衍生作品：
- Neroued/ninfer — 引擎本体（sm_120a / RTX 5090）
- Don-Chad/ninfer-3090 — SM86 兼容层与 Qwen3.8 运行时
- sergiuszm/ninfer-4090 — 本补丁的直接基线（sm_89）
- UDPSendToFailed/ninfer-4090 — E8 格点 KV 量化
上游源码不在此仓库内分发（见 .gitignore 的 src/）。请自行克隆基线后应用补丁。
改动清单见 NOTICE，符合 Apache-2.0 第 4(b) 条。