4. 协同计算模型

这一章，我们正式进入GPU的世界。说实话，很多做编解码优化的朋友，一开始都卡在CUDA的线程模型上。你想想看，CPU上我们习惯了串行思维，突然要面对成千上万个并发线程，确实需要一点时间适应。

我个人习惯把CUDA编程拆成两件事来理解：怎么组织线程，以及怎么用这些线程去处理数据。而NVIDIA的Video Codec SDK，说白了就是把编解码这种特定任务，封装成了可以直接调用的API，让你不用从零写shader。

4.1 CUDA线程层次结构：Grid、Block、Thread

先看这张图，它把CUDA的执行模型讲得很清楚。

这张图里，最外层是Grid，里面包含多个Block，每个Block又包含多个Thread。为什么搞这么复杂？因为GPU的硬件架构就是这样分层的——Thread对应一个CUDA核心，Block内的线程可以共享数据，而不同Block之间是独立的。

4.2 第一个CUDA Kernel：从Hello World到数据并行

写一个CUDA程序，其实就是写一个__global__函数，然后告诉GPU要启动多少个线程。来看个最简单的例子：

// 向量加法 kernel
__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

// 调用方式
int main() {
    // ... 分配内存、拷贝数据 ...
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    // ... 拷贝结果回CPU ...
}

这里有个细节我特别想强调：threadIdx.x是线程在Block内的索引，blockIdx.x是Block在Grid内的索引，blockDim.x是每个Block的线程数。三者组合起来，就能算出全局唯一的线程ID。

4.3 内存模型：搞清楚数据在哪

CUDA的内存模型，说白了就是搞清楚数据放在哪、谁可以访问。我见过太多人把全局内存当共享内存用，性能惨不忍睹。

内存类型	位置	访问速度	作用域	生命周期
Global Memory	显存	慢（~200周期）	所有线程	程序运行期间
Shared Memory	芯片内	快（~5周期）	单个Block	Block生命周期
Registers	寄存器	极快（1周期）	单个线程	线程生命周期
Constant Memory	显存（带缓存）	较快（~10周期）	所有线程（只读）	程序运行期间

为什么Shared Memory这么快？因为它就在SM（Streaming Multiprocessor）内部，相当于CPU的L1缓存。但容量有限——以RTX 4090为例，每个SM只有128KB Shared Memory，所以用的时候要精打细算。

4.4 NVIDIA Video Codec SDK：硬件编解码的钥匙

好了，有了CUDA基础，我们来看看怎么用它做编解码。NVIDIA从Kepler架构开始，就在GPU里集成了专门的硬件编解码单元——NVENC和NVDEC。Video Codec SDK就是操作这些硬件的API。

说实话，我第一次用这个SDK时，最困惑的是它的API设计。它不像CUDA那样直接写kernel，而是通过异步回调的方式工作。你提交一个编码任务，它完成后通知你，中间CPU可以干别的事。

4.5 NVENCODE API：硬件编码实战

用NVENC编码一个H.264视频，大致流程如下：

1. 创建编码器：调用NvEncCreateEncoder，指定编码格式、分辨率、码率等参数
2. 注册输入资源：把GPU显存里的YUV数据注册为输入表面
3. 提交编码任务：调用NvEncEncodeFrame，把输入表面交给编码器
4. 获取输出：在回调函数里拿到编码后的H.264/HEVC码流
5. 销毁编码器：释放资源

来看一个简化的代码片段：

// 初始化编码器参数
NV_ENC_INITIALIZE_PARAMS initParams = {0};
initParams.encodeWidth = 1920;
initParams.encodeHeight = 1080;
initParams.encodeGUID = NV_ENC_CODEC_H264_GUID;
initParams.presetGUID = NV_ENC_PRESET_P1_GUID;  // 低延迟模式

// 创建编码器
NvEncCreateEncoder(&hEncoder, &initParams);

// 注册输入表面
NV_ENC_INPUT_PTR inputSurface;
NvEncRegisterInputSurface(hEncoder, &inputDesc, &inputSurface);

// 编码一帧
NV_ENC_PIC_PARAMS picParams = {0};
picParams.inputBuffer = inputSurface;
picParams.bufferFmt = NV_ENC_BUFFER_FORMAT_NV12;
picParams.pictureStruct = NV_ENC_PIC_STRUCT_FRAME;
NvEncEncodeFrame(hEncoder, &picParams, &outputBitstream);

// 处理输出码流
processEncodedData(outputBitstream);

4.6 NVDECODE API：硬件解码的套路

解码和编码是对称的，但有个区别：解码器是被动的，你给它码流，它给你图像。流程如下：

1. 创建解码器：指定码流格式（H.264/HEVC/AV1等）
2. 提交码流数据：把压缩后的数据包传给解码器
3. 获取解码帧：从输出队列里取出YUV/RGB图像
4. 显示或处理：把图像送到显示器或传给下一个处理环节

这里有个坑：解码器内部有DPB（Decoded Picture Buffer），它会缓存多帧图像用于参考。如果你不按时取走解码帧，DPB满了就会阻塞解码。我曾经在项目里遇到解码卡顿，查了半天发现是取帧线程优先级太低，导致DPB溢出。

4.7 CPU与GPU协同：谁该做什么？

最后聊一个架构层面的问题。有了CUDA和Video Codec SDK，是不是所有事都该扔给GPU？

我的看法是：编解码本身交给硬件，但控制逻辑和预处理留在CPU。比如：

• CPU负责：解封装（demux）、参数解析、码流管理、错误恢复
• GPU负责：编解码核心计算、色彩空间转换、缩放、后处理

为什么这么分？因为CPU擅长处理分支逻辑和异常，而GPU擅长大规模并行计算。你想想看，如果让GPU去解析H.264的NAL单元，那得写多少if-else？而且GPU的warp divergence会严重拖慢性能。