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/PCIe32.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.