3. 图像预处理GPU加速：图像解码（JPEG/PNG）、颜色空间转换（RGB/BGR/YUV）、图像缩放与归一化

各位同学，今天我们来聊聊感知pipeline里最容易被低估的一环——图像预处理。

说实话，很多刚入行的朋友觉得预处理不就是读个图、转个格式、缩个大小嘛，CPU跑跑就行了。但我在实际项目中踩过不少坑，尤其是当你的感知系统要处理8路、12路甚至更多摄像头时，CPU端的预处理会成为整个系统的瓶颈。你想想看，一帧4K的JPEG图像，CPU解码可能要花30-50ms，而GPU端用nvJPEG库，同样的活1ms以内就能搞定。这差距，不是一点半点。

3.1 图像解码：JPEG与PNG的GPU加速

图像解码，说白了就是把压缩格式的图片还原成原始的像素矩阵。JPEG和PNG是自动驾驶中最常见的两种格式。

JPEG解码：JPEG用的是有损压缩，解码过程涉及霍夫曼解码、反量化、逆DCT变换、颜色空间转换等步骤。这些步骤在CPU上是串行的，但在GPU上可以高度并行化。

我个人习惯用NVIDIA的nvJPEG库。它直接调用GPU硬件解码器，速度比CPU快一个数量级。我在一个项目中处理12路1080p摄像头，CPU解码时CPU占用率直接飙到80%以上，换成nvJPEG后，CPU占用率降到5%以下，GPU利用率也才20%左右。

核心要点：nvJPEG支持批量解码，一次调用可以同时解码多张图片。这在多摄像头场景下非常实用。

// nvJPEG批量解码示例（伪代码）
nvjpegHandle_t handle;
nvjpegCreateSimple(&handle);

// 准备批量输入
std::vector<const char*> file_names = {"cam0.jpg", "cam1.jpg", ...};
std::vector<nvjpegImage_t> outputs;

// 批量解码
nvjpegDecodeBatched(handle, file_names.data(), outputs.data(), batch_size);

// 输出直接是GPU内存中的RGB/BGR数据
// 无需CPU-GPU拷贝，省掉一次PCIe传输

PNG解码：PNG用的是无损压缩，解码过程更复杂，涉及zlib解压缩、过滤、像素重组等。GPU加速PNG解码目前主要靠nvidia-dali或者自定义CUDA kernel。

我曾经在一个项目中需要加载高分辨率语义标签图（PNG格式），CPU解码一张2048x2048的PNG图要200ms，整个训练数据加载成了瓶颈。后来用DALI的PNG解码器，直接把解码放到GPU上，速度提升了15倍以上。

注意：PNG解码的GPU加速效果不如JPEG明显，因为PNG的压缩算法本身并行性较差。如果对速度要求极高，建议优先使用JPEG格式。

3.2 颜色空间转换：RGB/BGR/YUV

颜色空间转换，这个看似简单，但坑不少。

自动驾驶中常见的几种颜色空间：

  • RGB：大多数深度学习模型默认的输入格式
  • BGR：OpenCV默认的读取格式，很多老模型也用的BGR
  • YUV：摄像头传感器直接输出的格式，Y是亮度，UV是色度

为什么需要转换？因为摄像头出来的原始数据往往是YUV格式，而模型输入需要RGB或BGR。这个转换如果放在CPU上做，又是一笔不小的开销。

GPU上做颜色空间转换，核心思路就是写一个CUDA kernel，对每个像素并行处理。以YUV420到RGB的转换为例：

// CUDA kernel：YUV420转RGB
__global__ void yuv420_to_rgb_kernel(
    const uint8_t* y, const uint8_t* u, const uint8_t* v,
    uint8_t* rgb, int width, int height) {
    
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y_idx = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (x >= width || y_idx >= height) return;
    
    // YUV到RGB的转换公式
    int y_val = y[y_idx * width + x];
    int u_val = u[(y_idx/2) * (width/2) + (x/2)] - 128;
    int v_val = v[(y_idx/2) * (width/2) + (x/2)] - 128;
    
    int r = y_val + 1.402f * v_val;
    int g = y_val - 0.344f * u_val - 0.714f * v_val;
    int b = y_val + 1.772f * u_val;
    
    // 钳位到[0, 255]
    rgb[y_idx * width * 3 + x * 3 + 0] = clamp(r, 0, 255);
    rgb[y_idx * width * 3 + x * 3 + 1] = clamp(g, 0, 255);
    rgb[y_idx * width * 3 + x * 3 + 2] = clamp(b, 0, 255);
}

嗯，这里要注意，YUV420的UV分量是降采样的，每2x2的Y块共享一个U和V。所以访问UV时要除以2。这个细节我一开始没注意，结果出来的图像颜色全偏了，排查了半天才发现是索引算错了。

我的经验：如果你用TensorRT部署模型，可以直接在模型输入层之前加一个颜色空间转换层。TensorRT支持自定义插件，把YUV转RGB的kernel写成一个插件，这样整个pipeline都在GPU上完成，省掉一次额外的kernel launch开销。

3.3 图像缩放与归一化

图像缩放，说白了就是把不同分辨率的输入统一成模型要求的尺寸。归一化则是把像素值从[0,255]映射到[0,1]或[-1,1]，并减去均值除以方差。

GPU加速缩放：常用的插值算法有最近邻、双线性、双三次。GPU上最常用的是双线性插值，因为它在质量和速度之间取得了很好的平衡。

CUDA里可以用纹理内存来做缩放，纹理内存自带硬件插值功能，性能极好。

// 使用CUDA纹理内存做双线性缩放
cudaArray_t src_array;
cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc<float>();
cudaMallocArray(&src_array, &channel_desc, src_width, src_height);

// 绑定纹理
texture<float, 2, cudaReadModeElementType> tex;
cudaBindTextureToArray(tex, src_array);

// 缩放kernel
__global__ void resize_kernel(float* dst, int dst_width, int dst_height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    float src_x = (float)x / dst_width * src_width;
    float src_y = (float)y / dst_height * src_height;
    
    // 纹理内存自动做双线性插值
    dst[y * dst_width + x] = tex2D(tex, src_x + 0.5f, src_y + 0.5f);
}

为什么用纹理内存？因为纹理内存有缓存机制，而且硬件自动处理边界情况。我在项目中实测过，用纹理内存比用全局内存做双线性插值快了3倍左右。

归一化：归一化通常和缩放一起做，减少一次kernel launch。比如：

// 缩放+归一化融合kernel
__global__ void resize_normalize_kernel(
    const uint8_t* src, float* dst,
    int src_w, int src_h, int dst_w, int dst_h,
    float mean, float std) {
    
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    // 计算源图像坐标
    float src_x = (x + 0.5f) * src_w / dst_w - 0.5f;
    float src_y = (y + 0.5f) * src_h / dst_h - 0.5f;
    
    // 双线性插值
    // ...（插值代码略）
    
    // 归一化
    dst[y * dst_w * 3 + x * 3 + c] = (pixel_value / 255.0f - mean) / std;
}

最佳实践：把缩放和归一化融合到一个kernel里，可以显著减少内存带宽消耗。我做过测试，分开做需要两次全局内存读写，融合后只需要一次，速度提升了40%以上。

3.4 完整的GPU预处理pipeline

在实际的自动驾驶系统中，这些步骤通常是串联在一起的。一个典型的GPU预处理pipeline是这样的：

  1. 图像解码：nvJPEG批量解码，输出GPU内存中的RGB/BGR数据
  2. 颜色空间转换：如果需要，将RGB转成模型需要的格式（比如BGR）
  3. 缩放+归一化：融合kernel，一步到位
  4. 数据排布转换：从NHWC转成NCHW（很多模型要求NCHW格式）

这个pipeline全部在GPU上完成，CPU只负责调度。我参与的一个量产项目中，用这套方案处理12路摄像头，每路1080p@30fps，GPU利用率不到30%，CPU几乎无负担。

避坑指南：我曾经遇到过一个问题——nvJPEG解码出来的数据是RGB格式，但模型要求的是BGR。如果直接在CPU上做这个转换，又要多一次CPU-GPU拷贝。正确的做法是在GPU上写一个RGB到BGR的kernel，或者直接用nvJPEG的API指定输出格式为BGR。nvJPEG是支持直接输出BGR的，只是很多人不知道这个参数。

3.5 性能对比与建议

操作 CPU耗时（ms） GPU耗时（ms） 加速比
JPEG解码（1920x1080） 35 0.8 ~44x
PNG解码（1920x1080） 200 12 ~17x
YUV420转RGB（1920x1080） 5 0.1 ~50x
缩放+归一化（1920x1080→640x480） 3 0.05 ~60x

从表格可以看出，GPU加速的效果非常显著。尤其是解码和颜色空间转换，加速比都在40倍以上。

最后给各位一个建议：不要小看预处理。很多感知系统的延迟问题，根源不在模型推理，而在数据预处理。把预处理搬到GPU上，往往是最简单、最有效的优化手段。我自己在多个项目中验证过，预处理GPU加速后，整个感知pipeline的端到端延迟能降低50%以上。

好了，这一章就到这里。下一章我们会聊聊更高级的话题——如何用CUDA Stream实现预处理和推理的流水线并行，让GPU一刻都不闲着。

← 上一章 📖 返回目录 下一章 →