[Benchmark] 讀者問雙卡 AllReduce 吃多少 PCIe:只佔 8%,因為搬得少、跑得勤

前言

家裡網路慢的時候，你會先怪中華電信，很少有人第一時間想到是自己房間那台老舊的分享器。瓶頸的直覺幾乎都放錯位置，而且錯得很有信心。

上一篇講的是把第二張改裝 2080 Ti 接進來，用 llama.cpp 的 tensor parallel 疊上 MTP 投機解碼，同一顆 Qwen3.8-27B 從單卡 37.1 衝到雙卡 59.632 tok/s。文章發出去隔天，一位讀者寫信來，問了一件我自己完全沒想過要量的事：

I had a question, would it be possible for you run the same benchmarks while running nvidia-smi dmon -s t in another terminal? I am interested in the pcie bandwidth that is used for AllReduce.

他想知道兩張卡在對帳的時候，到底吃掉多少 PCIe 頻寬。

這是個好問題，因為多卡推論划不划算，關鍵就在兩張卡之間的溝通成本。但我第一個反應是「這題大概量不到」，而那個反應是錯的 —— 錯的方式還很有教育意義。

nvidia-smi dmon -s t 到底在看什麼

先把工具講清楚，這樣你在自己機器上也能照做。

nvidia-smi dmon 是裝置監看模式，每隔一段時間（預設一秒）印一行，一張卡一行，像 top 那樣一直刷。-s 選你要看哪一組指標：

旗標	看什麼
p	功耗與溫度
u	使用率（SM、記憶體、編解碼器）
c	運算與記憶體時脈
v	功耗與溫度的 violation
m	顯示記憶體與 BAR1
e	ECC 與 PCIe replay 錯誤
t	PCIe 收發吞吐

-s t 因此只有兩欄：rxpci（進）和 txpci（出），單位是 MB/s。

實際會用到的完整指令長這樣：

nvidia-smi dmon -s t -i 0,1 -d 1 -o T

-i 0,1 只看這兩張卡，-d 1 是取樣間隔一秒，-o T 會多印一欄時間戳 —— 最後這個很重要，等一下要拿它跟請求的起訖時間對齊，不然你分不出哪幾行是你的流量、哪幾行是別人的。

這條指令是唯讀的，不會動到任何東西，也不會拖慢跑分。

我以為這題量不到，因為我插著 NVLink

我的兩張 2080 Ti 之間有 NVLink 橋接器。nvidia-smi 認得它，狀態是 active：

$ nvidia-smi nvlink -s
GPU 0: NVIDIA GeForce RTX 2080 Ti
	 Link 0: 25.781 GB/s
	 Link 1: 25.781 GB/s

拓樸表也確認兩張卡是 NV2，也就是兩條 NVLink 綁在一起：

	GPU0	GPU1	GPU2
GPU0	 X 	NV2	NODE
GPU1	NV2	 X 	NODE

於是我的推理是：AllReduce 是兩張卡之間的溝通，NVLink 不就是為了這個而存在的嗎？那它應該走 NVLink，PCIe 計數器上會是一片安靜 —— Slava 想要的數字就量不到。

更糟的是，NVLink 那一側也看不到。GeForce 卡的 NVLink 吞吐計數器是被鎖住的：

$ nvidia-smi nvlink -gt d -i 0
GPU 0: NVIDIA GeForce RTX 2080 Ti
	 Link 0: Data Tx: N/A
	 Link 0: Data Rx: N/A

我這三張卡全部 N/A。這不是設定問題 —— NVIDIA 是逐型號決定支不支援的，沒有一條乾淨的「消費卡沒有、專業卡才有」的界線(RTX A6000 這種工作站卡就讀得到)。重點是：2080 Ti 讀不到。

照我當時的理解，這題是死的：流量在 PCIe 上看不到，在 NVLink 上也讀不到。

會這樣想其實不算離譜，你手上如果也插著橋接器，大概也會這樣推。問題是我沒有去查它到底走哪條路，只是照著「有 NVLink 就會用 NVLink」的常識往下推。

拆開原始碼：internal 這條路從頭到尾走主機記憶體

在放棄之前我去翻了 llama.cpp 的 CUDA 後端。答案就寫在 ggml/src/ggml-cuda/allreduce.cu 的開頭註解裡，而且寫得毫不含糊：

// Two reduction strategies are selected per call by tensor size:
//
//   * Chunked kernel path (small reductions): a single CUDA kernel both
//     stages data through pinned host memory and performs the local sum.
//     Cross-GPU synchronization happens *inside the kernel* (busy-wait on
//     a host-memory flag), which keeps launch overhead low for the
//     latency-sensitive token-generation case.
//
//   * Copy-engine path (large reductions): the transfer is split into
//     D2H + H2D cudaMemcpyAsync chunks driven by the GPU's copy engine,
//     followed by a small device-side add kernel.

兩條策略，兩條都經過 pinned host memory。小張量那條（也就是吐字時最常走的那條 hot path）由一顆 kernel 直接把資料搬到主機記憶體再算；大張量那條（prefill）拆成 D2H 加 H2D 兩段搬運。

大張量那條的實際程式碼是這兩行：

CUDA_CHECK(cudaMemcpyAsync(..., cudaMemcpyDeviceToHost, p->streams[i]));
CUDA_CHECK(cudaMemcpyAsync(
    p->dev_tmp[i] + offset, p->host_large[peer].host + offset, this_bytes,
    cudaMemcpyHostToDevice, p->streams[i]));

先寫下去，再從對方那份讀回來。中間那一站是主機記憶體，不是另一張卡。

而最有力的證據是同一個檔案再往下幾行的這句設計說明：

// atomicAdd_system() requires hostNativeAtomicSupported, which is unavailable
// on PCIe-attached consumer GPUs without NVLink, so the volatile path is the
// portable choice.

這條路徑是為了「沒有 NVLink 的消費卡」寫的。你有沒有插橋接器，它都走同一條。

先講清楚這篇量的是哪一條路

llama.cpp 的 AllReduce 不只一種實作。這篇從頭到尾量的是 GGML_CUDA_ALLREDUCE=internal 那條 —— 也就是上面拆出來、繞主機記憶體的那條，也是這台機器實際在跑的。

換成別的後端，行為可能完全不同，不能直接沿用下面的數字。 所以往下看之前，先確認你自己跑的是哪一條。

有一點也要講清楚，免得誤導：NVLink 並不是完全沒被啟用。llama.cpp 在初始化時確實會呼叫 cudaDeviceEnablePeerAccess，一般的跨卡張量搬運也走 cudaMemcpyPeerAsync，那些是會用到 NVLink 的。但每一層都要跑的那個 AllReduce 不走那條。橋接器不是完全沒在工作，只是每一層都要跑的那個 AllReduce 用不到它。

這麼一來，Slava 的問題不但量得到，而且 dmon -s t 剛好是最適合的工具 —— AllReduce 的每一個位元組都會出現在 PCIe 計數器上。

量一組數字，需要三個對照組

量到數字之後我自己更好奇一件事：那些流量裡，有多少真的是 AllReduce？直接開 dmon 分不出來 —— 主機與顯卡之間本來就一直有搬運。要把 AllReduce 單獨切出來，得有對照組。

所以我拆成四組，同一顆模型、同一組旗標，只差切分方式：

組	設定	有沒有 AllReduce
A	單卡 -sm none	沒有 —— 這是 PCIe 流量的最低基準
B	雙卡 -sm layer	沒有 —— 只在層邊界交接一次
C	雙卡 -sm tensor	有，每一層兩次
C'	C 再加上 MTP	有，每一層兩次

C 減 B 就是 AllReduce 的淨成本。 只看 C 的絕對值會把日常的主機搬運算進去，差值才乾淨。

模型用的是上一篇那顆 Huihui Qwen3.8-27B Q4（16.8 GB），剛好單卡塞得下，A 組才成立。四組統一 -c 16384、--parallel 1、KV 留 f16、不掛視覺模型，每組跑三次，每次量測前先靜置八秒確認基線是零。

量測期間我把線上正在跑的那個 instance 整個停掉。第一次隨手量的時候忘了這件事，結果抓到一波不是我發的一 GB/s 突波，那批數字整組不能用。

數字：AllReduce 把流量放大十幾倍，然後還是只用掉不到一成

四組各跑三次。rxpci 是每張卡的數字，不是兩張加起來：

組	decode tok/s	decode rx	prefill tok/s	prefill rx
單卡	24.1	31 MB/s	567.3	10 MB/s
雙卡 layer	25.6	22 MB/s	944.8	72 MB/s
雙卡 tensor	37.7	89 MB/s	844.9	1,045 MB/s
雙卡 tensor + MTP	44.3	146 MB/s	797.9	1,068 MB/s

把 tensor 減掉 layer，就是 AllReduce 自己的成本：

• 吐字時：每張卡 67 MB/s(tensor 的總流量是 layer 那組的 4.0 倍)
• prefill 時：每張卡 973 MB/s(總流量 14.5 倍)

放大十幾倍聽起來很兇。但這兩張卡的 PCIe 是 Gen3 x16，單向理論值 15,750 MB/s，所以：

吐字時用掉不到 1%，prefill 時用掉約 6.6%。把收發兩個方向都算進去，最忙的那一格是 8%。

計算方式：每個 rep 只取請求真正在跑的那段時間（用伺服器回報的起訖時刻對齊 dmon 的時間戳），在窗內取平均，三次之間再取中位數。

這裡用平均而不是中位數，是因為流量是爆發式的。舉個實際的例子，layer split 那組 prefill 的逐秒長這樣：13, 13, 12, 11, 13, 15, 13, 0, 1111, 10, 0, 70。看中位數只有 13 MB/s，但那個 1111 的秒真的搬了 1.1 GB。問「這條匯流排被用掉多少」要的是總搬運量除以時間，用中位數算會把爆發整個漏掉。

我第一版就是用中位數算的，得出 AllReduce 在 prefill 放大「80 倍」。那個數字是低估分母來的，錯的。

Slava 想知道 AllReduce 吃掉多少頻寬，答案是上面那些數字。但量完之後更有用的一句話是：這條匯流排連一成都沒用到。

換句話說，在這條路上，光為了雙卡去換 Gen4 主機板也換不到效能 —— 頻寬根本還沒吃滿。NVLink 橋接器同理，這條路徑不碰它。

⚠️ 但別把這句話推廣到別的 AllReduce 後端上，那不是這篇量的。

正式服役的那組組態,數字也一併量了

上面四組為了讓條件一致，context 壓到 16K、只開一個 slot、沒掛視覺模型。既然機器都已經淨空了，順手把上一篇那組真正在跑 production 的設定也照樣量一次：-sm tensor、262K context、--parallel 2、MTP n=3、掛 mmproj、KV 留 f16。

階段	tok/s	rx/卡	tx/卡	佔 Gen3 x16
decode	54.9	97–129 MB/s	244–379 MB/s	rx 0.6–0.8%，tx 1.5–2.4%
prefill	~801	1,132–1,166 MB/s	1,282–1,290 MB/s	rx 7.2–7.4%，tx 8.1–8.2%

整輪最忙的一格是 8.2%。 量的時候用的就是實際在服務的那組設定，不是為了實驗特別改過的版本，所以「PCIe 沒被吃滿」並不是只在實驗環境才看得到。

還有一個旁證可以確認這次確實跑出了同一組設定：兩張卡吃掉的顯示記憶體是 GPU0 18,890 MiB、GPU1 17,752 MiB，而上一篇發表的數字是 18,888 和 17,738。數字幾乎一模一樣。

decode 這次量到 54.9，比上一篇的 59.632 低了 8%。prompt 不同、量測窗口不同，我傾向當成兩次獨立量測而不是退步 —— 但也不打算把它凹成一致。

tensor parallel 是吐字的優化，prefill 反而變慢

這原本不在我的量測範圍內，是跑完對照組才意外看到的。

看 prefill 那一欄：最快的是雙卡 layer split 的 944.8 tok/s，不是 tensor 的 844.9。tensor parallel 在 prefill 這一段比單純的層切慢了大約 11%。

多一張卡幫忙，怎麼會反而變慢？道理其實不難懂。prefill 是一次餵一大把 token 進去，每張卡都有吃滿的工作量，本來就不缺平行度；tensor parallel 這時候多做的事只有「每一層都要對一次帳」，而那一趟就是上面量到的 1 GB/s 加上兩張卡互相等待。多付了同步成本，卻沒換到更多效能。

吐字就完全相反。一次只算一個 token，單張卡的運算單元根本吃不飽，把同一層攤開給兩張卡算才有意義 —— 所以 37.7 對 25.6，快了 47%。再疊上 MTP 就是 44.3。

所以「雙卡要用哪種切法」沒有單一答案，要看你的工作型態偏向哪一邊。 長 prompt 短回答（例如整份文件丟進去問一句話）偏 prefill，layer split 反而划算；長對話一直吐字則是 tensor parallel 的主場。

⚠️ 這裡的 44.3 tok/s 不能拿去跟上一篇的 59.632 直接比。這次為了讓四組條件一致，context 壓到 16K、只開一個 slot、prompt 也不同；上一篇是 262K context 的正式服役組態。兩次量測取的區間不同，這兩個數字不能直接比。

瓶頸不是頻寬，是往返次數

把數字換個算法會更清楚。tensor parallel 吐字時，每張卡每秒搬 89 MB，而這條匯流排每秒能搬 15,750 MB。如果卡住的真的是頻寬，速度理論上還有一百多倍的成長空間 —— 但顯然不是這樣。

時間不是花在「搬」，是花在「來回」。

每吐一個 token，模型的每一層都要做兩次 AllReduce。每一次都是：把自己算的那半寫進主機記憶體、等對方也寫完、再把對方那半讀回來。搬的位元組很少，但這個「等對方」的動作，一個 token 就要重複幾十次。

而且 llama.cpp 的小張量路徑是用 busy-wait 在等 —— kernel 裡面盯著主機記憶體上的一個旗標轉，直到對方的號碼牌變成自己等的那一個。原始碼裡的說明寫得很直白：這樣做是為了「keep launch overhead low for the latency-sensitive token-generation case」。整條路徑的設計目標從一開始就是壓低延遲，不是衝高頻寬。

真正能讓雙卡吐字更快的，是減少往返次數，不是加寬管路。MTP 就是這個道理 —— 一次前向算出好幾個 token，等於把同樣的往返成本攤到更多字上，44.3 對 37.7 那 17% 就是這樣來的。

進階

底下是完整的量測方法、原始數據，以及這次差點把對照組做壞的幾個坑。不讀不影響你照著前面的做。

四組的完整組態

四組共用這些旗標，一個字都不差：

-ngl 99 -c 16384 --parallel 1 --jinja -fa on --metrics -fit off
--chat-template-kwargs '{"enable_thinking":false}'
--temp 0.7 --top-p 0.80 --top-k 20 --min-p 0.0
--presence-penalty 1.5 --repeat-penalty 1.0

只有這裡不同：

# A  單卡
CUDA_VISIBLE_DEVICES=0                                    -sm none
# B  雙卡層切
CUDA_VISIBLE_DEVICES=0,1 GGML_CUDA_ALLREDUCE=internal      -sm layer
# C  雙卡 tensor
CUDA_VISIBLE_DEVICES=0,1 GGML_CUDA_ALLREDUCE=internal      -sm tensor
# C' 雙卡 tensor + MTP
CUDA_VISIBLE_DEVICES=0,1 GGML_CUDA_ALLREDUCE=internal      -sm tensor \
  --spec-type draft-mtp --spec-draft-n-max 3

模型是 Huihui 的 Qwen3.8-27B abliterated Q4_K（16.8 GB），KV 留 f16，不掛視覺模型。context 統一 16K —— 單卡塞不下 262K 的 KV，而 AllReduce 的搬運量只跟 hidden size 有關、跟 context 長度無關，所以壓低 context 不影響這篇的結論。

坑一：--fit 預設是開的，而 -sm tensor 不吃它

這個差一點就讓整組對照失效，而且不會有任何錯誤訊息。

llama.cpp 有個 --fit 功能，會自動調整你沒指定的參數，好讓模型塞得進顯示記憶體，預設是開的。而 -sm tensor 不支援它，開機時會直接跳過：

common_fit_params: failed to fit params to free device memory:
llama_params_fit is not implemented for SPLIT_MODE_TENSOR, abort

看出問題了嗎？A 組和 B 組會被自動調參，C 組和 C' 組不會。 四組就不再是同一組條件，而你完全看不出來 —— 數字照樣跑出來，也照樣很合理。

解法是四組全部明確寫上 -fit off。

一般化來說：做對照實驗時，任何會自動幫你調參數的旗標都要關掉。不是因為它做得不好，是因為它對不同的組做的決定不一樣。

坑二：--help 列得出 -sm row，但它已經被刪掉了

llama-server --help 目前長這樣：

-sm,   --split-mode {none,layer,row,tensor}

四個值都在。但 row 在 b10064 是死的 —— 上游 PR #24216（2026-07-06 merged，標題就叫「CUDA: remove -sm row, refactor cuBLAS」）把 CUDA 實作整段刪掉，開下去 model load 就死。上一篇有完整的死法。

我本來考慮拿 row 當第三個對照組。help 印得出來不等於有實作，這個坑值得記，因為它跟平常「查文件」的直覺是反的。

坑三：中位數把爆發丟掉了

這是我自己在分析階段踩的，前面已經提過，這裡補完整。

第一版分析我用中位數，算出 AllReduce 在 prefill 放大 80 倍。跟另一份獨立算出來的報告對照時差了將近十倍，回去看原始逐行才發現原因：

03:17:47  gpu0   13 MB/s
03:17:48  gpu0   13
03:17:49  gpu0   12
03:17:50  gpu0   11
03:17:51  gpu0   13
03:17:52  gpu0   15
03:17:54  gpu0   13
03:17:55  gpu0    0
03:17:56  gpu0 1111      ← 這一秒真的搬了 1.1 GB
03:17:57  gpu0   10
03:17:58  gpu0    0
03:17:59  gpu0   70

中位數 13，平均 110。這條匯流排在這 13 秒裡到底搬了多少？答案是照平均算的那個。中位數算出來的是「典型的一秒長怎樣」，但這篇要問的不是這個。

兩個統計量都沒有算錯 —— 是我選錯了指標，中位數回答的是另一個問題。

這批數字的限制

吐字階段的取樣數很少。dmon 最快就是一秒一筆，而每次吐字只跑三到七秒，所以每個 rep 只有 4 到 10 筆取樣。看得出來 —— C 組 GPU0 的三次分別是 51、175、以及中間值，離散得很開。

prefill 好很多，每個 rep 有 13 到 22 筆，三次之間也收斂（C 組 GPU0：1033、1176、1038）。

所以：prefill 那些數字可以拿去引用，吐字那些只能當量級參考。 要把吐字量準，得改用更細的取樣工具，dmon 的一秒解析度不夠。

不過這不影響主結論 —— 就算取吐字那組最大的一筆（C' 組 GPU1 的 209 MB/s），也只是天花板的 1.3%。

完整數據

每個 rep 在請求窗內取平均，三次取中位數。括號是三次之間的範圍。

組	階段	GPU	rx MB/s	tx MB/s	每 rep 取樣數
A	decode	0	31 (26–38)	23	10, 10, 6
A	prefill	0	10 (9–45)	6	21, 22, 22
B	decode	0	22 (21–26)	5	6, 8, 6
B	decode	1	22 (22–29)	34	6, 8, 6
B	prefill	0	33 (10–99)	3	13, 14, 13
B	prefill	1	110 (14–117)	108	13, 14, 13
C	decode	0	115 (51–175)	56	5, 4, 6
C	decode	1	63 (48–72)	50	5, 4, 6
C	prefill	0	1038 (1033–1176)	1265	14, 15, 15
C	prefill	1	1052 (1030–1087)	1240	14, 15, 15
C'	decode	0	97 (86–235)	154	4, 6, 5
C'	decode	1	195 (182–209)	245	4, 6, 5
C'	prefill	0	1071 (989–1169)	1118	15, 15, 14
C'	prefill	1	1064 (1048–1335)	1215	15, 15, 13

layer split 那一組的不對稱很明顯：GPU0 幾乎不送（tx 3–5 MB/s），GPU1 扛掉全部（tx 34–108）。那就是 pipeline 切法最明顯的特徵 —— 最後幾層跟輸出都在第二張卡上，資料是單向流過去的。tensor split 就對稱得多，兩張卡的收發都在同一個量級，因為它們每一層都在互相對帳。

NVLink 那一側，量不到

前面提過但值得單獨講：GeForce 卡的 NVLink 吞吐計數器是關的。

$ nvidia-smi nvlink -gt d -i 0
GPU 0: NVIDIA GeForce RTX 2080 Ti
	 Link 0: Data Tx: N/A
	 Link 0: Data Rx: N/A

我這三張卡全部 N/A。NVIDIA 沒有公布一條乾淨的產品線規則 —— 工作站等級的 RTX A6000 就讀得到，但 2080 Ti 讀不到。

所以如果 llama.cpp 的 AllReduce 真的走 NVLink，這篇文章就寫不出來了 —— PCIe 那側看不到，NVLink 那側讀不到，兩邊都是黑的。這題能回答，純粹是因為那條路徑不用 NVLink。

量測期間要把服役中的東西停掉

第一次隨手量的時候我忘了這件事，結果抓到這個：

02:41:41   0   963   367
02:41:41   1  1119   349
02:41:42   0  1390   316

那不是我發的請求 —— 是別的用戶端在打同一個伺服器。我自己的請求結束之後，也還有 150 MB/s 持續著。

正式量測時我把整個席位停掉，每次量之前先靜置八秒確認基線是零。四組十二次全部通過，沒有一次要作廢重跑。

如果你也在服役中的機器上量，這一步不能省。而且要驗的是「基線真的是零」，不是「我覺得應該沒人在用」。

這篇的收穫

讀者問了一個很具體的問題：AllReduce 吃掉多少 PCIe 頻寬。

量完才發現，我一開始問「吃掉多少頻寬」，其實已經先認定瓶頸在頻寬 —— 但最忙的一格也只有 8%。 用量最兇的那一格是天花板的 8%，其他都在 1% 上下。

三件可以帶走的事：

1. nvidia-smi dmon -s t 一行就能量，唯讀、不影響跑分。要對齊請求的起訖時間，不然你分不出哪些流量是自己的。
2. 雙卡切法沒有單一正解。 在這條路上，長 prompt 短回答偏 prefill，layer split 快 11%；長對話一直吐字才是 tensor parallel 的主場，快 47%。照工作型態選，不是設好之後就不管。
3. 同一項功能，背後可能是完全不同的實作。 這篇所有數字都出自 GGML_CUDA_ALLREDUCE=internal 那條路；llama.cpp 的 AllReduce 還有別的後端。看任何 benchmark 之前，先確認它量的是哪一條。

---

同系列其他文章：Part 14：兩張改裝 2080 Ti 玩 tensor parallel，Qwen3.8-27B 衝到 59.6 tok/s · Part 12：為什麼你的 4-bit 量化在 2080 Ti 上沒有比較快