第4章：CUDA流与并发

CUDA流，说白了就是GPU上的任务队列。我刚开始接触这个概念时，总觉得它有点抽象——不就是让GPU同时干几件事吗？后来真正调优时才发现，这里面的门道比想象中深得多。

4.1 CUDA流的基本概念

一个CUDA流，就是一系列按顺序执行的GPU操作。这些操作包括：

• 内核启动（Kernel launch）
• 内存拷贝（cudaMemcpy）
• 事件记录（Event record）

默认情况下，所有操作都在默认流（default stream）中执行。这意味着它们是串行的——一个做完，下一个才能开始。

我习惯把CUDA流想象成工厂里的多条流水线。每条流水线有自己的工人（SM资源）和物料（显存数据）。只要资源够用，多条流水线就能同时开工。

4.2 多流并发执行

多流并发的核心价值是什么？说白了就是让GPU别闲着。你想想看，如果只有一个流，那GPU在执行内核时，数据传输就得等着。反过来，数据传输时，计算单元又在空转。

多流并发的典型模式是这样的：

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

// 流1：传输数据到设备
cudaMemcpyAsync(d_data1, h_data1, size, cudaMemcpyHostToDevice, stream1);
// 流2：传输数据到设备
cudaMemcpyAsync(d_data2, h_data2, size, cudaMemcpyHostToDevice, stream2);

// 流1：执行内核
kernel1<<<grid, block, 0, stream1>>>(d_data1);
// 流2：执行内核
kernel2<<<grid, block, 0, stream2>>>(d_data2);

// 流1：传回结果
cudaMemcpyAsync(h_data1, d_data1, size, cudaMemcpyDeviceToHost, stream1);
// 流2：传回结果
cudaMemcpyAsync(h_data2, d_data2, size, cudaMemcpyDeviceToHost, stream2);

cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

这里有个关键点：cudaMemcpyAsync是异步的，它不会阻塞CPU。而cudaMemcpy是同步的，会阻塞直到传输完成。我在项目中遇到过有人把这两个搞混，结果多流并发完全没生效。

4.3 事件（Event）同步

多流并发虽然好，但有时候我们需要让不同流之间协调工作。比如流2要等流1算完某个中间结果才能开始。这时候就需要事件同步。

事件就像赛道上的接力棒。一个流记录事件，另一个流等待这个事件。用法很简单：

cudaEvent_t event;
cudaEventCreate(&event);

// 流1：执行内核并记录事件
kernel1<<<grid, block, 0, stream1>>>(d_data);
cudaEventRecord(event, stream1);

// 流2：等待事件，然后执行
cudaEventSynchronize(event);  // CPU等待
// 或者
cudaStreamWaitEvent(stream2, event, 0);  // 流2等待

kernel2<<<grid, block, 0, stream2>>>(d_data);

cudaEventDestroy(event);

我曾经在做一个多GPU协同的项目时，发现事件同步的粒度太粗了。每个事件都要等整个内核跑完，但有些中间结果其实可以更早传递。后来我改用细粒度的事件，把内核拆成多个小内核，每个小内核完成后记录事件。这样流之间的等待时间就短多了。

4.4 异步数据传输与计算重叠

这是多流并发的终极目标——让数据传输和计算同时进行。理想情况下，数据传输的延迟被计算完全隐藏。

实现计算重叠的经典模式是流水线（pipeline）：

// 假设有N个数据块需要处理
for (int i = 0; i < N; i++) {
    // 将第i块数据异步传输到设备
    cudaMemcpyAsync(d_data[i], h_data[i], blockSize, 
                    cudaMemcpyHostToDevice, stream[i % numStreams]);
    
    // 启动内核处理第i块数据
    kernel<<<grid, block, 0, stream[i % numStreams]>>>(d_data[i]);
    
    // 将第i块结果异步传回主机
    cudaMemcpyAsync(h_result[i], d_data[i], blockSize,
                    cudaMemcpyDeviceToHost, stream[i % numStreams]);
}

这里有个坑：流的数量不是越多越好。我记得有一次调优，把流从4个加到16个，性能反而下降了。为什么？因为GPU的硬件资源有限——DMA引擎、SM数量、显存带宽都是瓶颈。

一般来说，2到4个流就能达到不错的重叠效果。具体多少个最优，得看你的硬件和 workload。

4.5 知识体系总览

下面这张图总结了CUDA流与并发的核心逻辑。我建议你把它记在脑子里——每次调优时，对照这张图检查自己的实现。

4.6 避坑指南

最后，分享几个我踩过的坑：

• 流数量不是越多越好——我曾经在一个项目里用了32个流，结果性能还不如4个流。用profiler一看，大部分时间都花在流切换上了。
• 别忘了同步——多流并发时，如果某个流依赖另一个流的结果，一定要用事件同步。否则你会得到随机错误，而且很难排查。
• 异步传输必须用页锁定内存——这个我提过了，但还是要强调。用cudaMallocHost，别用malloc。
• 用profiler验证重叠效果——NVIDIA Nsight Systems可以直观地看到各个流的时间线。如果发现计算和传输没有重叠，说明你的实现有问题。

嗯，CUDA流与并发就讲到这里。记住，多流并发的本质是让GPU的各个硬件单元同时工作。理解了这个，你就能写出高效的GPU代码。