摘要：GB10 上 27B dense 模型 decode 卡 15 t/s 的根因是 GGML_CUDA_ENABLE_UNIFIED_MEMORY 让访存走 CPU 跨域慢路径（162 GB/s vs 237 GB/s）。删除该环境变量（设 0 无效，须整行删除）后 +16%；配合 MTP 投机解码 –spec-draft-n-max 4 累计 +57%（15.1 → 23.7 t/s）。
目录
- 一、问题：decode 卡在 15 t/s
- 二、核心发现：GB10 有两条访存路径，性能差 32%
- 实验设计
- 结论
- 排查你的 llama-server 走的哪条路
- 三、罪魁祸首：GGML_CUDA_ENABLE_UNIFIED_MEMORY=1
- ⚠️ 陷阱：设 =0 无效！
- 修复与效果
- 四、MTP 投机解码：这台机器的正确打开方式
- 五、踩坑实录（这三条价值最高）
- 坑 1：–backend-sampling 与 MTP 严重不兼容
- 坑 2：合并改参数导致无法归因
- 坑 3：nvidia-smi 在 GB10 上不能用来监控内存
- 六、-kvu：零代价修复长会话切换
- 问题现象
- 机制
- 修复
- 实测（同款请求，零性能代价）
- 七、最终配置与完整数据
- 调优全程数据
- 并发行为（提前认知）
- 八、编译配置（GB10 专项）
- 九、天花板与突破方向
- 十、调优清单（TL;DR）
适读人群：拥有 NVIDIA GB10 / DGX Spark 并在本地跑大模型的开发者 环境：Ubuntu 24.04 (aarch64) · CUDA 13.0 · llama.cpp b11310 · Qwen3.8-27B Q4_K_M（qwen35 hybrid 架构，16 层全注意力 + 48 层 GDN） 实测环境：121GB 统一内存，48 SM Blackwell (sm_121)
一、问题：decode 卡在 15 t/s
在 GB10 上跑 27B dense 模型，decode 稳定在 15 t/s 左右。按照带宽模型估算：
每 token 读全部权重 ≈ 16.8 GB
理论 decode = 总线带宽 ÷ 16.8 GB
如果总线是标称的 273 GB/s，应该能跑 ~16 t/s（无投机）。实测 12.7 t/s（llama-bench 无投机）→ 反推带宽 213 GB/s。看起来已经"贴墙"了？
但真正的问题藏在内存分配路径里。
二、核心发现：GB10 有两条访存路径，性能差 32%
实验设计
同一进程、同一 kernel，只变分配方式（两种负载状态各测一组）：
| cudaMalloc（GPU 侧分配） | 237 GB/s | 141 GB/s |
| cudaMallocManaged（统一内存） | 162 GB/s（−32%） | 91 GB/s |
| cudaMallocHost（CPU pinned） | — | 91 GB/s |
| malloc / mmap（CPU 侧） | — | 92 GB/s |
| managed + cudaMemPrefetchAsync | 162 GB/s ❌ prefetch 无效 | — |
| managed + cudaMemAdvise | 162 GB/s ❌ advise 无效 | — |
¹ 争用 = 有推理负载同时运行（分时复用下带宽折半，但相对关系不变）
无论哪种状态，managed ≈ CPU 侧分配（同一档），cudaMalloc 始终快 40-70%。
结论
Device 路径 (cudaMalloc): GPU 本地页表直连 → 237 GB/s
Host 跨域路径 (其余全部): 走 CPU 页表翻译 + 一致性协议 → 162 GB/s
cudaMallocManaged ≠ device 内存。它是"CPU 侧内存 + GPU 可见性"——在 GB10 上它没有获得任何 device 本地性。prefetch 和 advise 都救不回来（不是页面驻留问题，是访问路径本身）。
排查你的 llama-server 走的哪条路
# 看 nvidia-smi 的进程 GPU 内存：
nvidia-smi –query-compute-apps=pid,used_memory –format=csv,noheader
# 如果你的 llama-server 只显示 ~300 MiB（实际 RSS 90+ GB）→ 走的 managed 路径！
# nvidia-smi 在 GB10 上不记账 managed 分配，这是个"假象"指标
三、罪魁祸首：GGML_CUDA_ENABLE_UNIFIED_MEMORY=1
排查启动脚本发现：
# start-llama-server.sh 中:
export GGML_CUDA_ENABLE_UNIFIED_MEMORY="${GGML_CUDA_ENABLE_UNIFIED_MEMORY:-1}"
这个变量让 GGML 所有 device 分配走 cudaMallocManaged——整个推理的访存都在慢路径上。
⚠️ 陷阱：设 =0 无效！
GGML 源码判断的是变量是否存在，不看值：
// ggml-cuda.cu
if (getenv("GGML_CUDA_ENABLE_UNIFIED_MEMORY") != nullptr) {
err = cudaMallocManaged(ptr, size); // 任何值（包括"0"）都走这里
}
唯一正确做法：从脚本中整行删除。
修复与效果
# 1. 注释/删除启动脚本中的 export 行
# 2. 重启
./scripts/local/stop-llama-server.sh
./scripts/local/start-llama-server.sh
| tg（投机 decode） | 15.1 t/s | 17.5–18.0 t/s | +16~19% |
| T_step | 194 ms | 149 ms | −23% |
| PP（prefill） | 292 t/s | 309 t/s | +6% |
| GPU 功耗（忙态中位） | 47.4 W | 51.5 W | 计算密度↑（生效证据） |
| nvidia-smi 记账 | 305 MiB（假象） | 58 GB（真实） | 口径切换 |
验证脚本的核心逻辑（可直接复用）：
# 确认环境变量已移除
tr '\\0' '\\n' < /proc/$(pgrep -f llama-server)/environ | grep GGML_CUDA
# 确认 nvidia-smi 开始记账（>10GB）
nvidia-smi –query-compute-apps=used_memory –format=csv,noheader
四、MTP 投机解码：这台机器的正确打开方式
Qwen3.5 架构系列自带 MTP（Multi-Token Prediction）头，–spec-type draft-mtp 开箱即用：
一次 verify 读全部权重（16.8GB），产出 mean_len 个 token
tg = mean_len / T_step
实测健康指标：接受率 58–64%，mean_len 2.88–2.93。
默认 n_max=3 是个隐形截断（mean_len 上限 = n_max+1 = 4.00，实测 max 恰好 4.00）。
–spec-draft-n-max 4 实测显著提升（+15.6%）：
| tg（写作负载） | 20.5 t/s | 23.7 t/s | +15.6% |
| mean_len | 2.88 | 3.70 | +28% |
| 接受率 | ~0.60 | 0.67 | 更高 |
机制：写作上下文的 MTP 逐位接受率远比预期平缓（~0.90/0.85/0.80/0.75）， 第 4 位 draft 有 75% 接受率——加一位几乎是白送的。
⚠️ 测试方法坑：一次性短 prompt（如"写一个故事"）接受率天然低（~23%）， 不能代表真实写作负载。判定必须用真实多轮会话数据—— 我们最初的短 prompt 测试显示"下降"，写作负载实测才发现是 +15.6%。
n_max=7 可能再 +10%（逐位外推 mean_len ~4.4），未实测。
五、踩坑实录（这三条价值最高）
坑 1：–backend-sampling 与 MTP 严重不兼容
把采样搬到 GPU 听起来对（CPU 主线程 100% → 17%），实测灾难：
| tg | 17.3 t/s | 7.4 t/s（−57%） |
| PP | 155 t/s | 20 t/s（−87%） |
原因：MTP 每步做多次小采样（draft×n + verify + accept），GPU 采样路径引入每次同步往返，串行链放大 T_step。该特性标记 experimental，与 MTP 组合未经充分验证——不要在生产环境用。
坑 2：合并改参数导致无法归因
一次同时加 –backend-sampling -kvu → 出问题后分不清谁的责任，多花两次重启才定位。每个参数单独 A/B 验证再叠加。
坑 3：nvidia-smi 在 GB10 上不能用来监控内存
- managed 分配不记账（显示 305 MiB 假象）
- memory.total/used = N/A（无独立显存）
- 正确口径：进程 RSS（ps -o rss）或系统 free
六、-kvu：零代价修复长会话切换
问题现象
prompt state size 34826 MiB exceeds cache size limit 16384, skipping
长会话（100K+ token）切换时 state 存不下 → 完整重新 prefill（6-8 分钟）。93 次跳过。
机制
kv_unified=false 时 state 序列化 ≈ 68 KB/token × n_stream(=np)。np=3 时单满 slot state ~40GB，16GB cache 装不下。
修复
./scripts/local/start-llama-server.sh -kvu –cache-ram 24576 # 或 32768（内存富余时）
实测（同款请求，零性能代价）
| tg | 17.3 | 17.5（一致） |
| PP | 155 | 155.0（一致） |
| cache skip | 93 次 | 2 次（−98%） |
附带收益：slot 从硬均分 192K 变为共享池弹性（单序列可达模型原生 262K）。
七、最终配置与完整数据
#!/bin/bash
# GB10 llama.cpp 最优启动配置（实测验证）
# 1. 关键：不要设置 GGML_CUDA_ENABLE_UNIFIED_MEMORY（删掉 export 行！）
./build/bin/llama-server \\
-m /path/to/Qwen3.8-27B-Q4_K_M.gguf \\
–host 127.0.0.1 –port 8080 \\
-ngl 99 \\ # 全层上 GPU（CPU 侧带宽仅 76 GB/s）
-c 589824 -np 3 \\ # 3 槽共享 576K 池
-fa on \\ # Flash Attention
–spec-type draft-mtp \\ # MTP 投机（模型自带头）
–spec-draft-n-max 4 \\ # 加深一位 draft（+15%）
-kvu \\ # 统一 KV 池（cache 修复）
–cache-ram 24576 \\ # 24GB state 缓存
–jinja \\
–mmproj /path/to/mmproj.gguf
调优全程数据
| 原始（managed 路径） | 15.1 t/s | — | 基线 |
| +删环境变量（Device 路径） | 17.5 t/s | +16% | 访存路径切换 |
| +kvu（cache 修复） | 17.5 t/s | +16% | 零代价，skip −98% |
| +n_max=4（draft 加深） | 23.7 t/s | +57% | mean_len 2.88→3.70 |
以上均为写作负载统计中位（108 样本基线 → 3 样本 n_max=4）。 n_max=4 后的 tg_3s 峰值可达 27+ t/s。
并发行为（提前认知）
| 1 | 17.5 | 17.5 |
| 2 | ~11 | ~22（+26%） |
| 3 | ~7-8 | ~23 |
单会话慢 40% 是带宽共享的物理必然；总吞吐反而提升。延迟敏感就串行，吞吐敏感就并发。
八、编译配置（GB10 专项）
cmake -B build \\
-DGGML_CUDA=ON \\
-DCMAKE_CUDA_ARCHITECTURES=121a-real \\ # GB10 专属 (sm_121)
-DGGML_CPU_ARM_ARCH=armv9.2-a+dotprod+i8mm+bf16+sve2 \\
-DGGML_CUDA_GRAPHS=ON \\
-DCMAKE_BUILD_TYPE=Release
陷阱：GGML_NATIVE=ON 在此内核会退化为 -march=armv8-a（GCC 不识别 X925），dotprod/i8mm/sve 全关。必须显式指定 ARM_ARCH。
九、天花板与突破方向
dense 27B 在 GB10 的物理上限 ≈ 25-26 t/s（单序列投机）
↳ 带宽 237 GB/s ÷ 16.8 GB/token × mean_len
↳ vLLM 在同机型同模型收敛到同样区间（18-25 t/s）
要 2-3×：只有 MoE
↳ Qwen3.8-Flash-Next (125B/6B active): 每token读 ~3.5GB → 理论 60+ t/s
↳ 社区 GGUF 已有（IQ3_M 档 89.5GB 可装入 121GB）
↳ llama.cpp 已有 qwen4exp 架构支持
↳ 或 Nemotron-3-Nano (NVFP4 + vLLM)：论坛实测 56 t/s
十、调优清单（TL;DR）
✅ 必做（零成本，累计 +57%）:
1. 删除 GGML_CUDA_ENABLE_UNIFIED_MEMORY（整行！设0无效）→ +16%
2. -ngl 99（全层上 GPU）
3. –spec-type draft-mtp（模型自带头，开箱即用）
4. –spec-draft-n-max 4（加深 draft 一位）→ +15%
5. -kvu –cache-ram 24576（长会话切换修复，零代价）
6. -fa on
❌ 不要做:
1. –backend-sampling（与 MTP 灾难性不兼容）
2. 期望降 -c 提速（T_step 与上下文长度无关）
3. 用 nvidia-smi 监控 GB10 内存（用 RSS）
4. 设 GGML_CUDA_ENABLE_UNIFIED_MEMORY=0（等于开启）
⬜ 可选（待实测）:
1. –spec-draft-n-max 7（逐位外推或再 +10%，verify batch=8 仍在 GEMV）
2. 换 MoE 模型（2-3× 质变——唯一真正的数量级提升）
附录：完整启动脚本（脱敏版，基于源码构建）
环境：llama.cpp b11310+ (commit 0c1e57098) · CUDA 13.0 · GNU 13.3.0 · aarch64
1. 源码构建
git clone https://github.com/ggml-org/llama.cpp
cd llama.cpp
# GB10 专属编译参数
cmake -B build \\
-DGGML_CUDA=ON \\
-DCMAKE_CUDA_ARCHITECTURES=121a-real \\
-DGGML_CPU_ARM_ARCH=armv9.2-a+dotprod+i8mm+bf16+sve2 \\
-DGGML_CUDA_GRAPHS=ON \\
-DCMAKE_BUILD_TYPE=Release
ninja -C build llama-server
# 验证
./build/bin/llama-server –version
# 预期: version 0.5.0-dev (build 11312), built with GNU 13.3.0 for Linux aarch64
2. 启动脚本（88 行，含内存预检）
#!/usr/bin/env bash
# GB10/DGX Spark llama.cpp 启动脚本（实测最优配置，+55%）
#
# ⚠️ 核心优化: 不要设置 GGML_CUDA_ENABLE_UNIFIED_MEMORY（见文章 §3）
# GGML 只判断变量存在性（getenv != nullptr），设 =0 也等于开启！
set -euo pipefail
# ===== 用户配置 =====
MODELS_DIR="${MODELS_DIR:-$HOME/models}" # 模型目录
BUILD_DIR="${BUILD_DIR:-$(dirname "$0")/../build}"
LOG="${LOG:-$MODELS_DIR/llama-server.log}"
PORT="${PORT:-8080}"
# ===== 推理参数（实测最优组合）=====
NG=99 # 全层上 GPU（CPU 侧带宽仅 76 GB/s）
NP=3 # 并发槽数
CTX_PER_SLOT=192 # 每槽 192K → 总池 576K
CACHE_RAM=24576 # prompt state 缓存 24GB（配 -kvu）
SPEC_DRAFT_NMAX=4 # MTP draft 深度（实测 +14%）
# ===== 模型发现（自动选取唯一主模型）=====
mapfile -t models < <(find "$MODELS_DIR" -maxdepth 1 -name '*.gguf' ! -iname '*mmproj*' | sort)
if [[ ${#models[@]} -ne 1 ]]; then
echo "请指定模型（$MODELS_DIR 下有 ${#models[@]} 个）:" >&2
printf ' %s\\n' "${models[@]}" >&2; exit 1
fi
MODEL="${models[0]}"
# 视觉投影器（如果有）
MMPROJ_FILE=""
mapfile -t mm < <(find "$MODELS_DIR" -maxdepth 1 -iname '*mmproj*.gguf' 2>/dev/null | sort)
[[ ${#mm[@]} -eq 1 ]] && MMPROJ_FILE="${mm[0]}"
# ===== KV 内存预检 =====
KV_PER_TOK_KB=64 # 16层 full-attn × 4 KV头 × 256 dim × f16
CTX=$((NP * CTX_PER_SLOT * 1024))
KV_GB=$((CTX * KV_PER_TOK_KB / 1024 / 1024))
AVAIL_GB=$(awk '/MemAvailable/{print int($2/1024/1024)}' /proc/meminfo)
WEIGHTS_GB=$(du -g "$MODEL" | cut -f1)
echo "模型: $MODEL"
echo "配置: -ngl $NG -c $CTX -np $NP -fa on –jinja –spec-type draft-mtp"
echo "KV: ~${KV_GB}GB | 权重: ${WEIGHTS_GB}GB | 需求: ~$((KV_GB + WEIGHTS_GB + 6))GB | 可用: ${AVAIL_GB}GB"
if (( KV_GB + WEIGHTS_GB + 6 + 12 > AVAIL_GB )); then
echo "⚠️ 内存不足（建议留 12GB 系统余量）: 降 NP 或 CTX_PER_SLOT" >&2
exit 1
fi
# ===== 启动（注意: 无 GGML_CUDA_ENABLE_UNIFIED_MEMORY！）=====
BIN="$BUILD_DIR/bin/llama-server"
[[ -x "$BIN" ]] || { echo "找不到 $BIN，先构建" >&2; exit 1; }
MMPROJ_ARGS=()
[[ -n "$MMPROJ_FILE" ]] && MMPROJ_ARGS=(–mmproj "$MMPROJ_FILE")
nohup "$BIN" \\
-m "$MODEL" \\
–host 127.0.0.1 –port "$PORT" \\
-ngl "$NG" -c "$CTX" -np "$NP" \\
-fa on –jinja \\
–spec-type draft-mtp \\
–spec-draft-n-max "$SPEC_DRAFT_NMAX" \\
-kvu \\
–cache-ram "$CACHE_RAM" \\
"${MMPROJ_ARGS[@]}" \\
>> "$LOG" 2>&1 &
PID=$!
echo "$PID" > "$BUILD_DIR/llama-server.pid"
# 等就绪（大模型加载 1-5 分钟）
for _ in $(seq 1 600); do
kill -0 "$PID" 2>/dev/null || { echo "启动失败"; tail -20 "$LOG"; exit 1; }
curl -fsS "http://127.0.0.1:$PORT/health" >/dev/null 2>&1 && {
echo "✅ 就绪 (pid $PID): http://127.0.0.1:$PORT"
exit 0
}
sleep 1
done
echo "超时"; exit 1
3. 停止脚本
#!/usr/bin/env bash
PIDFILE="$(dirname "$0")/../build/llama-server.pid"
[[ -f "$PIDFILE" ]] && kill "$(cat "$PIDFILE")" && rm "$PIDFILE"
4. 快速验证
# 环境检查（决定性——确认无 unified memory 变量）
tr '\\0' '\\n' < /proc/$(pgrep -f llama-server)/environ | grep GGML_CUDA
# 预期: 无输出（有输出 = 优化未生效！）
# 记账检查
nvidia-smi –query-compute-apps=used_memory –format=csv,noheader
# 预期: >10GB（显示 ~300 MiB = 走了 managed 慢路径）
# 性能冒烟
curl -s http://127.0.0.1:8080/v1/chat/completions \\
-d '{"messages":[{"role":"user","content":"你好"}],"max_tokens":50}' \\
| python3 -c "import sys,json; print(json.load(sys.stdin)['timings'])"
