4、CUDA流与事件：异步操作、流的创建与同步、事件计时与同步
各位同学，今天我们来聊聊CUDA编程里一个非常核心的概念——流（Stream）和事件（Event）。
说实话，很多初学者刚开始接触CUDA时，总觉得GPU就是一块“大号计算卡”，把数据扔进去，等结果出来就行了。但实际项目中，你会发现GPU的潜力远不止于此。我个人习惯把GPU想象成一个“多流水线的工厂”，而流和事件，就是管理这些流水线的工具。
4.1 为什么需要异步操作？
先问大家一个问题：CPU和GPU之间是怎么配合的？
默认情况下，CPU发出一条指令给GPU，然后CPU就傻等着，直到GPU执行完。这叫同步操作。你想想看，CPU在等的时候，啥也干不了，这不就浪费了吗？
异步操作就是为了解决这个问题的。CPU发出指令后，不用等，直接去做别的事。GPU在后台默默算着，算完了通知CPU一声。
我在项目中遇到过这样一个场景：需要同时处理多路视频流。如果同步操作，CPU得等每一帧处理完才能发下一帧，帧率根本上不去。后来改成异步，CPU只管往GPU里塞数据，GPU自己排队处理，效率直接翻倍。
核心思想：异步操作让CPU和GPU可以“并行工作”，而不是“串行等待”。
4.2 流（Stream）的创建与使用
流是什么？说白了，流就是GPU上的一条“任务队列”。
默认情况下，所有操作都在默认流（default stream）里执行。但如果你只用一个流，那跟同步操作没啥区别。真正的威力在于——创建多个流，让它们并行执行。
4.2.1 创建流
创建流很简单，就两行代码：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);嗯，这里要注意：创建流本身是轻量级操作，但别创建太多。我见过有人一口气创建了1000个流，结果性能反而下降了。为什么呢？因为GPU的硬件资源是有限的，流太多反而会增加调度开销。
4.2.2 在流中执行操作
创建完流，怎么用？看这个例子：
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream1);
kernel<<<grid, block, 0, stream1>>>(d_data);
cudaMemcpyAsync(h_result, d_data, size, cudaMemcpyDeviceToHost, stream1);注意看，我用了 cudaMemcpyAsync 而不是 cudaMemcpy。前者是异步的，后者是同步的。还有内核启动的第四个参数，就是指定流。
我的经验：异步内存拷贝和内核启动要放在同一个流里，这样才能保证执行顺序。我曾经犯过一个错误，把拷贝放在流1，内核放在流2，结果数据还没拷贝完内核就开始算了，结果全错。
4.3 流的同步
多流并行虽然好，但有个问题：怎么知道流里的任务执行完了？
这就涉及到同步了。常用的同步方式有三种：
| 同步方式 | 作用 | 使用场景 |
|---|---|---|
cudaStreamSynchronize(stream) |
阻塞CPU，直到指定流中的所有任务完成 | 需要等待某个流的结果时 |
cudaDeviceSynchronize() |
阻塞CPU，直到设备上所有流都完成 | 程序结束前，确保所有任务完成 |
cudaStreamWaitEvent(stream, event) |
让一个流等待另一个流的事件 | 流间依赖控制 |
我个人习惯用 cudaStreamSynchronize，因为它更精细。比如我有三个流，我只关心流1的结果，那就只等流1，其他流继续跑。
避坑指南：我曾经在项目里滥用 cudaDeviceSynchronize，结果CPU频繁等待，性能还不如单流。记住：同步越少，并行度越高。
4.4 事件（Event）的计时与同步
事件是流的“好搭档”。它主要有两个用途：计时和同步。
4.4.1 用事件计时
想精确测量GPU上某段代码的执行时间？用事件就对了。
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, stream1);
// 执行你的内核
kernel<<<grid, block, 0, stream1>>>(d_data);
cudaEventRecord(stop, stream1);
cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("内核执行时间: %f ms\n", milliseconds);注意，cudaEventSynchronize 是必须的，否则时间还没记录完你就去读，结果不准。
我记得有一次做性能优化，用事件测出来某个内核花了12ms，但直觉告诉我应该更快。后来发现是内存拷贝和内核重叠了，事件记录的是“从记录点到结束点”的时间，而不是内核本身的执行时间。嗯，这个坑大家要注意。
4.4.2 用事件做流间同步
事件另一个重要用途是控制流之间的执行顺序。比如流2要等流1算完某个中间结果才能开始：
cudaEvent_t event;
cudaEventCreate(&event);
// 流1执行
kernel1<<<grid, block, 0, stream1>>>(d_data1);
cudaEventRecord(event, stream1);
// 流2等待流1的事件
cudaStreamWaitEvent(stream2, event);
kernel2<<<grid, block, 0, stream2>>>(d_data2);这样，流2会等到流1的 kernel1 执行完才开始。但注意，cudaStreamWaitEvent 是非阻塞的，CPU不会等，只是GPU内部会做同步。
总结一下：
- 流是任务队列，多流可以实现并行
- 同步要精细，别滥用全局同步
- 事件既能计时，又能做流间同步
好了，这一章的内容就到这里。下一章我们会讲CUDA内存管理，包括统一内存和页锁定内存，这些都是实际项目中经常用到的技巧。
公众号:蓝海资料掘金营，微信deep3321