異質計算叢集程式設計 練習題 (Practice - MPI Cluster Background, Stencil Example, and MPI Basics)

MPI 假設 distributed memory model:每個 process 有獨立位址空間,不能透過 shared memory 存取同一變數,只能靠送/收訊息 (message passing) 交換資料。程式模型是 SPMD(Single Program Multiple Data):所有 process 跑同一支程式,靠 rank 走不同控制流與函式呼叫路徑——對應 CUDA kernel 內所有 thread 跑同一個 kernel、靠 thread id 分支。

五大 API:MPI_Init(初始化)、MPI_Comm_rank(取本 process rank)、MPI_Comm_size(取總 process 數)、MPI_Abort(中止,錯誤碼非 0)、MPI_Finalize(釋放資源)。角色分工:rank np-1(最大 rank)保留當 data server 做 I/O;rank 0…np-2 共 np-1 個 compute process 算 stencil。注意傳給 compute 的 dimz 是 dimz/(np-1)。

MPI_Send 有 6 個參數 (buf, count, datatype, dest, tag, comm);MPI_Recv 有 7 個,多了最後的 MPI_Status*。第 4 參數:MPI_Send 是 dest(目的 rank),MPI_Recv 是 source(來源 rank)。不限 tag 用 MPI_ANY_TAG。傳輸位元組數 = count × sizeof(datatype)。

Halo cells = 計算邊界所需、來自鄰居 partition 的資料切片(某 partition 的 halo 即鄰居的 boundary)。Ghost cells = domain 最外緣、無真實鄰居處,如 convolution 邊界填 0,不需透過網路傳。25-point 每方向 4 鄰居 → 每側 4 片 halo。令 S = dimz/num_comp:internal 兩側皆有鄰居 → 收 dimx*dimy*(S+8);edge 僅單側 → 收 dimx*dimy*(S+4),缺側用 ghost。

天真做法讓系統永遠處於兩種模式之一:大家都在算時網路閒置,大家都在換 halo 時運算硬體閒置,兩者交替浪費。兩階段策略:stage 1 先算鄰居「下一步」需要的 boundary slices(用 stream0);stage 2 一邊把新 boundary 送出(D2H copy + MPI),一邊算剩下的 internal points(用 stream1)。理想時間 T_iter ≈ T_stage1 + max(T_internal_compute, T_halo_comm),只要內部運算夠久就能把通訊延遲完全藏住。

cudaMemcpy 底層用 DMA 硬體,DMA 在 host 端操作的是實體位址;但 OS 的 virtual memory 可能把 pageable 記憶體換出 (page out),實體位址被重指派 → DMA 資料損毀。cudaHostAlloc(..., cudaHostAllocDefault) 配置 pinned / page-locked memory,標記為不可換出,DMA 可安全直接搬運且免去「先複製到內部 pinned 暫存」的額外步驟。cudaMemcpyAsync 要求 host buffer 必須是 pinned,否則非同步語意不成立。釋放要用 cudaFreeHost(不是 free)。

Stream = 有序操作序列:同一 stream 內依放入順序循序執行;不同 stream 間無順序限制,可並行。把「短而急」的 boundary 運算 + 通訊鏈放 stream0、把「長」的 internal 運算放 stream1,硬體即可同時推進兩者達成重疊。cudaStreamSynchronize(s) 只等單一 stream s(確保 boundary 已 D2H 完成再做 MPI);cudaDeviceSynchronize() 等所有 device 活動(iteration 結尾確認 d_output 全就緒)。

MPI_Sendrecv = MPI_Send + MPI_Recv 合一,一次呼叫同時送出 boundary 與收進 halo,減少呼叫數並避免死結 (deadlock)。對 edge process,把缺少的鄰居設成 MPI_PROC_NULL(process 0 的 left_neighbor、process np-2 的 right_neighbor),MPI runtime 會自動跳過該方向的送/收,因此不需特別的 if-then-else 分支:left_neighbor = (pid>0)?(pid-1):MPI_PROC_NULL。

Collective 原語:barrier、broadcast (MPI_Bcast)、reduce (MPI_Reduce)、gather (MPI_Gather)、scatter (MPI_Scatter),全 communicator 共同參與。MPI_Barrier 確保所有 compute process 都收到輸入、就緒後才一起開始交換 halo,避免落後者拖累全體(類比 __syncthreads() 但作用於 process)。CUDA-aware MPI 讓 MPI 直接讀寫 GPU memory,移除 halo 交換的 D2H/H2D staging copy 與 pinned buffer(每 iteration 省 4 次 copy),只需把 device 位址傳給 MPI_Sendrecv。支援的實作:MVAPICH2、IBM Platform MPI、OpenMPI。

num_comp = 16,S = 2048/16 = 128 片,一片 = 64×64 = 4096 點,float = 4 bytes。
(a) 每 compute process 輸出 = 4096×128 = 524,288 點。
(b) halo:internal 兩側各 4 片 = 8×4096 = 32,768;edge 單側 4 片 = 4×4096 = 16,384。
(c) stage 1 boundary:internal 8 片 = 32,768;edge 4 片 = 16,384。
(d) stage 2 internal:internal (128-8)×4096 = 491,520;edge (128-4)×4096 = 507,904。
(e) stage 2 送出:internal 兩側各 16,384×4 = 65,536 bytes → 131,072 bytes;edge 單側 65,536 bytes。

(a) 每元素 = 4000 / 1000 = 4 bytes(即 MPI_FLOAT,單精度浮點)。傳輸位元組 = count × sizeof(datatype)。
(b) 「MPI_Recv 預設是 blocking」 為真——MPI_Recv 回傳即代表資料已收到。其餘皆錯:MPI 無 128-byte 下限;MPI 是 distributed memory,process 不共享變數、無 shared memory。

送完 P0(edge,只含右 halo)後,前進 dimx*dimy*(S-4):因 P1 的 partition 從 input + S 起,但須退回 4 片納入左 halo,故淨位移 S-4 片。internal 迴圈中每次前進 dimx*dimy*S(一個完整 partition)。差別在於:第一次位移要從「無左 halo 的 edge」切換到「有左 halo 的 internal」對齊;之後每個 internal 起點都同樣退 4 片,退的效果一致抵銷,故相鄰起點淨距仍是一個 partition (S 片)。

天真策略把執行硬性分成兩相:全算階段→網路閒置,全換 halo 階段→運算硬體閒置,兩者輪流空轉,T ≈ T_compute + T_comm(相加)。重疊策略先算 boundary(stage1),再用 stream1 算 internal、同時用 stream0 + MPI 交換 halo,使網路與運算同時忙碌,T ≈ T_stage1 + max(T_internal_compute, T_halo_comm)。完全藏住通訊的條件:T_halo_comm < T_internal_compute,即內部 slice 夠多(dimz-8 夠大)讓內部運算時間蓋過通訊時間。若內部 slice 太少 → 內部運算太短,網路又變成瓶頸、通訊外露,退化回接近天真策略。

傳統 staging MPI 每次 halo 交換需 4 次 copy:送出前 2 次 D2H(cudaMemcpyAsync 把 left/right boundary 搬到 pinned host bounce buffer),收到後 2 次 H2D(把 halo 搬回 device),外加每 compute process 4 個 cudaHostAlloc pinned buffer。CUDA-aware MPI 讓 MPI_Sendrecv 直接收送 device 位址(GPUDirect),移除全部 4 次 staging copy 與 pinned buffer,簡化主機端程式碼與記憶體佈局。但仍需保留 CUDA streams 與兩段式 kernel:CUDA-aware 移除的是 staging copy,不是 overlap 策略——還是要先算 boundary、再算 internal,才能在 halo 一算完就立刻開始跨節點通訊,維持計算/通訊重疊。

Pageable 路徑:DMA 只能在實體位址上搬,但 pageable 記憶體隨時可能被 OS 換出 → 為避免損毀,CUDA runtime 必須先把資料複製到內部 pinned 暫存再 DMA,多一次複製;且這個兩步流程被實作成 synchronous(host 必須等 copy 完才能繼續),所有複製被序列化。Pinned memory 標記為不可換出:DMA 可直接存取(免額外複製,解決速度)、且資料不會在搬運中被 page out(解決損毀),於是 cudaMemcpyAsync 的非同步語意才能成立,複製可與 kernel 並行。代價:pinned 記憶體不可換出,過度配置會壓縮系統可用實體記憶體、影響其他程序,因此只用於真正需要非同步傳輸的 bounce buffer。

核心觀念鏈:單卡 CUDA 放大成多節點叢集 → domain partitioning(沿 z 切,一片 z slice 連續)→ point-to-point 分發 partition + halo → 兩階段 + streams + pinned 讓計算與通訊重疊 → double buffering 換 buffer → collective / CUDA-aware MPI 簡化協作與資料搬移。

關鍵公式:S = dimz/num_comp、partition = dimx*dimy*S、一側 halo = 4*dimx*dimy、internal 收 (S+8)/edge 收 (S+4)、bytes = count*sizeof(datatype)、T_iter ≈ T_stage1 + max(T_internal, T_comm)。

CUDA ↔ MPI 對照:SPMD ↔ 同 kernel、MPI rank ↔ blockIdx*blockDim+threadIdx(一維)、MPI_Barrier ↔ __syncthreads()、send/recv ↔ global memory 讀寫(但 MPI 不共享)。

一句話總結:joint CUDA/MPI 把 stencil 從單卡擴展到叢集,靠 domain partitioning 切資料、halo exchange 換資料,並用 streams + pinned memory(或 CUDA-aware MPI)把通訊藏進運算。