1、CUDA流基础:CUDA流的概念、默认流与创建显式流、流的同步机制
各位同学,咱们今天来聊聊CUDA流。这东西说白了,就是GPU里的一条“任务流水线”。你想想看,CPU可以一边听歌一边写代码,靠的是多任务调度。GPU也一样,它需要同时处理多个计算任务,这时候就得靠CUDA流来帮忙了。
1.1 CUDA流的概念
CUDA流,官方定义是一系列按顺序执行的GPU操作。我个人习惯把它想象成一条“传送带”。你把任务放到传送带上,它就会按顺序执行。但如果你有多条传送带,它们就可以并行跑。
为什么会需要这个?我刚开始做CUDA开发时,遇到过一个大坑。那时候我写了一个程序,先做矩阵乘法,再做向量加法。结果发现,GPU一直在“摸鱼”——它得等矩阵乘法算完了,才开始向量加法。说白了,就是串行执行,浪费了GPU的并行能力。
后来我明白了,CUDA流就是用来解决这个问题的。它允许你把不相关的任务放到不同的流里,让它们同时执行。嗯,这里要注意:同一个流里的操作是严格按顺序的,但不同流之间可以乱序执行。
核心要点:CUDA流是GPU上的一组顺序操作序列。不同流之间可以并发执行,但同一流内的操作必须按提交顺序执行。
1.2 默认流与创建显式流
每个CUDA程序启动时,都会自带一个“默认流”。你如果不显式创建流,所有操作都会丢进这个默认流里。说白了,默认流就是一条“单车道”,所有车都得排队走。
我记得有一次,我接手一个老项目,里面所有核函数调用都没指定流。结果性能惨不忍睹。我当时就想:这不行,得给它开几条“快车道”。
创建显式流其实很简单,用 cudaStreamCreate() 就行。看代码:
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在stream1中执行核函数
kernel1<<<grid, block, 0, stream1>>>(d_data1);
// 在stream2中执行核函数
kernel2<<<grid, block, 0, stream2>>>(d_data2);
// 别忘了销毁流
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);
你看,核函数调用时,第四个参数就是流句柄。如果不传(或者传0),就默认用默认流。我个人建议:只要你的程序里有两个以上不相关的计算任务,就果断用显式流。
小技巧:创建流时可以用 cudaStreamCreateWithFlags() 设置非阻塞标志。这样默认流就不会被其他流阻塞了。我一般这么用:cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking)。
1.3 流的同步机制
流搞多了,同步就成了大问题。你想想看,如果stream1算完了结果,stream2要用这个结果,那stream2就得等着。怎么等?靠同步机制。
CUDA提供了几种同步方式:
- cudaDeviceSynchronize():阻塞CPU,直到所有流上的所有操作都完成。这招最暴力,但效率最低。我一般只在程序最后用一次。
- cudaStreamSynchronize(stream):只等某一个流完成。比上面那个温和一些。
- cudaEventSynchronize(event):用事件来做更精细的同步。这个我最常用。
我曾经在一个项目里,需要让stream2等stream1算完某个中间结果。一开始我用 cudaDeviceSynchronize(),结果CPU被卡死,GPU利用率也上不去。后来换成事件同步,性能直接翻倍。
事件同步的用法是这样的:
cudaEvent_t event;
cudaEventCreate(&event);
// stream1中记录事件
kernel1<<<grid, block, 0, stream1>>>(d_data);
cudaEventRecord(event, stream1);
// stream2等待事件完成
cudaStreamWaitEvent(stream2, event);
kernel2<<<grid, block, 0, stream2>>>(d_data);
cudaEventDestroy(event);
这里要注意:cudaEventRecord() 是在stream1里插入一个“标记”,cudaStreamWaitEvent() 是让stream2在这个标记完成前不要执行后续操作。说白了,就是给stream2设了个“红绿灯”。
避坑指南:我曾经犯过一个错误——在默认流里用事件去同步显式流。结果发现默认流会隐式阻塞所有显式流,导致同步失效。后来我查了文档才知道,默认流有“隐式同步”的特性。所以我的建议是:要么全用默认流,要么全用显式流,别混着用。
最后,我整理了一个简单的对比表,方便你理解不同同步方式的适用场景:
| 同步方式 | 作用范围 | 性能影响 | 推荐场景 |
|---|---|---|---|
| cudaDeviceSynchronize | 所有流 | 高(阻塞CPU) | 程序结束、调试 |
| cudaStreamSynchronize | 单个流 | 中 | 等待某个流完成 |
| cudaEventSynchronize | 事件点 | 低 | 流间依赖、性能优化 |
嗯,到这里,CUDA流的基础就讲完了。说白了,流就是让你把GPU的并行能力榨干的一个工具。下一节我们会聊更高级的流用法,比如流回调、MPS这些。到时候你会发现,流的玩法比你想的要多得多。