opencl => カーネルの基礎

前書き

このトピックでは、opencl用のカーネルを書く際の基本について説明します

グレースケールカーネル

グレースケール画像を生成するカーネルを構築します。それぞれのコンポーネントとRGBAのオーダーを使用して定義されたイメージデータを使用します。

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                                CLK_ADDRESS_CLAMP_TO_EDGE |
                                CLK_FILTER_LINEAR;

__kernel void Grayscale(__read_only image2d_t input, __write_only image2d_t output) {
    int2 gid = (int2)(get_global_id(0), get_global_id(1));
    int2 size = get_image_dim(input);

    if(all(gid < size)){
        uint4 pixel = read_imageui(input, sampler, gid);
        float4 color = convert_float4(pixel) / 255;
        color.xyz = 0.2126*color.x + 0.7152*color.y + 0.0722*color.z;
        pixel = convert_uint4_rte(color * 255);
        write_imageui(output, gid, pixel);
    }
}

今度は、そのコードをステップごとに歩いてみましょう。最初の行は、sampler_t型の__定数メモリ領域に変数を作成します。このサンプラーは、画像データへのアクセスをさらに指定するために使用されます。詳しくはKhronos Docsを参照してください。

カーネルを呼び出す前に入力をread_only、出力をwrite_onlyとして割り当てたので、ここでそれらの修飾子を追加します。

image2dとimage3dは常にグローバルメモリに割り当てられるため、ここでは__global修飾子を省略できます。

次に、グレースケールに変換するピクセルを定義するスレッドIDを取得します。また、スレッドが未割り当てメモリにアクセスしていないことを確認するためにサイズを問い合わせます。これは、あなたがそれを忘れた場合、あなたのカーネルをクラッシュさせます。

正当なスレッドであることを確認した後、入力画像からピクセルを読み込みます。次に、小数点以下の桁数の損失を避けるために浮動小数点数に変換し、計算を行い、逆変換して出力に書き出します。

カーネルスケルトン

そこにある最も単純なカーネルとそのバリエーションを見てみましょう

__kernel void myKernel() {
}

メインコードから起動できるカーネルは、__kernelキーワードで識別されます。カーネル関数は戻り値の型がvoidのみです。

__kernel void myKernel(float a, uint b, byte c) {

}

もちろん、カーネルとして公開されていない多くの関数を作成することができます。この場合、__kernel修飾子を省略することができます。

関数は、他のC / C ++関数のように変数を公開することがあります。唯一の違いは、メモリを参照したいときです。これは、引数であるかコードで使用されているかにかかわらず、すべてのポインタに適用されます。

float*  ptr;

実行中のスレッドのみがアクセスできるメモリ領域へのポインタです。実際には

__private float* ptr;

利用可能な4つの異なるメモリ領域修飾子があります。カーネルの中では、通常、心配する必要はありませんが、議論のためには、これは不可欠です。

__global：この修飾子は、グローバルメモリに配置されたポインタを参照します
__constant：定数メモリポインタを参照します。
__local：共有メモリポインタを参照します。
__private：ローカルメモリポインタを参照します。

さらに、メモリへのアクセス方法を定義することもできます

修飾子なし：読み書き
__read_only
__write_only

これらのフラグは、メモリバッファをホストに割り当てた方法と一致する必要があります。

カーネルID

データを適切に処理するために、各スレッドはスレッドブロック/グローバルスレッドプール内の位置を知る必要があります。これは

get_local_id($dim);
get_global_id($dim);

これらの2つの関数は、スレッドブロックまたはすべてのスレッドに対してスレッドの位置を返します。

get_working_dim();

カーネルが起動されたディメンションの合計数を取得します。

get_local_size($dim);
get_global_size($dim);

指定したディメンションのスレッドブロックまたはスレッドブロック内のスレッドの合計数を取得します。

警告：スレッドがデータサイズを超えていないことを常に確認してください。これは非常に起こりそうで、常にチェックする必要があります。

OpenCLのベクトル

それぞれの基本的なopencl型にはベクトルバージョンがあります。タイプの後に目的のコンポーネントの数を追加して、ベクタータイプを使用できます。サポートされているコンポーネントの数は2,3,4,8および16です.OpenCL 1.0には3つのコンポーネントはありません。

次の2つの方法でベクトルを初期化できます。

単一のスカラーを提供する
すべてのコンポーネントを満足させる

float4 a = (float4)(1); //a = (1, 1, 1, 1)

または

float4 b = (float4)(1, 2, 3, 4);
float4 c = (float4)(1, (float3)(2));

または構成要素の数を満たすベクトルの任意の他の組み合わせを含むことができる。ベクトルの要素にアクセスするには、さまざまな方法を使用できます。インデックスを使用するか、

a[0] = 2;

またはリテラルを使用します。リテラルの利点は、あなたが望む任意の方法でそれらを組み合わせることができることです。それをすぐに行います。すべてのベクターコンポーネントには、

a.s0 = 2; // same as a[0] = 2

複数のコンポーネントを組み合わせて新しいベクターにすることもできます

a.s02 = (float2)(0, 0); // same as  a[0] = 0; a[2] = 0; or even a.s20 = (float2)(0, 0)

必要な方法でコンポーネントの順序を変更することができます。

a.s1423 = a.s4132; // flip the vector

しかし、あなたは何かのようにすることはできません

a.s11 = ... // twice the same component is not possible

ベクトルコンポーネントにアクセスするためのいくつかの便利な省略形があります。以下の簡略表記は、サイズ2,4,8,16にのみ適用されます

a.hi //=a.s23 for vectors of size 4, a.4567 for size 8 and so on.
a.lo //=a.s01
a.even //=a.s02
a.odd //=a.13

ベクトル・サイズ2,3および4については、いくつかの追加の簡略表記

a.x //=a.s0
a.y //=a.s1
a.z //=a.s2
a.w //=a.s3

ガンマ補正カーネル

ガンマ補正カーネルを見てみましょう

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                                CLK_ADDRESS_CLAMP_TO_EDGE |
                                CLK_FILTER_LINEAR;

__kernel void Gamma(__read_only image2d_t input, __write_only image2d_t output, __constant float gamma) {
    int2 gid = (int2)(get_global_id(0), get_global_id(1));
    int2 size = get_image_dim(input);

    if(all(gid < size)){
        uint4 pixel = read_imageui(input, sampler, gid);
        float4 color = convert_float4(pixel) / 255;
        color = pow(color, (float4)(gamma));
        pixel = convert_uint4_rte(color * 255);
        write_imageui(output, gid, pixel);
    }
}

カーネルを呼び出す前に入力をread_only、出力をwrite_onlyとして割り当てたので、ここでそれらの修飾子を追加します。

image2dとimage3dは常にグローバルメモリに割り当てられるため、ここでは__global修飾子を省略できます。私たちのガンマ値は__定数メモリ内にあるので、それも指定します。

次に、ガンマ補正を行うピクセルを定義するスレッドIDを取得します。また、スレッドが未割り当てメモリにアクセスしていないことを確認するためにサイズを問い合わせます。これは、あなたがそれを忘れた場合、あなたのカーネルをクラッシュさせます。

Modified text is an extract of the original Stack Overflow Documentation

ライセンスを受けた CC BY-SA 3.0

所属していない Stack Overflow

opencl
カーネルの基礎

サーチ…

前書き

グレースケールカーネル

カーネルスケルトン

カーネルID

OpenCLのベクトル

ガンマ補正カーネル