CUDA Kernel实现P010LE转ARGB异常，求排查错误原因

我通过FFmpeg提取10位影片的P010LE格式帧，将帧数据加载到uint8_t类型的vector中，已验证数据无误。尝试编写CUDA Kernel将这些帧转换为ARGB格式，先仅处理亮度通道生成灰度图，但转换结果不符合预期（左侧为原图，右侧为转换后图），期望得到平滑的灰度图。以下是我的Kernel代码：

__global__ void P010leToArgbKernel(const uint8_t * P010le, uint8_t * argb, int width, int height)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < width && y < height)
    {
        int yIndexP010 = (y * width + x) * 2;

        uint16_t Y = *reinterpret_cast<const uint16_t*>(P010le + yIndexP010) & 0x3FF;
        uint8_t Y8 = static_cast<uint8_t>((Y * 255 + 511) / 1023); // Proper rounding

        int yIndex = (y * width + x) * 4;

        argb[yIndex + 0] = Y8;
        argb[yIndex + 1] = Y8;
        argb[yIndex + 2] = Y8;
        argb[yIndex + 3] = 255; // Alpha
    }
}

请问我哪里出错了？

问题分析与修复 #

你的Kernel核心问题是未对齐内存访问：CUDA设备对数据类型的内存访问有严格对齐要求，访问uint16_t类型数据时，地址必须是2字节对齐的。但你直接将uint8_t*强制转换为uint16_t*访问，当yIndexP010对应的地址不是2的整数倍时，会触发未对齐访问，导致读取到错误的16位数值，最终输出异常的灰度图。

另外，强制类型转换的方式在未对齐场景下会彻底破坏P010LE的小端字节序，进一步加剧数据错误。

修复方案 #

放弃强制类型转换，手动从uint8_t数组中拼接出10位的Y分量，规避内存对齐限制：

__global__ void P010leToArgbKernel(const uint8_t * P010le, uint8_t * argb, int width, int height)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < width && y < height)
    {
        int yIndexP010 = (y * width + x) * 2;

        // 手动拼接小端格式的16位数据，避免未对齐访问
        uint8_t y_lo = P010le[yIndexP010];
        uint8_t y_hi = P010le[yIndexP010 + 1];
        uint16_t Y = (static_cast<uint16_t>(y_hi) << 8) | y_lo;
        Y &= 0x3FF; // 保留低10位有效数据

        uint8_t Y8 = static_cast<uint8_t>((Y * 255 + 511) / 1023); // 四舍五入转换为8位灰度

        int yIndex = (y * width + x) * 4;
        argb[yIndex] = Y8;
        argb[yIndex + 1] = Y8;
        argb[yIndex + 2] = Y8;
        argb[yIndex + 3] = 255; // Alpha通道设为不透明
    }
}

补充说明 #

• 你原本的10位转8位灰度的计算逻辑(Y * 255 + 511) / 1023是正确的，通过加半最大值实现了四舍五入，保证了灰度转换的平滑性。
• 修复后的代码通过逐字节读取再拼接的方式，完全规避了CUDA的内存对齐限制，能正确解析P010LE格式的亮度数据。

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