跳转至内容
  • 版块
  • 最新
  • 标签
  • 热门
  • 用户
  • 群组
皮肤
  • 浅色
  • Brite
  • Cerulean
  • Cosmo
  • Flatly
  • Journal
  • Litera
  • Lumen
  • Lux
  • Materia
  • Minty
  • Morph
  • Pulse
  • Sandstone
  • Simplex
  • Sketchy
  • Spacelab
  • United
  • Yeti
  • Zephyr
  • 深色
  • Cyborg
  • Darkly
  • Quartz
  • Slate
  • Solar
  • Superhero
  • Vapor

  • 默认(LCZ-Blue)
  • 不使用皮肤
  • LCZ-Green
  • LCZ-Blue
  • LCZ-Black
折叠
品牌标识

抡锤者

首页 版块 标签
硬件
装备秀
AI 广场 App
yangkang5303Y

yangkang5303

@yangkang5303
体面乡绅
取消关注 关注
关于
帖子
7
主题
3
分享
0
群组
1
粉丝
0
关注
0

帖子

最新 最佳 有争议的

  • 实测报告:RTX 3080 20GB (GA102) 解锁前后 cuBLAS GEMM 对比
    yangkang5303Y yangkang5303
    AI硬件 rtx3080 nvidia

    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

  • 在 魔改 3080 20G上,运行当下流行的ninfer-3090代码,峰值70t/s,ctx70k
    yangkang5303Y yangkang5303
    LLM讨论区 rtx3080 cuda mtp

    对代码有一定的修改以适配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 有效控制。


  • [实测] RTX 3080 (魔改20G) 跑通 Ternary Bonsai 2 27B + rk4v4,256K 上下文,1000+ prefill 65t/s decode
    yangkang5303Y yangkang5303
    LLM讨论区 rtx3080 ninfer 量化

    在 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 是一件"开箱即用"的事:

    1. ninfer-all 的默认构建目标就是 86 —— 不需要 porting
    2. rk4v4 等 9 种 KV 格式上游全部实现,sm_86 不再被拒
    3. 混合注意力架构(只有 16/64 层占 KV)+ rk4v4,让 20GB 卡能开 256K 上下文
    4. --kv-headroom-mib 让桌面和模型和平共处
    5. 实测 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
  • 登录

  • 登录或注册以进行搜索。
  • 第一个帖子
    最后一个帖子
0
  • 版块
  • 最新
  • 标签
  • 热门
  • 用户
  • 群组