雙 7900XTX + SGLang 跑 Qwen3.8-27B 完整實戰：TP=2 由 ROCm 7.2 到 HiCache，decode 78 tok/s（附全套 benchmark）
-
跟住 https://lcz.me/topic/1532 同 https://lcz.me/topic/1567 兩位前輩嘅路線，我喺自己部機由零行咗一次完整 setup，順便做咗全套 benchmark 同 HiCache 逐出/還原實測。有幾個新發現（尤其 HiCache 嗰 part）想還返俾社群。
一、硬件
- CPU：EPYC 7K62 48C（洋垃圾）
- 兩張 Sapphire PULSE 7900 XTX 24G，PCIe 4.0 x16 + x16（lspci LnkSta 實測）— EPYC lanes 夠，TP all-reduce 冇瓶頸，比 X670E x8/x8 爽
- 125GB RAM、Ubuntu 24.04.4、kernel 7.0.0-31
- 機入面另外兩張 4080 SUPER 行緊另一個 FP8 instance，兩邊並行唔相撞（ROCm 只見 AMD 卡）
二、軟件棧同安裝
- ROCm 7.2.0 userspace：kernel 7.0 內核 amdgpu 直接用，唔使 dkms。注意 30.x 版 amdgpu-install 參數要等號：
sudo amdgpu-install -y --usecase=rocm --no-dkms --no-32- Python 3.12 venv + torch 鎖死 2.11.0+rocm7.2（whl/rocm7.2 index），任何 pip 操作之後 check 返 torch.version.hip 有值
- engine：StevenChenSE/sglang
gfx1100-support分支（f84475c），照 README：setup_rocm.py 編 19 個 kernels，成個 aot/python/sgl_kernel 目錄 copy 入 site-packages（要有 init.py 先至有 gptq_gemm） - 自訂 all-reduce：scripts/rdna_ar/rdna_ar_ext.py + make（kfd_event_age_fix.so 用 LD_PRELOAD）
三、兩個必改嘅源碼 patch（今次分享重點之一）
Patch 1：all-reduce wrapper 入面 copy_sys_kernel 用咗 CDNA3 先有嘅
__builtin_amdgcn_global_store_b128，gfx1100 編唔過（要 gfx940-insts）。改做 RDNA3 安全嘅 128-bit volatile store（方向同 farmer-node 喺 https://lcz.me/topic/1640 講嘅一致）：// scripts/rdna_ar/rdna_custom_all_reduce.cu - __builtin_amdgcn_global_store_b128((sgl_v4u_gptr)(dst + i), v, ""); + *(volatile sgl_v4u_gptr)(dst + i) = v; // RDNA3: no b128 builtinPatch 2：custom_all_reduce_hip.cuh 入面 RankData 嘅 const void* restrict ptrs[8] 會令 clang 生成隱式 copy assignment 失敗（std::fill 需要），剷走 restrict 就得（純優化提示，零語義影響）：
- const void* __restrict__ ptrs[8]; + const void* ptrs[8];改完用 PYTORCH_ROCM_ARCH=gfx1100 行 rdna_ar_ext.py 一發過。
四、模型
Vishva007/Qwen3.8-27B-W4A16-AutoRound-GPTQ（~19G，MTP 保持 bf16）。記得改 config.json：將 dynamic quantization rules 入面兩條
+:.*mtp.*規則換成一條排除規則-:.*mtp.*（bits 16、group 128），唔係嘅話 GPTQ loader 會喺 MTP tensor 上面炸。五、生產啟動配置（全套）
export SGL_DTYPE=bfloat16 SGLANG_RDNA_CUSTOM_AR=1 SGL_RDNA_NO_FUSED=1 \ SGL_RDNA_GEMMA_TRITON=1 SGL_RDNA_VLLM_VERIFY=1 export LD_PRELOAD=~/sglang-gfx1100/sglang/scripts/rdna_ar/vendor/kfd_event_age_fix.so:$LD_PRELOAD python3 -m sglang.launch_server \ --model-path ~/models/Qwen3.8-27B-W4A16 --host 0.0.0.0 --port 30001 \ --tp-size 2 --quantization gptq --dtype bfloat16 \ --mamba-ssm-dtype bfloat16 --kv-cache-dtype auto \ --attention-backend triton --triton-attention-num-kv-splits 16 \ --context-length 196608 --mem-fraction-static 0.91 \ --max-running-requests 4 --max-mamba-cache-size 20 \ --speculative-algorithm NEXTN --speculative-num-steps 3 \ --speculative-eagle-topk 1 --speculative-num-draft-tokens 4 \ --cuda-graph-bs-decode 1 2 4 --chunked-prefill-size 2048 \ --page-size 64 --schedule-policy lpm --sleep-on-idle \ --reasoning-parser qwen3 --tool-call-parser qwen3_coder --enable-metrics \ --enable-hierarchical-cache --hicache-ratio 2 \ --hicache-io-backend direct --hicache-write-policy write_through \ --default-chat-template-kwargs '{"reasoning_effort": "medium"}'KV pool 257,024 tokens（bf16），HiCache host pool 再加 514K（2×17.9G RAM + 1.6G mamba host）。
六、Benchmark（2026-09-15，temp 0，串流 + usage 計數）
Decode vs 深度（MTP-3 單流）： 0 → 78.4 tok/s（12.7ms/token）；16K → 71.7；32K → 64.9；65K → 61.3；131K → 42.0（23.8ms）
併發： 1u 67.6 ／ 2u 合計 118.9 ／ 4u 合計 176 tok/s
Prefill（冷啟動）： 4K 峰值 1,869 tok/s；16K 1,715；32K 1,477（TTFT 22s）；65K 1,136；130K 777。4×32K 同時只有 1,469 — prefill 係計算飽和，排隊唔會快。
快取： GPU radix 32K 冷 22s → 暖 0.41s（53.6×）；L2 還原見下節。
MTP： accept length 2.725（/metrics 嘅 spec_accept_length）。
同 2×4080S FP8 同日同方法對比： decode 78.4 vs 71.3（XTX +10%）；32K prefill 1,477 vs 1,914（4080S +29%）。同 #1532/#1640 嘅格局一致：XTX decode 勁、prefill 輸 NVIDIA 一截。
七、HiCache 實測（新發現）
背景：fork 對混合架構（48 層 Gated DeltaNet + 16 層 full attention）有寫 mamba host cache，但作者都話未測過。我嘅逐出測試（載入 A → 灌爆 257K pool 逼 A 被逐出 → 再問 A）：
- 備份管線通：write-through 持續 backup（hicache_backup_tokens_total 上升，零 drop）
- 純重複 prompt 觸發唔到還原（4.2s 全量 re-prefill）— 因為 mamba state 只會喺**請求完結點、prefill chunk 邊界（每 2048）、decode 追蹤點（每 256）**捐入 radix tree
- 續接形態（下一輪 = 上一輪完整序列 + 新內容，即 agent 天然用法）完美還原：7.6K context 被逐出後 0.21s 由 RAM 還原（20×）；chat template 重 render 有時令 token 序列喺 thinking 位分歧，退落最近 2048-grid checkpoint 都有 2.66s（仍慳 37%）
- 分支形態（同 doc 唔同問題）只有部分著數 — 呢個先係真正嘅 hybrid 缺口（對應上游 sgl-project/sglang#12826）
結論：開住佢。有效容量 257K GPU + 514K host ≈ 771K，長 session agent 逐出後秒級恢復。io-backend 一定要 direct（kernel backend 對 hybrid mamba 會炸，sglang #24121）。
八、陷阱清單
- 千祈唔好數 SSE chunk 做 benchmark：MTP 下一個 chunk 載 2.6 個 token（= accept length），數 chunk 會低估 2.6 倍，要用 usage 計數
- 唔好 pkill -f sglang.launch_server — 會撈埋 Docker 入面其他 instance，用 port 搵 PID 精準殺
- 永遠唔好 rocm-smi --gpureset（gfx1100 鎖 PCIe root port，要冷開機）
- pip 依賴要手動補：orjson、gguf、sentencepiece、dill、xgrammar、python-multipart、prometheus_client、openai、apache-tvm-ffi==0.1.11…
- fork 冇 --max-consecutive-prefill-batches（上游 PR #34058 未併入），併發 prefill 打斷 decode 嘅卡頓暫時無解
致謝
StevenChenSE 嘅 fork、farmer-node 嘅 all-reduce 修法方向、以及 #1532 #1567 #1640 #1674 嘅前輩。mamba 分支還原嗰 part 如果有人焗到更深，歡迎交流。
-
跟住 https://lcz.me/topic/1532 同 https://lcz.me/topic/1567 兩位前輩嘅路線，我喺自己部機由零行咗一次完整 setup，順便做咗全套 benchmark 同 HiCache 逐出/還原實測。有幾個新發現（尤其 HiCache 嗰 part）想還返俾社群。
一、硬件
- CPU：EPYC 7K62 48C（洋垃圾）
- 兩張 Sapphire PULSE 7900 XTX 24G，PCIe 4.0 x16 + x16（lspci LnkSta 實測）— EPYC lanes 夠，TP all-reduce 冇瓶頸，比 X670E x8/x8 爽
- 125GB RAM、Ubuntu 24.04.4、kernel 7.0.0-31
- 機入面另外兩張 4080 SUPER 行緊另一個 FP8 instance，兩邊並行唔相撞（ROCm 只見 AMD 卡）
二、軟件棧同安裝
- ROCm 7.2.0 userspace：kernel 7.0 內核 amdgpu 直接用，唔使 dkms。注意 30.x 版 amdgpu-install 參數要等號：
sudo amdgpu-install -y --usecase=rocm --no-dkms --no-32- Python 3.12 venv + torch 鎖死 2.11.0+rocm7.2（whl/rocm7.2 index），任何 pip 操作之後 check 返 torch.version.hip 有值
- engine：StevenChenSE/sglang
gfx1100-support分支（f84475c），照 README：setup_rocm.py 編 19 個 kernels，成個 aot/python/sgl_kernel 目錄 copy 入 site-packages（要有 init.py 先至有 gptq_gemm） - 自訂 all-reduce：scripts/rdna_ar/rdna_ar_ext.py + make（kfd_event_age_fix.so 用 LD_PRELOAD）
三、兩個必改嘅源碼 patch（今次分享重點之一）
Patch 1：all-reduce wrapper 入面 copy_sys_kernel 用咗 CDNA3 先有嘅
__builtin_amdgcn_global_store_b128，gfx1100 編唔過（要 gfx940-insts）。改做 RDNA3 安全嘅 128-bit volatile store（方向同 farmer-node 喺 https://lcz.me/topic/1640 講嘅一致）：// scripts/rdna_ar/rdna_custom_all_reduce.cu - __builtin_amdgcn_global_store_b128((sgl_v4u_gptr)(dst + i), v, ""); + *(volatile sgl_v4u_gptr)(dst + i) = v; // RDNA3: no b128 builtinPatch 2：custom_all_reduce_hip.cuh 入面 RankData 嘅 const void* restrict ptrs[8] 會令 clang 生成隱式 copy assignment 失敗（std::fill 需要），剷走 restrict 就得（純優化提示，零語義影響）：
- const void* __restrict__ ptrs[8]; + const void* ptrs[8];改完用 PYTORCH_ROCM_ARCH=gfx1100 行 rdna_ar_ext.py 一發過。
四、模型
Vishva007/Qwen3.8-27B-W4A16-AutoRound-GPTQ（~19G，MTP 保持 bf16）。記得改 config.json：將 dynamic quantization rules 入面兩條
+:.*mtp.*規則換成一條排除規則-:.*mtp.*（bits 16、group 128），唔係嘅話 GPTQ loader 會喺 MTP tensor 上面炸。五、生產啟動配置（全套）
export SGL_DTYPE=bfloat16 SGLANG_RDNA_CUSTOM_AR=1 SGL_RDNA_NO_FUSED=1 \ SGL_RDNA_GEMMA_TRITON=1 SGL_RDNA_VLLM_VERIFY=1 export LD_PRELOAD=~/sglang-gfx1100/sglang/scripts/rdna_ar/vendor/kfd_event_age_fix.so:$LD_PRELOAD python3 -m sglang.launch_server \ --model-path ~/models/Qwen3.8-27B-W4A16 --host 0.0.0.0 --port 30001 \ --tp-size 2 --quantization gptq --dtype bfloat16 \ --mamba-ssm-dtype bfloat16 --kv-cache-dtype auto \ --attention-backend triton --triton-attention-num-kv-splits 16 \ --context-length 196608 --mem-fraction-static 0.91 \ --max-running-requests 4 --max-mamba-cache-size 20 \ --speculative-algorithm NEXTN --speculative-num-steps 3 \ --speculative-eagle-topk 1 --speculative-num-draft-tokens 4 \ --cuda-graph-bs-decode 1 2 4 --chunked-prefill-size 2048 \ --page-size 64 --schedule-policy lpm --sleep-on-idle \ --reasoning-parser qwen3 --tool-call-parser qwen3_coder --enable-metrics \ --enable-hierarchical-cache --hicache-ratio 2 \ --hicache-io-backend direct --hicache-write-policy write_through \ --default-chat-template-kwargs '{"reasoning_effort": "medium"}'KV pool 257,024 tokens（bf16），HiCache host pool 再加 514K（2×17.9G RAM + 1.6G mamba host）。
六、Benchmark（2026-09-15，temp 0，串流 + usage 計數）
Decode vs 深度（MTP-3 單流）： 0 → 78.4 tok/s（12.7ms/token）；16K → 71.7；32K → 64.9；65K → 61.3；131K → 42.0（23.8ms）
併發： 1u 67.6 ／ 2u 合計 118.9 ／ 4u 合計 176 tok/s
Prefill（冷啟動）： 4K 峰值 1,869 tok/s；16K 1,715；32K 1,477（TTFT 22s）；65K 1,136；130K 777。4×32K 同時只有 1,469 — prefill 係計算飽和，排隊唔會快。
快取： GPU radix 32K 冷 22s → 暖 0.41s（53.6×）；L2 還原見下節。
MTP： accept length 2.725（/metrics 嘅 spec_accept_length）。
同 2×4080S FP8 同日同方法對比： decode 78.4 vs 71.3（XTX +10%）；32K prefill 1,477 vs 1,914（4080S +29%）。同 #1532/#1640 嘅格局一致：XTX decode 勁、prefill 輸 NVIDIA 一截。
七、HiCache 實測（新發現）
背景：fork 對混合架構（48 層 Gated DeltaNet + 16 層 full attention）有寫 mamba host cache，但作者都話未測過。我嘅逐出測試（載入 A → 灌爆 257K pool 逼 A 被逐出 → 再問 A）：
- 備份管線通：write-through 持續 backup（hicache_backup_tokens_total 上升，零 drop）
- 純重複 prompt 觸發唔到還原（4.2s 全量 re-prefill）— 因為 mamba state 只會喺**請求完結點、prefill chunk 邊界（每 2048）、decode 追蹤點（每 256）**捐入 radix tree
- 續接形態（下一輪 = 上一輪完整序列 + 新內容，即 agent 天然用法）完美還原：7.6K context 被逐出後 0.21s 由 RAM 還原（20×）；chat template 重 render 有時令 token 序列喺 thinking 位分歧，退落最近 2048-grid checkpoint 都有 2.66s（仍慳 37%）
- 分支形態（同 doc 唔同問題）只有部分著數 — 呢個先係真正嘅 hybrid 缺口（對應上游 sgl-project/sglang#12826）
結論：開住佢。有效容量 257K GPU + 514K host ≈ 771K，長 session agent 逐出後秒級恢復。io-backend 一定要 direct（kernel backend 對 hybrid mamba 會炸，sglang #24121）。
八、陷阱清單
- 千祈唔好數 SSE chunk 做 benchmark：MTP 下一個 chunk 載 2.6 個 token（= accept length），數 chunk 會低估 2.6 倍，要用 usage 計數
- 唔好 pkill -f sglang.launch_server — 會撈埋 Docker 入面其他 instance，用 port 搵 PID 精準殺
- 永遠唔好 rocm-smi --gpureset（gfx1100 鎖 PCIe root port，要冷開機）
- pip 依賴要手動補：orjson、gguf、sentencepiece、dill、xgrammar、python-multipart、prometheus_client、openai、apache-tvm-ffi==0.1.11…
- fork 冇 --max-consecutive-prefill-batches（上游 PR #34058 未併入），併發 prefill 打斷 decode 嘅卡頓暫時無解
致謝
StevenChenSE 嘅 fork、farmer-node 嘅 all-reduce 修法方向、以及 #1532 #1567 #1640 #1674 嘅前輩。mamba 分支還原嗰 part 如果有人焗到更深，歡迎交流。
干货很足，尤其 HiCache 那段。补几条：
-
混合架构的 checkpoint 粒度是理解「純重複 prompt 還原唔到」的关键：48 层 GDN 的 mamba state 只在请求结束点、prefill chunk 边界、decode 追踪点落 radix tree。纯重复 prompt 若 token 串和上次完全一致，本应命中 GPU radix；你看到全量 re-prefill，多半是命中节点上没挂 mamba state，或 chat template 重 render 让 thinking 段 token 分叉。建议打日志看实际 match length，再决定是否在序列末尾加一个尾 token 把它变成「续接」形态——agent 负载本来就是续接，所以能用。
-
L2 还原本建议按命中长度分层报。短前缀走 GPU radix（0.41s 那个），真正有价值的是 host pool 还原本。7.6K 的 0.21s 要标清是纯 host 命中还是 GPU+host 混合，否则和 2.66s 的 2048-grid 回退不好比。
-
write_through 建议和 write_back 对一下首轮 prefill 吞吐，长 prefill 上写放大能吃回你 1,477 tok/s 的一部分。direct backend 对 hybrid 是必须的，这点同意。
-
把 MTP accept length 2.725 和「不要数 SSE chunk」写在一起很关键，多数人 benchmark 就死在这。78.4 是 MTP 生效后的端到端，和单 forward 对比时要注明。
问一个：L2 还原到卡的路径是 pinned + PCIe 拷回，还是走了别的通道？EPYC 7K62 八通道 DDR4 带宽不缺，0.21s 是不是已经贴到 PCIe 上限，量一下 host 读带宽就知道还有没有空间。
4080S prefill 领先 29% 不意外（FP8 + tile 调度），XTX 靠 decode 和显存容量打，方向是对的。
-
@Geekyang 我放在辦公室
-
@franklee006 我也是，放办公室。满载的时候我测试过，整机功耗992W，真是高。
@懒人烘培 我設了300W 單卡上限
-
，
T terry 固定了此主题
-
，系统 取消固定了此主题
