MPI Collective Communication 與 CUDA-aware MPI

#advanced-practice #mpi-cluster #collective-communication #cuda-aware-mpi #barrier-synchronization

重點總覽 (Overview)

本章收尾涵蓋兩個主題：Collective communication（群體通訊原語）與 CUDA-aware MPI（讓 MPI 直接讀寫 GPU 記憶體）。前者把「多對多」的協作抽象成單一最佳化函式；後者移除 halo 交換時的 host staging copy，簡化主機端程式碼。

主題	關鍵函式 / 機制	解決的問題	代價 / 限制
Barrier	`MPI_Barrier`	全體同步，確保大家都就緒再交換資料	最慢的 process 拖累全體
Broadcast	`MPI_Bcast`	root 把同一份資料送給所有 process	需指定 root rank
Reduce	`MPI_Reduce`	把各 process 的值用運算子 (sum/max…) 匯總到 root	只有 root 拿到結果
Gather	`MPI_Gather`	把各 process 的分段資料收集到 root	root 需有完整緩衝區
Scatter	`MPI_Scatter`	root 把大陣列切片分發給各 process	與 Gather 互為反向
CUDA-aware MPI	傳 device pointer 給 `MPI_Send/Recv`	移除 D2H/H2D staging copy 與 pinned buffer	需 MVAPICH2 / OpenMPI / IBM Platform MPI

Tip

Collective 函式由 MPI runtime 開發者與系統廠商高度最佳化，通常比自己用一堆 MPI_Send/MPI_Recv 組合更快、可讀性與生產力也更高。能用 collective 就別手刻。

群體通訊 (Collective Communication)

定義：collective communication 是一群 process（由 communicator 界定，通常為 MPI_COMM_WORLD）全部共同參與的通訊操作，而非 point-to-point 的一對一。第 03 篇用過的 MPI_Barrier 就是其中一種。

最常用的五種 collective 原語：

原語	語意	root?	CUDA 對應概念
`MPI_Barrier`	所有 process 互相等待，全到齊才放行	無	`__syncthreads()`（block 內同步）
`MPI_Bcast`	root 的同一份資料複製給全體	有	constant memory 廣播 / 同一值寫入所有 thread
`MPI_Reduce`	全體的值經運算子匯總到 root（`MPI_SUM`, `MPI_MAX`…）	有	parallel reduction tree
`MPI_Gather`	各 process 的片段收集成 root 的大陣列	有	多 block 結果寫回 global memory
`MPI_Scatter`	root 的大陣列切片分發給各 process	有	grid 把資料分給各 block

Barrier 是最常用的一個。在 stencil 範例中，barrier 確保所有 compute process 都收到輸入資料、準備就緒後，才一起開始交換 halo —— 避免少數落後的 process 在資料交換階段拖累全體。

Reduce (MPI_SUM) 樹狀匯總到 root                Broadcast 樹狀散播自 root
   P0   P1   P2   P3                              root(P0)
    \   /     \   /                               /      \
     v0+v1     v2+v3                            P1        P2
        \       /                               / \      /  \
         v0+v1+v2+v3   --> root              (data) (data)(data)...

Important

本章 stencil 範例的「data server 收集輸出」雖然是用 for 迴圈逐一 MPI_Recv（見第 02 篇），但這個 pattern 正是 gather；實務上可直接用 MPI_Gather 取代迴圈，更簡潔也更快。同理，分發 domain partition 的動作概念上接近 scatter（但因每段大小含 halo 而不等長，範例才用顯式 MPI_Send）。

Warning

Collective 函式要求 communicator 內所有 process 都呼叫到同一個 collective，否則會死結（deadlock）。不要在條件分支裡只讓部分 process 進入 MPI_Barrier/MPI_Bcast。

CUDA-aware MPI

動機：現代 MPI 實作「知道」CUDA 程式模型，能直接在不同節點的 GPU memory 之間傳訊息，免去傳統流程的 host staging。支援的實作：MVAPICH2、IBM Platform MPI、OpenMPI。

資料路徑對比 (Data Path)

傳統 MPI（非 CUDA-aware）— 每次 halo 交換需 4 次 staging copy
  GPU mem ──cudaMemcpyAsync(D2H)──▶ pinned host buf ──MPI_Send──▶ 網路
                                                                    │
  GPU mem ◀──cudaMemcpyAsync(H2D)── pinned host buf ◀──MPI_Recv──◀─┘

CUDA-aware MPI — device pointer 直接進出網路（GPUDirect RDMA）
  GPU mem ──MPI_Send(d_ptr)──▶ 網路 ──▶ remote GPU mem
  GPU mem ◀──MPI_Recv(d_ptr)──◀ 網路 ◀── (remote)

面向	傳統 staging MPI	CUDA-aware MPI
pinned host buffer (`cudaHostAlloc`)	需要（left/right boundary + halo）	不需要（可刪 Fig.20.11 lines 21–24）
D2H copy (送出前)	`cudaMemcpyAsync` ×2	移除（刪 Fig.20.15 lines 42–43）
H2D copy (收到後)	`cudaMemcpyAsync` ×2	移除（刪 Fig.20.15 lines 47–48）
`MPI_Sendrecv` 參數	host buffer 位址	device 位址（`d_output + offset`）
主機端程式碼	較複雜、記憶體佈局多一層	簡化

改寫後的 halo 交換

直接把 device 位址傳給 MPI_Sendrecv，由 CUDA-aware runtime 直接更新 GPU memory：

// CUDA-aware 版本：MPI 直接讀寫 device memory，不再經過 host bounce buffer
MPI_Sendrecv(d_output + num_halo_points,      num_halo_points, MPI_FLOAT,
             left_neighbor,  i,
             d_output + left_halo_offset,      num_halo_points, MPI_FLOAT,
             right_neighbor, i, MPI_COMM_WORLD, &status);

MPI_Sendrecv(d_output + right_stage1_offset,  num_halo_points, MPI_FLOAT,
             right_neighbor, i,
             d_output + right_halo_offset,     num_halo_points, MPI_FLOAT,
             left_neighbor,  i, MPI_COMM_WORLD, &status);

Tip

即使用了 CUDA-aware MPI，仍需保留 CUDA streams 與兩段式 kernel（先算 boundary slices、再算 internal points），這樣才能在 halo 一算完就立刻開始跨節點通訊，維持計算/通訊重疊。CUDA-aware 移除的是 staging copy，不是 overlap 策略。詳見 20-Heterogeneous-Computing-Cluster/03-Overlapping-Computation-and-Communication。

Important

不只 halo 交換可受惠：初始輸入分發與最終輸出收集的 memory copy 也能省去，做法是讓 data server 直接對 compute process 的 GPU memory 收送 input/output。

量化：每次 halo 交換省下的 copy

staging copies per iteration = \underset{D 2 H}{\underset{⏟}{2}} + \underset{H 2 D}{\underset{⏟}{2}} = 4 \overset{CUDA-aware}{\to} 0

外加每個 compute process 省下 4 個 pinned buffer 配置（cudaHostAlloc）。

本章總結 (Summary)

joint CUDA/MPI 程式設計的核心觀念，幾乎都能在 CUDA 找到對應：

MPI 概念	說明	CUDA 對應
SPMD	所有 process 跑同一支程式	所有 thread 跑同一個 kernel
MPI rank (pid)	唯一識別 process，用來分流角色	`blockIdx*blockDim+threadIdx` 全域索引
角色分化	data server vs compute process 走不同控制流	thread 依 id 做不同工作
Barrier	`MPI_Barrier` 全體同步	`__syncthreads()` block 內同步
重疊計算與通訊	CUDA streams + async copy	同（kernel/copy 併發）
CUDA-aware MPI	直接交換 device memory	簡化資料搬移

Tip

作者的結語：把一種平行模型（CUDA）學透，就能快速上手其他模型（MPI）—— 因為 SPMD、rank、barrier 等大觀念是共通的。

考試/面試重點 (Exam / Test Patterns)

情境 / 關鍵字	答案 / 技巧
「barrier、broadcast、reduce、gather、scatter 是哪類通訊？」	Collective communication（群體通訊），全 communicator 共同參與
「最常用的 collective 是哪個？」	`MPI_Barrier`（確保全體就緒再互動）
「為何優先用 collective 而非手刻 send/recv？」	runtime/廠商高度最佳化 → 效能、可讀性、生產力皆較佳
「CUDA-aware MPI 移除了什麼？」	halo 交換的 D2H/H2D staging copy 與 pinned buffer；MPI 直接讀寫 device 位址
「哪些 MPI 實作支援 CUDA-aware？」	MVAPICH2、IBM Platform MPI、OpenMPI
「用了 CUDA-aware 還需要 streams 嗎？」	需要 —— 仍用兩段式 kernel 來重疊計算與通訊
「`MPI_Send(ptr, 1000, MPI_FLOAT,…)` 傳了 4000 bytes，每元素幾 bytes？」	4000/1000 = 4 bytes（MPI_FLOAT）
「MPI_Recv 預設是 blocking 嗎？」	是（blocking by default）；MPI process 不共享變數、無 shared memory
「MPI rank 對應 CUDA 的什麼？」	thread 全域索引 `blockIdx.x*blockDim.x+threadIdx.x`（但 rank 是一維）
「collective 用在條件分支內的風險？」	部分 process 未呼叫 → deadlock

重點總覽 (Overview)

群體通訊 (Collective Communication)

CUDA-aware MPI

資料路徑對比 (Data Path)

改寫後的 halo 交換

量化：每次 halo 交換省下的 copy

本章總結 (Summary)

考試/面試重點 (Exam / Test Patterns)

Related Notes