2、CUDA流(Stream)基础：CUDA流的概念、默认流与非默认流、流的同步与异步、流与事件
2.1 CUDA流到底是什么？
咱们先聊聊CUDA流的概念。说白了，CUDA流就是一个操作队列。你把kernel启动、内存拷贝这些操作扔进去，它们就会按顺序执行。但不同流之间的操作，是可以并发的。
我记得刚接触CUDA时，以为GPU就是个黑盒子，一次只能干一件事。后来才发现，GPU内部有多个引擎——比如拷贝引擎和计算引擎是独立的。你想想看，如果数据拷贝和计算能重叠，那性能不就翻倍了吗？
嗯，这里要注意：流本身不提供任何数据保护机制。如果你在流A里写数据，在流B里读同一块内存，那结果就是未定义的。我早期一个项目就踩过这个坑，排查了两天才发现是流之间没做同步。
核心要点：流是GPU上的命令队列，不同流之间可以并发执行，但同一流内的操作严格按序执行。
2.2 默认流与非默认流
每个CUDA程序启动时，都会有一个默认流。在旧版CUDA中，默认流叫NULL流。你如果不显式创建流，所有操作都跑在默认流里。
我个人习惯是：只要涉及多流，就显式创建非默认流。为什么？因为默认流在某些场景下会有隐式同步行为，容易让人迷惑。
来看代码：
// 创建非默认流
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在流中启动kernel
myKernel<<<grid, block, 0, stream1>>>(d_data1);
myKernel<<<grid, block, 0, stream2>>>(d_data2);
// 别忘了销毁
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);这里有个细节：kernel启动时的第四个参数就是流句柄。传0表示使用默认流。我建议你养成习惯，哪怕只用单流，也显式创建一个。这样代码迁移到多流场景时，改动最小。
小技巧：可以用 cudaStreamCreateWithFlags 创建非阻塞流，避免默认流的隐式同步。这在多线程场景下特别有用。
2.3 流的同步与异步
CUDA操作默认是异步的。你调用 cudaMemcpyAsync 或者启动kernel，CPU会立刻返回，不会等GPU干完活。但异步不代表不排队——操作还是按顺序进入流的队列。
那怎么知道GPU干完了？需要同步。同步方式有几种：
| 同步方式 | 作用范围 | 说明 |
|---|---|---|
cudaDeviceSynchronize() |
整个设备 | 阻塞CPU直到设备上所有流都完成 |
cudaStreamSynchronize(stream) |
单个流 | 只等待指定流完成 |
cudaEventSynchronize(event) |
单个事件 | 等待事件被记录 |
我曾经在项目中遇到一个性能问题：CPU端频繁调用 cudaDeviceSynchronize()，导致GPU利用率上不去。后来改成用事件做细粒度同步，吞吐量直接提升了30%。
你想想看，如果每次启动kernel都等GPU干完，那跟串行有啥区别？异步的精髓就是：让CPU和GPU各干各的，只在必要时才同步。
避坑指南：我曾经在循环里忘了同步，结果CPU端读取了还没拷贝完成的数据。排查时发现数据时对时错，非常诡异。记住：异步操作后，如果CPU要访问GPU数据，必须先同步。
2.4 流与事件
事件（Event）是CUDA提供的一种轻量级同步工具。你可以把它想象成一个标记点——在流里插入一个事件，然后在另一个流里等待这个事件。
事件最常见的用途有两个：
- 计时：记录kernel或拷贝操作的执行时间
- 流间同步：让流B等流A干完某件事再继续
看个计时的例子：
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, stream1);
myKernel<<<grid, block, 0, stream1>>>(d_data);
cudaEventRecord(stop, stream1);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);
printf("Kernel耗时: %f ms\n", ms);
cudaEventDestroy(start);
cudaEventDestroy(stop);这里有个坑：cudaEventElapsedTime 只能在同步之后调用，否则结果未定义。我刚开始用时就犯过这个错，测出来的时间忽大忽小，后来才发现是没等事件完成。
事件同步的另一个典型场景是流水线：
// 流A做计算，流B做拷贝
cudaEvent_t event;
cudaEventCreate(&event);
cudaMemcpyAsync(d_in, h_in, size, cudaMemcpyHostToDevice, streamA);
computeKernel<<<grid, block, 0, streamA>>>(d_in, d_out);
cudaEventRecord(event, streamA);
// 流B等待流A完成计算
cudaStreamWaitEvent(streamB, event, 0);
cudaMemcpyAsync(h_out, d_out, size, cudaMemcpyDeviceToHost, streamB);你想想看，如果没有事件，流B怎么知道流A算完了？用 cudaStreamSynchronize 会阻塞CPU，用事件则只阻塞流B，CPU还能继续干别的活。
关键理解：事件是流间同步的轻量级方案。它不阻塞CPU，只阻塞指定流。这在构建复杂流水线时非常有用。
2.5 实战中的流管理
说了这么多理论，咱们聊聊实际项目中怎么用。我个人习惯把流分成三类：
- 计算流：专门跑kernel，一般2-4个
- 拷贝流：专门做H2D和D2H，1-2个
- 控制流：处理同步和事件，通常用默认流
为什么要分开？因为拷贝和计算可以重叠。你让拷贝流在后台搬数据，计算流在前台做运算，GPU的利用率就能拉满。
我曾经优化过一个视频处理管线，最初所有操作都在一个流里，GPU利用率只有40%。后来拆成3个流——一个做解码拷贝，一个做推理计算，一个做后处理拷贝。利用率直接飙到85%以上。
嗯，这里要提醒一句：流不是越多越好。GPU的硬件并发能力有限，一般4-8个流就够用了。开太多流反而会增加调度开销。
个人经验：调试多流程序时，可以用 nsys 或 nvvp 查看时间轴。如果发现流之间有空隙，说明同步点设置不合理。我每次优化多流代码，第一件事就是看profiler的时间线。
最后总结一下：CUDA流是并发执行的基础。理解默认流和非默认流的区别，掌握同步与异步的时机，用好事件做流间协调——这三件事做好了，你的GPU程序性能就不会差。