4、CUDA流与并发执行：Stream的概念、默认流与非默认流、Stream同步、Stream优先级、使用Stream实现Kernel执行与数据传输的重叠
各位同学，咱们今天来聊聊CUDA编程里一个特别重要的概念——Stream（流）。
说实话，我刚开始接触CUDA的时候，觉得这玩意儿不就是把数据扔给GPU算嘛，有啥复杂的？后来在实际项目中，发现GPU利用率老是上不去，显存带宽也跑不满，才意识到自己根本没用好Stream。说白了，Stream就是让你手里的GPU真正忙起来的关键。
4.1 什么是CUDA Stream？
Stream，你可以把它理解成GPU上的一条任务流水线。所有提交到同一个Stream里的操作（比如Kernel执行、数据拷贝），都会按照你提交的顺序，一个接一个地执行。
但不同Stream里的操作呢？它们可以乱序执行，甚至可以同时执行——这就是并发。
我打个比方：你开了一家奶茶店，只有一个店员（GPU）。Stream就是不同的订单队列。如果你把所有订单都塞进同一个队列，那店员只能一杯一杯地做。但如果你把订单分到两个队列，一个做奶茶，一个做果茶，那店员就可以穿插着做，效率自然就上来了。
核心要点：Stream是CUDA实现并发执行的基本单位。每个Stream维护一个独立的操作序列，不同Stream之间的操作可以并发执行。
4.2 默认流 vs 非默认流
CUDA里有两种流：默认流（也叫NULL流）和非默认流（你显式创建的流）。
默认流是每个CUDA上下文自带的，你如果不显式创建Stream，所有操作都会跑到默认流里。嗯，这里要注意：默认流的行为在CUDA 6.0前后有过一次大改。
- CUDA 6.0之前：默认流是阻塞流。意思是，如果你在默认流里提交了一个操作，它会阻塞其他所有流里的操作，直到它执行完。
- CUDA 6.0之后：默认流变成了非阻塞流。它不再阻塞其他流了，但为了兼容旧代码，你可以通过编译选项切换回旧行为。
我个人习惯是：永远不要依赖默认流来做并发。显式创建非默认流，代码更清晰，行为也更可控。
创建非默认流的API很简单：
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 使用流
kernel1<<<grid, block, 0, stream1>>>(...);
kernel2<<<grid, block, 0, stream2>>>(...);
// 销毁流
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);你看，创建、使用、销毁，三步走。Kernel启动时，第四个参数就是Stream句柄。不传的话，默认就是NULL流。
4.3 Stream同步
并发虽好，但同步不能少。你想想看，如果Stream A在写数据，Stream B在读同一块数据，那不乱套了？
CUDA提供了几种同步方式：
| 同步方式 | API | 说明 |
|---|---|---|
| 设备同步 | cudaDeviceSynchronize() | 阻塞CPU，直到设备上所有Stream的所有操作完成 |
| 流同步 | cudaStreamSynchronize(stream) | 阻塞CPU，直到指定Stream的所有操作完成 |
| 事件同步 | cudaEventSynchronize(event) | 阻塞CPU，直到指定事件被记录 |
| 流等待事件 | cudaStreamWaitEvent(stream, event) | 让一个Stream等待另一个Stream的事件 |
我在项目中遇到过一个问题：两个Stream之间需要协调，但我不想让CPU等着。这时候cudaStreamWaitEvent就派上用场了。它让Stream之间可以互相等待，而CPU该干嘛干嘛。
小技巧：能用cudaStreamSynchronize就别用cudaDeviceSynchronize。后者太粗暴了，会把所有Stream都停下来，影响并发效率。
4.4 Stream优先级
CUDA支持给Stream设置优先级。高优先级的Stream里的操作，会优先被调度执行。
这个功能在CUDA 4.0之后就有了。用法如下：
int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
cudaStream_t high_pri_stream, low_pri_stream;
cudaStreamCreateWithPriority(&high_pri_stream, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&low_pri_stream, cudaStreamNonBlocking, priority_low);priority_high是最高优先级（数值最小），priority_low是最低优先级（数值最大）。
我曾经在一个实时渲染项目里用过这个特性。渲染管线里的Kernel必须尽快执行，不然画面会卡顿。而一些后处理Kernel可以慢一点。把渲染Kernel放到高优先级Stream里，后处理放到低优先级Stream里，效果立竿见影。
注意：不是所有GPU都支持Stream优先级。你可以用cudaDeviceGetStreamPriorityRange查询一下。如果返回的priority_low == priority_high，说明不支持。
4.5 使用Stream实现Kernel执行与数据传输的重叠
好了，前面铺垫了这么多，终于到实战环节了。Stream最经典的应用场景，就是让Kernel计算和数据传输重叠起来。
你想想看，传统的做法是：
- 把数据从CPU拷贝到GPU（cudaMemcpy）
- 在GPU上执行Kernel
- 把结果从GPU拷贝回CPU
这三步是串行的。GPU在拷贝数据的时候，计算单元闲着；GPU在计算的时候，拷贝引擎闲着。浪费啊！
用Stream可以这样优化：
// 假设有4个数据块
const int N = 4;
float *d_data[N], *h_data[N];
cudaStream_t streams[N];
// 创建流和分配内存
for (int i = 0; i < N; i++) {
cudaStreamCreate(&streams[i]);
cudaMalloc(&d_data[i], size);
// h_data[i] 假设已经准备好
}
// 流水线执行
for (int i = 0; i < N; i++) {
// 异步拷贝到GPU
cudaMemcpyAsync(d_data[i], h_data[i], size,
cudaMemcpyHostToDevice, streams[i]);
// 启动Kernel
kernel<<<grid, block, 0, streams[i]>>>(d_data[i]);
// 异步拷贝回CPU
cudaMemcpyAsync(h_data[i], d_data[i], size,
cudaMemcpyDeviceToHost, streams[i]);
}
// 等待所有完成
cudaDeviceSynchronize();
// 清理
for (int i = 0; i < N; i++) {
cudaStreamDestroy(streams[i]);
}这段代码的精髓在于：每个Stream内部的操作是串行的，但不同Stream之间是并发的。当Stream 0在计算时，Stream 1可能正在拷贝数据。这样一来，计算和传输就重叠起来了。
我曾经在一个视频处理项目里用这个技巧，把整体吞吐量提升了将近2倍。当然，前提是你的GPU有足够多的拷贝引擎和计算单元。你可以用cudaDeviceGetAttribute查询一下设备属性。
关键点：要实现重叠，必须使用cudaMemcpyAsync（异步拷贝），不能使用cudaMemcpy（同步拷贝）。同步拷贝会阻塞当前Stream，破坏并发性。
另外，别忘了使用固定内存（pinned memory）。异步拷贝需要主机内存是页锁定的，否则会退化成同步拷贝。用cudaMallocHost分配就行。
嗯，说到这里，我想起一个坑。有一次我写代码，发现重叠效果不明显，查了半天才发现是主机内存没固定。换成cudaMallocHost之后，性能立马就上来了。所以，这个细节一定要注意。
4.6 本章小结
Stream是CUDA并发编程的基石。用好Stream，你的GPU才能真正发挥出它的潜力。
- 默认流方便但不灵活，建议显式创建非默认流
- 同步方式要选对，别动不动就cudaDeviceSynchronize
- Stream优先级在实时场景下很有用
- 计算与传输重叠是Stream最经典的应用，记得用固定内存和异步拷贝
下一章，我们会深入讨论事件（Event）和更高级的同步技巧。到时候，你会发现Stream和Event配合起来，能玩出更多花样。