2、CUDA流与并发：CUDA流的概念、默认流与非默认流、流同步机制、流回调函数

各位同学，咱们今天聊聊CUDA流。说实话，我刚接触CUDA那会儿，觉得流这玩意儿可有可无。不就是把任务排个队嘛，有啥好学的？直到我在一个实时渲染项目里被性能瓶颈卡了整整两周……嗯，从那以后，我再也不敢小看流了。

2.1 CUDA流到底是什么？

说白了，CUDA流就是一个操作队列。你把kernel启动、内存拷贝这些操作丢进去，GPU就按顺序执行。但关键来了——不同流之间的操作是可以并发的。

我习惯把流想象成工厂里的流水线。一条流水线只能顺序干活，但多条流水线可以同时开工。你想想看，如果你的GPU有80个SM（流多处理器），却只用一条流水线，那多浪费啊。

2.2 默认流与非默认流

每个CUDA程序启动时，系统会自动给你创建一个默认流。你平时写的那些不带流的kernel调用，其实都跑在默认流里。

// 默认流调用
kernel<<<grid, block>>>(d_data);

// 等价于
kernel<<<grid, block, 0, 0>>>(d_data);  // 0就是默认流

非默认流需要你手动创建。我个人习惯用cudaStreamCreate()，简单直接。

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

// 在指定流中启动kernel
kernel1<<<grid, block, 0, stream1>>>(d_data1);
kernel2<<<grid, block, 0, stream2>>>(d_data2);

这里有个坑，我曾经踩过——默认流在某些情况下会阻塞其他流。在CUDA 6.0之前，默认流是所谓的「阻塞流」，它会跟所有其他流同步。后来引入了「每线程默认流」的概念，才解决了这个问题。

2.3 流同步机制

流同步，说白了就是让CPU等GPU干完活。但怎么等、等多久，这里面学问大了。

最粗暴的方式是cudaDeviceSynchronize()——等所有流都完事。但这种方式效率很低，你想想看，如果只有一条流在干活，其他流早就闲着了，你还得等最慢的那个。

更好的做法是用cudaStreamSynchronize()，只等特定的一条流。

cudaStreamSynchronize(stream1);  // 只等stream1
// 此时stream2可能还在跑，你可以继续做其他CPU工作

还有一种更精细的同步方式——事件（Event）。我特别喜欢用事件来做流间同步。

cudaEvent_t event;
cudaEventCreate(&event);

// 在stream1中记录事件
cudaEventRecord(event, stream1);

// 在stream2中等待事件
cudaStreamWaitEvent(stream2, event, 0);

// 这样stream2就会等stream1干完某件事后才开始

我在做多GPU协同计算时，经常用事件来协调不同GPU上的流。比如GPU0算完中间结果后，通过事件通知GPU1开始消费。这种方式比全局同步高效得多。

2.4 流回调函数

流回调是个很有意思的特性。它允许你在流中所有操作完成后，自动执行一个CPU函数。

void CUDART_CB myCallback(cudaStream_t stream, cudaError_t status, void *data) {
    printf("流 %p 执行完毕，状态: %d\n", stream, status);
    // 这里可以做清理工作，或者触发后续操作
}

// 注册回调
cudaStreamAddCallback(stream, myCallback, NULL, 0);

嗯，这里要注意——回调函数是在CPU上执行的，但由GPU驱动触发。这意味着你不能在回调里调用CUDA API（除了少数几个）。我曾经犯过这个错，在回调里又启动了一个kernel，结果程序直接崩了。

回调的典型应用场景是什么？我举两个例子：

• 异步数据后处理：GPU算完结果后，自动触发CPU进行数据整理或写入磁盘
• 流水线控制：一条流完成后，自动通知下一条流开始工作

// 新接口方式
void CUDART_CB myHostFunc(void *data) {
    printf("主机函数被调用\n");
}

cudaLaunchHostFunc(stream, myHostFunc, NULL);

2.5 实战经验总结

说了这么多，我给大家总结几条实战经验：

1. 流的数量不是越多越好。我见过有人创建了64个流，结果性能反而下降了。一般来说，流的数量跟GPU的硬件队列数相关，4-8个流通常就够了。
2. 注意内存传输和kernel的搭配。用流做并发时，最好让内存传输和计算重叠。比如stream1在算数据，stream2在拷数据。
3. 调试多流程序很痛苦。我建议先用cuda-gdb或Nsight Systems做可视化分析，看看各个流的时间线，一眼就能看出哪里没并发起来。

最后说一句，CUDA流这东西，光看书是学不会的。你得动手写代码，跑起来，用工具看性能数据。我第一次成功让两个流并发跑起来时，那个兴奋劲儿，到现在还记得。好了，下一章咱们聊聊多GPU编程，那才是真正的大场面。