4. CUDA基础与编解码API：从线程模型到硬件加速
这一章，我们正式进入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之间是独立的。
核心概念速记：
- Thread：最小的执行单元，每个线程执行相同的kernel函数
- Block：线程的集合，内部可通过Shared Memory通信
- Grid：Block的集合，一个kernel启动对应一个Grid
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。
我的经验：刚开始写CUDA时，我总忘记做边界检查（if (idx < N)）。有一次处理非对齐的数据，结果读到了非法地址，程序直接崩了。从那以后，我每个kernel都会加边界判断，哪怕数据量刚好是线程数的整数倍。
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，所以用的时候要精打细算。
避坑指南：我曾经在项目里把整个帧的数据都塞进Shared Memory，结果编译报错说"too many resources requested"。后来才意识到，Shared Memory是每个Block独享的，不是整个Grid共享的。正确的做法是：把频繁访问的小块数据（比如宏块、滤波窗口）放到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视频，大致流程如下：
- 创建编码器：调用
NvEncCreateEncoder，指定编码格式、分辨率、码率等参数 - 注册输入资源：把GPU显存里的YUV数据注册为输入表面
- 提交编码任务：调用
NvEncEncodeFrame，把输入表面交给编码器 - 获取输出：在回调函数里拿到编码后的H.264/HEVC码流
- 销毁编码器：释放资源
来看一个简化的代码片段：
// 初始化编码器参数
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);我的建议：编码参数里的presetGUID很关键。P1（最快）适合实时推流，P7（最高质量）适合离线转码。我在做直播项目时，用的就是P1，延迟能控制在10ms以内。但如果是做视频编辑的导出，我会用P6或P7，画质差距肉眼可见。
4.6 NVDECODE API：硬件解码的套路
解码和编码是对称的，但有个区别：解码器是被动的，你给它码流，它给你图像。流程如下：
- 创建解码器：指定码流格式（H.264/HEVC/AV1等）
- 提交码流数据：把压缩后的数据包传给解码器
- 获取解码帧：从输出队列里取出YUV/RGB图像
- 显示或处理：把图像送到显示器或传给下一个处理环节
这里有个坑：解码器内部有DPB（Decoded Picture Buffer），它会缓存多帧图像用于参考。如果你不按时取走解码帧，DPB满了就会阻塞解码。我曾经在项目里遇到解码卡顿，查了半天发现是取帧线程优先级太低，导致DPB溢出。
性能关键点：
- 解码器和编码器都支持异步操作，建议用单独的线程管理
- 输入输出表面尽量用CUDA Array，避免CPU-GPU间拷贝
- 对于4K/8K视频，注意显存带宽，别让解码器和渲染管线抢带宽
4.7 CPU与GPU协同：谁该做什么？
最后聊一个架构层面的问题。有了CUDA和Video Codec SDK，是不是所有事都该扔给GPU？
我的看法是：编解码本身交给硬件，但控制逻辑和预处理留在CPU。比如：
- CPU负责：解封装（demux）、参数解析、码流管理、错误恢复
- GPU负责：编解码核心计算、色彩空间转换、缩放、后处理
为什么这么分？因为CPU擅长处理分支逻辑和异常，而GPU擅长大规模并行计算。你想想看，如果让GPU去解析H.264的NAL单元，那得写多少if-else？而且GPU的warp divergence会严重拖慢性能。
避坑指南：我曾经尝试在GPU上做码流解析，结果性能还不如CPU单线程。后来老老实实把解析放回CPU，GPU只做编解码，整体吞吐量反而提升了30%。记住：不要为了用GPU而用GPU。