CPU-GPU协同工作模型：异构计算的核心

好，咱们进入第三章。前两章我们把CPU和GPU各自的编解码能力摸了个底。现在要聊的，才是真正拉开性能差距的地方——怎么让这两个家伙好好配合。

我个人习惯把CPU比作项目经理，GPU比作施工队。项目经理脑子快，擅长处理各种突发状况和复杂逻辑；施工队力气大，能同时干一堆重复的体力活。但问题来了——项目经理和施工队不在一个办公室，他们之间传递图纸和材料，是有成本的。

Host-Device模型：谁是指挥官？

在异构计算里，CPU所在的系统内存侧叫Host（主机端），GPU及其显存叫Device（设备端）。这个模型很简单：

• Host：负责控制流、复杂逻辑、文件I/O、内存管理
• Device：负责大规模并行计算、数据吞吐

说白了，CPU是大脑，GPU是肌肉。大脑不能直接指挥肌肉里的每根纤维，它只需要下达指令：「把这块视频帧给我处理了」。剩下的，GPU自己搞定。

我在项目中遇到过不少新手，上来就把所有逻辑往GPU里塞，结果性能反而更差。为什么？因为GPU不擅长分支判断、链表操作、递归调用。你让施工队去画图纸，那不是找罪受吗？

数据搬运：被低估的开销

好，现在问题来了。CPU和GPU有各自独立的内存空间。CPU的数据在DDR里，GPU的数据在显存里。它们之间要交换数据，得走PCIe总线。

这个搬运成本有多高？我给你算笔账：

传输方向	带宽（PCIe 3.0 x16）	延迟	传输4K帧（~8MB）耗时
Host → Device	~12 GB/s	~5-10 μs	~0.7 ms
Device → Host	~12 GB/s	~5-10 μs	~0.7 ms
GPU内核计算（4K帧）	—	—	~1-3 ms

看到了吗？光搬运一帧数据，就要花掉将近1毫秒。如果一帧的处理总共才3毫秒，搬运就占了三分之一。这还没算上多次搬运的情况。

异步传输：别让GPU干等着

既然搬运这么贵，怎么优化？答案是异步传输。

传统的做法是同步的：

cudaMemcpy(dst, src, size, cudaMemcpyHostToDevice);  // 阻塞，等搬运完
kernel<<<grid, block>>>(...);                      // 再开始计算
cudaMemcpy(dst, src, size, cudaMemcpyDeviceToHost);  // 等计算完，再搬回来

这就像你让施工队等着，等图纸从总部快递到了再开工。快递到了，施工队干活，干完了再等下一批图纸。中间全是空闲时间。

异步传输就不一样了：

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

// 在stream1里搬运数据
cudaMemcpyAsync(dst1, src1, size, cudaMemcpyHostToDevice, stream1);
// 在stream2里执行计算（用上一批已经搬好的数据）
kernel<<<grid, block, 0, stream2>>>(dst2, ...);
// 在stream1里搬回结果
cudaMemcpyAsync(src1, dst1, size, cudaMemcpyDeviceToHost, stream1);

说白了，就是让搬运和计算重叠起来。GPU在算第N帧的时候，CPU已经在搬第N+1帧的数据了。这样总时间就不再是「搬运+计算」，而是max(搬运, 计算)。

流水线设计：让每一毫秒都不浪费

异步传输是基础，流水线设计才是真正的杀手锏。你想想看，如果能把整个编解码流程拆成多个阶段，每个阶段都在不同的数据上并行工作，那效率得多高？

一个典型的编解码流水线长这样：

阶段1：CPU读取文件 → 阶段2：CPU预处理（解析头信息、分片）
阶段3：CPU→GPU搬运数据
阶段4：GPU编解码核心计算
阶段5：GPU→CPU搬运结果
阶段6：CPU后处理（封装、写文件）

如果串行执行，处理N帧的时间就是 N × (t1 + t2 + ... + t6)。但如果用流水线：

• CPU在处理第3帧的读取时，GPU正在算第2帧，同时第1帧的结果正在被搬回
• 理想情况下，总时间 ≈ N × max(各阶段耗时) + 启动开销

我画个图帮你理解：

你看，三个Stream各干各的，互不干扰。GPU在算第2帧的时候，CPU已经把第3帧搬进去了，同时第1帧的结果正在被搬回来。每一毫秒都在干活，没有空闲。

实现流水线的关键点

光有图还不够，实际编码时要注意几个坑：

1. 固定内存（Pinned Memory）：cudaMemcpyAsync必须用cudaHostAlloc分配的内存，否则DMA控制器没法直接访问，会退化成同步拷贝。我见过有人调了半天异步没效果，最后发现是内存没固定。
2. Stream数量不是越多越好：Stream太多会导致GPU上下文切换开销增大。我个人习惯用2-4个Stream，具体要看你的编解码流程拆成了几个阶段。
3. 依赖关系要理清：如果第N帧的计算依赖第N-1帧的结果（比如参考帧），那就不能简单地把所有帧都扔进流水线。需要用cudaStreamSynchronize或事件机制来同步。
4. CPU端别闲着：流水线里CPU不只是搬运数据，它还可以做预处理和后处理。比如在GPU算第N帧的时候，CPU可以解析第N+1帧的NAL单元头信息。