1. 为什么需要StreamGPU执行模型下的性能黑洞1.1 等等GPU不是已经很快了吗很多刚接触CUDA的开发者都有过这样的体验花了一周时间把算法改成kernel跑起来一看确实比CPU快了不少但仔细一测GPU利用率只有三成显存带宽没吃满SM有一大半在闲着。这时候你通常会怀疑kernel写得不够好于是开始调block数量、调共享内存、调访存模式折腾一圈下来提升有限。其实问题可能根本不在kernel内部而在kernel之间——你压根没有让GPU真正“忙起来”。CUDA Stream流就是解决这个问题的核心机制。它不是一个新算法也不是某种黑魔法而是一条“任务流水线”的概念GPU上的所有操作kernel计算、显存拷贝、事件记录都会被放入某个流中排队执行。如果你只会用默认流那所有操作都挤在同一条队伍里一个跑完才能轮到下一个。而如果你擅长用多个流就能让互不依赖的计算和拷贝同时进行把GPU那张“多核大网”真正铺开。我见过太多人跳过Stream直接学什么tensor core、warp shuffle结果连最基本的并行度都没榨干净。这一篇我先把Stream的底层逻辑和基础API讲透下一篇再聊事件同步和多流协作的高级玩法。1.2 默认流的“串行陷阱”一次真实的性能浪费先看一个非常典型的例子。假设你在做视频帧处理每一帧都要先从CPU拷贝到GPUH2D在GPU上做一次滤波kernel再把结果拷回CPUD2H。初学者往往会这样写for (int i 0; i frameCount; i) { cudaMemcpy(d_input i * frameSize, h_input i * frameSize, frameBytes, cudaMemcpyHostToDevice); myFilterKernelblocks, threads(d_input i * frameSize, d_output i * frameSize, frameSize); cudaMemcpy(h_output i * frameSize, d_output i * frameSize, frameBytes, cudaMemcpyDeviceToHost); }这段代码在默认流里执行时完整的时间线是拷贝、计算、拷贝、计算……所有操作严格串行。但是仔细想想GPU引擎其实分两大类拷贝引擎Copy Engine负责DMA传输计算引擎SM负责跑kernel。我只需要让第i帧的kernel在执行时第i1帧的H2D拷贝同时进行两个引擎就能并行工作。仅仅这一步帧处理吞吐量就能明显提升。实测过一个1920×1080灰度图的滤波流程单流版本的H2D拷贝约0.6mskernel约1.2msD2H约0.6ms单帧总耗时接近2.4ms用双流把拷贝和计算重叠后吞吐能提升30%~40%。这还只是最简单的场景如果kernel足够长重叠收益会更明显。所以我说理解Stream的第一个价值就是让你意识到“GPU不是只能做一件事”。2. Stream基础API创建、使用、销毁的一整套规矩2.1 cudaStreamCreate与cudaStreamDestroy的细节Stream的API极其简洁创建和销毁就是两个函数cudaStream_t stream; cudaError_t err cudaStreamCreate(stream); // 使用 cudaStreamDestroy(stream);但细节藏在“使用”里。创建流之后所有能接受流参数的CUDA操作都可以指定到这个流上。最常见的三类是kernel启动myKernelgrid, block, sharedMemSize, stream(args...)异步拷贝cudaMemcpyAsync(dst, src, size, kind, stream)事件操作cudaEventRecord(event, stream)这里有个很多人忽略的关键点cudaMemcpyAsync才是异步版本普通的cudaMemcpy即使传了流参数也依然是同步行为实际上cudaMemcpy不接受流参数编译器会报错。你要做流水线重叠必须用Async版本。另外一个常用的创建方式是带标志位cudaStreamCreateWithFlags(stream, cudaStreamNonBlocking);默认创建的流跟默认流即stream 0之间存在隐式同步关系而cudaStreamNonBlocking可以让这个新流不被默认流“管住”。具体的同步语义我建议你先记在脑子里默认流是一条“老大流”普通的命名流在默认流面前是要主动让路的。这个规则的细节我在第4节会详细讲。2.2 三种必须掌握的流相关操作除了创建销毁日常用得最多的还有三个操作cudaStreamSynchronize(stream); // 阻塞CPU直到该流所有操作完成 cudaStreamQuery(stream); // 非阻塞查询返回cudaSuccess表示完成 cudaStreamWaitEvent(stream, event); // 让流等待某个事件cudaStreamSynchronize在多流场景下是调试利器——你可以在程序关键节点确认某个流真的算完了。但别在性能敏感的代码里到处撒同步点否则前面做的并行重叠就全白费了。cudaStreamQuery适合做“轮询”场景。比如你在等一个流算完但又不想彻底阻塞主线程可以每隔一小段时间查询一次状态。cudaStreamWaitEvent则是跨流依赖的核心如果流B的kernel需要用到流A的结果就让B等待A上的某个事件。这比直接cudaDeviceSynchronize精度高得多不会把无关的流也一起锁住是精细调度必不可少的工具。2.3 流的生命周期与资源管理流不是用完就自动销毁的。它在创建时会占用GPU上下文资源包括内部的命令缓冲区、事件对象等。如果一个程序反复创建和销毁大量流你会发现CUDA context的内存占用持续波动甚至触发驱动层的资源回收开销。我自己的习惯是流的数量按需创建最常用的是2~4个超过16个流的场景很少长期运行的服务型程序应该在初始化阶段就建好流运行期复用而不是每帧都新建。提示调用cudaStreamDestroy时如果流内还有未完成的操作CUDA会自动阻塞直到操作完成再销毁。所以没必要在销毁前手动cudaStreamSynchronize但如果你希望“销毁但不等待”那需要用cudaStreamDestroy以外的机制来管理这个比较复杂一般场景用不到。3. 多Stream并发的硬件真相调度器、SM占用与背压3.1 GPU内部到底是怎么调度多个流的很多人以为“多流并发”就是GPU把多个kernel同时塞进所有SM里公平地运行这个理解不完全对。GPU的硬件调度器Work Distributor负责把线程块thread block分发到各个SM上。当一个流中的kernel启动时它的一堆线程块会进入调度队列如果此时还有其他流的kernel也在排队调度器会尽可能把不同流的线程块混合分发到不同SM上。问题来了一个SM能同时运行多少个线程块取决于每个线程块的资源占用寄存器、共享内存以及SM本身的硬件限制。比如一个SM最多能驻留2048个线程如果你的kernel每个block用了512个线程那最多驻留4个block如果每个block又用了大量共享内存可能连2个block都放不下。SM资源被第一个kernel占满后第二个kernel的线程块只能排队等待。所以多流并发能不能生效本质上是看“资源还有没有空位”。在实际项目中我曾用Nsight Compute看过一个占用率76%的kernel它把SM的线程槽位占了大半结果我开了4个流并发效率依然很差——因为线程块根本插不进去。反而把kernel的block数调小比如从256线程改成128线程让多个kernel的block能混布在同一批SM上并发度才真正提上来。3.2 “背压”现象为什么流开多了反而更慢Stream的并发不是无限制的。当大量流同时启动kernel时GPU的工作队列会堆积调度器要在多个流之间切换、分配资源这本身有开销。表现就是流从4个增加到16个吞吐先升后降或者总延迟明显变长。从原理上讲这是“背压”backpressure机制在起作用当任务提交速度超过GPU执行速度时提交端会被阻塞。CUDA驱动的API提交本身有锁和队列管理开销流太多会导致CPU端提交kernel的时间变长反而盖过了GPU并行带来的收益。实测数据我做过一个矩阵乘法的多流测试数据分块后分别放到2、4、8、16个流里执行。2个流时吞吐提升最明显约1.7倍4个流进一步小幅提升8个流基本持平16个流出现了轻微下降。结论很明确流的数量不是越多越好要结合你的kernel规模、SM资源占用和硬件代数来测试。我一般建议从2和4开始试别一上来就8个流。3.3 用Nsight工具验证并发是否真的发生很多人在自己的机器上写了多流代码跑完发现性能没变化就开始怀疑Stream有没有用。我建议不要靠感觉直接用工具看。在Nsight Systems里你可以清楚看到每个流上kernel的时间轴如果多个流的kernel在时间轴上重叠说明并发真的发生了如果它们首尾相接说明资源不够或者同步没解除。这个工具比任何文字解释都直观。另外Nsight Compute里有个“SOL”Streaming Occupation Limit的指标直接显示当前kernel的线程块在SM上的驻留上限。SOL低于50%时多流并发的空间很小SOL很低说明SM资源被某个kernel的单个block占得太狠需要调整block尺寸和资源占用。4. 事件Event与细粒度同步把调度控制权抓在手里4.1 默认流与命名流的同步关系先暂时回到第2节提到的隐式同步。为什么默认流这么特殊因为CUDA为了兼容早期代码给默认流设计了一套“过度保守”的同步规则默认流上的操作会等待此前所有命名流上的操作完成同时其后所有命名流的操作也会等待默认流的操作完成。简单说默认流和命名流之间有一条隐形的“全等栅栏”谁都要等对方。这就是很多多流程序的隐形性能杀手——你在命名流里辛辛苦苦做了并行分解但中间某个地方调了默认流的kernel或cudaMemcpy它运行在默认流上把整个并行时间线全部打断了。我第2节提到的cudaStreamNonBlocking标志就是为了让命名流不跟默认流产生隐式同步。创建流时加上这个标志命名流就和默认流彻底解耦完全按你设定的依赖走。在现代CUDA代码里我强烈建议统一用cudaStreamNonBlocking创建流除非你明确需要老式同步语义。4.2 事件的本质GPU时间线上的里程碑Event事件可以理解为“GPU执行时间线上的一个标记点”。你在某个流上cudaEventRecord(event, stream)就是在该流的操作队列末尾插一个标记当GPU执行到这个标记时event被置为“已完成”。CPU可以查询event状态或者让其他流cudaStreamWaitEvent等待它。用event做跨流同步比cudaDeviceSynchronize精细得多。举个例子流A做完预处理流B和流C都需要它的结果但流B和流C之间没有依赖。你可以cudaEvent_t preDone; cudaEventCreate(preDone); // 流A执行预处理后记录事件 cudaEventRecord(preDone, streamA); // 让B和C等待事件而不是等待所有设备工作完成 cudaStreamWaitEvent(streamB, preDone); cudaStreamWaitEvent(streamC, preDone); // 此刻B和C可以并行执行A自身后续操作也可以继续如果不需要等待B/C这种“点对点”的同步方式非常强大它把“全局同步”拆成了“局部依赖”整个流水线的并行窗口一下就打开了。4.3 时间测量用事件替代cudaDeviceSynchronize事件还有一个特别重要的用途——精确测量kernel耗时。很多人用cudaDeviceSynchronize加CPU计时器这会包含内核启动的开销和CPU同步等待时间测出来偏大且不稳定。正确做法是cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start, stream); myKernelblocks, threads, 0, stream(); cudaEventRecord(stop, stream); cudaEventSynchronize(stop); float ms 0; cudaEventElapsedTime(ms, start, stop);这里的事件计时是在GPU时间线上测量的不会把CPU端的启动延迟算进去。排流水线时我通常会在每个流的“总入口”和“总出口”各放一个事件测出来的数值才是真实的任务耗时。5. 实战多Stream流水线让图像处理吞吐翻倍5.1 场景拆解数据分块与任务流水理论说太多了直接上一个完整的实战例子。假设我们要对一个由1000帧组成的图像序列做灰度化加高斯模糊。每帧处理流程分三段H2D把帧数据从CPU内存拷贝到GPU显存Kernel在GPU上执行灰度化模糊的kernelD2H把处理结果拷贝回CPU内存如果不使用流时间线是“拷-算-拷-算”严格串行。如果使用2个流我们可以把帧分成两组流0处理偶数帧流1处理奇数帧。由于帧之间没有数据依赖两个流上的kernel可以并发执行同时流0的kernel运行时流1的H2D拷贝引擎可以在另一个引擎上并行执行。5.2 双流实现的完整代码代码的核心并不复杂const int numStreams 2; cudaStream_t streams[numStreams]; for (int i 0; i numStreams; i) { cudaStreamCreateWithFlags(streams[i], cudaStreamNonBlocking); } for (int i 0; i frameCount; i) { int sid i % numStreams; int idx i / numStreams; char* d_in d_input sid * batchFrameBytes idx * frameBytes; char* d_out d_output sid * batchFrameBytes idx * frameBytes; const char* h_in h_input i * frameBytes; char* h_out h_output i * frameBytes; cudaMemcpyAsync(d_in, h_in, frameBytes, cudaMemcpyHostToDevice, streams[sid]); processFrameKernelblocks, threads, 0, streams[sid](d_in, d_out, frameSize); cudaMemcpyAsync(h_out, d_out, frameBytes, cudaMemcpyDeviceToHost, streams[sid]); } for (int i 0; i numStreams; i) { cudaStreamSynchronize(streams[i]); }注意这里给每个流单独分配了输入输出缓冲区的不同区域sid * batchFrameBytes idx * frameBytes避免不同流同时读写同一块显存。这一点非常重要——如果两个流共用一个缓冲区cudaMemcpyAsync只是发起异步传输实际传输可能在kernel运行时才发生数据竞争就会悄悄出现。5.3 实测数据性能提升与瓶颈分析我在一张RTX 3090上跑过这个例子图像尺寸1024×1024灰度化高斯模糊kernel约0.9msH2D拷贝约0.3msD2H拷贝约0.3ms。单流版本的每帧总耗时约1.5ms双流版本的理论下限是max(总拷贝时间, 总kernel时间)而不是三者之和实测帧间吞吐提升了约38%接近理论重叠极限。如果kernel本身的SM占用率偏高例如用了大量共享内存双流的提升会变小但通常依然优于单流。想进一步压榨性能可以用4个流但要注意数据分块的大小块太小会导致kernel启动开销占比升高。6. 实测中的坑与性能调优来自项目现场的踩坑笔记6.1 坑一cudaMemcpyAsync不是魔法别乱用cudaMemcpyAsync虽然叫异步但它只是把拷贝操作提交到指定流上真正的数据传输不一定和计算重叠。它有三个前提使用cudaMemcpyAsync不是cudaMemcpy传输的源和目标必须是“可分页内存”或“映射内存”其中可分页内存特指用cudaMallocHost分配的锁页内存pinned memory普通malloc分配的内存无法真正异步化拷贝的方向和设备支持情况要满足条件我见过太多人直接对普通malloc数组调cudaMemcpyAsync结果性能没变化还以为是流没用。要发挥异步拷贝的威力一定要用cudaMallocHost或cudaHostAlloc分配主机端内存否则异步调用会退化成同步传输。6.2 坑二数据依赖没理清并行变乱序多流并发的另一大坑是数据依赖。比如流B的kernel需要流A算出的中间结果但你忘了加cudaStreamWaitEvent结果流B在流A算完之前就开始读数据得到的是垃圾值。这种Bug很难复现因为GPU调度时机不固定可能跑100次才崩一次。我的排查习惯是一旦怀疑多流数据竞争先砍到单流跑一遍确认结果正确再一个流一个流地加回来每加一个流就检查一次结果。如果某个流加进来后结果开始错乱重点查它和上游流的数据依赖以及事件等待是否齐全。6.3 调优技巧流优先级、MPS与硬件差异最后分享几个压箱底的经验。第一CUDA流本身支持优先级cudaStreamCreateWithPriority可以给不同流设置不同优先级。实测中把延迟敏感的小kernel放在高优先级流把大kernel放在低优先级流交互场景的卡顿明显减少。第二如果你在数据中心GPU上跑多进程或多任务可以考虑开MPSMulti-Process Service。MPS能让多个进程的kernel共享GPU资源相当于把多流概念扩展到进程级别。但MPS配置复杂单机单卡场景用多流就够了。第三GPU架构对流的并发能力影响很大。Turing及之后架构的Hyper-Q增强了并发kernel的调度能力Ampere更进一步优化了SM资源感知。老架构上流并发效果差是正常的不必怀疑自己写错。我自己的切身体会是多流编程最难的并不是API本身——API就那几个难的是建立“GPU是一个并行流水线工厂”的思维模式。每次写kernel之前先画一条任务时间线标出哪里有等待、哪里可以重叠然后才动手写代码。这套方法比盲目堆流数量有效得多。 
长河君 · 2026/10/3 17:23:38 · 阅读 1.2k CUDA Stream深入解析：多流并发让GPU性能翻倍
1. 为什么需要StreamGPU执行模型下的性能黑洞1.1 等等GPU不是已经很快了吗很多刚接触CUDA的开发者都有过这样的体验花了一周时间把算法改成kernel跑起来一看确实比CPU快了不少但仔细一测GPU利用率只有三成显存带宽没吃满SM有一大半在闲着。这时候你通常会怀疑kernel写得不够好于是开始调block数量、调共享内存、调访存模式折腾一圈下来提升有限。其实问题可能根本不在kernel内部而在kernel之间——你压根没有让GPU真正“忙起来”。CUDA Stream流就是解决这个问题的核心机制。它不是一个新算法也不是某种黑魔法而是一条“任务流水线”的概念GPU上的所有操作kernel计算、显存拷贝、事件记录都会被放入某个流中排队执行。如果你只会用默认流那所有操作都挤在同一条队伍里一个跑完才能轮到下一个。而如果你擅长用多个流就能让互不依赖的计算和拷贝同时进行把GPU那张“多核大网”真正铺开。我见过太多人跳过Stream直接学什么tensor core、warp shuffle结果连最基本的并行度都没榨干净。这一篇我先把Stream的底层逻辑和基础API讲透下一篇再聊事件同步和多流协作的高级玩法。1.2 默认流的“串行陷阱”一次真实的性能浪费先看一个非常典型的例子。假设你在做视频帧处理每一帧都要先从CPU拷贝到GPUH2D在GPU上做一次滤波kernel再把结果拷回CPUD2H。初学者往往会这样写for (int i 0; i frameCount; i) { cudaMemcpy(d_input i * frameSize, h_input i * frameSize, frameBytes, cudaMemcpyHostToDevice); myFilterKernelblocks, threads(d_input i * frameSize, d_output i * frameSize, frameSize); cudaMemcpy(h_output i * frameSize, d_output i * frameSize, frameBytes, cudaMemcpyDeviceToHost); }这段代码在默认流里执行时完整的时间线是拷贝、计算、拷贝、计算……所有操作严格串行。但是仔细想想GPU引擎其实分两大类拷贝引擎Copy Engine负责DMA传输计算引擎SM负责跑kernel。我只需要让第i帧的kernel在执行时第i1帧的H2D拷贝同时进行两个引擎就能并行工作。仅仅这一步帧处理吞吐量就能明显提升。实测过一个1920×1080灰度图的滤波流程单流版本的H2D拷贝约0.6mskernel约1.2msD2H约0.6ms单帧总耗时接近2.4ms用双流把拷贝和计算重叠后吞吐能提升30%~40%。这还只是最简单的场景如果kernel足够长重叠收益会更明显。所以我说理解Stream的第一个价值就是让你意识到“GPU不是只能做一件事”。2. Stream基础API创建、使用、销毁的一整套规矩2.1 cudaStreamCreate与cudaStreamDestroy的细节Stream的API极其简洁创建和销毁就是两个函数cudaStream_t stream; cudaError_t err cudaStreamCreate(stream); // 使用 cudaStreamDestroy(stream);但细节藏在“使用”里。创建流之后所有能接受流参数的CUDA操作都可以指定到这个流上。最常见的三类是kernel启动myKernelgrid, block, sharedMemSize, stream(args...)异步拷贝cudaMemcpyAsync(dst, src, size, kind, stream)事件操作cudaEventRecord(event, stream)这里有个很多人忽略的关键点cudaMemcpyAsync才是异步版本普通的cudaMemcpy即使传了流参数也依然是同步行为实际上cudaMemcpy不接受流参数编译器会报错。你要做流水线重叠必须用Async版本。另外一个常用的创建方式是带标志位cudaStreamCreateWithFlags(stream, cudaStreamNonBlocking);默认创建的流跟默认流即stream 0之间存在隐式同步关系而cudaStreamNonBlocking可以让这个新流不被默认流“管住”。具体的同步语义我建议你先记在脑子里默认流是一条“老大流”普通的命名流在默认流面前是要主动让路的。这个规则的细节我在第4节会详细讲。2.2 三种必须掌握的流相关操作除了创建销毁日常用得最多的还有三个操作cudaStreamSynchronize(stream); // 阻塞CPU直到该流所有操作完成 cudaStreamQuery(stream); // 非阻塞查询返回cudaSuccess表示完成 cudaStreamWaitEvent(stream, event); // 让流等待某个事件cudaStreamSynchronize在多流场景下是调试利器——你可以在程序关键节点确认某个流真的算完了。但别在性能敏感的代码里到处撒同步点否则前面做的并行重叠就全白费了。cudaStreamQuery适合做“轮询”场景。比如你在等一个流算完但又不想彻底阻塞主线程可以每隔一小段时间查询一次状态。cudaStreamWaitEvent则是跨流依赖的核心如果流B的kernel需要用到流A的结果就让B等待A上的某个事件。这比直接cudaDeviceSynchronize精度高得多不会把无关的流也一起锁住是精细调度必不可少的工具。2.3 流的生命周期与资源管理流不是用完就自动销毁的。它在创建时会占用GPU上下文资源包括内部的命令缓冲区、事件对象等。如果一个程序反复创建和销毁大量流你会发现CUDA context的内存占用持续波动甚至触发驱动层的资源回收开销。我自己的习惯是流的数量按需创建最常用的是2~4个超过16个流的场景很少长期运行的服务型程序应该在初始化阶段就建好流运行期复用而不是每帧都新建。提示调用cudaStreamDestroy时如果流内还有未完成的操作CUDA会自动阻塞直到操作完成再销毁。所以没必要在销毁前手动cudaStreamSynchronize但如果你希望“销毁但不等待”那需要用cudaStreamDestroy以外的机制来管理这个比较复杂一般场景用不到。3. 多Stream并发的硬件真相调度器、SM占用与背压3.1 GPU内部到底是怎么调度多个流的很多人以为“多流并发”就是GPU把多个kernel同时塞进所有SM里公平地运行这个理解不完全对。GPU的硬件调度器Work Distributor负责把线程块thread block分发到各个SM上。当一个流中的kernel启动时它的一堆线程块会进入调度队列如果此时还有其他流的kernel也在排队调度器会尽可能把不同流的线程块混合分发到不同SM上。问题来了一个SM能同时运行多少个线程块取决于每个线程块的资源占用寄存器、共享内存以及SM本身的硬件限制。比如一个SM最多能驻留2048个线程如果你的kernel每个block用了512个线程那最多驻留4个block如果每个block又用了大量共享内存可能连2个block都放不下。SM资源被第一个kernel占满后第二个kernel的线程块只能排队等待。所以多流并发能不能生效本质上是看“资源还有没有空位”。在实际项目中我曾用Nsight Compute看过一个占用率76%的kernel它把SM的线程槽位占了大半结果我开了4个流并发效率依然很差——因为线程块根本插不进去。反而把kernel的block数调小比如从256线程改成128线程让多个kernel的block能混布在同一批SM上并发度才真正提上来。3.2 “背压”现象为什么流开多了反而更慢Stream的并发不是无限制的。当大量流同时启动kernel时GPU的工作队列会堆积调度器要在多个流之间切换、分配资源这本身有开销。表现就是流从4个增加到16个吞吐先升后降或者总延迟明显变长。从原理上讲这是“背压”backpressure机制在起作用当任务提交速度超过GPU执行速度时提交端会被阻塞。CUDA驱动的API提交本身有锁和队列管理开销流太多会导致CPU端提交kernel的时间变长反而盖过了GPU并行带来的收益。实测数据我做过一个矩阵乘法的多流测试数据分块后分别放到2、4、8、16个流里执行。2个流时吞吐提升最明显约1.7倍4个流进一步小幅提升8个流基本持平16个流出现了轻微下降。结论很明确流的数量不是越多越好要结合你的kernel规模、SM资源占用和硬件代数来测试。我一般建议从2和4开始试别一上来就8个流。3.3 用Nsight工具验证并发是否真的发生很多人在自己的机器上写了多流代码跑完发现性能没变化就开始怀疑Stream有没有用。我建议不要靠感觉直接用工具看。在Nsight Systems里你可以清楚看到每个流上kernel的时间轴如果多个流的kernel在时间轴上重叠说明并发真的发生了如果它们首尾相接说明资源不够或者同步没解除。这个工具比任何文字解释都直观。另外Nsight Compute里有个“SOL”Streaming Occupation Limit的指标直接显示当前kernel的线程块在SM上的驻留上限。SOL低于50%时多流并发的空间很小SOL很低说明SM资源被某个kernel的单个block占得太狠需要调整block尺寸和资源占用。4. 事件Event与细粒度同步把调度控制权抓在手里4.1 默认流与命名流的同步关系先暂时回到第2节提到的隐式同步。为什么默认流这么特殊因为CUDA为了兼容早期代码给默认流设计了一套“过度保守”的同步规则默认流上的操作会等待此前所有命名流上的操作完成同时其后所有命名流的操作也会等待默认流的操作完成。简单说默认流和命名流之间有一条隐形的“全等栅栏”谁都要等对方。这就是很多多流程序的隐形性能杀手——你在命名流里辛辛苦苦做了并行分解但中间某个地方调了默认流的kernel或cudaMemcpy它运行在默认流上把整个并行时间线全部打断了。我第2节提到的cudaStreamNonBlocking标志就是为了让命名流不跟默认流产生隐式同步。创建流时加上这个标志命名流就和默认流彻底解耦完全按你设定的依赖走。在现代CUDA代码里我强烈建议统一用cudaStreamNonBlocking创建流除非你明确需要老式同步语义。4.2 事件的本质GPU时间线上的里程碑Event事件可以理解为“GPU执行时间线上的一个标记点”。你在某个流上cudaEventRecord(event, stream)就是在该流的操作队列末尾插一个标记当GPU执行到这个标记时event被置为“已完成”。CPU可以查询event状态或者让其他流cudaStreamWaitEvent等待它。用event做跨流同步比cudaDeviceSynchronize精细得多。举个例子流A做完预处理流B和流C都需要它的结果但流B和流C之间没有依赖。你可以cudaEvent_t preDone; cudaEventCreate(preDone); // 流A执行预处理后记录事件 cudaEventRecord(preDone, streamA); // 让B和C等待事件而不是等待所有设备工作完成 cudaStreamWaitEvent(streamB, preDone); cudaStreamWaitEvent(streamC, preDone); // 此刻B和C可以并行执行A自身后续操作也可以继续如果不需要等待B/C这种“点对点”的同步方式非常强大它把“全局同步”拆成了“局部依赖”整个流水线的并行窗口一下就打开了。4.3 时间测量用事件替代cudaDeviceSynchronize事件还有一个特别重要的用途——精确测量kernel耗时。很多人用cudaDeviceSynchronize加CPU计时器这会包含内核启动的开销和CPU同步等待时间测出来偏大且不稳定。正确做法是cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start, stream); myKernelblocks, threads, 0, stream(); cudaEventRecord(stop, stream); cudaEventSynchronize(stop); float ms 0; cudaEventElapsedTime(ms, start, stop);这里的事件计时是在GPU时间线上测量的不会把CPU端的启动延迟算进去。排流水线时我通常会在每个流的“总入口”和“总出口”各放一个事件测出来的数值才是真实的任务耗时。5. 实战多Stream流水线让图像处理吞吐翻倍5.1 场景拆解数据分块与任务流水理论说太多了直接上一个完整的实战例子。假设我们要对一个由1000帧组成的图像序列做灰度化加高斯模糊。每帧处理流程分三段H2D把帧数据从CPU内存拷贝到GPU显存Kernel在GPU上执行灰度化模糊的kernelD2H把处理结果拷贝回CPU内存如果不使用流时间线是“拷-算-拷-算”严格串行。如果使用2个流我们可以把帧分成两组流0处理偶数帧流1处理奇数帧。由于帧之间没有数据依赖两个流上的kernel可以并发执行同时流0的kernel运行时流1的H2D拷贝引擎可以在另一个引擎上并行执行。5.2 双流实现的完整代码代码的核心并不复杂const int numStreams 2; cudaStream_t streams[numStreams]; for (int i 0; i numStreams; i) { cudaStreamCreateWithFlags(streams[i], cudaStreamNonBlocking); } for (int i 0; i frameCount; i) { int sid i % numStreams; int idx i / numStreams; char* d_in d_input sid * batchFrameBytes idx * frameBytes; char* d_out d_output sid * batchFrameBytes idx * frameBytes; const char* h_in h_input i * frameBytes; char* h_out h_output i * frameBytes; cudaMemcpyAsync(d_in, h_in, frameBytes, cudaMemcpyHostToDevice, streams[sid]); processFrameKernelblocks, threads, 0, streams[sid](d_in, d_out, frameSize); cudaMemcpyAsync(h_out, d_out, frameBytes, cudaMemcpyDeviceToHost, streams[sid]); } for (int i 0; i numStreams; i) { cudaStreamSynchronize(streams[i]); }注意这里给每个流单独分配了输入输出缓冲区的不同区域sid * batchFrameBytes idx * frameBytes避免不同流同时读写同一块显存。这一点非常重要——如果两个流共用一个缓冲区cudaMemcpyAsync只是发起异步传输实际传输可能在kernel运行时才发生数据竞争就会悄悄出现。5.3 实测数据性能提升与瓶颈分析我在一张RTX 3090上跑过这个例子图像尺寸1024×1024灰度化高斯模糊kernel约0.9msH2D拷贝约0.3msD2H拷贝约0.3ms。单流版本的每帧总耗时约1.5ms双流版本的理论下限是max(总拷贝时间, 总kernel时间)而不是三者之和实测帧间吞吐提升了约38%接近理论重叠极限。如果kernel本身的SM占用率偏高例如用了大量共享内存双流的提升会变小但通常依然优于单流。想进一步压榨性能可以用4个流但要注意数据分块的大小块太小会导致kernel启动开销占比升高。6. 实测中的坑与性能调优来自项目现场的踩坑笔记6.1 坑一cudaMemcpyAsync不是魔法别乱用cudaMemcpyAsync虽然叫异步但它只是把拷贝操作提交到指定流上真正的数据传输不一定和计算重叠。它有三个前提使用cudaMemcpyAsync不是cudaMemcpy传输的源和目标必须是“可分页内存”或“映射内存”其中可分页内存特指用cudaMallocHost分配的锁页内存pinned memory普通malloc分配的内存无法真正异步化拷贝的方向和设备支持情况要满足条件我见过太多人直接对普通malloc数组调cudaMemcpyAsync结果性能没变化还以为是流没用。要发挥异步拷贝的威力一定要用cudaMallocHost或cudaHostAlloc分配主机端内存否则异步调用会退化成同步传输。6.2 坑二数据依赖没理清并行变乱序多流并发的另一大坑是数据依赖。比如流B的kernel需要流A算出的中间结果但你忘了加cudaStreamWaitEvent结果流B在流A算完之前就开始读数据得到的是垃圾值。这种Bug很难复现因为GPU调度时机不固定可能跑100次才崩一次。我的排查习惯是一旦怀疑多流数据竞争先砍到单流跑一遍确认结果正确再一个流一个流地加回来每加一个流就检查一次结果。如果某个流加进来后结果开始错乱重点查它和上游流的数据依赖以及事件等待是否齐全。6.3 调优技巧流优先级、MPS与硬件差异最后分享几个压箱底的经验。第一CUDA流本身支持优先级cudaStreamCreateWithPriority可以给不同流设置不同优先级。实测中把延迟敏感的小kernel放在高优先级流把大kernel放在低优先级流交互场景的卡顿明显减少。第二如果你在数据中心GPU上跑多进程或多任务可以考虑开MPSMulti-Process Service。MPS能让多个进程的kernel共享GPU资源相当于把多流概念扩展到进程级别。但MPS配置复杂单机单卡场景用多流就够了。第三GPU架构对流的并发能力影响很大。Turing及之后架构的Hyper-Q增强了并发kernel的调度能力Ampere更进一步优化了SM资源感知。老架构上流并发效果差是正常的不必怀疑自己写错。我自己的切身体会是多流编程最难的并不是API本身——API就那几个难的是建立“GPU是一个并行流水线工厂”的思维模式。每次写kernel之前先画一条任务时间线标出哪里有等待、哪里可以重叠然后才动手写代码。这套方法比盲目堆流数量有效得多。