5

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 ビット符号なし整数です (最初の 1 つは 4 つの Y サンプルを表し、2 番目はは 4 U サンプルを表し、最後のものは 4 V サンプルを表します)。

私の質問は、これらの 3 つの 32 ビット整数を 12 の 8 ビット整数から直接取得する方法です。これらの 3 つの 32 ビット整数を取得する式はありますか、または古い式を使用して 12 個の 8 ビット整数の結果 (4 Y、4 U、4 V) を取得し、3 つの 32 ビット整数をビットで構築する必要があります。 -賢明な操作？

4

3 に答える 3

10

この質問は 2 年前に出されたものですが、動作するコードがここで役立つと思います。8 ビット値に直接アクセスするときのパフォーマンスの低下に関する最初の懸念に関しては、可能であれば 32 ビット直接アクセスを実行することをお勧めします。

少し前に、次の OpenCL カーネルを開発して使用し、libx264 の入力として ARGB (典型的な Windows ビットマップ ピクセル レイアウト) を y プレーン (フル サイズ)、u/v-ハーフプレーン (1/4 サイズ) メモリ レイアウトに変換しました。エンコーディング。

__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 を使用するなど)。

于 2013-07-30T20:36:22.123 に答える
2

このような？プラットフォームが int3 を使用できない場合は、int4 を使用してください。また、5 ピクセルを int16 にパックできるため、メモリ帯域幅の 1/4 ではなく 1/16 を浪費します。

__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;
}
于 2011-02-14T22:09:34.223 に答える
2

opencl 仕様に伴い、データ型 int3 は存在しません。

123ページ：

サポートされている n の値は、2、4、8、および 16 です...

カーネル変数rgbでRはG、、、、Bおよびyuv少なくとも__private int4.

OpenCL 1.1 では、typenwhereのサポートが追加されましたn = 3。ただし、使用しないことを強くお勧めします。異なるベンダーの実装には異なるバグがあり、何も節約できません.

于 2011-02-28T01:01:52.127 に答える