rigayaの日記兼メモ帳

今回はconvolution計算を例に、imageNd(image1d/image2d/image3d)の使い方を確認した。

---

簡単な例として、2次元畳みこみ演算を考える。

float *input;  //入力配列
float *output; //出力配列
int width, pitch, height; //2Dサイズ
//重みは適当
const float weight[3][3] = {
    { 1.0f, 2.0f, 1.0f },
    { 2.0f, 4.0f, 2.0f },
    { 1.0f, 2.0f, 1.0f },
};
for (int y = 0; y < height; y++) {
    for (int x = 0; x < width; x++) {
        float sump = 0.0f;
        float sumw = 0.0f;
        for (int j = -1; j <= 1; j++) {
            int yj = CLAMP(y+j, 0, height-1);
            for (int i = -1; i <= 1; i++) {
                int xi = CLAMP(x+i, 0, width-1);
                float w = weight[j+1][i+1];
                sump += w * input[yj * pitch + xi];
                sumw += w;
            }
        }
        output[y * pitch + x] = sump / sumw;
    }
}

これをGPUのkernelに置き換えると、

__kernel void convolve(
    __global float* pIn,
    __global float* pOut,
    int width, int pitch, int height) {
    const int x = get_global_id(0);
    const int y = get_global_id(1);

    if (x < width && y < height) {
        const float weight[] = {
            1.0f, 2.0f, 1.0f,
            2.0f, 4.0f, 2.0f,
            1.0f, 2.0f, 1.0f,
        };
        float sump = 0.0f;
        float sumw = 0.0f;
        int iw = 0;
        for (int j = -1; j <= 1; j++) {
            int yj = clamp(y+j, 0, height-1);
            for (int i = -1; i <= 1; i++) {
                int xi = clamp(x+i, 0, width-1);
                float w = weight[iw];
                sumw += w;
                sump += w * pIn[yj * pitch + xi];
                iw++;
            }
        }
        pOut[y * pitch + x] = sump / sumw;
    }
}

本来は重みも引数で渡すべきだけど、今回はそこは本題ではないので、気にしないことにする。

まあ、この畳みこみは非常に簡単な計算で、自分とその周辺について重みをかけて足しこむというものである。



このとき、入力フレームの画素は、重複して何度もロードされることになる。こういうアクセスパターンの場合、OpenCLのimageを使うとテクスチャキャッシュが効いて速い…と聞いたことがあるので、とりあえずやってみた。

---

image2dは、cl_image_format構造体とcl_image_desc構造体に適切にパラメータを設定して、clCreateImage()によって作ることができる。

ホスト(CPU)のポインタからimageNd(image1d/image2d/image3d)を作る場合には、clCreateImageによって直接imageNdを作る方法と、一度OpenCLバッファを経由してからclCreateImageで作る方法がある。ただ、後者の方法で作成しないとZeroCopyにならず、転送が無駄に発生してしまう点に注意する必要がある。(わりと罠)

直接image2dを作る
ホスト側のポインタをclCreateImageのhost_ptr引数に指定している。ZeroCopyにならない。

//パラメータを適切に設定する
cl_image_format format;
format.image_channel_order = CL_R;         //チャンネル数
format.image_channel_data_type = CL_FLOAT; //データ型

cl_image_desc img_desc;
img_desc.image_type = CL_MEM_OBJECT_IMAGE2D; //2D
img_desc.image_width = arrayWidth;   //サイズ
img_desc.image_height = arrayHeight; //サイズ
img_desc.image_depth = 0;            //サイズ
img_desc.image_array_size = 0;
img_desc.image_row_pitch = arrayPitch * sizeof(cl_float);
img_desc.image_slice_pitch = 0;
img_desc.num_mip_levels = 0;
img_desc.num_samples = 0;
img_desc.buffer = 0;
img_desc.mem_object = 0;

ocl->srcImg = clCreateImage(ocl->context,
    CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
    &format, &img_desc,
    input, //ホスト側のポインタ
    &err);

一度OpenCLバッファを経由する
この方法はこちらで触れられている。

一度OpenCLバッファを作成し、そのバッファをcl_image_desc構造体のmem_objectに指定してclCreateImage()を呼ぶ。この場合は、もとのOpenCLバッファがZeroCopyなら、作成したimage2dもZeroCopyとなる。

//これまでと同じようにOpenCLバッファを作成する
cl_uint nSize = sizeof(cl_float) * arrayPitch * arrayHeight;
ocl->srcMem = clCreateBuffer(ocl->context,
    CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
    nSize,
    input, //ホスト側のポインタ
    &err);

//パラメータを適切に設定する
cl_image_format format;
format.image_channel_order = CL_R;         //チャンネル数
format.image_channel_data_type = CL_FLOAT; //データ型

cl_image_desc img_desc;
img_desc.image_type = CL_MEM_OBJECT_IMAGE2D; //2D
img_desc.image_width = arrayWidth;   //サイズ
img_desc.image_height = arrayHeight; //サイズ
img_desc.image_depth = 0;
img_desc.image_array_size = 0;
img_desc.image_row_pitch = arrayPitch * sizeof(cl_float);
img_desc.image_slice_pitch = 0;
img_desc.num_mip_levels = 0;
img_desc.num_samples = 0;
img_desc.buffer = 0;
img_desc.mem_object = ocl->srcMem; //作成したOpenCLバッファをここに指定する

ocl->srcImg = clCreateImage(ocl->context,
    CL_MEM_READ_ONLY,
    &format, &img_desc,
    nullptr,
    &err);

作成したimageの解放は、OpenCLバッファと同じように行えばよい。

err = clReleaseMemObject(srcImg);

cl_image_format構造体で、使用するimageのformat(チャンネルのフォーマットとデータ型)を指定したが、imageで使用可能なformatは様々なものが用意されている。

まず、チャンネル数については、下記から選択できるが、フォーマットによっては特定のデータ型との組み合わせのみ可能なものがある。

フォーマット	データ型の制約
CL_R CL_Rx CL_A
CL_INTENSITY	CL_UNORM_INT8, CL_UNORM_INT16, CL_SNORM_INT8, CL_SNORM_INT16, CL_HALF_FLOAT, CL_FLOAT
CL_LUMINANCE	CL_UNORM_INT8, CL_UNORM_INT16, CL_SNORM_INT8, CL_SNORM_INT16, CL_HALF_FLOAT, CL_FLOAT
CL_DEPTH	CL_UNORM_INT16, CL_FLOAT
CL_RG CL_RGx CL_RA
CL_RGB CL_RGBx	CL_UNORM_SHORT_565, CL_UNORM_SHORT_555, CL_UNORM_INT101010
CL_RGBA
CL_sRGB CL_sRGBx CL_sRGBA CL_sBGRA	CL_UNORM_INT8
CL_ARGB CL_BGRA CL_ABGR	CL_UNORM_INT8, CL_SNORM_INT8, CL_SIGNED_INT8, CL_UNSIGNED_INT8
CL_DEPTH_STENCIL	CL_UNORM_INT24, CL_FLOAT

また、データ型は以下から選ぶことになる。選んだデータ型によってkernelコードで使用すべき関数と、kernelで取り出せる値のデータ型が変わる。また、データ型によっては、特定のフォーマットとの組み合わせのみ可能なものがある。

データ型	kernelで取り出せるデータ型	formatの制約
CL_SNORM_INT8	規格化された符号付き8bit整数
CL_SNORM_INT16	規格化された符号付き16bit整数
CL_UNORM_INT8	規格化された符号なし8bit整数
CL_UNORM_INT16	規格化された符号なし16bit整数
CL_UNORM_SHORT_565	規格化された 5-6-5 3-channel RGB	CL_RGB, CL_RGBx
CL_UNORM_SHORT_555	規格化された x-5-5-5 4-channel xRGB	CL_RGB, CL_RGBx
CL_UNORM_INT_101010	規格化された x-10-10-10 4-channel xRGB	CL_RGB, CL_RGBx
CL_SIGNED_INT8	符号付き8bit整数
CL_SIGNED_INT16	符号付き16bit整数
CL_SIGNED_INT32	符号付き32bit整数
CL_UNSIGNED_INT8	符号なし8bit整数
CL_UNSIGNED_INT16	符号なし16bit整数
CL_UNSIGNED_INT32	符号なし32bit整数
CL_HALF_FLOAT	半精度浮動小数点
CL_FLOAT	単精度浮動小数点
CL_UNORM_INT24	規格化された符号なし24bit整数

imageNdをカーネルで使う際には、すこしコードを変える必要があり、
・引数としては、imageNd_t型
・取り出す際にread_imagef(floatで取り出す場合)、read_imagei/read_imageui(整数型で取り出す場合)を使用
の2点変更する必要がある。

read_imageの取り出し方にはさまざまな方法があるが、まずデータ型によって使用すべき関数が決まっている。

使用すべき関数	データ型	値域
read_imagef	CL_UNORM_INT8 CL_UNORM_INT16	規格化された値 [-1.0 … 1.0]
read_imagef	CL_SNORM_INT8 CL_SNORM_INT16	規格化された値 [0.0 … 1.0]
read_imagef	CL_HALF_FLOAT CL_FLOAT	入力と同じ値
read_imagei	CL_SIGNED_INT8 CL_SIGNED_INT16 CL_SIGNED_INT32	入力と同じ値
read_imageui	CL_UNSIGNED_INT8 CL_UNSIGNED_INT16 CL_UNSIGNED_INT32	入力と同じ値

座標値の指定は、image1d/image2dであればint2型/float2型で、image3dであればint4型/float4型で指定する。

座標の指定方法と値の取り出し方について、3つのオプションがあり、これを指定することができる。

・取り出す座標を指定する際に、0.0～1.0に規格化された座標を用いるかどうか
　- CLK_NORMALIZED_COORDS_TRUE (規格化された浮動小数点の座標で指定)
　- CLK_NORMALIZED_COORDS_FALSE (通常のインデックスで指定)

・領域外の座標を指定した場合にどのように処理するか
　- CLK_ADDRESS_NONE (なにもしない > どんな値が取得されるかは未定義)
　- CLK_ADDRESS_CLAMP (imageの境界の色が採用される)
　- CLK_ADDRESS_CLAMP_TO_EDGE (imageの端の値が採用される)
　- CLK_ADDRESS_REPEAT (imageの反対側の端から繰り返す)
　- CLK_ADDRESS_MIRRORED_REPEAT (imageを折り返して繰り返す)

・座標値に合わせて値の補間を行うか
　- CLK_FILTER_NEAREST (補間を行わず、最も近い座標の値をそのまま使用する)
　- CLK_FILTER_LINEAR (線形補間)

また、関数によりimageNdから取得された値は、float4, half4, int4, short4, char4などの4要素のベクトル型であるが、imageのフォーマット(チャンネル)によって、何番目の要素にどのデータが入っているかが異なる。

チャンネル	data4の中身
CL_R	(r, 0, 0, 1)
CL_A	(0, 0, 0, a)
CL_RG	(r, g, 0, 1)
CL_RA	(r, 0, 0, a)
CL_RGB	(r, g, b, 1)
CL_RGBA CL_BGRA CL_ARGB	(r, g, b, a)
CL_INTENSITY	(I, I, I, I)
CL_LUMINANCE	(L, L, L, 1)

これらを踏まえて、kernelを書き直すと、以下のようになる。

__kernel void convolve(
    __read_only image2d_t imgIn,
    __global float* pOut,
    int width, int pitch, int height) {
    const int x = get_global_id(0);
    const int y = get_global_id(1);

    if (x < width && y < height) {
        const float weight[] = {
            1.0f, 2.0f, 1.0f,
            2.0f, 4.0f, 2.0f,
            1.0f, 2.0f, 1.0f,
        };
        float sump = 0.0f;
        float sumw = 0.0f;
        int iw = 0;
        for (int j = -1; j <= 1; j++) {
            int yj = clamp(y+j, 0, height-1);
            for (int i = -1; i <= 1; i++) {
                int xi = clamp(x+i, 0, width-1);
                float w = weight[iw];
                sumw += w;
                sump += w * read_imagef(imgIn, (int2)(xi,yj)).x;
                iw++;
            }
        }
        pOut[y * pitch + x] = sump / sumw;
    }
}

kernelのほうは引数を変更したが、ホスト(CPU)からkernel引数を渡す場合は特に変更は必要なく、普通にclSetKernelArgで渡せばよい。

clSetKernelArg(ocl->kernel, //対象のkernel
    0,                     //第0引数
    sizeof(cl_mem),        //引数のサイズ
    (void *)&ocl->srcImg); //引数へのポインタ

とすればよい。

---

計算速度の確認

image2dを使用しない場合


image2dを使用した場合


というわけでkernelの実行時間をみると1.91ms → 2.94msと、image2dを使うことで逆に遅くなってしまった。原因は残念ながらよくわからない…。コードの書き方の問題だろうか…?

まあ、とりあえずOpenCLのimageの使い方の確認ができたのと、(コードの書き方が悪いのかもしれないが)遅くなることもある、ということがわかった。

コードはこちら。

---

続き>>

---

OpenCLメモリスト
OpenCLメモ (その1) - かんたんな計算
OpenCLメモ (その2) - Intel GPUの構造
OpenCLメモ (その3) - work sizeの調整
OpenCLメモ (その4) - 転送(コピー)の排除 (USE_HOST_PTR)
OpenCLメモ (その5) - 転送(コピー)の排除 (SVM : OpenCL 2.0)
OpenCLメモ (その6) - imageの使用 (いまここ)
OpenCLメモ (その7) [終] - reductionとshared local memory(SLM)の使用

---