NVIDIA CUDA NV12(YUV)转RGB内核逻辑与色度复用问题咨询

修改NVIDIA Video SDK提供的原生YUV转RGB CUDA内核时，需要厘清部分设计逻辑，涉及的内核代码如下：

template<class YuvUnitx2, class Rgb, class RgbIntx2>
__global__ static void YuvToRgbKernel(uint8_t* pYuv, int nYuvPitch, uint8_t* pRgb, int nRgbPitch, int nWidth, int nHeight) {

int x = (threadIdx.x + blockIdx.x * blockDim.x) * 2;
int y = (threadIdx.y + blockIdx.y * blockDim.y) * 2;
if (x + 1 >= nWidth || y + 1 >= nHeight) {
    return;
}

uint8_t* pSrc = pYuv + x * sizeof(YuvUnitx2) / 2 + y * nYuvPitch;
uint8_t* pDst = pRgb + x * sizeof(Rgb) + y * nRgbPitch;

YuvUnitx2 l0 = *(YuvUnitx2*)pSrc;
YuvUnitx2 l1 = *(YuvUnitx2*)(pSrc + nYuvPitch);
YuvUnitx2 ch = *(YuvUnitx2*)(pSrc + (nHeight - y / 2) * nYuvPitch);

//YuvToRgbForPixel - returns rgba encoded in uint32_t (.d)

*(RgbIntx2*)pDst = RgbIntx2{
    YuvToRgbForPixel<Rgb>(l0.x, ch.x, ch.y).d,
    YuvToRgbForPixel<Rgb>(l0.y, ch.x, ch.y).d,
};
*(RgbIntx2*)(pDst + nRgbPitch) = RgbIntx2{
    YuvToRgbForPixel<Rgb>(l1.x, ch.x, ch.y).d,
    YuvToRgbForPixel<Rgb>(l1.y, ch.x, ch.y).d,
};
}

提出的两点基础假设如下：

• 假设1：NV12格式包含两个平面，第1个平面存储Luma（亮度）分量，第2个平面存储交错排列的Chroma（色度）分量
• 假设2：该内核设计为单次处理写入4个像素
  核心疑问：若假设2成立，为何所有4个像素都使用相同的ch（色度）值？若假设2不成立，请详细解释该内核的实际执行逻辑。

假设正确性判定 #

两点假设全部成立：

1. 对NV12格式的描述完全准确：NV12是典型的4:2:0采样YUV格式，亮度平面分辨率与画面全分辨率一致，UV色度平面的水平、垂直分辨率均为亮度平面的1/2，且U、V分量以交错方式存储。
2. 单线程处理4像素的判断也完全准确：每个线程计算初始坐标时，x、y方向都做了乘2偏移，最终覆盖(x,y)、(x+1,y)、(x,y+1)、(x+1,y+1)组成的2×2像素块，单线程完成这4个像素的转换写入。

4个像素共用同一组色度值的原因 #

这不是代码逻辑问题，完全是NV12格式本身的4:2:0采样规则决定的：

NV12的4:2:0采样规则明确规定，每相邻的2×2亮度像素块，共用同一组U、V色度值。

代码中取色度值的逻辑也完全匹配这个规则：

• 色度平面位于亮度平面之后，在YUV缓冲区中的起始偏移为nHeight * nYuvPitch（亮度平面共nHeight行，每行步长为nYuvPitch）
• 色度平面的垂直坐标取y/2，对应垂直方向每2行亮度匹配1行色度的规则
• 读取到的ch.x是当前2×2像素块共用的U分量，ch.y是共用的V分量
  代码给4个像素传入相同的ch.x、ch.y做色彩转换，完全符合NV12的格式规范。

代码设计细节补充 #

• 模板参数YuvUnitx2一般定义为ushort2或等价的2字节聚合类型，一次可以读取2个连续的8位Y/U/V值，读2行亮度数据就能凑齐2×2块的4个亮度值，配合单线程处理4像素的逻辑做了向量化访存优化，比逐字节读写效率高很多
• 模板参数RgbIntx2一般定义为uint2类型，一次可以存储2个32位的RGBA像素值，刚好对应一行2个像素的输出，写2次就完成2行共4个像素的输出写入
• 边界判断逻辑x + 1 >= nWidth || y + 1 >= nHeight也和2×2块的处理逻辑匹配，只要块的右边界或下边界超出画面范围就直接返回，避免越界访存。

内容的提问来源于stack exchange，提问作者Michael IV