4. 单GPU任务调度：流(Stream)的概念、并发执行、同步机制、优先级调度

好，咱们今天聊聊单GPU上的任务调度。说白了，就是怎么让GPU这个“超级打工人”高效干活，别闲着。

很多初学者觉得，把数据扔给GPU，它自己就会安排得明明白白。嗯，理想很丰满，现实很骨感。默认情况下，GPU确实会按顺序执行你发的指令，但这往往不是最高效的方式。我刚开始做GPU开发时，就吃过这个亏——明明GPU算力很强，但利用率就是上不去，一看profiling结果，好家伙，大部分时间都在等数据搬运。

4.1 流(Stream)的概念

什么是流？你可以把它想象成一条“任务流水线”。

在CUDA里，流是一个操作序列，这些操作会按你提交的顺序在GPU上执行。每个流里的操作是串行的，但不同流之间的操作可以并发执行。

我个人的习惯是，把流比作“车道”。默认情况下，所有任务都在同一条车道上跑，后面的车必须等前面的车走完。但如果你开辟了多条车道（多个流），那不同车道的车就可以并排跑了。

核心要点：流是GPU任务调度的基本单位。每个流维护着自己的执行顺序，流之间相互独立。

创建流很简单，API调用就完事了：

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

// 使用流执行任务
kernel1<<<grid, block, 0, stream1>>>(...);
kernel2<<<grid, block, 0, stream2>>>(...);

// 记得销毁
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

注意那个第四个参数 stream1，它就是指定这个kernel在哪个流里执行。如果不指定，默认就是 0 流，也就是默认流。

4.2 并发执行

并发执行是流的核心价值。为什么需要并发？

你想想看，GPU里有计算单元、有显存、有数据传输总线。如果只用一个流，那这些资源经常是“轮流上班”——计算的时候总线闲着，传数据的时候计算单元闲着。这多浪费啊！

用多个流，就可以让计算和数据传输重叠起来。比如：

cudaStream_t streamA, streamB;
cudaStreamCreate(&streamA);
cudaStreamCreate(&streamB);

// 流A：搬运数据到GPU
cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, streamA);

// 流B：执行计算（假设数据已经在GPU里了）
kernelB<<<grid, block, 0, streamB>>>(d_B);

// 流A：把结果拷回来
cudaMemcpyAsync(h_A, d_A, size, cudaMemcpyDeviceToHost, streamA);

这里 cudaMemcpyAsync 是异步版本，它不会阻塞CPU。流A在搬运数据的同时，流B可以执行计算。只要它们操作的是不同的数据，就不会冲突。

避坑指南：我曾经犯过一个错误——在同一个流里既做数据传输又做计算，以为异步就万事大吉了。结果发现，同一个流里的操作还是串行的！异步只是说CPU不等待，但GPU内部还是按顺序来。所以，想并发，必须用不同的流。

并发执行能带来多大的提升？我做过一个图像处理项目，用4个流把数据加载、预处理、推理、后处理流水线化，吞吐量直接翻了3倍多。当然，具体收益取决于你的任务特性和硬件能力。

4.3 同步机制

并发虽好，但也不能乱来。有时候，一个流必须等另一个流干完某些事才能开始。这时候就需要同步机制了。

CUDA提供了几种同步方式：

cudaDeviceSynchronize()：阻塞CPU，直到GPU上所有流的所有操作都完成。简单粗暴，但效率低。
cudaStreamSynchronize(stream)：只等待指定流完成。更精细一些。
cudaEvent：事件机制，可以实现流之间的精确同步。

事件同步是我用得最多的。举个例子：

cudaEvent_t event;
cudaEventCreate(&event);

// 流A做完某件事后，记录事件
kernelA<<<grid, block, 0, streamA>>>(...);
cudaEventRecord(event, streamA);

// 流B等待这个事件
cudaStreamWaitEvent(streamB, event, 0);
kernelB<<<grid, block, 0, streamB>>>(...);  // 等流A的kernelA完成后再执行

这里 cudaEventRecord 会在流A执行完 kernelA 后记录一个“标记”。cudaStreamWaitEvent 让流B停下来，直到这个标记出现。说白了，就是流B在说：“兄弟，你那个活干完了我再动手。”

注意：事件同步是GPU级别的，不会阻塞CPU。CPU可以继续提交其他任务。这一点和 cudaStreamSynchronize 不同，后者会阻塞CPU直到流完成。

我记得有一次调试一个复杂的多流程序，发现结果总是不对。查了半天，原来是忘了加 cudaStreamWaitEvent，导致流B在流A还没写完数据就开始读了。嗯，从那以后，我对同步就特别敏感——该等的时候一定要等，不该等的时候别乱等。

4.4 优先级调度

流还支持优先级设置。高优先级的流会优先获得GPU资源。

这个功能在什么场景下有用？我举个例子：

假设你有一个实时推理任务，要求低延迟；同时还有一些后台的数据预处理任务，不着急。这时候，你可以把推理任务放在高优先级流里，预处理放在低优先级流里。GPU会优先执行高优先级流里的操作，确保推理的实时性。

创建带优先级的流：

int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);

cudaStream_t high_prio_stream, low_prio_stream;
cudaStreamCreateWithPriority(&high_prio_stream, cudaStreamDefault, priority_high);
cudaStreamCreateWithPriority(&low_prio_stream, cudaStreamDefault, priority_low);

cudaDeviceGetStreamPriorityRange 会返回当前设备支持的优先级范围。priority_high 是最高优先级（数值上可能更小），priority_low 是最低优先级。

优先级	典型用途	注意事项
高优先级	实时推理、交互式应用	不要滥用，否则低优先级流可能饿死
默认优先级	大多数通用计算任务	平衡性能和公平性
低优先级	后台批处理、数据预处理	延迟可以接受，但别让GPU闲着

我的经验：优先级调度不是万能的。它只影响GPU内部的计算资源分配，不影响数据传输。而且，不同GPU对优先级的支持程度不一样。我建议你先用默认优先级跑通，再根据profiling结果决定是否需要调整优先级。

最后，我用一张图来总结单GPU任务调度的核心逻辑：

从这张图可以看出，CPU提交任务到不同的流，流内部的操作串行执行，流之间可以并发。通过事件同步，我们可以控制流之间的执行顺序。最终，GPU根据流的优先级来调度计算资源。

嗯，单GPU任务调度就讲到这里。流的概念、并发执行、同步机制、优先级调度，这四个点环环相扣。你只要记住一句话：流是并发的载体，同步是秩序的保证，优先级是资源的倾斜。把这个搞明白了，单GPU的调度你就入门了。

最后一个小建议：刚开始用流的时候，别贪多。先试试2个流，看看性能有没有提升。用NVIDIA Nsight Systems或者Nsight Compute看看profiling结果，确认计算和传输是否真的重叠了。我见过太多人一上来就搞8个流，结果性能反而下降了——因为流的管理开销超过了并发收益。

想要更多资料，联系可可

📚 课程目录

4. 单GPU任务调度：流(Stream)的概念、并发执行、同步机制、优先级调度

4.1 流(Stream)的概念

4.2 并发执行

4.3 同步机制

4.4 优先级调度