Multi-GPU Programming

Often one host thread (or process) per GPU.

32.1 Multiple devices in one process

Often one host thread (or process) per GPU.

int nDev; cudaGetDeviceCount(&nDev);
std::vector<float*> d_parts(nDev);
for (int d = 0; d < nDev; ++d) {
    cudaSetDevice(d);                             // subsequent calls target device d
    cudaMalloc(&d_parts[d], chunkBytes);
    cudaMemcpyAsync(d_parts[d], h_in + d * chunk, chunkBytes, cudaMemcpyHostToDevice);
    kernel<<<grid, block>>>(d_parts[d], chunk);   // runs concurrently on each GPU
}
for (int d = 0; d < nDev; ++d) { cudaSetDevice(d); cudaDeviceSynchronize(); }

32.2 Peer-to-peer (P2P)

With P2P enabled, kernels on one GPU can dereference pointers to another GPU's memory (UVA).

int canAccess;
cudaDeviceCanAccessPeer(&canAccess, 0, 1);
cudaSetDevice(0); cudaDeviceEnablePeerAccess(1, 0);   // GPU 0 can access GPU 1's memory
cudaMemcpyPeerAsync(d_dst_on1, 1, d_src_on0, 0, bytes, stream);   // direct copy over NVLink/PCIe

32.3 NCCL — collective communication

Collectives: AllReduce, Broadcast, Reduce, AllGather, ReduceScatter, AllToAll, send/recv. Topology-aware (NVLink, NVSwitch, InfiniBand); the backbone of distributed deep learning (PyTorch DDP/FSDP).

#include <nccl.h>
ncclComm_t comms[4]; int devs[4] = {0, 1, 2, 3};
ncclCommInitAll(comms, 4, devs);
ncclGroupStart();
for (int i = 0; i < 4; ++i) {
    cudaSetDevice(i);
    ncclAllReduce(sendbuf[i], recvbuf[i], count, ncclFloat, ncclSum, comms[i], streams[i]);
}
ncclGroupEnd();

32.4 Multi-node: CUDA-aware MPI & NVSHMEM

GPUDirect RDMA moves data NIC↔GPU without CPU copies. NVSHMEM provides a partitioned global address space with GPU-initiated communication from inside kernels.

// With CUDA-aware MPI (Open MPI/MVAPICH2/HPC-X built with CUDA), pass device pointers directly:
MPI_Allreduce(MPI_IN_PLACE, d_buf, n, MPI_FLOAT, MPI_SUM, MPI_COMM_WORLD);

32.5 Decomposition strategies

Data parallelism (split data, replicate model), domain decomposition with halo exchange (stencils, CFD), model/tensor/pipeline parallelism (large neural networks), task parallelism.