3、默认流与自定义流:默认流的阻塞特性与自定义流的非阻塞优势
说到CUDA流,我不得不先提一个常见的误解。很多初学者以为,把数据拷贝到GPU、执行核函数、再拷回来,这三步是自动并行执行的。其实不然。默认流(也叫NULL流)有个让人头疼的特性——它会阻塞其他流。
说白了,默认流就像一个霸道总裁。它不干完活,别人别想动。我早期做项目时就被这个特性坑过,当时一个简单的图像处理程序,GPU利用率死活上不去。后来才发现,是默认流在作祟。
3.1 默认流的阻塞特性
默认流是CUDA自动创建的流。你如果不显式创建流,所有操作都默认在这个流里执行。它的阻塞特性体现在哪里?
- 隐式同步:默认流中的操作会与所有其他流同步
- 串行执行:默认流内的操作严格按顺序执行
- 阻塞其他流:默认流启动的操作,会等待其他流完成;反之亦然
核心问题:默认流的存在,让多流并行变得几乎不可能。你想想看,只要有一个操作在默认流里,其他流就得等着。
为什么会这样?这是CUDA设计的历史遗留问题。为了兼容老代码,默认流保留了这种阻塞行为。我个人习惯是,新项目一律不用默认流。
3.2 自定义流的非阻塞优势
自定义流就灵活多了。你可以创建多个流,每个流独立执行自己的任务。流与流之间默认不阻塞,这才是真正的异步。
我记得有一次优化一个视频处理管线。输入是4K视频流,需要做色彩空间转换、降噪、缩放三个步骤。如果用默认流,只能串行做。换成三个自定义流后,吞吐量直接翻倍。
创建和使用自定义流
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 异步拷贝到设备
cudaMemcpyAsync(d_data1, h_data1, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_data2, h_data2, size, cudaMemcpyHostToDevice, stream2);
// 核函数在不同流中执行
kernel1<<<grid, block, 0, stream1>>>(d_data1);
kernel2<<<grid, block, 0, stream2>>>(d_data2);
// 异步拷贝回主机
cudaMemcpyAsync(h_result1, d_data1, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(h_result2, d_data2, size, cudaMemcpyDeviceToHost, stream2);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);
注意看,每个操作都指定了流参数。这样GPU就能同时处理两个数据流。嗯,这里要注意:cudaMemcpyAsync需要主机内存是页锁定内存(pinned memory),否则会退化成同步拷贝。
3.3 默认流 vs 自定义流:对比表格
| 特性 | 默认流 | 自定义流 |
|---|---|---|
| 创建方式 | 自动创建 | cudaStreamCreate |
| 阻塞行为 | 阻塞其他流 | 非阻塞(默认) |
| 同步方式 | 隐式同步 | 显式同步(cudaStreamSynchronize) |
| 并行能力 | 无 | 支持多流并行 |
| 适用场景 | 简单串行任务 | 复杂管线、多数据流 |
3.4 避坑指南:我曾经踩过的坑
我曾经在一个金融计算项目中,用默认流跑蒙特卡洛模拟。数据量很大,我天真地以为GPU会自动并行。结果性能惨不忍睹。排查了两天,才发现是默认流把所有操作串行化了。
后来改成自定义流,每个模拟路径一个流。配合cudaStreamWaitEvent做细粒度同步,性能提升了4倍。这里分享几个经验:
- 别混用默认流和自定义流:一旦混用,默认流的阻塞特性会拖慢所有流
- 流不是越多越好:GPU硬件能同时执行的流数量有限,一般4-8个就够
- 注意流同步开销:频繁同步会抵消并行收益
小技巧:可以用cudaStreamNonBlocking标志创建自定义流,这样它连默认流都不阻塞。但慎用,容易引发数据竞争。
3.5 知识体系图:流的工作机制
这张图很直观。默认流就像单车道,所有车都得排队。自定义流就像多车道,车可以并行跑。但注意,默认流这个「霸道总裁」一旦出现,其他车道也得等它先走。
警告:不要以为创建了自定义流就万事大吉。如果核函数之间没有数据依赖,并行没问题。但如果有依赖关系,必须用cudaStreamWaitEvent或cudaEventSynchronize做同步。否则数据错乱,调试起来很痛苦。
最后说一句:默认流不是不能用。如果你的任务很简单,比如就一个核函数,用默认流反而省事。但一旦涉及多数据流、管线化处理,请果断抛弃默认流,拥抱自定义流。这是我从无数次性能优化中总结出来的经验。