OpenCL 中的快速 RGB => YUV 转换答案

【问题标题】：Fast RGB => YUV conversion in OpenCLOpenCL 中的快速 RGB => YUV 转换
【发布时间】：2011-02-12 17:36:34
【问题描述】：

我知道以下公式可用于将 RGB 图像转换为 YUV 图像。下式中，R、G、B、Y、U、V均为8位无符号整数，中间值为16位无符号整数。

Y = ( (  66 * R + 129 * G +  25 * B + 128) >> 8) +  16  
U = ( ( -38 * R -  74 * G + 112 * B + 128) >> 8) + 128  
V = ( ( 112 * R -  94 * G -  18 * B + 128) >> 8) + 128

但在 OpenCL 中使用该公式时，情况就不同了。
1. 8 位内存写访问是一个可选扩展，这意味着一些 OpenCL 实现可能不支持它。
2. 即使支持上述扩展，与32位写访问相比，它的速度要慢得多。

为了得到更好的性能，每4个像素会同时处理，所以输入是12个8位整数，输出是3个32位无符号整数（第一个代表4个Y样本，第二个代表 4 U 样本，最后一个代表 4 V 样本）。

我的问题是如何直接从 12 个 8 位整数中得到这 3 个 32 位整数？是否有公式可以得到这 3 个 32 位整数，或者我只需要使用旧公式得到 12 个 8 位整数结果（4 Y，4 U，4 V）并用位构造 3 个 32 位整数- 明智的操作？

【问题讨论】：

标签： alignment rgb opencl yuv

【解决方案1】：

尽管这个问题是在 2 年前提出的，但我认为一些工作代码在这里会有所帮助。考虑到直接访问 8 位值时性能不佳的最初担忧，最好尽可能执行 32 位直接访问。

前段时间，我开发并使用以下 OpenCL 内核将 ARGB（典型的 windows 位图像素布局）转换为 y 平面（全尺寸）、u/v 半平面（四分之一尺寸）内存布局，如下所示libx264 编码的输入。

__kernel void ARGB2YUV ( 
                            __global  unsigned int * sourceImage,
                            __global unsigned int * destImage,
            unsigned int srcHeight,
            unsigned int srcWidth,
            unsigned int yuvStride // must be srcWidth/4 since we pack 4 pixels into 1 Y-unit (with 4 y-pixels)
            )
{
    int i,j;
    unsigned int RGBs [ 4 ];
    unsigned int posSrc, RGB, Value4 = 0, Value, yuvStrideHalf, srcHeightHalf, yPlaneOffset, posOffset;
    unsigned char red, green, blue;

    unsigned int posX = get_global_id(0);
    unsigned int posY = get_global_id(1);

    if ( posX < yuvStride ) {
        // Y plane - pack 4 y's within each work item
        if ( posY >= srcHeight )
            return;

        posSrc = (posY * srcWidth) + (posX * 4);

        RGBs [ 0 ] = sourceImage [ posSrc ];
        RGBs [ 1 ] = sourceImage [ posSrc + 1 ];
        RGBs [ 2 ] = sourceImage [ posSrc + 2 ];
        RGBs [ 3 ] = sourceImage [ posSrc + 3 ];

        for ( i=0; i<4; i++ ) {
            RGB = RGBs [ i ];

            blue = RGB & 0xff; green = (RGB >> 8) & 0xff; red = (RGB >> 16) & 0xff;

            Value = ( ( 66 * red + 129 * green + 25 * blue ) >> 8 ) + 16;
            Value4 |= (Value << (i * 8));
        }

        destImage [ (posY * yuvStride) + posX ] = Value4;
        return;
    }

    posX -= yuvStride;
    yuvStrideHalf = yuvStride >> 1;

    // U plane - pack 4 u's within each work item
    if ( posX >= yuvStrideHalf )
        return;

    srcHeightHalf = srcHeight >> 1; 
    if ( posY < srcHeightHalf ) {
        posSrc = ((posY * 2) * srcWidth) + (posX * 8);

        RGBs [ 0 ] = sourceImage [ posSrc ];
        RGBs [ 1 ] = sourceImage [ posSrc + 2 ];
        RGBs [ 2 ] = sourceImage [ posSrc + 4 ];
        RGBs [ 3 ] = sourceImage [ posSrc + 6 ];

        for ( i=0; i<4; i++ ) {
            RGB = RGBs [ i ];

            blue = RGB & 0xff; green = (RGB >> 8) & 0xff; red = (RGB >> 16) & 0xff;
            Value = ( ( -38 * red + -74 * green + 112 * blue ) >> 8 ) + 128;
            Value4 |= (Value << (i * 8));
        }
        yPlaneOffset = yuvStride * srcHeight;
        posOffset = (posY * yuvStrideHalf) + posX;
        destImage [ yPlaneOffset + posOffset ] = Value4;
        return;
    }

    posY -= srcHeightHalf;
    if ( posY >= srcHeightHalf )
        return;

    // V plane - pack 4 v's within each work item
    posSrc = ((posY * 2) * srcWidth) + (posX * 8);

    RGBs [ 0 ] = sourceImage [ posSrc ];
    RGBs [ 1 ] = sourceImage [ posSrc + 2 ];
    RGBs [ 2 ] = sourceImage [ posSrc + 4 ];
    RGBs [ 3 ] = sourceImage [ posSrc + 6 ];

    for ( i=0; i<4; i++ ) {
        RGB = RGBs [ i ];

        blue = RGB & 0xff; green = (RGB >> 8) & 0xff; red = (RGB >> 16) & 0xff;

        Value = ( ( 112 * red + -94 * green + -18 * blue ) >> 8 ) + 128;
        Value4 |= (Value << (i * 8));
    }

    yPlaneOffset = yuvStride * srcHeight;
    posOffset = (posY * yuvStrideHalf) + posX;

    destImage [ yPlaneOffset + (yPlaneOffset >> 2) + posOffset ] = Value4;
    return;
}

此代码仅执行全局 32 位内存访问，而在每个工作项中进行 8 位处理。

哦..以及调用内核的正确代码

unsigned int width = 1024;
unsigned int height = 768;

unsigned int frameSize = width * height;
const unsigned int argbSize = frameSize * 4; // ARGB pixels

const unsigned int yuvSize = frameSize + (frameSize >> 1); // Y,U,V planes

const unsigned int yuvStride = width >> 2; // since we pack 4 RGBs into "one" YYYY

// Allocates ARGB buffer
ocl_rgb_buffer = clCreateBuffer ( context, CL_MEM_READ_WRITE, argbSize, 0, &error );
// ... error handling ...

ocl_yuv_buffer = clCreateBuffer ( context, CL_MEM_READ_WRITE, yuvSize, 0, &error );
// ... error handling ...

error = clSetKernelArg  ( kernel, 0, sizeof(cl_mem), &ocl_rgb_buffer );
error |= clSetKernelArg ( kernel, 1, sizeof(cl_mem), &ocl_yuv_buffer );

error |= clSetKernelArg ( kernel, 2, sizeof(unsigned int), &height);
error |= clSetKernelArg ( kernel, 3, sizeof(unsigned int), &width);

error |= clSetKernelArg ( kernel, 4, sizeof(unsigned int), &yuvStride);
// ... error handling ...

const size_t local_ws[] = { 16, 16 };
const size_t global_ws[] = { yuvStride + (yuvStride >> 1), height };

error = clEnqueueNDRangeKernel ( queue, kernel, 2, NULL, global_ws, local_ws, 0, NULL, NULL );
// ... error handling ...

注意：查看工作项计算。需要添加一些额外的代码（例如，使用 mod 以添加足够的备用项）以确保工作项大小适合本地工作大小。

【讨论】：

【解决方案2】：

像这样？除非您的平台可以使用 int3，否则请使用 int4。您还可以将 5 个像素打包到 int16 中，这样您就浪费了 1/16 而不是 1/4 的内存带宽。

__kernel void rgb2yuv( __global int3* input, __global int3* output){


rgb = input[get_global_id(0)];
R = rgb.x;
G = rgb.y;
B = rgb.z;    

yuv.x = ( (  66 * R + 129 * G +  25 * B + 128) >> 8) +  16; 
yuv.y = ( ( -38 * R -  74 * G + 112 * B + 128) >> 8) + 128; 
yuv.z = ( ( 112 * R -  94 * G -  18 * B + 128) >> 8) + 128;

output[get_global_id(0)] = yuv;
}

【讨论】：

【解决方案3】：

与 opencl specification 一起，数据类型 int3 不存在。

第 123 页：

支持的 n 值为 2、4、8 和 16...

在您的内核变量中，rgb、R、G、B 和 yuv 至少应为 __private int4。

OpenCL 1.1 添加了对typen 的支持，其中n = 3。但是，我强烈建议您不要使用它。不同的供应商实现有不同的错误，它并没有为您节省任何东西。

【讨论】：