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 知识体系图:流的工作机制

主机(CPU) 设备(GPU) 默认流(阻塞) 操作1 → 操作2 → 操作3 串行执行,阻塞其他流 自定义流1 操作A → 操作B → 操作C 自定义流2 操作X → 操作Y → 操作Z 流1和流2可以并行执行 默认流会阻塞流1和流2 默认流(阻塞) 自定义流(非阻塞)

这张图很直观。默认流就像单车道,所有车都得排队。自定义流就像多车道,车可以并行跑。但注意,默认流这个「霸道总裁」一旦出现,其他车道也得等它先走。

警告:不要以为创建了自定义流就万事大吉。如果核函数之间没有数据依赖,并行没问题。但如果有依赖关系,必须用cudaStreamWaitEventcudaEventSynchronize做同步。否则数据错乱,调试起来很痛苦。

最后说一句:默认流不是不能用。如果你的任务很简单,比如就一个核函数,用默认流反而省事。但一旦涉及多数据流、管线化处理,请果断抛弃默认流,拥抱自定义流。这是我从无数次性能优化中总结出来的经验。


专注资料整理