4、CUDA流与异步操作：让GPU真正忙起来
说实话，很多刚接触GPU编程的朋友，最容易犯的一个错误就是——把GPU当CPU用。什么意思呢？就是一条指令执行完，再执行下一条。GPU的并行能力，其实被浪费了大半。
今天我们要聊的CUDA流，就是解决这个问题的核心工具。我个人习惯把CUDA流比作「GPU上的传送带」。你可以同时开多条传送带，每条传送带独立干活，互不干扰。
4.1 CUDA流的概念
CUDA流，说白了就是一个操作队列。你把kernel启动、内存拷贝这些操作丢进队列里，GPU按顺序执行。但不同流之间的操作，是可以并发执行的。
默认情况下，所有操作都在默认流（NULL流）里。这就像只有一条传送带，所有任务排着队来。你想想看，GPU那么多计算单元，就干等着？太浪费了。
核心要点：CUDA流是实现GPU并发执行的关键机制。每个流内部是串行的，但不同流之间可以并行。
4.2 创建与销毁流
创建流很简单，API就两个：cudaStreamCreate 和 cudaStreamDestroy。我在项目中遇到过一种情况：有人创建了几百个流，结果性能反而下降了。为什么？因为流太多，调度开销反而成了瓶颈。
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在流中启动kernel
myKernel<<<grid, block, 0, stream1>>>(d_data);
myKernel<<<grid, block, 0, stream2>>>(d_data);
// 销毁流
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);我的建议：流的数量一般不超过GPU的硬件队列数。对于现代GPU，4-8个流通常就够用了。别贪多。
4.3 异步内存拷贝
这是很多人的知识盲区。常规的 cudaMemcpy 是同步的，它会阻塞CPU。而 cudaMemcpyAsync 是异步的，它把拷贝操作丢进流里，CPU立刻返回继续干活。
// 同步拷贝 - CPU会等
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
// 异步拷贝 - CPU不等待
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream1);嗯，这里要注意：异步拷贝要求主机内存必须是页锁定内存（pinned memory）。否则，拷贝会退化成同步操作。我曾经因为这个坑，排查了整整一个下午。
避坑指南：我曾经在项目里用 cudaMemcpyAsync 但没分配pinned memory，结果性能一点没提升。后来才发现，CUDA在背后默默做了同步。记住：异步拷贝必须搭配 cudaHostAlloc 或 cudaMallocHost 分配的主机内存。
4.4 多流并发执行
多流并发，才是CUDA流的精髓。你可以让一个流做计算，另一个流做数据拷贝，两者同时进行。这叫做「计算-拷贝重叠」。
// 流1：做计算
kernelA<<<grid, block, 0, stream1>>>(d_data1);
// 流2：同时拷贝数据
cudaMemcpyAsync(d_data2, h_data2, size, cudaMemcpyHostToDevice, stream2);
// 流3：再做另一个计算
kernelB<<<grid, block, 0, stream3>>>(d_data3);你想想看，如果这三个操作都在默认流里，那就得排队：先拷贝，再算A，再算B。现在呢？拷贝和计算同时进行，时间直接砍半。
| 场景 | 默认流（串行） | 多流（并发） |
|---|---|---|
| 数据拷贝 + 计算 | 拷贝 → 计算 → 总时间 = T_copy + T_compute | 拷贝 || 计算 → 总时间 ≈ max(T_copy, T_compute) |
| 多个小kernel | 依次执行 | 可同时执行（取决于硬件资源） |
4.5 事件（Event）与同步
多流并发是好，但问题来了：怎么保证执行顺序？比如，流2需要等流1算完才能开始。这时候就需要事件（Event）了。
事件就像赛道上的「计时点」和「信号灯」。你可以记录某个流执行到哪一步了，也可以让另一个流停下来等。
cudaEvent_t event;
cudaEventCreate(&event);
// 流1执行完kernel后，记录事件
kernelA<<<grid, block, 0, stream1>>>(d_data);
cudaEventRecord(event, stream1);
// 流2等待事件完成后再执行
cudaStreamWaitEvent(stream2, event, 0);
kernelB<<<grid, block, 0, stream2>>>(d_data);
// 测量时间
cudaEventSynchronize(event);
float elapsed;
cudaEventElapsedTime(&elapsed, startEvent, event);关键理解：cudaEventRecord 只是「记录」事件，不阻塞任何流。cudaStreamWaitEvent 才是让流等待的指令。这两个配合使用，就能实现流之间的精确同步。
我记得有一次做视频处理管线，需要解码、预处理、推理、后处理四个阶段。每个阶段一个流，用事件串联。结果呢？整个管线的吞吐量提升了3倍。说白了，就是让GPU的每个部分都别闲着。
4.6 实战建议
最后，我总结几条实战经验：
- 先分析瓶颈：你的程序是计算密集还是IO密集？如果是计算密集，多流帮助不大；如果是IO密集，多流是神器。
- 别滥用事件：事件同步是有开销的。我曾经见过有人每个kernel之间都插事件，结果同步开销比计算还大。
- 用工具验证：NVIDIA Nsight Systems可以直观看到各个流的时间线。我每次优化完都会跑一遍，看看流之间有没有真正的重叠。
- 注意默认流阻塞：默认流有个「暗坑」——它会阻塞其他流。所以，一旦用了多流，就尽量别碰默认流。
小技巧：如果你不确定自己的多流优化是否有效，可以先用 cudaStreamSynchronize(stream) 改成同步版本，对比一下性能。如果异步版本更快，说明优化有效；如果差不多甚至更慢，那就要检查是不是有隐式的同步点了。
好了，CUDA流的内容就这些。说白了，它就是让GPU从「单线程」变成「多线程」的关键。下一章我们会聊更高级的CUDA图（Graph），那又是另一个维度的优化了。