把只支持 5090 的推理引擎跑在 4090/4080 上，性能全面超过 SGLang
-
在 Ada 上跑 NInfer：RTX 4090 + 4080 双卡部署记
NInfer 是一个为 RTX 5090 单卡硬编码的手写 C++/CUDA 推理引擎，CMake 第 9 行就拒绝任何非
sm_120a架构。借助一个被上游拒收的社区移植分支，我们把它跑在了 RTX 4090 与 4080 上，两张卡都全面超过原先的 SGLang 部署。日期 2026-09-15 主机 Ubuntu 24.04.4 · 内核 7.0.0-31-generic 模型 Qwen3.8-27B-Uncensored-NInfer · groupwise-int · 18.21 GB 状态 已接入生产客户端 RTX 4090 RTX 4080 --- --- --- 显存 49,140 MiB 32,760 MiB 端口 / model-id :8001 / 4090:8002 / 4080并发 3 2 KV 池 786,432 350,208 可同时吃满 262K 的通道 3 1 decode 单路 115.5 tok/s 88.2 tok/s decode 满并发 207.5 tok/s（3 路） 128.1 tok/s（2 路） prefill（长上下文） 1,068 tok/s 701 tok/s 权重 / 启动后空闲 16.9 / 1.98 GiB 16.9 / 1.13 GiB 启动耗时 12.0 s 12.0 s
一、环境与构建
上游
Neroued/ninfer的 CMake 对非sm_120a直接FATAL_ERROR，官方只支持 RTX 5090。下面用的是社区的 Ada 移植分支memoriaru/ninfer的sm89-ada-port（1 个 commit、18 个文件，上游以「难以维护」为由拒收，已落后上游 93 个 commit，想跟进得自己 rebase）。组件 版本 / 值 OS / 内核 Ubuntu 24.04.4 LTS · 7.0.0-31-generic GPU 0 RTX 4090 · 49,140 MiB · UUID GPU-aaaaaaaa-…（nvidia-smi -L查）GPU 1 RTX 4080 · 32,760 MiB · UUID GPU-bbbbbbbb-…驱动 580.159.04（只支持到 CUDA 13.0） CUDA 工具链 13.1.115（NInfer 硬要求 ≥ 13.1） 源码 memoriaru/ninfer@baf9feb3e41039c7ce92323b9f4f0f732e120866二进制 ninfer-serve177 MiB模型 fullmetaljackass/Qwen3.8-27B-Uncensored-NInfer· 18.21 GB · groupwise-int驱动版本差一档，但能跑。 驱动 580.159.04 报告支持到 CUDA 13.0，而引擎要求工具链 ≥ 13.1。靠 CUDA minor version compatibility——13.1 编出的二进制能跑在支持 13.0 的驱动上——实测无问题，全程零运行期错误。不需要升级驱动。
# 1. 工具链（与系统原有 CUDA 12.8 并存，不动 /usr/local/cuda 符号链接） sudo apt-get install -y cuda-toolkit-13-1 sudo apt-get install -y pkg-config libavformat-dev libavcodec-dev \ libavutil-dev libswscale-dev libcurl4-openssl-dev # 2. 源码 git clone --depth 1 --branch sm89-ada-port \ https://github.com/memoriaru/ninfer.git src # 3. 配置 —— 打印 "arch=89 NVFP4=OFF" 即为正确 cmake -S src -B build -G Ninja -DCMAKE_BUILD_TYPE=Release \ -DCMAKE_CUDA_ARCHITECTURES=89 \ -DCMAKE_CUDA_COMPILER=/usr/local/cuda-13.1/bin/nvcc \ -DCUDAToolkit_ROOT=/usr/local/cuda-13.1 \ -DNINFER_BUILD_APPS=ON # 4. 编译 —— 288 步，零错误 ninja -C build -j6FFmpeg 那几个 dev 包是视频/多模态路径的硬依赖，缺了 CMake 直接失败。系统里有运行时库不够，必须有头文件和
.pc。Ubuntu 24.04 自带的 FFmpeg 6.1.1 正好满足版本下限（libavformat 60 / libavcodec 60 / libavutil 58 / libswscale 7）。编译耗时最久的单个文件是
src/ops/linear/w8/w8_small_t.cu——正是移植里被改造的那个（136 个 kernel 实例从静态共享内存挪到动态）。它编过了，就说明那套绕开 48 KB 限制的办法在 Ada 上成立。
二、跑的是什么模型
下载地址：https://huggingface.co/fullmetaljackass/Qwen3.8-27B-Uncensored-NInfer
格式 / 体积 .ninfer单文件 · 18.21 GB ·weights_id: groupwise-int基座 orcarouter/Qwen3.8-27B-Uncensored←Qwen/Qwen3.8-27B能力 文本 + 视觉 + MTP 投机解码 · 上下文 262,144 · 中英 许可 Apache-2.0 架构是 Qwen3.8-27B：64 层 dense，其中 16 层 full-attention + 48 层 linear-attention（GDN）——这个混合结构决定了 KV 只要 32.4 KB/token。
三点要注意：
- 仓库里只有 3 个文件（
.gitattributes/README.md/.ninfer），没有 safetensors、config.json、tokenizer——全部内嵌在单文件里。这个格式只能给 NInfer 用，喂不了 vLLM / SGLang / llama.cpp。 - 标签里写着
RTX-5090：打包者的目标平台是 Blackwell。Ada 能跑是本文移植的结果，不是发布者的承诺。 - ★必须选
groupwise-int的 artifact。同一个模型的 NVFP4 版本在 Ada 上永远不能用（FP4 张量核是 Blackwell 硬件）。判别方法：.ninfer文件开头是明文 JSON manifest，用 HTTP Range 取前 1.5 MB 就能读出identity.weights_id，不必下整个 18 GB。
三、RTX 4090：NInfer 对比 SGLang
同一天、同一张卡、同一段 prompt、同一套计时方法，512 token 输出，关闭 thinking。这是本文里最干净的一组对照。
项目 SGLang 0.5.19 NInfer 变化 模型 twolven abliterated W4-AWQ · 19 GB Uncensored groupwise-int · 18.2 GB 换了基座 KV 数据类型 fp8_e4m3 int8 — KV 池 451,959 786,432 +74% 单请求上下文 262,144 262,144 持平 并发上限 2 3 +1 可同时吃满 262K 的通道 1 3 ×3 投机解码 NEXTN 树形 topk4/steps5/draft16 MTP 线性 draft3 + lm-head-draft 降级为线性 启动耗时 ~82 s 12.0 s −85% 启动后空闲显存 4.16 GB 1.98 GiB 更紧 decode 单路 101.4 115.5 +14% decode 双路合计 173.2 178.3 +3% decode 三路合计 配置上限 2，未测 207.5 — prefill（长上下文） ~1,115 @207K 1,068 @249,952 基本持平 一个漂亮的整除：
--max-concurrency 3配--kv-capacity auto，引擎解出的 KV 池是 786,432 = 3 × 262,144，页表12,288 / 12,288也刚好用满。3 是「每一路都能吃满满上下文」的最大档，4 路需要 1,048,576，差 26 万装不下。对比之下 SGLang 的 451,959 连两路满上下文（524,288）都装不下——第二路会在准入阶段被
AddReqResult.NO_TOKEN直接拒绝。两处口径说明：
- SGLang 的三路那一栏不是「测了跑不动」，是配置上做不到——
--max-running-requests 2写死了。短请求场景下它未必输，把上限调到 3 再测大概率也能跑。NInfer 站得住的优势是长上下文：786,432 的池子能让三路同时吃满 262K。 - SGLang 那个 prefill 数字采集时有其它客户端并发流量，被污染过，可比性一般。
四、RTX 4080：NInfer 对比 SGLang
4080 上原本跑的是另一个模型（cyankiwi AWQ-INT4），所以这组对照不如 4090 那组严格，配置差异已逐项列出。
项目 SGLang 0.5.19 NInfer 变化 模型 cyankiwi AWQ-INT4 Uncensored groupwise-int 换了模型 KV 池 151,419 350,208 +131% 单请求上下文 147,456 262,144 +78% 并发上限 1 2 +1 投机解码 NEXTN 链式 topk1/steps5/draft6 MTP 线性 draft3 + lm-head-draft — 启动后空闲显存 2.97 GB 1.13 GiB 更紧 decode 单路 未测该模型 88.2 见下 decode 双路合计 并发上限 1 128.1 — prefill（长上下文） 未测 701 @200,932 — 4080 的速度对照要打折看。 被替换掉的那个 SGLang 服务（cyankiwi AWQ-INT4）我们没有测过它的 tok/s。同一张卡上另一个模型 twolven-AWQ 的历史记录是 56–64 tok/s（峰值 68，2026-08-21，上下文 131,072）。拿它和 NInfer 的 88.2 比是 +40~55%——但那是不同模型、不同日期、不同上下文长度的数字，只能当指示性参考，不能当严格 A/B。4090 那一组才是同日同条件的干净对照。
五、启动器与服务单元
两张卡用同一个二进制、同一个 artifact，靠
CUDA_VISIBLE_DEVICES绑 UUID 区分，端口与--model-id沿用原 SGLang 的值，客户端零改动。/data/launchers/start-ninfer-4090.sh#!/bin/bash export CUDA_VISIBLE_DEVICES=GPU-aaaaaaaa-aaaa-aaaa-aaaa-aaaaaaaaaaaa export LD_LIBRARY_PATH=/usr/local/cuda-13.1/lib64:$LD_LIBRARY_PATH exec /data/ninfer-build/build/apps/ninfer-serve \ /data/ninfer-models/qwen3_8_27b_uncensored.ninfer \ --host 0.0.0.0 --port 8001 --model-id 4090 \ --max-context 262144 --kv-capacity auto --max-concurrency 3 \ --kv-dtype int8 --vision \ --spec mtp --draft-tokens 3 --lm-head-draft \ --pending-timeout-ms 600000 \ --request-log-jsonl /data/ninfer-build/reqlog-4090.jsonl \ --log-stats-interval-ms 5000/data/launchers/start-ninfer-4080.sh#!/bin/bash export CUDA_VISIBLE_DEVICES=GPU-bbbbbbbb-bbbb-bbbb-bbbb-bbbbbbbbbbbb export LD_LIBRARY_PATH=/usr/local/cuda-13.1/lib64:$LD_LIBRARY_PATH exec /data/ninfer-build/build/apps/ninfer-serve \ /data/ninfer-models/qwen3_8_27b_uncensored.ninfer \ --host 0.0.0.0 --port 8002 --model-id 4080 \ --max-context 262144 --kv-capacity auto --max-concurrency 2 \ --kv-dtype int8 --vision \ --spec mtp --draft-tokens 3 --lm-head-draft \ --pending-timeout-ms 600000 \ --request-log-jsonl /data/ninfer-build/reqlog-4080.jsonl \ --log-stats-interval-ms 5000参数逐项说明
参数 作用与取值理由 --model-id覆盖 artifact 自报的 identity.model_id。沿用原 SGLang 的served-model-name（4090 / 4080），客户端配置一行都不用改--max-context单条序列的逻辑上限。超出会明确报 context_length_exceeded而不是静默截断--kv-capacity auto按「权重之后的剩余显存 − 1 GiB headroom」反推最大合法容量。所以启动后空闲显存必然落在 1 GiB 出头，那是设计值不是意外 --max-concurrency启动时固定，运行期不可变，引擎硬上限是 1–8。4090 取 3 是为了保证每路都能吃满 262,144 --kv-dtype int8可选 bf16 / int8 / fp8 / nvfp4 / k8v4。nvfp4 在 Ada 上不可用--vision必须显式加，否则不加载视觉的 GPU 分配，多模态请求会失败 --spec mtp --draft-tokens 3MTP 投机解码，draft window 范围 0–5。作者全部基准都用 3 --lm-head-draft启用优化过的 proposal head。该 artifact 内含 text/draft_head张量--pending-timeout-ms默认只有 30,000（30 秒），覆盖 preparation + 媒体获取 + 引擎 FIFO 准入的全过程。长 prompt 场景必须调大，理由见第七节 --request-log-jsonl机器可读的测量日志。含 prefix_cache_hit_tokens与prefix_reuse_path，是排查缓存的唯一直接手段。纯元数据，不含 prompt 文本或图片/etc/systemd/system/ninfer-4090.service（4080 的同构，仅 Conflicts 与 ExecStart 不同）[Unit] Description=NInfer 4090 Qwen3.8-27B-Uncensored groupwise-int+MTP+Vision After=network-online.target Wants=network-online.target Conflicts=sglang-4090-twolven.service sglang-4090-heretic.service StartLimitIntervalSec=0 [Service] Type=simple User=<你的用户名> ExecStart=/data/launchers/start-ninfer-4090.sh Restart=always RestartSec=10 TimeoutStartSec=600 KillSignal=SIGTERM TimeoutStopSec=60 StandardOutput=journal StandardError=journal [Install] WantedBy=multi-user.targetConflicts=让新旧引擎互斥：启动任意一个会自动停掉另一个，回滚就是一条systemctl start。旧的 SGLang launcher 与 unit 全部保留，只是disable。
六、启动日志：容量账怎么读
RTX 4090
engine ready | qwen3.8-27b/groupwise-int | total 12.0s | weights 16.9 GiB capacity | KV 786,432 tokens, int8, auto | pages 12,288/12,288 | runtime 28.2 GiB | free 1.98 GiB context cache | 3 active + 3 cached device states | host 8 states, 8.00 GiB KV | private 6 | shared 4 | anchors 2 media | 12 preprocess workers | cache 1.00 GiB | live 2.00 GiB listening on http://0.0.0.0:8001 | model 4090 | auth disabledRTX 4080
engine ready | qwen3.8-27b/groupwise-int | total 12.0s | weights 16.9 GiB capacity | KV 350,208 tokens, int8, auto | pages 5,472/8,192 | runtime 13.3 GiB | free 1.13 GiB context cache | 2 active + 2 cached device states | host 8 states, 8.00 GiB KV | private 4 | shared 4 | anchors 2 media | 12 preprocess workers | cache 1.00 GiB | live 2.00 GiB listening on http://0.0.0.0:8002 | model 4080 | auth disabled- device states 是 GDN recurrent state 槽位，投机解码回滚要用。这个模型一份 state 是
48 GDN 层 × 48 value heads × 128×128 FP32 ≈ 144 MiB，默认 active 与 cached 各等于并发数，所以每加一路并发约多吃 288 MiB。 - host 8 states / 8.00 GiB KV 在主机内存里，不占显存，是 Device/Host 两级缓存的第二级。长上下文前缀被挤出显存时落到这里，比重算快得多。
- pages 12,288/12,288：4090 的页表刚好用满，4080 的
5,472/8,192还有余量——那是--max-context 262144预留的寻址空间，不是浪费。
七、上线之后才暴露的三件事
跑通只是开始。下面三条在基准测试里完全看不到，只有真实客户端接上去才会撞上，也是这次部署里最值钱的经验。
7.1 前缀缓存决定一切：226 秒对 0.4 秒
把
--request-log-jsonl打开后，request_done事件里有prefix_cache_hit_tokens和prefix_reuse_path两个字段。同一个会话连续几轮的真实记录：消息数 prompt 命中 命中率 真算 reuse_path TTFT 174 107,124 0 0.0% 107,124 root226.2s 176 107,451 107,429 100.0% 22 private_endpoint0.4s 114 171,957 0 0.0% 171,957 root145.0s 116 173,779 173,721 100.0% 58 private_endpoint0.6s 118 174,291 174,125 99.9% 166 private_endpoint0.6s 17.4 万 token 的 prompt，第二轮只真算了 58 个 token，TTFT 从 145 秒降到 0.6 秒。 首轮走
root（全量重算）是不可避免的冷启动成本，之后每轮几乎免费。直接推论：别轻易重启。 每次重启引擎都会清空前缀缓存，所有活跃会话的下一轮都要重付一次全量 prefill。两个 17 万 token 的会话同时冷启动 = 大约 4 分钟不可用。调参时把改动攒到一起，一次重启改完。
7.2 prefill 是串行的，
--max-concurrency给的是 decode 通道throughput | 5.0s | prefill 1.43k tok/s | decode 4.2 tok/s | running 2 (prefill 1, decode-ready 1)注意
prefill 1：并发设成 3，同时在跑的请求可以有 3 个，但任一时刻只有一条在 prefill。两个客户端同时发大请求时，第二个的 TTFT = 自己的 prefill + 前面那个的全部 prefill。实测一条请求queue 2m 20.9s，纯粹是在等前一条 prefill 完。加并发数解决不了这个——瓶颈是 prefill 吞吐（本机 1.2~1.4k tok/s），不是通道数。
7.3
--pending-timeout-ms默认 30 秒，长 prompt 场景必炸req#104 failed during generation | HTTP 503 | request queue timeout # 等了 30 秒 req#105 failed during generation | HTTP 503 | request queue timeout # 等了 30 秒 req#106 failed during generation | HTTP 503 | request queue timeout # 等了 31 秒这个计时从 preparation 之前就开始，覆盖媒体获取和引擎 FIFO 准入的全过程。只要有一条大请求在 prefill（十几万 token = 2 分钟以上），排在它后面的一律 30 秒判死。改成 600,000（10 分钟）后，那条等了 2 分 21 秒的请求正常完成了。
这三条是连环的：串行 prefill 造成长队列 → 30 秒默认超时把队列里的人杀掉 → 而如果缓存命中，这些长 prefill 根本不会发生。先保证缓存能用，再把超时调大兜底。
八、视觉的隐藏上限：32,768 merged token
这条在任何文档的显眼处都找不到，但多图会话一定会撞上。接上真实客户端后出现的第二个 400：
{"code":"media_budget_exceeded","message":"vision raw patches exceed processor budget"} req#78 rejected during prepare | HTTP 400 | media budget exceeded | messages 162 | media 20 | tools 27根因是两个编译期常量（
prepared_prompt.h）：inline constexpr std::uint64_t kMaximumPromptVisionTokens = 32'768; // 整请求 inline constexpr std::uint64_t kMaximumVisionItemTokens = 16'384; // 单张图 inline constexpr std::uint64_t kRawPatchesPerVisionToken = 4; // frontend.cpp vision_tokens = min(options.max_context, kMaximumPromptVisionTokens); // → 32,768注意那个
min()：纯文本能到--max-context的 262,144，但一旦带图，视觉部分的上限就是 32,768。 这两个值没有任何命令行参数可以调。换算（
patch_size 16+merge_size 2，即每个 merged token = 32×32 像素）：图片 像素 raw patches 视觉 token 单请求可容纳 3840×2160 8.29 M 32,400 8,100 4 张 1920×1080 2.07 M 8,100 2,025 16 张 1280×720 0.92 M 3,600 900 36 张 640×360 0.23 M 900 225 145 张 一个攒了 20 张 1080p 截图的会话 =
20 × 2,025 = 40,500 > 32,768，整条请求被拒。一个说明打包意图的细节。 artifact 内嵌的
preprocessor_config.json里size.longest_edge = 16,777,216像素，而kMaximumVisionItemTokens × 4 × 256 = 16,777,216——两个数字完全相等。打包者是刻意把单图上限顶满的：这个 artifact 是为「一张超大图」调的，不是为「很多张图」调的。对照之下，NInfer 引擎结构体里image_max_pixels的默认值只有1024×1024。怎么修：改客户端，别改 artifact
image_max_pixels读自 artifact 内嵌的longest_edge，而smart_resize_image()对超限图片是等比缩小而不是拒绝（processor.cpp:85）。所以把这个值调小确实能让更多图片装进预算——但那会让所有客户端的单图精度一起下降。更好的做法是在客户端侧限制。 我们的客户端（DeepSeek Harness）提供了 provider 级参数，逐图缩放 + 超预算时把最旧的图换成文字占位符：
requestImagePixelBudget: 1048576 # 逐图 ≤1 Mpx，1080p 缩到 ~1365×768，每张 2,025 → 1,024 token requestImageMaxBytes: 393216 # 逐图编码目标 384 KB maxRequestImageBytes: 12582912 # 整请求 12 MB，超出替换最旧的图改完之后同样 20 张图的请求顺利通过。关键是这样做只影响那一个客户端——我们另一条直连 API 的截图判读链路仍然拿到未缩放的 1080p 原图（实测
prompt 2,061，正好等于 2,025 个未缩放的视觉 token）。单位错配是根本原因。 客户端按字节限额，引擎按像素/patches限额。截图压缩率极高，20 MB 的字节额度能装下远超 32,768 视觉 token 的内容，所以字节限额永远不会先触发。要对症就得用像素单位的那个参数。
九、能力与质量验证
三项在 Ada 上此前无人验证过的能力，加一组质量回归。
项目 结果 视觉（sm_89 首次跑通） 通过 640×360 图，四行文字逐字正确 + 形状识别正确，0.6–1.4 s。issue #189 的移植作者只测了文本负载，ViT 路径在 Ada 上没有任何先例 tool call 双向通过 tools+tool_choice:auto→finish_reason: tool_calls与正确的 JSON 参数；role:tool结果回传后正确续写长上下文 通过 4090 吃下 249,952 token（上限的 95%）；4080 吃下 200,932 token 超长保护 正确 明确报 context_length_exceeded，不静默截断智商 12 题 12 / 12（开 thinking），与原模型历史满分一致 拒答 8 题 0 / 8 完全去审 MTP 接受率 4090 53.5–68.1%；4080 67.3–73.5%（强内容相关，见下） 一条容易踩的测量陷阱。 同一套智商题，关 thinking 只有 9/12，开 thinking 是 12/12。错的三题全是数理推理（时钟夹角、鸡兔同笼、星期推算）。测模型智力必须开 thinking，否则会系统性低估。
MTP 接受率与内容类型强相关。引用作者在 5090 上的分类基准（groupwise-int，draft=3）：Structured 89.5%、Code 76.3%、Translation 74.9%，而 Story 只有 37.9%。对应的 decode 是 224.4 / 200.3 / 198.1 / 130.4 tok/s——同一套配置，写作类比结构化输出慢了 42%。所以单一场景的 tok/s 数字意义有限。
十、性能天花板在哪
拿作者在 RTX 5090 上的官方基准做硬件缩放换算，两条线独立验证。
指标 5090（作者） 4090（实测） 硬件比 换算预期 偏差 decode · Code 类 200.3 115.5 带宽 1.78× 112.5 +2.7% prefill · ~26 万 token 1,609.7 1,068 算力 ≈1.5× ~1,073 −0.5% 两条线都精确落在硬件缩放预期上，误差都在 3% 以内。结论是当前配置已经跑在硬件极限，不存在「参数没调好」的空间——能动的只有「每 token 读多少字节」和「每轮吐多少 token」这两件事。
顺带解释一个反常现象：issue #189 的移植作者在 RTX 4090 上只跑出 85.4 / 88.8 tok/s，反而不如我们在更弱的 4080 上跑出的 88.2。原因有二——他跑在 Windows/WDDM（自述有 3.4 GB 桌面合成器基线），且 MTP 接受率只有 53.5–58.3%。Linux 上不要拿那个数字当 Ada 的上限。
还没实施的三条杠杆
方向 机制与预期 风险 --kv-dtype k8v4K 8bit + V 4bit，每 token KV 字节降 25%。4090 上 786,432 ÷ 0.75 = 1,048,576 = 4 × 262,144——并发可升到 4 且每路仍吃满 262K。满上下文时 KV 约占每步流量 1/3，砍 V 一半 → 总流量 −12%4-bit V cache 无任何公开质量数据，作者全部基准都用 int8 --draft-tokens 3 → 4/5verify 那一步无论窗口多长权重只读一次，窗口越长＝每次权重读取吐越多 token。按 Code 类接受率 76.3% 推，draft=4 可到 1+4×0.76 ≈ 4.04tokens/round，约 +23%。作者从未发布过窗口扫描数据强内容相关，写作类（接受率 37.9%）可能反而变慢 并发 → 8 作者 5090 数据：C=1 185.8 → C=2 247.0（1.33×）→ C=4 309.5（1.67×）→ C=8 535.0（2.88×）。我们 C=3 已达 1.80× 每路上下文压到约 9.8 万，与「保满 262K」冲突
十一、踩过的坑
- 首次启动必定 FAILURE 一次。 日志是
ERROR cannot bind 0.0.0.0:8001——Conflicts=停旧服务与新服务抢端口之间有竞态。Restart=always10 秒后重试即成功，不是故障。开机场景不会出现（旧服务已 disable，没有竞态对象）。 pkill -f ninfer-serve会杀掉执行它的 shell 自己，因为那条命令的 command line 里也含这个字符串。用pkill -f "[n]infer-serve"。- 通过 ssh 用
nohup … &起守护进程不可靠——日志文件都没被创建。直接上 systemd。 - 三个默认值在长 prompt 场景下都偏小：
--pending-timeout-ms30 秒（排队必炸）、视觉预算 32,768 merged token（多图会话必炸）、以及客户端侧按字节而非像素限额（永远不先触发）。这三条在单请求基准里全都看不见。 - 先开
--request-log-jsonl再排障。prefix_reuse_path一个字段就能区分「缓存没命中」和「缓存不工作」——前者是冷启动，后者才是 bug。我们一度以为是后者，日志证明是前者。 - 不必为了验证格式下载整个 18 GB。
.ninfer是私有单文件格式，但开头就是明文 JSON manifest。用 HTTP Range 取前 1.5 MB 就能读出identity.weights_id，并数清vision/（333 个张量）和mtp/（12 个张量）命名空间是否存在。 - 在中国大陆需要走 HF 镜像。 本机
huggingface.co被 DNS 污染（解析到 Facebook 的 IPv6 段2a03:2880:…，curl 卡在 SYN-SENT）。配HF_ENDPOINT=https://hf-mirror.com后正常。 - 装 CUDA 前先做 apt 模拟。
apt-get install -s数Remv行确认为 0 再执行。实测cuda-toolkit-13-1是 63 新装 + 1 升级（cuda-toolkit-config-common纯配置包）、零删除，不碰驱动、dkms、libnvidia-*。装完要复核/usr/local/cuda符号链接没有被改走。
十二、值不值得
值得，但要清楚自己押上了什么。
收益是真实的：4090 上 KV 池 +74%、并发从 2 到 3 且每路都能吃满 262K、单路 decode +14%、启动快 85%；4080 上 KV +131%、上下文 +78%。视觉和 tool call 都完整可用。
代价同样真实：两张卡现在跑同一个引擎、同一个 artifact，而这个引擎的 Ada 支持来自一个 0 star、单 commit、上游明确拒收、已落后 93 commit 的 fork。 引擎与模型的多样性都没了，它出问题两张卡一起受影响。上游每前进一步，这个分支的 rebase 成本就高一分。
我们的对冲是保留完整回滚路径——旧的 SGLang launcher 和 unit 一个都没删，只是
disable，Conflicts=会在启动旧服务时自动停掉 NInfer：sudo systemctl start sglang-4090-twolven # 4090 回退 sudo systemctl start sglang-4080-awq-int4 # 4080 回退如果只带走一条：先确认前缀缓存在你的客户端上真的命中（开
--request-log-jsonl看prefix_reuse_path）。命中与不命中之间是 226 秒和 0.4 秒的差距，比任何参数调优都重要一个数量级。如果你手上是 5090，这些移植工作一样都不用做，直接用上游；如果是 Ada 消费卡，上面这条路是通的，但请把它当成一个需要你自己维护的分支，而不是一个能长期依赖的上游产品。
全部数值为 2026-09-15 单机实测。5090 数字引自 NInfer 仓库
docs/performance.md。
硬件 RTX 4090 48 GB + RTX 4080 32 GB · Ubuntu 24.04.4 · CUDA 13.1.115 · 驱动 580.159.04 - 仓库里只有 3 个文件（
-
-
奇怪了
4090 48gb 居然跟 3090 x2 差不多？
这样看来 3090 还真是性比价 之王 -
在 Ada 上跑 NInfer：RTX 4090 + 4080 双卡部署记
NInfer 是一个为 RTX 5090 单卡硬编码的手写 C++/CUDA 推理引擎，CMake 第 9 行就拒绝任何非
sm_120a架构。借助一个被上游拒收的社区移植分支，我们把它跑在了 RTX 4090 与 4080 上，两张卡都全面超过原先的 SGLang 部署。日期 2026-09-15 主机 Ubuntu 24.04.4 · 内核 7.0.0-31-generic 模型 Qwen3.8-27B-Uncensored-NInfer · groupwise-int · 18.21 GB 状态 已接入生产客户端 RTX 4090 RTX 4080 --- --- --- 显存 49,140 MiB 32,760 MiB 端口 / model-id :8001 / 4090:8002 / 4080并发 3 2 KV 池 786,432 350,208 可同时吃满 262K 的通道 3 1 decode 单路 115.5 tok/s 88.2 tok/s decode 满并发 207.5 tok/s（3 路） 128.1 tok/s（2 路） prefill（长上下文） 1,068 tok/s 701 tok/s 权重 / 启动后空闲 16.9 / 1.98 GiB 16.9 / 1.13 GiB 启动耗时 12.0 s 12.0 s
一、环境与构建
上游
Neroued/ninfer的 CMake 对非sm_120a直接FATAL_ERROR，官方只支持 RTX 5090。下面用的是社区的 Ada 移植分支memoriaru/ninfer的sm89-ada-port（1 个 commit、18 个文件，上游以「难以维护」为由拒收，已落后上游 93 个 commit，想跟进得自己 rebase）。组件 版本 / 值 OS / 内核 Ubuntu 24.04.4 LTS · 7.0.0-31-generic GPU 0 RTX 4090 · 49,140 MiB · UUID GPU-aaaaaaaa-…（nvidia-smi -L查）GPU 1 RTX 4080 · 32,760 MiB · UUID GPU-bbbbbbbb-…驱动 580.159.04（只支持到 CUDA 13.0） CUDA 工具链 13.1.115（NInfer 硬要求 ≥ 13.1） 源码 memoriaru/ninfer@baf9feb3e41039c7ce92323b9f4f0f732e120866二进制 ninfer-serve177 MiB模型 fullmetaljackass/Qwen3.8-27B-Uncensored-NInfer· 18.21 GB · groupwise-int驱动版本差一档，但能跑。 驱动 580.159.04 报告支持到 CUDA 13.0，而引擎要求工具链 ≥ 13.1。靠 CUDA minor version compatibility——13.1 编出的二进制能跑在支持 13.0 的驱动上——实测无问题，全程零运行期错误。不需要升级驱动。
# 1. 工具链（与系统原有 CUDA 12.8 并存，不动 /usr/local/cuda 符号链接） sudo apt-get install -y cuda-toolkit-13-1 sudo apt-get install -y pkg-config libavformat-dev libavcodec-dev \ libavutil-dev libswscale-dev libcurl4-openssl-dev # 2. 源码 git clone --depth 1 --branch sm89-ada-port \ https://github.com/memoriaru/ninfer.git src # 3. 配置 —— 打印 "arch=89 NVFP4=OFF" 即为正确 cmake -S src -B build -G Ninja -DCMAKE_BUILD_TYPE=Release \ -DCMAKE_CUDA_ARCHITECTURES=89 \ -DCMAKE_CUDA_COMPILER=/usr/local/cuda-13.1/bin/nvcc \ -DCUDAToolkit_ROOT=/usr/local/cuda-13.1 \ -DNINFER_BUILD_APPS=ON # 4. 编译 —— 288 步，零错误 ninja -C build -j6FFmpeg 那几个 dev 包是视频/多模态路径的硬依赖，缺了 CMake 直接失败。系统里有运行时库不够，必须有头文件和
.pc。Ubuntu 24.04 自带的 FFmpeg 6.1.1 正好满足版本下限（libavformat 60 / libavcodec 60 / libavutil 58 / libswscale 7）。编译耗时最久的单个文件是
src/ops/linear/w8/w8_small_t.cu——正是移植里被改造的那个（136 个 kernel 实例从静态共享内存挪到动态）。它编过了，就说明那套绕开 48 KB 限制的办法在 Ada 上成立。
二、跑的是什么模型
下载地址：https://huggingface.co/fullmetaljackass/Qwen3.8-27B-Uncensored-NInfer
格式 / 体积 .ninfer单文件 · 18.21 GB ·weights_id: groupwise-int基座 orcarouter/Qwen3.8-27B-Uncensored←Qwen/Qwen3.8-27B能力 文本 + 视觉 + MTP 投机解码 · 上下文 262,144 · 中英 许可 Apache-2.0 架构是 Qwen3.8-27B：64 层 dense，其中 16 层 full-attention + 48 层 linear-attention（GDN）——这个混合结构决定了 KV 只要 32.4 KB/token。
三点要注意：
- 仓库里只有 3 个文件（
.gitattributes/README.md/.ninfer），没有 safetensors、config.json、tokenizer——全部内嵌在单文件里。这个格式只能给 NInfer 用，喂不了 vLLM / SGLang / llama.cpp。 - 标签里写着
RTX-5090：打包者的目标平台是 Blackwell。Ada 能跑是本文移植的结果，不是发布者的承诺。 - ★必须选
groupwise-int的 artifact。同一个模型的 NVFP4 版本在 Ada 上永远不能用（FP4 张量核是 Blackwell 硬件）。判别方法：.ninfer文件开头是明文 JSON manifest，用 HTTP Range 取前 1.5 MB 就能读出identity.weights_id，不必下整个 18 GB。
三、RTX 4090：NInfer 对比 SGLang
同一天、同一张卡、同一段 prompt、同一套计时方法，512 token 输出，关闭 thinking。这是本文里最干净的一组对照。
项目 SGLang 0.5.19 NInfer 变化 模型 twolven abliterated W4-AWQ · 19 GB Uncensored groupwise-int · 18.2 GB 换了基座 KV 数据类型 fp8_e4m3 int8 — KV 池 451,959 786,432 +74% 单请求上下文 262,144 262,144 持平 并发上限 2 3 +1 可同时吃满 262K 的通道 1 3 ×3 投机解码 NEXTN 树形 topk4/steps5/draft16 MTP 线性 draft3 + lm-head-draft 降级为线性 启动耗时 ~82 s 12.0 s −85% 启动后空闲显存 4.16 GB 1.98 GiB 更紧 decode 单路 101.4 115.5 +14% decode 双路合计 173.2 178.3 +3% decode 三路合计 配置上限 2，未测 207.5 — prefill（长上下文） ~1,115 @207K 1,068 @249,952 基本持平 一个漂亮的整除：
--max-concurrency 3配--kv-capacity auto，引擎解出的 KV 池是 786,432 = 3 × 262,144，页表12,288 / 12,288也刚好用满。3 是「每一路都能吃满满上下文」的最大档，4 路需要 1,048,576，差 26 万装不下。对比之下 SGLang 的 451,959 连两路满上下文（524,288）都装不下——第二路会在准入阶段被
AddReqResult.NO_TOKEN直接拒绝。两处口径说明：
- SGLang 的三路那一栏不是「测了跑不动」，是配置上做不到——
--max-running-requests 2写死了。短请求场景下它未必输，把上限调到 3 再测大概率也能跑。NInfer 站得住的优势是长上下文：786,432 的池子能让三路同时吃满 262K。 - SGLang 那个 prefill 数字采集时有其它客户端并发流量，被污染过，可比性一般。
四、RTX 4080：NInfer 对比 SGLang
4080 上原本跑的是另一个模型（cyankiwi AWQ-INT4），所以这组对照不如 4090 那组严格，配置差异已逐项列出。
项目 SGLang 0.5.19 NInfer 变化 模型 cyankiwi AWQ-INT4 Uncensored groupwise-int 换了模型 KV 池 151,419 350,208 +131% 单请求上下文 147,456 262,144 +78% 并发上限 1 2 +1 投机解码 NEXTN 链式 topk1/steps5/draft6 MTP 线性 draft3 + lm-head-draft — 启动后空闲显存 2.97 GB 1.13 GiB 更紧 decode 单路 未测该模型 88.2 见下 decode 双路合计 并发上限 1 128.1 — prefill（长上下文） 未测 701 @200,932 — 4080 的速度对照要打折看。 被替换掉的那个 SGLang 服务（cyankiwi AWQ-INT4）我们没有测过它的 tok/s。同一张卡上另一个模型 twolven-AWQ 的历史记录是 56–64 tok/s（峰值 68，2026-08-21，上下文 131,072）。拿它和 NInfer 的 88.2 比是 +40~55%——但那是不同模型、不同日期、不同上下文长度的数字，只能当指示性参考，不能当严格 A/B。4090 那一组才是同日同条件的干净对照。
五、启动器与服务单元
两张卡用同一个二进制、同一个 artifact，靠
CUDA_VISIBLE_DEVICES绑 UUID 区分，端口与--model-id沿用原 SGLang 的值，客户端零改动。/data/launchers/start-ninfer-4090.sh#!/bin/bash export CUDA_VISIBLE_DEVICES=GPU-aaaaaaaa-aaaa-aaaa-aaaa-aaaaaaaaaaaa export LD_LIBRARY_PATH=/usr/local/cuda-13.1/lib64:$LD_LIBRARY_PATH exec /data/ninfer-build/build/apps/ninfer-serve \ /data/ninfer-models/qwen3_8_27b_uncensored.ninfer \ --host 0.0.0.0 --port 8001 --model-id 4090 \ --max-context 262144 --kv-capacity auto --max-concurrency 3 \ --kv-dtype int8 --vision \ --spec mtp --draft-tokens 3 --lm-head-draft \ --pending-timeout-ms 600000 \ --request-log-jsonl /data/ninfer-build/reqlog-4090.jsonl \ --log-stats-interval-ms 5000/data/launchers/start-ninfer-4080.sh#!/bin/bash export CUDA_VISIBLE_DEVICES=GPU-bbbbbbbb-bbbb-bbbb-bbbb-bbbbbbbbbbbb export LD_LIBRARY_PATH=/usr/local/cuda-13.1/lib64:$LD_LIBRARY_PATH exec /data/ninfer-build/build/apps/ninfer-serve \ /data/ninfer-models/qwen3_8_27b_uncensored.ninfer \ --host 0.0.0.0 --port 8002 --model-id 4080 \ --max-context 262144 --kv-capacity auto --max-concurrency 2 \ --kv-dtype int8 --vision \ --spec mtp --draft-tokens 3 --lm-head-draft \ --pending-timeout-ms 600000 \ --request-log-jsonl /data/ninfer-build/reqlog-4080.jsonl \ --log-stats-interval-ms 5000参数逐项说明
参数 作用与取值理由 --model-id覆盖 artifact 自报的 identity.model_id。沿用原 SGLang 的served-model-name（4090 / 4080），客户端配置一行都不用改--max-context单条序列的逻辑上限。超出会明确报 context_length_exceeded而不是静默截断--kv-capacity auto按「权重之后的剩余显存 − 1 GiB headroom」反推最大合法容量。所以启动后空闲显存必然落在 1 GiB 出头，那是设计值不是意外 --max-concurrency启动时固定，运行期不可变，引擎硬上限是 1–8。4090 取 3 是为了保证每路都能吃满 262,144 --kv-dtype int8可选 bf16 / int8 / fp8 / nvfp4 / k8v4。nvfp4 在 Ada 上不可用--vision必须显式加，否则不加载视觉的 GPU 分配，多模态请求会失败 --spec mtp --draft-tokens 3MTP 投机解码，draft window 范围 0–5。作者全部基准都用 3 --lm-head-draft启用优化过的 proposal head。该 artifact 内含 text/draft_head张量--pending-timeout-ms默认只有 30,000（30 秒），覆盖 preparation + 媒体获取 + 引擎 FIFO 准入的全过程。长 prompt 场景必须调大，理由见第七节 --request-log-jsonl机器可读的测量日志。含 prefix_cache_hit_tokens与prefix_reuse_path，是排查缓存的唯一直接手段。纯元数据，不含 prompt 文本或图片/etc/systemd/system/ninfer-4090.service（4080 的同构，仅 Conflicts 与 ExecStart 不同）[Unit] Description=NInfer 4090 Qwen3.8-27B-Uncensored groupwise-int+MTP+Vision After=network-online.target Wants=network-online.target Conflicts=sglang-4090-twolven.service sglang-4090-heretic.service StartLimitIntervalSec=0 [Service] Type=simple User=<你的用户名> ExecStart=/data/launchers/start-ninfer-4090.sh Restart=always RestartSec=10 TimeoutStartSec=600 KillSignal=SIGTERM TimeoutStopSec=60 StandardOutput=journal StandardError=journal [Install] WantedBy=multi-user.targetConflicts=让新旧引擎互斥：启动任意一个会自动停掉另一个，回滚就是一条systemctl start。旧的 SGLang launcher 与 unit 全部保留，只是disable。
六、启动日志：容量账怎么读
RTX 4090
engine ready | qwen3.8-27b/groupwise-int | total 12.0s | weights 16.9 GiB capacity | KV 786,432 tokens, int8, auto | pages 12,288/12,288 | runtime 28.2 GiB | free 1.98 GiB context cache | 3 active + 3 cached device states | host 8 states, 8.00 GiB KV | private 6 | shared 4 | anchors 2 media | 12 preprocess workers | cache 1.00 GiB | live 2.00 GiB listening on http://0.0.0.0:8001 | model 4090 | auth disabledRTX 4080
engine ready | qwen3.8-27b/groupwise-int | total 12.0s | weights 16.9 GiB capacity | KV 350,208 tokens, int8, auto | pages 5,472/8,192 | runtime 13.3 GiB | free 1.13 GiB context cache | 2 active + 2 cached device states | host 8 states, 8.00 GiB KV | private 4 | shared 4 | anchors 2 media | 12 preprocess workers | cache 1.00 GiB | live 2.00 GiB listening on http://0.0.0.0:8002 | model 4080 | auth disabled- device states 是 GDN recurrent state 槽位，投机解码回滚要用。这个模型一份 state 是
48 GDN 层 × 48 value heads × 128×128 FP32 ≈ 144 MiB，默认 active 与 cached 各等于并发数，所以每加一路并发约多吃 288 MiB。 - host 8 states / 8.00 GiB KV 在主机内存里，不占显存，是 Device/Host 两级缓存的第二级。长上下文前缀被挤出显存时落到这里，比重算快得多。
- pages 12,288/12,288：4090 的页表刚好用满，4080 的
5,472/8,192还有余量——那是--max-context 262144预留的寻址空间，不是浪费。
七、上线之后才暴露的三件事
跑通只是开始。下面三条在基准测试里完全看不到，只有真实客户端接上去才会撞上，也是这次部署里最值钱的经验。
7.1 前缀缓存决定一切：226 秒对 0.4 秒
把
--request-log-jsonl打开后，request_done事件里有prefix_cache_hit_tokens和prefix_reuse_path两个字段。同一个会话连续几轮的真实记录：消息数 prompt 命中 命中率 真算 reuse_path TTFT 174 107,124 0 0.0% 107,124 root226.2s 176 107,451 107,429 100.0% 22 private_endpoint0.4s 114 171,957 0 0.0% 171,957 root145.0s 116 173,779 173,721 100.0% 58 private_endpoint0.6s 118 174,291 174,125 99.9% 166 private_endpoint0.6s 17.4 万 token 的 prompt，第二轮只真算了 58 个 token，TTFT 从 145 秒降到 0.6 秒。 首轮走
root（全量重算）是不可避免的冷启动成本，之后每轮几乎免费。直接推论：别轻易重启。 每次重启引擎都会清空前缀缓存，所有活跃会话的下一轮都要重付一次全量 prefill。两个 17 万 token 的会话同时冷启动 = 大约 4 分钟不可用。调参时把改动攒到一起，一次重启改完。
7.2 prefill 是串行的，
--max-concurrency给的是 decode 通道throughput | 5.0s | prefill 1.43k tok/s | decode 4.2 tok/s | running 2 (prefill 1, decode-ready 1)注意
prefill 1：并发设成 3，同时在跑的请求可以有 3 个，但任一时刻只有一条在 prefill。两个客户端同时发大请求时，第二个的 TTFT = 自己的 prefill + 前面那个的全部 prefill。实测一条请求queue 2m 20.9s，纯粹是在等前一条 prefill 完。加并发数解决不了这个——瓶颈是 prefill 吞吐（本机 1.2~1.4k tok/s），不是通道数。
7.3
--pending-timeout-ms默认 30 秒，长 prompt 场景必炸req#104 failed during generation | HTTP 503 | request queue timeout # 等了 30 秒 req#105 failed during generation | HTTP 503 | request queue timeout # 等了 30 秒 req#106 failed during generation | HTTP 503 | request queue timeout # 等了 31 秒这个计时从 preparation 之前就开始，覆盖媒体获取和引擎 FIFO 准入的全过程。只要有一条大请求在 prefill（十几万 token = 2 分钟以上），排在它后面的一律 30 秒判死。改成 600,000（10 分钟）后，那条等了 2 分 21 秒的请求正常完成了。
这三条是连环的：串行 prefill 造成长队列 → 30 秒默认超时把队列里的人杀掉 → 而如果缓存命中，这些长 prefill 根本不会发生。先保证缓存能用，再把超时调大兜底。
八、视觉的隐藏上限：32,768 merged token
这条在任何文档的显眼处都找不到，但多图会话一定会撞上。接上真实客户端后出现的第二个 400：
{"code":"media_budget_exceeded","message":"vision raw patches exceed processor budget"} req#78 rejected during prepare | HTTP 400 | media budget exceeded | messages 162 | media 20 | tools 27根因是两个编译期常量（
prepared_prompt.h）：inline constexpr std::uint64_t kMaximumPromptVisionTokens = 32'768; // 整请求 inline constexpr std::uint64_t kMaximumVisionItemTokens = 16'384; // 单张图 inline constexpr std::uint64_t kRawPatchesPerVisionToken = 4; // frontend.cpp vision_tokens = min(options.max_context, kMaximumPromptVisionTokens); // → 32,768注意那个
min()：纯文本能到--max-context的 262,144，但一旦带图，视觉部分的上限就是 32,768。 这两个值没有任何命令行参数可以调。换算（
patch_size 16+merge_size 2，即每个 merged token = 32×32 像素）：图片 像素 raw patches 视觉 token 单请求可容纳 3840×2160 8.29 M 32,400 8,100 4 张 1920×1080 2.07 M 8,100 2,025 16 张 1280×720 0.92 M 3,600 900 36 张 640×360 0.23 M 900 225 145 张 一个攒了 20 张 1080p 截图的会话 =
20 × 2,025 = 40,500 > 32,768，整条请求被拒。一个说明打包意图的细节。 artifact 内嵌的
preprocessor_config.json里size.longest_edge = 16,777,216像素，而kMaximumVisionItemTokens × 4 × 256 = 16,777,216——两个数字完全相等。打包者是刻意把单图上限顶满的：这个 artifact 是为「一张超大图」调的，不是为「很多张图」调的。对照之下，NInfer 引擎结构体里image_max_pixels的默认值只有1024×1024。怎么修：改客户端，别改 artifact
image_max_pixels读自 artifact 内嵌的longest_edge，而smart_resize_image()对超限图片是等比缩小而不是拒绝（processor.cpp:85）。所以把这个值调小确实能让更多图片装进预算——但那会让所有客户端的单图精度一起下降。更好的做法是在客户端侧限制。 我们的客户端（DeepSeek Harness）提供了 provider 级参数，逐图缩放 + 超预算时把最旧的图换成文字占位符：
requestImagePixelBudget: 1048576 # 逐图 ≤1 Mpx，1080p 缩到 ~1365×768，每张 2,025 → 1,024 token requestImageMaxBytes: 393216 # 逐图编码目标 384 KB maxRequestImageBytes: 12582912 # 整请求 12 MB，超出替换最旧的图改完之后同样 20 张图的请求顺利通过。关键是这样做只影响那一个客户端——我们另一条直连 API 的截图判读链路仍然拿到未缩放的 1080p 原图（实测
prompt 2,061，正好等于 2,025 个未缩放的视觉 token）。单位错配是根本原因。 客户端按字节限额，引擎按像素/patches限额。截图压缩率极高，20 MB 的字节额度能装下远超 32,768 视觉 token 的内容，所以字节限额永远不会先触发。要对症就得用像素单位的那个参数。
九、能力与质量验证
三项在 Ada 上此前无人验证过的能力，加一组质量回归。
项目 结果 视觉（sm_89 首次跑通） 通过 640×360 图，四行文字逐字正确 + 形状识别正确，0.6–1.4 s。issue #189 的移植作者只测了文本负载，ViT 路径在 Ada 上没有任何先例 tool call 双向通过 tools+tool_choice:auto→finish_reason: tool_calls与正确的 JSON 参数；role:tool结果回传后正确续写长上下文 通过 4090 吃下 249,952 token（上限的 95%）；4080 吃下 200,932 token 超长保护 正确 明确报 context_length_exceeded，不静默截断智商 12 题 12 / 12（开 thinking），与原模型历史满分一致 拒答 8 题 0 / 8 完全去审 MTP 接受率 4090 53.5–68.1%；4080 67.3–73.5%（强内容相关，见下） 一条容易踩的测量陷阱。 同一套智商题，关 thinking 只有 9/12，开 thinking 是 12/12。错的三题全是数理推理（时钟夹角、鸡兔同笼、星期推算）。测模型智力必须开 thinking，否则会系统性低估。
MTP 接受率与内容类型强相关。引用作者在 5090 上的分类基准（groupwise-int，draft=3）：Structured 89.5%、Code 76.3%、Translation 74.9%，而 Story 只有 37.9%。对应的 decode 是 224.4 / 200.3 / 198.1 / 130.4 tok/s——同一套配置，写作类比结构化输出慢了 42%。所以单一场景的 tok/s 数字意义有限。
十、性能天花板在哪
拿作者在 RTX 5090 上的官方基准做硬件缩放换算，两条线独立验证。
指标 5090（作者） 4090（实测） 硬件比 换算预期 偏差 decode · Code 类 200.3 115.5 带宽 1.78× 112.5 +2.7% prefill · ~26 万 token 1,609.7 1,068 算力 ≈1.5× ~1,073 −0.5% 两条线都精确落在硬件缩放预期上，误差都在 3% 以内。结论是当前配置已经跑在硬件极限，不存在「参数没调好」的空间——能动的只有「每 token 读多少字节」和「每轮吐多少 token」这两件事。
顺带解释一个反常现象：issue #189 的移植作者在 RTX 4090 上只跑出 85.4 / 88.8 tok/s，反而不如我们在更弱的 4080 上跑出的 88.2。原因有二——他跑在 Windows/WDDM（自述有 3.4 GB 桌面合成器基线），且 MTP 接受率只有 53.5–58.3%。Linux 上不要拿那个数字当 Ada 的上限。
还没实施的三条杠杆
方向 机制与预期 风险 --kv-dtype k8v4K 8bit + V 4bit，每 token KV 字节降 25%。4090 上 786,432 ÷ 0.75 = 1,048,576 = 4 × 262,144——并发可升到 4 且每路仍吃满 262K。满上下文时 KV 约占每步流量 1/3，砍 V 一半 → 总流量 −12%4-bit V cache 无任何公开质量数据，作者全部基准都用 int8 --draft-tokens 3 → 4/5verify 那一步无论窗口多长权重只读一次，窗口越长＝每次权重读取吐越多 token。按 Code 类接受率 76.3% 推，draft=4 可到 1+4×0.76 ≈ 4.04tokens/round，约 +23%。作者从未发布过窗口扫描数据强内容相关，写作类（接受率 37.9%）可能反而变慢 并发 → 8 作者 5090 数据：C=1 185.8 → C=2 247.0（1.33×）→ C=4 309.5（1.67×）→ C=8 535.0（2.88×）。我们 C=3 已达 1.80× 每路上下文压到约 9.8 万，与「保满 262K」冲突
十一、踩过的坑
- 首次启动必定 FAILURE 一次。 日志是
ERROR cannot bind 0.0.0.0:8001——Conflicts=停旧服务与新服务抢端口之间有竞态。Restart=always10 秒后重试即成功，不是故障。开机场景不会出现（旧服务已 disable，没有竞态对象）。 pkill -f ninfer-serve会杀掉执行它的 shell 自己，因为那条命令的 command line 里也含这个字符串。用pkill -f "[n]infer-serve"。- 通过 ssh 用
nohup … &起守护进程不可靠——日志文件都没被创建。直接上 systemd。 - 三个默认值在长 prompt 场景下都偏小：
--pending-timeout-ms30 秒（排队必炸）、视觉预算 32,768 merged token（多图会话必炸）、以及客户端侧按字节而非像素限额（永远不先触发）。这三条在单请求基准里全都看不见。 - 先开
--request-log-jsonl再排障。prefix_reuse_path一个字段就能区分「缓存没命中」和「缓存不工作」——前者是冷启动，后者才是 bug。我们一度以为是后者，日志证明是前者。 - 不必为了验证格式下载整个 18 GB。
.ninfer是私有单文件格式，但开头就是明文 JSON manifest。用 HTTP Range 取前 1.5 MB 就能读出identity.weights_id，并数清vision/（333 个张量）和mtp/（12 个张量）命名空间是否存在。 - 在中国大陆需要走 HF 镜像。 本机
huggingface.co被 DNS 污染（解析到 Facebook 的 IPv6 段2a03:2880:…，curl 卡在 SYN-SENT）。配HF_ENDPOINT=https://hf-mirror.com后正常。 - 装 CUDA 前先做 apt 模拟。
apt-get install -s数Remv行确认为 0 再执行。实测cuda-toolkit-13-1是 63 新装 + 1 升级（cuda-toolkit-config-common纯配置包）、零删除，不碰驱动、dkms、libnvidia-*。装完要复核/usr/local/cuda符号链接没有被改走。
十二、值不值得
值得，但要清楚自己押上了什么。
收益是真实的：4090 上 KV 池 +74%、并发从 2 到 3 且每路都能吃满 262K、单路 decode +14%、启动快 85%；4080 上 KV +131%、上下文 +78%。视觉和 tool call 都完整可用。
代价同样真实：两张卡现在跑同一个引擎、同一个 artifact，而这个引擎的 Ada 支持来自一个 0 star、单 commit、上游明确拒收、已落后 93 commit 的 fork。 引擎与模型的多样性都没了，它出问题两张卡一起受影响。上游每前进一步，这个分支的 rebase 成本就高一分。
我们的对冲是保留完整回滚路径——旧的 SGLang launcher 和 unit 一个都没删，只是
disable，Conflicts=会在启动旧服务时自动停掉 NInfer：sudo systemctl start sglang-4090-twolven # 4090 回退 sudo systemctl start sglang-4080-awq-int4 # 4080 回退如果只带走一条：先确认前缀缓存在你的客户端上真的命中（开
--request-log-jsonl看prefix_reuse_path）。命中与不命中之间是 226 秒和 0.4 秒的差距，比任何参数调优都重要一个数量级。如果你手上是 5090，这些移植工作一样都不用做，直接用上游；如果是 Ada 消费卡，上面这条路是通的，但请把它当成一个需要你自己维护的分支，而不是一个能长期依赖的上游产品。
全部数值为 2026-09-15 单机实测。5090 数字引自 NInfer 仓库
docs/performance.md。
硬件 RTX 4090 48 GB + RTX 4080 32 GB · Ubuntu 24.04.4 · CUDA 13.1.115 · 驱动 580.159.04@michael-zhou 移植到 sm_89 后有个关键点值得写清楚：5090 原生的 FP4 / tcgen05 路径在 Ada 上没了。你们是把 sm_120a 的 kernel 换成 INT8/FP8 反量化，还是重写成了 groupwise-int kernel？这决定了它能快多久——如果只是把 5090 的 kernel 用兼容层糊过去、实际跑 FP16/INT8，那「全面超过 SGLang」更多是单路 decode 的 kernel 调优，不代表算力效率真追上了。
想核对几点：
- 4090 48G 和 4080 32G 的 KV 池数字（786,432 / 350,208）对应的 block size 和 KV 量化是什么？能吃满 262K 的通道数差距主要来自显存，不是引擎。
- 4080 是 16G 原生魔改成 32G，长时间满载的显存稳定性怎么样，有没有 memtest / 长时间压测过。
- prefill 1068 / 701 tok/s 是 batch=1 还是并发下测的？agent 场景更看多路 prefill 总吞吐，和 SGLang 的 chunked prefill 对比才有意义。
这套 kernel 要是稳定，单卡聊天确实很香；双卡、多模型、agent 还是 SGLang 的主场。
- 仓库里只有 3 个文件（
