計算與通訊重疊 (CUDA Streams 與 Pinned Memory)

重點總覽 (Overview)

目標:讓通訊網路 (communication network) 與運算硬體 (computation hardware) 同時忙碌,而不是交替閒置。核心手法是把每次 iteration 的工作切成 stage 1 (boundary) 與 stage 2 (internal),並用 CUDA streams + pinned memory 讓 halo 交換與內部運算重疊。

為何需要重疊 (The Two-Mode Problem)

最直觀的做法:整個 partition 算一步 → 交換 halo → 重複。問題是系統永遠只在兩種模式之一:

天真做法 (serialized):
  ┌── compute (整個 partition) ──┐                     ┌── compute ──┐
  │   network idle              │   ┌── halo exch ──┐  │  net idle   │ ...
  └─────────────────────────────┘   │ compute idle │  └─────────────┘
                                     └──────────────┘
時間 ───────────────────────────────────────────────────────────────►

• 模式一:大家都在算 → 網路閒置。
• 模式二:大家都在換 halo → 運算硬體閒置。

兩階段策略 (Two-Stage Boundary / Internal Kernels)

每個 compute process 在 device memory 中的 partition 佈局(沿 z 切割,每片 = dimx*dimy):

device memory (per partition), 共 dimz+8 片:
 offset(片): 0        4            4           dimz       dimz+4    dimz+8
            │ L halo │ L boundary │  internal  │ R bndry │ R halo │
            │ (虛線) │  (彩色)    │ (dimz-8)   │ (彩色)  │ (虛線) │
            └────────┴────────────┴────────────┴─────────┴────────┘
  num_halo_points = 4*dimx*dimy  (一邊 4 片 halo 的元素數)

 stage1 left  : 用 offset 0          的 12 片 → 算出 4 片 L boundary
 stage1 right : 用 offset dimz-4     的 12 片 → 算出 4 片 R boundary
 stage2       : 用 offset 4(skip L halo)    → 算 (dimz-8) 內部點 (stream1)

stage 1:先算鄰居「下一步」需要的 boundary slice(edge process 只算一邊)。
stage 2:一邊把新 boundary 送出去(D2H copy + MPI),一邊算剩下的 internal 點。

// 重疊用的 offset (Fig. 20.14)
int left_halo_offset    = 0;
int right_halo_offset   = dimx * dimy * (4 + dimz);
int left_stage1_offset  = 0;
int right_stage1_offset = dimx * dimy * (dimz - 4);
int stage2_offset       = num_halo_points;          // = 4*dimx*dimy

MPI_Barrier(MPI_COMM_WORLD);                          // 全部就緒才開始
for (int i = 0; i < nreps; i++) {
  // stage 1: 先算鄰居要的 boundary，放 stream0，兩個 kernel 循序
  call_stencil_kernel(d_output + left_stage1_offset,
                      d_input  + left_stage1_offset, dimx, dimy, 12, stream0);
  call_stencil_kernel(d_output + right_stage1_offset,
                      d_input  + right_stage1_offset, dimx, dimy, 12, stream0);
  // stage 2: 內部點放 stream1，可與 stream0 + 通訊並行
  call_stencil_kernel(d_output + stage2_offset,
                      d_input  + stage2_offset, dimx, dimy, dimz, stream1);
  /* ... 見下方 halo 交換 ... */
}

Pinned Memory 與 DMA (Pinned Memory, DMA, cudaHostAlloc)

cudaMemcpy 底層用 DMA (direct memory access) 硬體搬資料,DMA 在 host 端操作的是實體位址 (physical address)。但 OS 的 virtual memory 會把分頁 (page) 換出 (page out) 到磁碟,實體位址可能被重新指派 → DMA 資料被覆寫、損毀。

一般 malloc 記憶體 (pageable):
  host buf ──(CUDA 先複製)──► pinned 暫存 ──(DMA)──► device   ← 兩步、且同步阻塞!

pinned 記憶體 (cudaHostAlloc):
  host buf ─────────────────(DMA 直接搬)──────────────► device   ← 一步、可非同步

• Pinned / page-locked memory:標記為「不可換出」,DMA 可安全直接存取。
• 一般 malloc 緩衝區 → CUDA runtime 須先複製到內部 pinned 暫存再 DMA,多一次複製且強迫 synchronous(host 必須等 copy 完才能繼續),所有複製被序列化。
• 用 cudaHostAlloc 配置的 buffer 稱為 bounce buffer(彈跳緩衝),作為 device ↔ 遠端 MPI process 之間的暫存站。

// 為 left/right boundary(要送出)與 left/right halo(要收進)配置 pinned buffer
cudaHostAlloc((void**)&h_left_boundary,  num_halo_bytes, cudaHostAllocDefault);
cudaHostAlloc((void**)&h_right_boundary, num_halo_bytes, cudaHostAllocDefault);
cudaHostAlloc((void**)&h_left_halo,      num_halo_bytes, cudaHostAllocDefault);
cudaHostAlloc((void**)&h_right_halo,     num_halo_bytes, cudaHostAllocDefault);

非同步複製與 Streams (cudaMemcpyAsync & CUDA Streams)

Stream = 一條有序的操作序列。把 cudaMemcpyAsync 或 kernel launch 指定到某 stream:

• 同一 stream 內:依放入順序循序執行。
• 不同 stream 間:無順序限制,可並行。

stream0:  [stage1 L]→[stage1 R]→[D2H boundary]──(MPI)──[H2D halo]
stream1:  ─────────────[ stage2 internal compute (最久) ]─────────
                         ▲ 兩條 stream 並行 → 通訊被內部運算藏住

cudaStream_t stream0, stream1;
cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1);
// ...
// boundary 算完後,D2H 複製到 pinned host(同在 stream0，自動等 stage1 kernel)
cudaMemcpyAsync(h_left_boundary,  d_output + num_halo_points,
                num_halo_bytes, cudaMemcpyDeviceToHost, stream0);
cudaMemcpyAsync(h_right_boundary, d_output + right_stage1_offset + num_halo_points,
                num_halo_bytes, cudaMemcpyDeviceToHost, stream0);
cudaStreamSynchronize(stream0);   // 只等 stream0：確保 boundary 已在 host

Halo 交換與 MPI_Sendrecv (Halo Exchange)

MPI_Sendrecv = MPI_Send + MPI_Recv 合一,一次呼叫同時送出與收進,減少呼叫數、避免死結。模式:所有人都往左鄰居送 boundary,等同所有人都從右鄰居收 halo。

int left_neighbor  = (pid > 0)      ? (pid - 1) : MPI_PROC_NULL;
int right_neighbor = (pid < np - 2) ? (pid + 1) : MPI_PROC_NULL;
// ...
// 送 L boundary 給左鄰、從右鄰收 R halo
MPI_Sendrecv(h_left_boundary,  num_halo_points, MPI_FLOAT, left_neighbor,  i,
             h_right_halo,      num_halo_points, MPI_FLOAT, right_neighbor, i,
             MPI_COMM_WORLD, &status);
// 送 R boundary 給右鄰、從左鄰收 L halo
MPI_Sendrecv(h_right_boundary, num_halo_points, MPI_FLOAT, right_neighbor, i,
             h_left_halo,      num_halo_points, MPI_FLOAT, left_neighbor,  i,
             MPI_COMM_WORLD, &status);
// 收到的 halo 搬回 device（stream0，與 stream1 的 internal kernel 並行）
cudaMemcpyAsync(d_output + left_halo_offset,  h_left_halo,  num_halo_bytes,
                cudaMemcpyHostToDevice, stream0);
cudaMemcpyAsync(d_output + right_halo_offset, h_right_halo, num_halo_bytes,
                cudaMemcpyHostToDevice, stream0);

MPI_Barrier 與 Double Buffering (Barrier & Double Buffering)

MPI_Barrier:跨 process 的 barrier,類比 CUDA 的 __syncthreads()(但作用於 process)。所有指定 process 互相等待,全部抵達才放行。

• iteration 開始前的 barrier:確保所有 compute node 都收到輸入再一起開始換資料,避免少數落後者拖累全體。
• 結尾的 barrier(MPI_Barrier):等所有 process 算完、送出 output 給 data server。

Double buffering:每步結束後交換 d_input 與 d_output 指標,本步輸出 → 下步輸入,避免讀寫同一緩衝。

  cudaDeviceSynchronize();        // 等所有 kernel + copy 完成，d_output 全就緒
  float *temp = d_output;         // double-buffer swap
  d_output = d_input;
  d_input  = temp;
} // end for (nreps)

時間軸(一個 iteration,重疊後):
 stream0: │stg1 L│stg1 R│D2H│        │H2D L│H2D R│
                          └─sync─►│MPI_Sendrecv│─►
 stream1: │────────── stage2 internal compute (藏住通訊) ──────────│
                                                          → cudaDeviceSync → swap

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

Related Notes

• 20-Heterogeneous-Computing-Cluster/01-MPI-Cluster-Background-and-Basics
• 20-Heterogeneous-Computing-Cluster/02-MPI-Point-to-Point-Communication
• 20-Heterogeneous-Computing-Cluster/04-Collective-Communication-and-CUDA-Aware-MPI
• 08-Stencil/02-Shared-Memory-Tiling-for-Stencil
• 07-Convolution/03-Tiled-Convolution-and-Halo-Handling
• 22-Advanced-Practices-And-Future-Evolution/02-Kernel-Execution-Control