第四章：CUDA流与事件——异步操作、流调度、事件同步与性能分析
各位同学，欢迎来到第四章。这一章我们聊点真正能让GPU跑起来的东西——CUDA流与事件。
说实话，很多初学者写CUDA程序，都是“串行思维”：把数据拷过去，算完，再拷回来。GPU利用率低得可怜。我刚开始做AI加速器那会儿，也犯过这毛病。后来被一个老工程师点醒：“你让GPU闲着，它可不会感谢你。”
嗯，这一章我们就来解决这个问题。
4.1 什么是CUDA流？
CUDA流，说白了就是一个操作队列。你把kernel启动、内存拷贝这些操作扔进队列里，GPU按顺序执行。但不同流之间，可以并行执行。
你想想看，如果只有一个流，那所有操作都是串行的。但如果你有多个流，GPU就能同时干好几件事。比如一个流在算数据，另一个流在拷数据，互不干扰。
我个人习惯把流想象成“GPU上的多个工人”。每个工人有自己的任务清单，工人之间可以同时干活。但要注意，工人多了也会打架——资源竞争。
核心概念：流是GPU上的命令序列。同一流内严格顺序执行，不同流之间可以乱序或并行执行。
4.2 流的创建与使用
创建流很简单，就两行代码：
cudaStream_t stream;
cudaStreamCreate(&stream);用完之后记得销毁：
cudaStreamDestroy(stream);那怎么用呢？看个例子：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在stream1中启动kernel
kernel1<<<grid, block, 0, stream1>>>(d_data1);
// 在stream2中启动kernel
kernel2<<<grid, block, 0, stream2>>>(d_data2);
// 在stream1中拷贝数据
cudaMemcpyAsync(h_data1, d_data1, size, cudaMemcpyDeviceToHost, stream1);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);注意看，kernel1和kernel2是在不同流里启动的。如果GPU资源够，它们可以同时执行。但cudaMemcpyAsync这个函数，我建议你记住它——异步拷贝，不阻塞CPU。
个人经验：我在做视频编解码加速时，把解码、预处理、推理分别放到三个流里。解码流负责从显存读数据，预处理流做图像缩放，推理流跑模型。三个流流水线作业，吞吐量直接翻倍。
4.3 默认流与NULL流
你可能注意到了，之前我们写kernel时没指定流，那就是在用默认流。默认流也叫NULL流，写法是0或者nullptr。
但这里有个坑——默认流的行为在CUDA 6.0之后变了。老版本里，默认流会阻塞其他所有流。新版本里，每个线程有自己的默认流，互不干扰。
我曾经在迁移老代码时踩过这个坑。一个项目从CUDA 5.0升级到8.0，结果性能反而下降了。查了半天，发现是默认流的行为变了，导致并行度降低。
避坑指南：如果你在维护老代码，一定要确认默认流的行为。建议显式创建流，别依赖默认流。这样代码更清晰，也更容易移植。
4.4 流调度与硬件机制
流是怎么被GPU执行的？这涉及到硬件调度机制。
GPU内部有多个硬件队列（Hardware Work Queue）。CUDA驱动会把流里的操作映射到这些硬件队列上。但注意，不是每个流对应一个硬件队列。实际上，硬件队列数量有限，通常是几十个。
调度策略是这样的：
- 同一流内的操作，按顺序提交到同一个硬件队列
- 不同流的操作，可以提交到不同硬件队列
- 硬件队列之间可以并行执行，但受限于GPU资源
为什么会这样设计？因为硬件队列是有限的，但流可以创建很多个。驱动会做负载均衡，把流分配到可用的硬件队列上。
我记得有一次做性能调优，创建了32个流，结果发现性能还不如8个流。后来分析发现，硬件队列只有16个，多余的流反而增加了调度开销。
经验之谈：流的数量不是越多越好。一般建议流数量不超过硬件队列数的2倍。具体数值可以通过cudaDeviceGetAttribute查询。
4.5 CUDA事件：精确的时间测量
事件（Event）是CUDA提供的时间戳机制。它可以精确测量GPU上操作的执行时间。
为什么不用CPU计时？因为CPU计时只能测到启动和结束的时间，但kernel是异步执行的。CPU计时会把启动开销也算进去，不准确。
事件的使用方法：
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 milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("Kernel耗时: %f ms\n", milliseconds);
cudaEventDestroy(start);
cudaEventDestroy(stop);注意，cudaEventRecord是异步的，它只是把事件插入到流里。真正的时间测量是在GPU上完成的，所以精度很高。
小技巧：我习惯在性能调优时，用事件测量每个阶段的耗时。比如数据拷贝、kernel计算、结果回传。这样能快速定位瓶颈在哪。
4.6 事件同步：流之间的协调
事件不仅能计时，还能做同步。比如，你希望stream2等stream1算完了再开始，就可以用事件：
cudaEvent_t event;
cudaEventCreate(&event);
// stream1中记录事件
kernel1<<<grid, block, 0, stream1>>>(d_data);
cudaEventRecord(event, stream1);
// stream2等待事件
cudaStreamWaitEvent(stream2, event, 0);
// stream2中的操作会等stream1完成后再执行
kernel2<<<grid, block, 0, stream2>>>(d_data);
cudaEventDestroy(event);这个机制非常有用。比如在流水线中，下游流必须等上游流处理完才能开始。用事件同步，比用cudaDeviceSynchronize高效得多。
我曾经在做一个实时推理系统时，用事件同步实现了三级流水线：输入流、计算流、输出流。每个流之间用事件同步，延迟控制在毫秒级。
注意：cudaStreamWaitEvent不会阻塞CPU。它只是告诉GPU：这个流要等事件完成才能继续。CPU可以继续做其他事情。
4.7 性能分析：用Nsight看流的行为
理论讲完了，我们来点实战。怎么知道你的流用得好不好？用Nsight Systems。
Nsight Systems可以直观地看到每个流的时间线。你会看到：
- 哪些流在并行执行
- 哪些操作在等待
- GPU利用率是多少
我建议你养成一个习惯：每次写完CUDA程序，都用Nsight跑一遍。看看GPU有没有闲着，流之间有没有不必要的等待。
举个例子，如果你看到时间线上有大片空白，说明GPU在空转。这时候就要检查：是不是同步太多了？是不是数据拷贝和计算没重叠？
性能调优口诀：让GPU忙起来，让数据跑起来，让流并行起来。
4.8 实战：流水线优化案例
最后，我分享一个实际案例。这是一个图像处理流水线：
- 从CPU拷贝图像到GPU
- 在GPU上做预处理（归一化、增强）
- 运行推理模型
- 把结果拷回CPU
优化前，我用一个流，顺序执行。GPU利用率只有40%。
优化后，我用了三个流：
- 流1：负责数据拷贝
- 流2：负责预处理
- 流3：负责推理
用事件同步保证顺序：流1完成 -> 流2开始 -> 流3开始。但流1可以继续拷贝下一批数据。
结果GPU利用率提升到85%，吞吐量翻了一倍多。
个人建议：做流水线优化时，先画出数据流图。标清楚哪些操作可以并行，哪些必须串行。然后根据依赖关系分配流。
好了，这一章的内容就到这里。流和事件是CUDA编程的核心技能，掌握了它们，你才能真正发挥GPU的并行能力。下一章我们聊内存优化，那又是一个大坑，到时候见。