抡锤者

yangkang5303

@zrlhk 有道理

yangkang5303

LLM 推理实测(真实模型对照 + 锁回 A/B)

用两个真实模型(同机)、锁回 vs 解锁两次启动实测(prefill 全部冷启动,0 前缀缓存命中):

bf16 稠密模型:Qwen/Qwen2.5-3B-Instruct,torch bf16 + transformers,prefill 8192 tokens + 贪心 decode 128 tokens(精确 token 计数)
三值量化模型:Ternary-Bonsai-2-27B(ninfer-serve,三值权重 × int8 激活整数内核),prefill 18369 tokens,引擎 /metrics 计数器口径

bf16 稠密模型(权重矩阵 = bf16 GEMM,解锁受益路径)

指标	锁定	解锁	提升
Prefill @8192	7,136 t/s	10,579 t/s	+48%
Decode @8192	46.97 t/s	47.18 t/s	~0%

同机 cuBLAS bf16 GEMM(8192³)锁回重测:63.2 TFLOPS(与首轮锁定基线 64.1 一致)→ 解锁 106.8,+69%。模型 prefill 提升低于内核(注意力/归一化/内核启动开销占一部分),方向与量级一致。

三值量化模型(整数内核 + 带宽受限,对照组)

指标	基线(解锁前记录)	解锁后实测	差异
Prefill @~20K	~1998 t/s	2162 / 2169 t/s	+8%(噪声)
Decode @20K	63.0 t/s	70.6 / 70.8 t/s	~0%

解读

	bf16 稠密 3B	三值 27B
计算路径	bf16 GEMM(FFMA/FMLA32)	整数三值 × int8(IMLA)
prefill	+48%	无
decode	~0%(显存带宽瓶颈)	~0%(同)

→ bf16 稠密模型的 prefill(首个 token / 大批量)吃到约 1.5x;decode tok/s 两模型都不变——那是带宽的活,解锁管不到。 三值量化恰好绕开了解锁唯一的受益路径。

yangkang5303

https://lcz.me/topic/2488/22 中没有对魔改20G 3080进行测试，我来补充

v67/PLM 早启动解锁(gpu-unlock 1.0.0)实测。4 种精度 × 3 种规模(4096/8192/16384,各 5 预热 + 12 次计时取中位数),解锁前后各 12 组有效测量,波动 <4%。

环境

项目	值
GPU	NVIDIA GeForce RTX 3080 20GB(改装版,GA102,10DE:2206,68 SM)
驱动	610.57.04(open kernel modules,DKMS)
CUDA	13.3.1
系统	Omarchy 4.0.4(Arch 系),内核 7.2.5-3-omarchy
引导	Limine + UKI,内核参数 `gpu_unlock.targets=all`
功耗上限	320 W(两轮一致)
解锁日志	`pci 0000:07:00.0: gpu_unlock: unlocked`
基准工具	自编译 cuBLAS 13 GEMM(`nvcc -arch=sm_86`,5 预热 + 12 次计时取中位数)

方法

直接调用 cuBLAS 13 的 GEMM(单卡,FP32 累加语义):

FP32:cublasSgemm(CUDA 核 FFMA)
TF32:cublasGemmEx + CUBLAS_COMPUTE_32F_FAST_TF32(张量核)
FP16:cublasGemmEx + CUBLAS_COMPUTE_16F(fp16 输入/输出,内部 FP32 累加)
BF16:cublasGemmEx + BF16 输入 + CUBLAS_COMPUTE_32F(张量核 BF16 路径,FP32 输出)

矩阵规模 N×N,N = 4096 / 8192 / 16384,报告 TFLOPS 为各次计时中位数。测量前有小规模矩阵正确性校验,四种精度全部通过。

结果(中位数 TFLOPS)

精度	规模	解锁前	解锁后	加速
FP32 (SGEMM)	4096	21.6	21.5	1.00x
	8192	21.2	21.2	1.00x
	16384	21.3	21.3	1.00x
TF32	4096	31.0	58.3	1.88x
	8192	31.6	52.2	1.65x
	16384	31.4	53.4	1.70x
FP16	4096	111.7	111.5	1.00x
	8192	107.6	109.1	1.01x
	16384	110.5	111.0	1.00x
BF16	4096	61.6	112.9	1.83x
	8192	64.1	106.8	1.67x
	16384	63.9	105.3	1.65x

各精度平均加速

精度	平均加速
FP32 (SGEMM)	1.00x
TF32	1.74x
FP16	1.01x
BF16	1.72x

解读

解锁生效:主打的 BF16 / TF32 GEMM 分别获得 +65% ~ +88% 提升(BF16 4096³ 从 61.6 → 112.9 TFLOPS,TF32 从 31.0 → 58.3 TFLOPS),与官方报告的"接近 2 倍"一致。
️ FP16 / FP32 基本无变化:这两个精度在解锁前已顶到本卡的性能平台，也是原版3080解锁后的结果（神奇：魔改卡3080默认解锁fp16算力了）。(FP16 ≈ 110 TFLOPS,FP32 ≈ 21.3 TFLOPS)。解锁后 BF16 恰好抬升到与 FP16 相同的平台值(~110 TFLOPS),说明平台由 320 W 功耗上限 / 其他瓶颈决定,而非 FMA 发放速率。
实用收益:BF16 训练/推理、TF32 混合精度矩阵运算受益最大;纯 FP32 与 FP16 负载在你这张卡上不受影响(功耗余量更大的 3090 Ti 上 FP16 也有 +35%,见上游报告)。

参考

项目:Zagitta/gpu-unlock(v67/PLM 早驱动解锁,module GPL-2.0)
同型号先例:Issue #10 — Successful unlock on 3080 20GB
3080 Ti 数据:Issue #2 — BF16 71.6→114,TF32 35.7→65.3

yangkang5303

补充一段真实 agent 负载下的实测，供参考。

场景：本地个人 agent（Hermes）连续 8 轮工具调用，每轮约 20 个 tools，上下文从 20K 累积到 34K。

结果：
Decode 稳定在 70–72 t/s，上下文从 20K 涨到 34K 几乎不衰减 —— rk4v4 的 KV 压缩有效抑制了注意力随长度的成本增长。
前缀缓存命中率随会话爬升：88.5% → 91.1% → 93.1% → 94.9%。
TTFT 从首轮的 9.8s 降到稳态 1.1–1.5s，正是缓存命中带来的直接收益。
Prefill 稳定 1.7–2.25k t/s，排队延迟仅 10–26 ms，主机 CPU 占用 5–8%。

两处配置修正：
--structured-output 是 agent 前端的刚需 —— tool calling 依赖结构化输出，缺了会直接 HTTP 400。

原用 --kv-headroom-mib 2048 时出现警告：CUDA graphs used 34.7 MiB but the KV sizing allowed 12.0 MiB。调到 3072 后警告消失，且 KV 容量完整保住 262,144，剩余显存反而略增（7.00 → 7.12 GiB）。

yangkang5303

在 RTX 3080 20GB 上，用 iamwavecut/ninfer-all
（master d1a0e07）原生构建跑通了 Ternary Bonsai 2 27B，使用 rk4v4 4-bit KV 量化，
完整 262,144 token 上下文，无需任何代码移植。

零编译改动、零源码补丁，CMAKE_CUDA_ARCHITECTURES=86 就是上游默认值。

背景

WaveCut/Ternary-Bonsai-2-27B-NInfer-v3 的模型卡列出了 RTX 3090 / 4090 / 5090 的数据，
但 RTX 3080 在 docs/gpus.md 里是未测试（）状态，而且卡片写的是 12/10GB SKU。

另一个已知限制：在非 ninfer-all 的 NInfer 分支上，rk8v4 在 sm_86 上被硬编码拒绝，
rk4v4 甚至不存在。所以很多人（包括我）一开始以为 3080 需要自己 porting。

实测结论：不需要。 ninfer-all 上游已经把 4-bit KV 全部实现好，并且
"accepted on every build target"。

环境

项目	值
GPU	NVIDIA GeForce RTX 3080 20480 MiB（20GB SKU，GA102）
Compute Capability	8.6
驱动	610.57.04（CUDA 13.3 toolkit）
OS	Omarchy（Arch 衍生）+ 内核 6.x
桌面	Hyprland（同时运行，见下文显存部分）
引擎	`ninfer-serve 0.12.0-rtx3090+d1a0e07`
构建目标	`CMAKE_CUDA_ARCHITECTURES=86`（上游默认）

验证是真·原生 sm_86

cuobjdump --list-elf 显示二进制里每个 cubin 都是 sm_86 SASS，不是 PTX JIT：

ELF file 1: device_calibration.sm_86.cubin
ELF file 2: ninfer-serve.2.sm_86.cubin
ELF file 7: device_tier.sm_86.cubin
ELF file 8: hyper_connection.sm_86.cubin

编译 flag：

--generate-code=arch=compute_86,code=[compute_86,sm_86]

转换

从 PrismML 官方的 Ternary-Bonsai-2-27B-PQ2_0.gguf（6.71 GB，GGUF type 142 / PQ2_0）
转换成 NInfer v3 artifact：

python3 -m tools.convert \
  --model /path/to/Qwen3.8-27B-config \
  --recipe bonsai2_27b_ternary \
  --source ternary=Ternary-Bonsai-2-27B-PQ2_0.gguf \
  --components text \
  --name bonsai2-27b \
  --device cpu \
  --out bonsai2-27b.ninfer

结果：778 objects，771 个张量，70.4 秒，7.20 GB。

注：--model 需要 Qwen3.8-27B 的 config.json + tokenizer 资源（不需要权重，
recipe 只读 model.config，所有张量都从 GGUF 来）。

启动

ninfer-serve bonsai2-27b.ninfer \
  --model-id bonsai2-27b \
  --host 127.0.0.1 --port 1234 \
  --max-context 262144 \
  --kv-capacity auto \
  --kv-headroom-mib 2048 \
  --kv-dtype rk4v4 \
  --gdn-state-fp16 \
  --no-thinking

关键点是 --kv-capacity auto + --kv-headroom-mib 2048：引擎先加载权重、实测剩余显存，
再选最大合法 KV 容量，并强制预留 2 GiB。这是桌面和 CUDA 抢显存问题的正解（见下）。

实测性能

合成基准（单请求）

项目	数值
Prefill @20K	1998 t/s
Decode @20K	63.0 t/s
Decode 短上下文	70.5 t/s

真实负载（28 次请求，约 12 分钟连续使用）

项目	数值
Prefill（平均）	1571 t/s
Decode（平均）	64.1 t/s
累计 prefill	70,174 tokens
累计 decode	10,497 tokens
失败请求	0
拒绝请求	0
引擎恢复事件	0

Prefill 从 1998 降到 1571 是正常的：真实负载里混入了更长、更碎的前缀，
且 95% 的输入被前缀缓存吃掉（见下），剩余重算部分占比小、统计噪声大。

前缀缓存命中率 95.1%

这是个人助理场景最有价值的数据：

前缀缓存命中: 1,353,465 tokens   ← 直接复用，不重算
实际重算:        70,174 tokens
─────────────────────────────────
总输入:       1,423,639 tokens
节省:              95.1%

超过 95% 的输入 token 直接从 context cache 恢复，不经过 prefill。

/slots 也能看到多级 checkpoint 在复用：

{"frontier": 15977, "session_digest": "0b6afd8c..."}
{"frontier": 18602, "session_digest": "0fda02d0..."}
{"frontier": 18962, "session_digest": "706f16d4..."}

显存：rk4v4 是决定性的

这是 3080 20GB 能开 256K 的核心原因。

KV 理论占用（16 层全注意力 / 共 64 层）

该模型是 混合架构（full_attention_interval = 4），64 层里只有 16 层有 KV cache：

KV 格式	256K 占用
F16	16.0 GiB
Q8	8.0 GiB
rk4v4	4.36 GiB

实测整机占用

nvidia-smi: 13165 MiB used / 20480 MiB total   (剩 6865 MiB)

分解（--max-context 262144 --kv-capacity auto）：

项目	占用
模型权重	6.70 GiB
rk4v4 KV（动态分配）	—
运行时 / workspace / CUDA Graph	—
合计	12.86 GiB
剩余	6.70 GiB

同时跑桌面 + 256K 上下文，还剩 6.7 GiB 空闲，桌面完全安全。

和 llama.cpp (Prism fork) 的对比

顺手对比了 PrismML 官方 fork（PrismML-Eng/llama.cpp，分支 prism），
它原生支持同一份 GGUF（type 142）：

指标	llama.cpp (prism) + Q8 KV	ninfer-all + rk4v4
上下文	64,000	262,144（4×）
KV @256K	8.0 GiB	4.36 GiB（-45%）
Prefill @20K	~1,050 t/s	1,998 t/s（1.9×）
Decode @20K	~45 t/s	63 t/s（1.4×）
Decode 短	65.5 t/s	70.5 t/s

llama.cpp 那条路也完全可用，而且胜在零转换（直接读 GGUF）。
但长上下文场景下 ninfer-all 优势明显。

一些说明 / 未做的事

RTX 3080 未在上游 profile 列表，所以首次启动可能触发设备校准（20 秒 ~ 3.5 分钟），
结果会缓存。我这次没遇到明显延迟，可能命中了同 SM 数的 profile。
没有用投机解码。模型卡上 3090 的 201 t/s 是 DFlash2 (5 drafts) 的成绩，
需要额外的 MTP / DFlash2 组件文件，我没有下载。纯 decode 是 ~64-71 t/s。
Vision 未启用（--components text），Bonsai 的视觉塔是独立组件。
这是纯文本个人助理用途的配置。

结论

RTX 3080 (sm_86) 上跑 Ternary Bonsai 2 27B 是一件"开箱即用"的事：

ninfer-all 的默认构建目标就是 86 —— 不需要 porting
rk4v4 等 9 种 KV 格式上游全部实现，sm_86 不再被拒
混合注意力架构（只有 16/64 层占 KV）+ rk4v4，让 20GB 卡能开 256K 上下文
--kv-headroom-mib 让桌面和模型和平共处
实测 prefill ~1571-1998 t/s、decode ~64-71 t/s、前缀缓存命中 95%+、零失败

感谢 upstream 的工作 —— 尤其是把 4-bit KV 从"sm_86 拒绝"变成"全平台可用"这件事。

参考

模型：https://huggingface.co/WaveCut/Ternary-Bonsai-2-27B-NInfer-v3
引擎：https://github.com/iamwavecut/ninfer-all （master d1a0e07）
PrismML llama.cpp fork：https://github.com/PrismML-Eng/llama.cpp （分支 prism）
上游构建文档：docs/build-linux.md、docs/cli.md#context-and-memory

yangkang5303

已经通过打补丁在3080上跑了70t/s 70k ctx（看我的帖子）。RotorQuant rk8v4 KV 模式解决之后，可以进一步增大上下文。

yangkang5303

对代码有一定的修改以适配RTX3080 代码仓库

====================

修复的 bug

文件	问题	修复
`src/ops/gdn_gating_proj/bf16/bf16_gdn_gating_proj_plan.cpp`	协作式 MMA 路由上界按 82 SM 硬编码,68 SM 上 prefill 宽 768 时 144 CTA 超出 136 驻留预算,模型加载即抛 `candidate is not legal for exact problem`	路由上界改为按运行时 `device_sm_count()` 动态推导
`src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_plan.cpp`	A16 融合 swiglu 仅注册到 T≤16,prefill 宽 >16 即抛异常	在 token 轴上按 32 分块调用小-T 核(镜像 FP8 A16 路线)
`src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_small_t.cu`	launcher 只注册 2..16	扩展到 2..32,与线性小-T 生产 schedule 一致

多角度优化结果

指标	nvfp4 A16(原)	groupwise-int(推荐)	提升
权重显存	16.03 GiB	15.25 GiB	−0.8 GiB
prefill(长 prompt)	19.4 tok/s	558–636 tok/s	~29×
decode(贪心)	11.7 tok/s	33.7 tok/s	2.9×
decode(MTP3)	33.0 tok/s	69.5 tok/s	2.1×

投机解码:MTP3 最优(87–91% 接受率);draft=5 反降至 67.6 tok/s(接受率 73%)。

KV 量化

选项	payload(2048 tok)	可用性
bf16	128 MiB	可用
int8-group64	66 MiB	可用(推荐,显存最低)
fp8	—	sm_86 拒绝(需 Blackwell mma.f8f6f4)
rk8v4(q4 V)	—	拒绝(未移植到 `kv_cache_append` Op)

70K 上下文长任务测试

配置:groupwise-int + --kv-dtype int8 + --spec mtp --draft-tokens 3,prompt 69,770 token。

指标	结果
prompt tokens	69,770
prefill	636 tok/s(109.7 s)
decode(MTP3)	62.97 tok/s(91% 接受率)
decode(贪心 baseline)	28.66 tok/s
KV payload	2.41 GiB
free after startup	625 MiB

对比 20-token 上下文:MTP3 从 69.5 降到 62.97 tok/s,仅 ~9% 退化,注意力随上下文增长的成本被 int8 KV 有效控制。

抡锤者

yangkang5303

帖子