3、多GPU编程模型：单进程多线程（SPMD）、多进程多GPU、CUDA Stream与Event同步

聊到多GPU编程，很多人第一反应就是“把数据分一分，每个GPU算一块”。嗯，道理是这个道理，但实际落地的时候，你会发现事情远没那么简单。我这些年调过不少多GPU的坑，说白了，核心就三个问题：谁来管任务？数据怎么分？怎么同步？

今天咱们就把这三个问题掰开揉碎，看看SPMD、多进程多GPU、还有Stream和Event这些家伙到底怎么用。

3.1 单进程多线程（SPMD）模型

SPMD，全称是Single Program Multiple Data。名字挺唬人，其实意思就是：所有GPU跑同一份代码，但处理不同的数据。

我个人习惯用CUDA里的SPMD模式。你想想看，一个进程里开多个线程，每个线程绑定一个GPU。代码逻辑完全一样，只是通过cudaSetDevice()来切换当前设备。

举个例子，我曾在项目中做过一个4卡的数据并行训练。代码结构大概是这样的：

// 伪代码：SPMD风格的多GPU控制
int nGPUs = 4;
std::vector<std::thread> threads(nGPUs);

for (int i = 0; i < nGPUs; ++i) {
    threads[i] = std::thread([i]() {
        cudaSetDevice(i);  // 每个线程绑定一个GPU
        // 每个GPU处理自己的数据分片
        kernel<<<grid, block, 0, stream[i]>>>(data + i * chunkSize);
    });
}

for (auto& t : threads) t.join();

这里有个坑，我曾经踩过：线程之间不能共享CUDA context。每个线程必须独立调用cudaSetDevice，否则会报错。嗯，这个细节很多人容易忽略。

3.2 多进程多GPU模型

SPMD虽然好用，但有个硬伤——一个进程能管理的GPU数量有限。而且进程内线程调度有开销。这时候，多进程多GPU模型就派上用场了。

说白了，就是每个GPU跑一个独立的进程。进程之间通过MPI或者NCCL通信。这种方式在超算集群里特别常见。

我记得有一次做千卡规模的训练，用的就是MPI+多进程模型。每个MPI rank绑定一个GPU，代码写起来反而比SPMD更清爽：

// 伪代码：MPI多进程多GPU
int rank, nRanks;
MPI_Init(&argc, &argv);
MPI_Comm_rank(MPI_COMM_WORLD, &rank);
MPI_Comm_size(MPI_COMM_WORLD, &nRanks);

cudaSetDevice(rank % nGPUsPerNode);  // 每个进程绑定一个GPU

// 每个进程处理自己的数据
kernel<<<grid, block>>>(data + rank * chunkSize);

// 跨GPU通信
MPI_Allreduce(sendbuf, recvbuf, count, MPI_FLOAT, MPI_SUM, MPI_COMM_WORLD);

你可能会问：那SPMD和多进程到底怎么选？我个人的经验是：

场景	推荐模型	原因
单机多卡（2-8卡）	SPMD	共享内存方便，线程间通信开销低
多机多卡（集群）	多进程+MPI	天然支持跨节点通信，容错性好
需要细粒度同步	SPMD	线程间同步比进程间快得多
需要隔离性	多进程	一个进程挂了不影响其他进程

3.3 CUDA Stream与Event同步

好了，模型选好了，数据也分好了。但还有个关键问题——同步。你想想看，多个GPU同时干活，怎么知道对方算完了？怎么保证数据依赖正确？

CUDA Stream和Event就是干这个的。

3.3.1 CUDA Stream：流水线并行

Stream说白了就是一个操作队列。同一个Stream里的操作按顺序执行，不同Stream之间可以并行。我习惯把Stream想象成工厂里的流水线——每条线独立干活，互不干扰。

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 两个Stream可以并行执行
kernel1<<<grid, block, 0, stream1>>>(data1);
kernel2<<<grid, block, 0, stream2>>>(data2);

// 等待所有Stream完成
cudaDeviceSynchronize();

这里有个细节：默认Stream（stream 0）是同步的。如果你不显式创建Stream，所有操作都塞进默认Stream，那就退化成串行了。我曾经见过一个团队，代码里没写Stream参数，结果4卡跑得比单卡还慢……嗯，就是因为默认Stream把并行给扼杀了。

3.3.2 Event：细粒度同步

Stream解决了并行问题，但同步还得靠Event。Event就像赛道上的计时器——你可以在某个Stream里插入一个Event，然后在另一个Stream里等这个Event完成。

cudaEvent_t event;
cudaEventCreate(&event);

// Stream A：计算+记录事件
kernelA<<<grid, block, 0, streamA>>>(data);
cudaEventRecord(event, streamA);

// Stream B：等待事件完成后再执行
cudaStreamWaitEvent(streamB, event, 0);
kernelB<<<grid, block, 0, streamB>>>(data);

你想想看，这有什么用？最常见的场景就是数据依赖。比如GPU 0算完中间结果，GPU 1才能开始算。用Event就能精确控制这个“等”的时机，而不是粗暴地cudaDeviceSynchronize()把整个设备都停掉。

3.3.3 多GPU场景下的Stream与Event

当Stream和Event遇上多GPU，事情就更有意思了。不同GPU之间的Stream可以通过跨设备Event来同步。不过要注意，跨设备Event需要启用cudaEventInterprocess标志。

// GPU 0上的Stream
cudaSetDevice(0);
cudaEventCreateWithFlags(&event0, cudaEventInterprocess | cudaEventDisableTiming);
kernel0<<<grid, block, 0, stream0>>>(data0);
cudaEventRecord(event0, stream0);

// GPU 1上的Stream等待GPU 0
cudaSetDevice(1);
cudaStreamWaitEvent(stream1, event0, 0);
kernel1<<<grid, block, 0, stream1>>>(data1);

我曾经在项目中用这个模式做pipeline并行，效果非常好。每个GPU只等它依赖的上一个GPU，而不是等所有GPU。这样就把同步开销降到了最低。

3.4 三种模型的对比与选择

说了这么多，咱们来总结一下。这三种模型各有千秋，没有银弹。我个人的选择逻辑是这样的：

• 单机小规模（2-4卡）：SPMD+Stream，简单直接，调试方便
• 单机大规模（4-8卡）：多进程+MPI，扩展性好，适合生产环境
• 多机集群：多进程+MPI+NCCL，这是标配
• 需要极致性能：SPMD+多Stream+Event细粒度同步

你想想看，其实核心就一句话：把并行度做上去，把同步开销降下来。Stream和Event就是帮你实现这个目标的工具。

嗯，这一章的内容就到这里。多GPU编程模型是后续NVSHMEM的基础，理解透了，后面学起来会轻松很多。

---