多 GPU 程式設計 (Multi-GPU Programming)

重點總覽

多 GPU 程式設計透過聚合多顆 GPU 的算力、記憶體容量與頻寬，突破單一 GPU 的問題規模與效能上限。基本流程為：枚舉裝置 → 選擇裝置並建立 context → 分配資料、啟動 kernel → 透過 peer-to-peer 傳輸或存取進行通訊與收集結果。

多 GPU 程式設計總覽

CUDA 透過 host API、driver 基礎設施與 GPU 硬體技術支援多 GPU 程式設計，核心能力包含：

• Host thread CUDA context management：應用需同時管理多個 active context。
• Unified memory addressing：系統內所有 processor 共用 unified virtual address space。
• Peer-to-peer bulk memory transfers：GPU 間的大量記憶體搬移。
• Fine-grained peer-to-peer load/store：kernel 直接以指標存取對方 device memory。
• 上層抽象與系統軟體：CUDA IPC、用 NCCL 做 parallel reduction、用 NVLink / GPUDirect RDMA 配合 NVSHMEM、MPI 做通訊。

常見的多 GPU 編程映射方式（依平行度與既有程式結構選擇）：

[單 host thread] ──► 驅動多顆 GPU
[多 host threads] ──► 每個 thread 各驅動一顆 GPU
[多個單執行緒 process] ──► 每個 process 各驅動一顆 GPU
[多執行緒多 process] ──► 每 thread 各驅動一顆 GPU
[Multi-node NVLink cluster] ──► 跨節點 OS instance 由 thread/process 驅動

3.4.1.1 Device Enumeration（裝置枚舉）

使用前先查詢可用的 CUDA 裝置數量、逐一枚舉並查詢其硬體屬性。

int deviceCount;
cudaGetDeviceCount(&deviceCount);
int device;
for (device = 0; device < deviceCount; ++device) {
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, device);
    printf("Device %d has compute capability %d.%d.\n",
           device, deviceProp.major, deviceProp.minor);
}

重點：cudaGetDeviceCount() 回傳裝置總數；對每個 device index 用 cudaGetDeviceProperties() 填入 cudaDeviceProp，其中 major/minor 即 compute capability。

• 應依硬體屬性、CPU affinity、與對 peer 的連通性來選擇要使用哪些裝置。
• 枚舉是建立 context 之前的第一步。

3.4.1.2 Device Selection（裝置選擇與 current device 語意）

Host thread 可在任意時刻呼叫 cudaSetDevice() 切換目前操作的裝置。

size_t size = 1024 * sizeof(float);
cudaSetDevice(0);              // Set device 0 as current
float* p0;
cudaMalloc(&p0, size);         // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0);   // Launch kernel on device 0
cudaSetDevice(1);              // Set device 1 as current
float* p1;
cudaMalloc(&p1, size);         // Allocate memory on device 1
MyKernel<<<1000, 128>>>(p1);   // Launch kernel on device 1

關鍵語意：current device 決定後續操作的歸屬。

• Device memory 分配 (cudaMalloc) 與 kernel launch 都作用在 current device 上。
• Stream 與 event 在建立時即與 current device 產生關聯。
• 在 host thread 第一次呼叫 cudaSetDevice() 之前，current device 預設為 device 0。

        cudaSetDevice(0)                cudaSetDevice(1)
            │                               │
   current=0│  cudaMalloc → p0@dev0   current=1│  cudaMalloc → p1@dev1
            │  launch     → @dev0             │  launch     → @dev1

3.4.1.3 Multi-Device Stream / Event / Memory Copy 行為

跨裝置時，stream 與 event 的「綁定裝置」會決定操作成敗。

cudaSetDevice(0);
cudaStream_t s0;
cudaStreamCreate(&s0);              // stream s0 綁定 device 0
MyKernel<<<100, 64, 0, s0>>>();     // OK：current=0，s0 屬 device 0
cudaSetDevice(1);
cudaStream_t s1;
cudaStreamCreate(&s1);              // stream s1 綁定 device 1
MyKernel<<<100, 64, 0, s1>>>();     // OK：current=1，s1 屬 device 1
// 失敗：current=1，但 s0 綁定的是 device 0
MyKernel<<<100, 64, 0, s0>>>();

重點：kernel launch 必須送往綁定 current device 的 stream，否則 launch 失敗。

各 API 的跨裝置行為整理：

3.4.2.1 Peer-to-Peer Memory Transfers（P2P 記憶體傳輸）

CUDA 可在裝置間搬移記憶體，並在 P2P access 可行時利用專屬 copy engine 與 NVLink 硬體最大化效能。

• 使用 cudaMemcpy 搭配 cudaMemcpyDeviceToDevice 或 cudaMemcpyDefault 可做裝置間 copy。
• 否則須用 cudaMemcpyPeer() / cudaMemcpyPeerAsync() / cudaMemcpy3DPeer() / cudaMemcpy3DPeerAsync()。

cudaSetDevice(1);
cudaMemcpyPeer(p1, 1, p0, 0, size); // 將 device 0 的 p0 複製到 device 1 的 p1
MyKernel<<<1000, 128>>>(p1);        // Launch kernel on device 1

cudaMemcpyPeer(dst, dstDevice, src, srcDevice, size) 明確指定來源/目的裝置編號。

implicit NULL stream 的跨裝置 copy 具有同步語意：

其他命令(任一裝置) ──┐
                     ▼  copy 不會開始，直到先前所有命令完成
              [ 跨裝置 copy ](NULL stream)
                     │  copy 完成前，之後送往任一裝置的命令都不能開始
                     ▼
   copy 後續命令(任一裝置) ───►

• 同步點 (1)：copy 不會開始，直到先前送往任一裝置的所有命令都完成。
• 同步點 (2)：copy 須跑完，之後送往任一裝置的命令才能開始。
• 但若是在非 NULL stream 的 async copy，仍依一般 stream 行為，可與其他 stream 的 copy/kernel 重疊。

3.4.2.2 Peer-to-Peer Memory Access（P2P 記憶體存取）

依系統的 PCIe / NVLink 拓樸，一個裝置上執行的 kernel 可以直接 dereference 指向另一裝置記憶體的指標。

• 用 cudaDeviceCanAccessPeer() 查詢兩裝置間是否支援 P2P access（回傳 true 才支援）。
• 用 cudaDeviceEnablePeerAccess() 啟用（由 current device 指向某 peer device）。
• 因兩裝置共用 unified virtual address space，同一個指標可同時定址兩裝置的記憶體。

cudaSetDevice(0);
float* p0;
cudaMalloc(&p0, size);
MyKernel<<<1000, 128>>>(p0);       // Launch kernel on device 0

cudaSetDevice(1);
cudaDeviceEnablePeerAccess(0, 0);  // device 1 啟用對 device 0 的 peer access
// 此 kernel 在 device 1 執行，可存取 device 0 上位於 p0 的記憶體
MyKernel<<<1000, 128>>>(p0);

   ┌─────────┐  cudaDeviceEnablePeerAccess(0,0)  ┌─────────┐
   │ Device 1│ ───────────────(NVLink/PCIe)────► │ Device 0│
   │ kernel  │   deref p0 (UVA 同一指標)          │  p0 mem │
   └─────────┘                                   └─────────┘

3.4.2.3 Peer-to-Peer Memory Consistency（P2P 一致性）

跨多裝置分散的 grid 中，並行執行的 thread 之間必須以同步操作強制記憶體存取的順序與正確性。

• 跨裝置同步的 thread 運作於 thread_scope_system 同步範圍。
• 對應地，這些記憶體操作落在 thread_scope_system 記憶體同步 domain。
• CUDA atomic 函式可對 peer device memory 中的物件做 read-modify-write，但僅限只有單一 GPU 存取該物件時。peer atomicity 的需求與限制見 CUDA memory model 的 atomicity 要求。

3.4.2.4 Multi-Device Managed Memory（多裝置 managed memory）

• Managed memory 可在具 P2P 支援的多 GPU 系統上使用。
• 並行 multi-device managed memory 存取的詳細需求，以及 GPU-exclusive 存取 managed memory 的 API，屬 Unified/Managed Memory 的 Multi-GPU 主題。

3.4.2.5 Host IOMMU / PCI ACS / VMs 的影響

硬體與虛擬化設定會直接影響 P2P 是否可用與效能。

考試/測驗重點

Related Notes

• 03-Advanced-CUDA/03-Batched-Transfers-and-Env-Vars
• 03-Advanced-CUDA/05-Thread-Scopes-and-Scoped-Atomics
• 03-Advanced-CUDA/08-CUDA-Driver-API
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 02-Programming-GPUs/14-Async-Streams-and-Events
• 02-Programming-GPUs/16-Unified-and-System-Memory
• 03-Advanced-CUDA/Practice-Advanced-CUDA
• 00-Dashboard/Exam-Traps