CUDA流（Stream）基础：概念、默认流与非默认流、创建与销毁

大家好，欢迎来到《计算与通信重叠技巧实战课程》的第三讲。今天我们来聊聊CUDA流——这个在GPU编程里看似基础、实则暗藏玄机的概念。

说实话，我刚开始接触CUDA时，对流的概念也是一知半解。觉得不就是把任务排个队嘛，有啥好研究的？直到有一次，我写了一个图像处理管线，数据拷贝和核函数执行串行跑，性能死活上不去。后来一位老同事看了一眼，说：“你试试用流。” 嗯，从那以后，我才真正理解了流的价值。

什么是CUDA流？

简单来说，CUDA流就是一个操作队列。你把一系列GPU操作（比如内存拷贝、核函数启动）按顺序丢进去，GPU就会按这个顺序执行。不同流之间的操作，则可以并发执行。

你想想看，如果没有流，所有操作都挤在一条单行道上。有了流，就像把单车道变成了多车道。不同的任务可以并行跑，互不干扰。

默认流 vs 非默认流

每个CUDA程序启动时，都会自动创建一个默认流（也叫NULL流）。你平时写的cudaMemcpy、kernel<<>>，如果不指定流，默认就丢进这个默认流里。

这里有个坑，我必须要提醒你：默认流的行为在CUDA 6.0前后发生了变化。

• CUDA 6.0之前：默认流是阻塞流。它会阻塞其他流中的操作，说白了就是“我干活的时候你们都等着”。
• CUDA 6.0之后：默认流变成了非阻塞流（per-thread default stream）。每个线程有自己的默认流，互不干扰。

我个人习惯，在写新代码时，永远显式创建非默认流。为什么？因为依赖默认流的行为，容易写出“看起来对、但换个环境就崩”的代码。我曾经在一个项目里，就因为默认流的阻塞特性，导致两个本该并行的核函数串行执行，排查了整整两天。

流的创建与销毁

创建流很简单，就一个API：

cudaStream_t stream;
cudaError_t err = cudaStreamCreate(&stream);

销毁流也很直接：

cudaError_t err = cudaStreamDestroy(stream);

嗯，这里有几个细节要注意：

• 创建流时，可以指定标志位来控制流的行为。比如cudaStreamNonBlocking，让这个流不阻塞默认流。
• 销毁流时，如果流里还有未完成的操作，cudaStreamDestroy会立即返回，不会等待操作完成。但GPU会继续执行这些操作，直到完成。如果你在销毁后立即分配新的资源，可能会出问题。

我建议的实践是：

1. 在程序初始化时创建所有需要的流。
2. 在程序退出前，先同步所有流，再销毁。
3. 不要频繁创建和销毁流，开销不小。

流的使用示例

光说不练假把式。我们来看一个实际例子：

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 在stream1中执行数据拷贝和核函数
cudaMemcpyAsync(d_data1, h_data1, size, cudaMemcpyHostToDevice, stream1);
kernel1<<<grid, block, 0, stream1>>>(d_data1);

// 在stream2中执行另一组操作
cudaMemcpyAsync(d_data2, h_data2, size, cudaMemcpyHostToDevice, stream2);
kernel2<<<grid, block, 0, stream2>>>(d_data2);

// 等待所有流完成
cudaDeviceSynchronize();

cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

你看，cudaMemcpyAsync和核函数启动都指定了流参数。这样，stream1和stream2里的操作就可以并发执行了。数据拷贝和计算可以重叠，效率自然就上去了。

流的同步机制

流之间需要同步时，可以用cudaStreamSynchronize等待某个流完成，或者用cudaEvent实现更精细的同步。

举个例子：

cudaEvent_t event;
cudaEventCreate(&event);

// stream1做完后，记录事件
kernel1<<<grid, block, 0, stream1>>>(...);
cudaEventRecord(event, stream1);

// stream2等待事件，再执行
cudaStreamWaitEvent(stream2, event, 0);
kernel2<<<grid, block, 0, stream2>>>(...);

这样，stream2里的kernel2就会等stream1里的kernel1完成后再执行。但stream2里的其他操作（比如前面的数据拷贝）不受影响。这种细粒度同步，是优化性能的关键。

知识体系总览

下面这张图，帮你梳理本章的核心逻辑：

这张图展示了CUDA流的五个核心维度：概念、默认流与非默认流、创建与销毁、同步机制、以及最终的实战目标——计算与通信重叠。每个维度都是我们后续课程的基础。

总结

好了，这一讲的内容就到这里。总结一下：

• CUDA流是操作队列，让GPU操作可以并发执行。
• 默认流有历史遗留问题，建议显式创建非默认流。
• 创建和销毁流很简单，但要注意生命周期管理。
• 同步可以用cudaStreamSynchronize或cudaEvent，后者更灵活。

下一讲，我们会深入探讨如何用流实现真正的计算与通信重叠。到时候，我会分享一个我实际优化过的案例，从2倍加速到5倍加速的过程。敬请期待。