4. 任务提交与分发
好，咱们接着聊。上一节我们把Kernel函数写好了，也配置好了执行参数。那问题来了——这个Kernel到底是怎么从CPU手里，跑到GPU上去执行的？
说白了，CPU和GPU是两个独立的处理器。它们有各自的内存空间，各自的指令流水线。想让GPU干活，CPU得把任务“递”过去。这个递的过程，就是我们今天要聊的核心。
4.1 任务提交的“快递”模型
我个人习惯把任务提交想象成一个快递发货的过程：
- 你（CPU）：寄件人，负责打包Kernel函数和参数
- Command Queue（命令队列）：快递柜，任务先放进去排队
- GPU：收件人，从队列里取任务执行
嗯，这个比喻虽然简单，但核心逻辑是对的。CPU不会直接把Kernel扔给GPU，而是通过一个中间结构——Command Queue。
核心要点：CPU提交任务到Command Queue，GPU从Queue中取任务执行。两者通过这个队列解耦，实现异步执行。
4.2 Command Queue：任务的“传送带”
Command Queue是什么？你可以把它理解成一个FIFO（先进先出）的缓冲区。CPU往里面塞命令，GPU从另一头取命令。
我在项目中遇到过一个问题：有一次调试一个图像处理管线，发现GPU利用率一直上不去。后来一查，原来是Command Queue里塞了太多同步指令，导致GPU频繁等待。嗯，这个坑后面会细讲。
Command Queue里装的命令类型包括：
- Kernel执行命令：启动一个Grid的线程
- 内存拷贝命令：Host到Device，Device到Host
- 同步命令：比如Event、Barrier
- 资源绑定命令：设置常量内存、纹理等
每个命令在Queue里都是一个“Command Packet”。GPU的硬件调度器会逐个解析并执行。
4.3 Stream：任务的“独立流水线”
好，Queue有了。但问题来了——如果所有任务都塞到同一个Queue里，那GPU只能串行执行。你想想看，这多浪费啊？
所以，现代GPU引入了Stream的概念。Stream本质上就是一条独立的命令流水线。每个Stream有自己的Command Queue，Stream之间的任务可以并行执行。
我的建议：刚开始学的时候，把Stream想象成“车道”。一个Stream是一条车道，多个Stream就是多条车道。车（任务）可以在不同车道上并行跑。
来看一个典型的CUDA代码示例：
// 创建两个Stream
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在stream1上执行Kernel A
kernel_A<<<grid, block, 0, stream1>>>(d_data);
// 在stream2上执行Kernel B
kernel_B<<<grid, block, 0, stream2>>>(d_data);
// 同步等待所有Stream完成
cudaDeviceSynchronize();
// 清理
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);注意看<<<grid, block, 0, stream1>>>这个语法。第四个参数就是Stream的句柄。如果不指定，默认使用默认Stream（也叫NULL Stream）。
4.4 默认Stream vs 非默认Stream
这里有个容易踩的坑。默认Stream的行为在不同GPU架构下不一样：
| 特性 | 默认Stream | 非默认Stream |
|---|---|---|
| 同步行为 | 老架构下会阻塞其他Stream | 相互独立，不阻塞 |
| 创建方式 | 自动创建 | cudaStreamCreate() |
| 适用场景 | 简单串行任务 | 需要并行的复杂管线 |
| 性能风险 | 可能造成隐式同步 | 需要手动管理依赖 |
我曾经踩过的坑：有一次在Pascal架构的GPU上写代码，用了默认Stream和自定义Stream混用。结果发现性能还不如全用默认Stream。查了半天，原来是默认Stream在Pascal上会隐式同步所有其他Stream。后来我改成全部用非默认Stream，性能直接翻倍。
4.5 任务提交的完整流程
好，我们把整个流程串起来。从CPU提交Kernel到GPU执行，大致经历这么几步：
- CPU端准备：把Kernel参数打包，放到寄存器或内存中
- 驱动层处理：CUDA Runtime把Kernel调用转换成Command Packet
- 入队：Command Packet被Push到指定Stream的Command Queue
- 硬件调度：GPU的Work Distributor从Queue中取出Packet
- 线程块分发：Work Distributor把线程块分配到SM（流多处理器）
- 执行：SM上的Warp Scheduler开始调度线程束执行
下面这张图展示了这个流程：
4.6 异步提交与同步点
这里有个关键点：CPU提交Kernel到Command Queue是异步的。也就是说，CPU把命令丢进Queue后，不会等GPU执行完，而是立刻返回继续干自己的事。
为什么会这样？因为如果CPU要等GPU执行完再继续，那CPU就闲着没事干了。你想想看，CPU和GPU并行工作，才能最大化吞吐量。
但异步也带来一个问题——你怎么知道GPU执行完了？这时候就需要同步点：
- cudaDeviceSynchronize()：阻塞CPU，直到所有Stream上的所有任务完成
- cudaStreamSynchronize(stream)：只等待指定Stream完成
- cudaEventSynchronize(event)：等待某个Event触发
我的习惯：在性能调优时，我尽量少用全局同步（cudaDeviceSynchronize），而是用Event来做细粒度的同步控制。这样可以减少不必要的等待。
4.7 多Stream的实战技巧
最后，分享几个我在项目中总结的多Stream使用技巧：
- Stream数量不是越多越好：一般2-4个Stream就能把GPU喂饱。太多Stream反而增加调度开销。
- 注意Stream间的依赖：如果Kernel B依赖Kernel A的结果，记得用Event做同步，别让它们乱跑。
- 内存拷贝和Kernel执行可以重叠：把内存拷贝放到一个Stream，Kernel执行放到另一个Stream，让DMA和计算单元同时工作。
- 避免隐式同步：比如在默认Stream上操作，或者跨Stream访问统一内存，都可能触发隐式同步，打乱你的并行计划。
嗯，关于任务提交和分发，今天就聊这么多。记住一句话：Command Queue是任务的传送带，Stream是独立的流水线。理解了这两个概念，你就掌握了CPU和GPU之间通信的核心机制。