4. CUDA流与事件：流的概念与创建、流同步、事件与计时、流与并发执行、默认流与非默认流

各位同学，今天我们来聊聊CUDA编程里一个非常核心，但也容易让人迷糊的概念——流（Stream）和事件（Event）。

说实话，我刚开始接触CUDA的时候，觉得流这玩意儿可有可无。不就是把任务排个队嘛，默认流不香吗？直到我在一个项目里，需要同时处理多路视频流，CPU端还得忙着做预处理，结果发现GPU利用率死活上不去。嗯，那时候我才意识到，不懂流，你根本玩不转真正的并发。

这一章，我们就把它彻底讲透。

4.1 流的概念：GPU里的任务流水线

你可以把流想象成一条任务流水线。GPU会按照你提交任务的顺序，一条一条地执行。但如果你只有一条流水线（默认流），那所有任务都得排队等着。

流的本质是什么？说白了，它是一个操作序列。这些操作包括：

• 内存拷贝（Host to Device / Device to Host）
• Kernel 启动
• 事件记录

同一个流里的操作，是严格顺序执行的。但不同流之间的操作，则可以乱序、并发执行。这就是提升性能的关键。

4.2 流的创建与销毁

创建流非常简单。API就两个：

cudaStream_t stream;
cudaError_t cudaStreamCreate(cudaStream_t *pStream);

销毁流：

cudaError_t cudaStreamDestroy(cudaStream_t stream);

我个人习惯在程序初始化的时候，就把需要用到的流全部创建好。比如：

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

// ... 使用流 ...

cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

这里有个坑，我曾经踩过：不要频繁创建和销毁流。创建流本身有开销，如果你在循环里反复创建销毁，性能反而会下降。我建议在程序启动时一次性创建，结束时统一销毁。

4.3 默认流与非默认流

每个CUDA程序都有一个默认流，也叫NULL流。你如果不显式指定流，所有操作都默认提交到默认流里。

默认流有个特殊行为：它会阻塞其他所有流。什么意思呢？

• 默认流里的操作，必须等所有其他流里的操作完成，才能开始执行。
• 反过来，其他流里的操作，也必须等默认流里的操作完成，才能开始执行。

这其实是个历史遗留问题。早期的CUDA为了兼容性，设计了这种“隐式同步”。但这也意味着，如果你混合使用默认流和非默认流，并发效果会大打折扣。

创建非默认流，就是上面说的 cudaStreamCreate。然后提交操作时，把流作为参数传进去：

kernel_name<<<grid, block, 0, stream>>>(args...);
cudaMemcpyAsync(dst, src, size, direction, stream);

注意，这里用的是 cudaMemcpyAsync，它是异步的，不会阻塞CPU。而 cudaMemcpy 是同步的，会阻塞CPU直到拷贝完成。

4.4 流同步：别让CPU跑得太快

流是异步的，这意味着你提交任务后，CPU会立刻返回，继续执行后面的代码。但有时候，你需要等待流里的所有操作都完成，才能进行下一步。比如，你要把计算结果从GPU拷回CPU。

这时候就需要流同步。

最常用的API是：

cudaError_t cudaStreamSynchronize(cudaStream_t stream);

它会阻塞CPU，直到指定流里的所有操作都完成。

举个例子：

cudaStream_t stream;
cudaStreamCreate(&stream);

// 异步启动Kernel
myKernel<<<grid, block, 0, stream>>>(d_data);

// CPU可以在这里做其他事情
doSomeCPUWork();

// 等Kernel执行完
cudaStreamSynchronize(stream);

// 现在可以安全地拷贝结果了
cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);

除了 cudaStreamSynchronize，还有几个同步方式：

• cudaDeviceSynchronize()：等待所有流里的所有操作完成。这个比较粗暴，一般用于程序结束前的清理。
• cudaStreamWaitEvent()：让一个流等待某个事件。这个我们后面会讲。

4.5 事件（Event）：精确计时与流间同步

事件是CUDA提供的一个轻量级标记。你可以把它插到流里的某个位置，然后查询它是否已经执行完毕。

事件有两个主要用途：

1. 精确计时：测量Kernel或内存拷贝的执行时间。
2. 流间同步：让一个流等待另一个流里的某个点。

4.5.1 事件的创建与使用

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

// 在流里记录开始事件
cudaEventRecord(start, stream);

// 执行Kernel
myKernel<<<grid, block, 0, stream>>>(d_data);

// 记录结束事件
cudaEventRecord(stop, stream);

// 等待事件完成
cudaEventSynchronize(stop);

// 计算耗时（毫秒）
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);

printf("Kernel耗时: %f ms\n", elapsedTime);

cudaEventDestroy(start);
cudaEventDestroy(stop);

这里要注意：cudaEventRecord 是异步的，它只是把事件插入到流的队列里。cudaEventSynchronize 会阻塞CPU，直到该事件被GPU处理完毕。

4.5.2 流间同步：让流学会等待

有时候，流A里的某个操作，需要等流B里的某个操作完成才能开始。这时候就可以用 cudaStreamWaitEvent。

cudaEvent_t event;
cudaEventCreate(&event);

// 在流B里记录事件
cudaEventRecord(event, streamB);

// 让流A等待这个事件
cudaStreamWaitEvent(streamA, event, 0);

// 现在流A里的操作，会在流B的事件之后执行
kernelA<<<grid, block, 0, streamA>>>(d_data);

这个机制非常有用。比如，流B负责从CPU拷贝数据到GPU，流A负责计算。你可以在流B拷贝完成后记录一个事件，然后让流A等待这个事件。这样就能保证数据就绪后再计算，同时又不阻塞CPU。

4.6 流与并发执行：让GPU忙起来

流的终极目标，就是实现并发执行。但并发不是你想并，想并就能并。它有几个前提条件：

• 硬件资源充足：比如SM（流多处理器）数量、共享内存大小、寄存器数量等。
• 操作类型不同：比如一个流在做内存拷贝，另一个流在做计算，它们可以并发。但如果两个流都在做计算，且SM资源不够，那就会串行化。
• 没有数据依赖：流之间不能有读写冲突。

我画了一张图，帮你理解流的并发模型：

从图上可以看到，三个流里的操作在时间上是重叠的。这就是并发执行的效果。如果只用默认流，这些操作会串行执行，总时间会大大增加。

4.7 实战建议：如何用好流

讲了这么多，我总结几条实战经验：

1. 尽量使用非默认流。默认流会阻塞其他流，影响并发。
2. 合理划分任务。把没有数据依赖的任务放到不同流里。比如，数据预处理、计算、结果回传，可以分别放到三个流里。
3. 不要过度创建流。流的数量不是越多越好。一般来说，2-4个流就够用了。太多流反而会增加调度开销。
4. 用事件做精确计时。别再用CPU计时了，不准。
5. 注意异步API的使用。比如 cudaMemcpyAsync 和 cudaMemcpy 的区别，搞错了会导致隐式同步。

好了，这一章的内容就到这里。流和事件是CUDA并发编程的基石。理解了它们，你才能真正发挥GPU的潜力。下一章我们会深入讨论CUDA内存模型，到时候你会看到，流和内存管理是如何配合的。

---