rigayaの日記兼メモ帳

---

配列の総和の計算は、CPUだと非常に単純で、

float sum = 0.0f;
for (int i = 0; i < nSize; i++) {
    sum += ptrArray[i];
}

とまあ、これだけなのだが、このままのループの形だと並列性がないので、GPUでやろうとするとやけに面倒だったりする。

こういったいわゆるreductionの計算は、なんとか並列性をひねりだし、work group内で共有できるメモリ(shared local memory, SLM)を使うと効率的に計算できることが知られている。

計算の仕方としては、もう図にしたほうが早いので、下のような感じ。(図はwork groupあたり16 work itemの例)
CUDAでもお馴染みの方法だったりする(最近はshuffleとかも使うらしいが)

まあ、ひとつづつ足すのではなく、なるべく並列に計算できるところを見つけて、GPUを活かそう、というもの。無限に演算器があるなら、計算のオーダーをO(n)からO(log(n))に落とすイメージ。

通常、GPUでは異なるスレッドの計算結果を得ることはできないが、__localメモリを使えば、work group内の計算結果を取得できることを利用している。



このようにkernelを1回起動することで、work group内のwork item分の和を取ることができる。つまり、work groupあたり16work itemの設定なら、1回で要素数を1/16にすることができ、これを繰り返すことで総和を得ることができる。もちろん、work groupあたりのwork item数が大きいほど一気に要素数を減らせるので、実際にはwork groupあたり16work itemなどにする。

実際のkernelコードはこんな感じ。

__kernel void reduce_add(__global float* pIn, __global float* pOut, int nSize) {
    const int lid = get_local_id(0);
    const int gid = get_global_id(0);
    __local float shared[GROUP_SIZE];

    shared[lid] = (gid < nSize) ? pIn[gid] : 0;
    barrier(CLK_LOCAL_MEM_FENCE);
    for (int offset = get_local_size(0) >> 1;
        offset > 0; offset >>= 1) {
        if (lid < offset) {
            shared[lid] += shared[lid + offset];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    if (lid == 0) {
        pOut[get_group_id(0)] = shared[0];
    }
}

shared local memoryは、kernelコードでは__localで取ることができる。

__localに書き込んだ値を他のスレッドで使う場合、かならず

barrier(CLK_LOCAL_MEM_FENCE);

によりwork group内のスレッドの同期をとる必要がある。work group内のスレッドはすべてが同時に実行されるわけではないので、きちんとすべてのスレッドが書き込みを終えてから次に進むようにしないと、おかしな値を使用してしまう可能性があるからである。

実際に実行している様子はこちら。



よく見えないけど、reduction_addが2回実行されていて4M要素の総和を取っている。(4M → 16K → 64)

---

この方法、GPUでもある程度高速に総和を取れてよいのだけど、加算の計算順序が変わるので、計算結果がCPUと多少変わるのが面倒なところ。浮動小数点演算は数学の計算と違って可換ではないのだ…。

コードはこちら。

---

OpenCLは勉強自体はこれまで少しづつやっていたことだが、時間がなかったのでこれまでまとめたりできなかった。この3連休でやっとまとめることができたので、これできっともう忘れないだろう(棒)

だいたい通常の計算と、image、__localの使い方がわかったし、これでやっとOpenCLの基本の基本ぐらいはわかったように思う。おおまかにCUDAとの対応関係もつかめたし、多少はCUDAの知識も使えそうな感じ。

あとは、実際にやりたいこと次第、になるのだと思う。

---

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)の使用 (いまここ)

---

---

簡単な例として、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)の使用

---

---

まず、選択したデバイスがSVM(Shared Virtual Memory)をサポートしているか確認する必要がある。こんな感じで調べられる。

cl_device_svm_capabilities caps;
err = clGetDeviceInfo(ocl->device, CL_DEVICE_SVM_CAPABILITIES,
    sizeof(cl_device_svm_capabilities), &caps, NULL);
if (CL_SUCCESS != err) {
	return err;
}
ocl->svmCoarse     = 0!=(caps & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER);
ocl->svmFineBuffer = 0!=(caps & CL_DEVICE_SVM_FINE_GRAIN_BUFFER);
ocl->svmFineSystem = 0!=(caps & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM);

SVMには3種類あって、順番により強力になる。

・SVM Coarse Grain Buffer
ｃｌSVMAllocという専用の関数を使って確保する。ホストから触るときには、USE_HOST_PTRなどの時と同様にMap / Unmapによる同期が必要なSVM。

・SVM Fine Grain Buffer
ｃｌSVMAllocという専用の関数を使って確保する。ホストから触るときでも、Map / Unmap不要(常に同期がとられる)。

・SVM Fine Grain System
普通のホスト(CPU)メモリのポインタをデバイス(GPU)と共有できる。

現時点でIntel HD Graphics 530 (i7 6700K)がサポートするのはFine Grain Bufferまでとなっている。

---

SVM Coarse Grain Buffer

---

まず、_aligned_mallocの代わりにclSVMAllocでメモリを確保する。

cl_uint nSize = sizeof(cl_int) * arrayWidth * arrayHeight;
cl_int* inputA  = (cl_int*)clSVMAlloc(
    ocl.context,      //関連付けるcontext
    CL_MEM_READ_ONLY, //kernelからの使用方法
    nSize,            //確保するサイズ
    0);               //アライメント(0 = 自動)

cl_int* inputB  = (cl_int*)clSVMAlloc(
    ocl.context, CL_MEM_READ_ONLY, nSize, 0);
cl_int* outputC = (cl_int*)clSVMAlloc(
    ocl.context, CL_MEM_READ_ONLY, nSize, 0);

これだけで、_aligned_mallocとclCreateBufferが同時にできたような感じ。

clSVMAllocしたメモリはclSVMFreeで解放すればよい。

    clSVMFree(ocl.context, inputA);
    clSVMFree(ocl.context, inputB);
    clSVMFree(ocl.context, outputC);

ホスト(CPU)側からこの領域を操作するには、USE_HOST_PTRの時同様、Map / Unampが必要になる。clEnqueueSVMMap() / clEnqueueSVMUnmap()を使う。

//対応するホスト側のポインタを取得する
err = clEnqueueSVMMap(
    ocl.commandQueue, //投入キュー
    CL_FALSE,         //終了までブロックするか -> しない
    CL_MAP_WRITE,     //CPUが書き込むためにMapする
                      //(読み込みならCL_MAP_READ)
                      //(両方ならCL_MAP_READ | CL_MAP_WRITE)
    inputA,           //マップするポインタ
    sizeof(cl_uint) * arrayWidth * arrayHeight, //マップするサイズ
    0,      //この関数が待機すべきeventの数
    NULL,   //この関数が待機すべき関数のリストへのポインタ
    NULL);  //この関数の返すevent

//終了を待機
err = clFinish(ocl.commandQueue);

//ホスト(CPU)から操作
//(...略...)

err = clEnqueueSVMUnmap(
    ocl.commandQueue, //投入キュー
    inputA,     //対象のポインタ
    0,          //この関数が待機すべきeventの数
    NULL,       //この関数が待機すべき関数のリストへのポインタ
    NULL);      //この関数の返すevent

//終了を待機
err = clFinish(ocl.commandQueue);

//デバイス(GPU) kernelを実行
//(...略...)

このSVMのポインタをkernelに渡す場合、clSetKernelArgSVMPointer()という専用関数を使う必要がある。

err = clSetKernelArgSVMPointer(ocl->kernel, 0, inputA);
err = clSetKernelArgSVMPointer(ocl->kernel, 1, inputB);
err = clSetKernelArgSVMPointer(ocl->kernel, 2, outputC);

この後の流れは同じ。

確認すると、きちんと転送がなくなっていることがわかる。



"Add"kernelの実行時間は1.741ms。



ちゃんと動くコードはこちら。

---

SVM Fine Grain Buffer

---

clSVMAllocで指定するflagにCL_MEM_SVM_FINE_GRAIN_BUFFERをつけるだけで使用できる。

cl_uint nSize = sizeof(cl_int) * arrayWidth * arrayHeight;
cl_int* inputA  = (cl_int*)clSVMAlloc(
    ocl.context,      //関連付けるcontext
    CL_MEM_READ_ONLY | CL_MEM_SVM_FINE_GRAIN_BUFFER,
    nSize,            //確保するサイズ
    0);               //アライメント(0 = 自動)

cl_int* inputB  = (cl_int*)clSVMAlloc(
    ocl.context, CL_MEM_READ_ONLY | CL_MEM_SVM_FINE_GRAIN_BUFFER, nSize, 0);
cl_int* outputC = (cl_int*)clSVMAlloc(
    ocl.context, CL_MEM_READ_ONLY | CL_MEM_SVM_FINE_GRAIN_BUFFER, nSize, 0);

転送は全く発生せず計算できている。



"Add"kernelの実行時間は2.067msとCoarse Grain Bufferの時と比べて増加してしまっている。CPUと常に同期をとる分、ややkernelの実行が遅くなるようだ。

ちゃんと動くコードはこちら。

---

SVM Fine Grain System

---

clSVMAllocしたポインタだけでなく、どんなCPUのポインタも共有できるという、最強のSVM。

だが、Skylakeではまだ使用できないのだった…。

---

続き>>

---

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)の使用

---

---

USE_HOST_PTRを使用する場合の手順はこんな感じ。

まず、CPU側のメモリ確保で2つの点で注意する。
・4096バイトでアラインされている
・バッファのサイズは64バイトの倍数

これは、_aligned_malloc()を使ってアライン済みアドレスを取得することで解決できる。

cl_uint optimizedSize
  = ceil_int(sizeof(cl_int) * arrayWidth * arrayHeight, 64);
cl_int* inputA  = (cl_int*)_aligned_malloc(optimizedSize, 4096);
cl_int* inputB  = (cl_int*)_aligned_malloc(optimizedSize, 4096);
cl_int* outputC = (cl_int*)_aligned_malloc(optimizedSize, 4096);

そして、このポインタを使って、OpenCLバッファを作成する。

int CreateBufferArguments(
    ocl_args_d_t *ocl,
    cl_int* inputA,  //CPU側の入力配列へのポインタ 
    cl_int* inputB,  //CPU側の入力配列へのポインタ 
    cl_int* outputC, //CPU側の出力配列へのポインタ
    cl_uint width, cl_uint height //配列のサイズ
) {
    cl_int err = CL_SUCCESS;
    cl_uint optimizedSize
      = ceil_int(sizeof(cl_int) * arrayWidth * arrayHeight, 64);
    ocl->srcA = clCreateBuffer(ocl->context,
	    //CL_MEM_USE_HOST_PTRを追加
        CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
        optimizedSize, //配列のサイズ
        inputA, //対応するCPUメモリへのポインタ
        &err);  //エラー情報を受け取る

    ocl->srcB = clCreateBuffer(ocl->context,
	    //CL_MEM_USE_HOST_PTRを追加
        CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
        optimizedSize, //配列のサイズ
        inputB, //対応するCPUメモリへのポインタ
        &err);  //エラー情報を受け取る

    ocl->dstMem = clCreateBuffer(ocl->context,
	    //CL_MEM_USE_HOST_PTRを追加
        CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
        optimizedSize, //配列のサイズ
        outputC, //対応するCPUメモリへのポインタ
        &err);   //エラー情報を受け取る

    return err;
}

このように作成した領域について、ホスト(CPU)側から操作する場合には、clEnqueueMapBuffer()を呼び出し、操作し終わったら、clEnqueueUnmapMemObject()を呼び出す必要がある。(これでCPU-GPU間の同期がとられる)

//対応するホスト側のポインタを取得する
cl_int *ptrMappedA = (cl_int *)clEnqueueMapBuffer(
    ocl.commandQueue, //投入キュー
    ocl.srcA,         //対象のOpenCLバッファ
    CL_FALSE,         //終了までブロックするか -> しない
    CL_MAP_WRITE,     //CPUが書き込むためにMapする
                      //(読み込みならCL_MAP_READ)
                      //(両方ならCL_MAP_READ | CL_MAP_WRITE)
    0,                //オフセット
    sizeof(cl_uint) * arrayWidth * arrayHeight, //マップするサイズ
    0,      //この関数が待機すべきeventの数
    NULL,   //この関数が待機すべき関数のリストへのポインタ
    NULL,   //この関数の返すevent
    &err);

//終了を待機
err = clFinish(ocl.commandQueue);

//ホスト(CPU)から操作
//(...略...)

err = clEnqueueUnmapMemObject(
    ocl.commandQueue, //投入キュー
    ocl.srcA,   //対象のOpenCLバッファ
    ptrMappedA, //取得したホスト側のポインタ
    0,          //この関数が待機すべきeventの数
    NULL,       //この関数が待機すべき関数のリストへのポインタ
    NULL);      //この関数の返すevent

//終了を待機
err = clFinish(ocl.commandQueue);

//デバイス(GPU) kernelを実行
//(...略...)

Map / Unmapが少し面倒であるが、このようにすることで転送が完全に省略される。clEnqueueUnmapMemObject()での転送時間が1usととても短くなり、事実上転送は行われていないことがわかる。




これに対し、アライメントが取れていないアドレスをclCreateBufferに渡すと、clEnqueueUnmapMemObject()で転送が起こってしまい、1.3ms前後かかってしまっている。

---

ZeroCopyを使用するには、今回のように、自分でCPU側のメモリを_aligned_malloc()で確保し、そのアドレスからUSE_HOST_PTR付きのclCreateBuffer()でOpenCLバッファを作る方法のほかに、ALLOC_HOST_PTR付きのclCreateBuffer()でOpenCLバッファを作る方法があり、この場合はOpenCL側でメモリを確保してくれる。

    ocl->srcA = clCreateBuffer(ocl->context,
	    //CL_MEM_ALLOC_HOST_PTRを追加
        CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR,
        optimizedSize, //配列のサイズ
        inputA, //OpenCL側で確保してくれるので不要
        &err);  //エラー情報を受け取る

こちらのほうが使いやすい場合もあるかもしれない。

---

今回のちゃんと動くコードはこちら。

---

続き>>

---

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)の使用

---

---

OpenCLのkernelがどのように実行されているのか確認したい。

まあ、vtuneがあれば悩まなくていいのだが、さすがにあれはお値段がお高すぎる。とりあえずいろいろ調べると、簡単にはIntel Graphics Performance Analyzersを使えばよいようだ。

(なんか昔も調べた気がするので、たぶん調べるの2回目…)

まず、Intel Graphics Monitorを起動する。「Analyze Application」から、起動する実行ファイルのパスを指定する。



正しくプロファイルを取得するには、トレースする時間を適切に指定する必要がある。[Managae Profiles...]から、[Tracing]タブのTracing Durationを調整する。



以下のような順序になるよう調整しないと、うまくトレースが取得できなかった。

1. 起動
2. OpenCLが実行され、完了
3. 起動してから、Tracing Durationの時間が経過し、トレースが終了
4. プログラムが終了

特にはまるのは、3と4が逆転してしまうと、トレースが取得できなかったことで、プログラムのほうにOpenCL実行後、

Sleep(5000);

を入れて調整した。なんかもっといいやり方がありそうだけど…。

さて、取得したトレースはIntel Platform Analyzerで見ることができて、こんな感じ。



キューにタスクが積まれ、転送→転送→Add→転送の順序で実行している様子とその時間が確認できる。

---

さて、この前のkernelは、以下のような単純な配列加算を行うものだった。

__kernel void Add(__global int* pA, __global int* pB, __global int* pC) {
    const int x     = get_global_id(0); //0次元目についてglobal idを取得
    const int y     = get_global_id(1); //1次元目についてglobal idを取得
    const int width = get_global_size(0);

    const int id = y * width + x;

    pC[id] = pA[id] + pB[id];
}

このkernelを以下のように実行していた。

//work sizeの指定
//ここでは1要素に対して1 work item
//またグループあたり1 work item (実は効率的でない)
//width * heightの2次元でwork itemを作成
size_t globalWorkSize[2] = { width, height };
size_t localWorkSize[2] = { 1, 1 };

//タスクをキューに積む
clEnqueueNDRangeKernel(
    ocl->commandQueue, //タスクを投入するキュー
    ocl->kernel,       //実行するカーネル
    2,                 //work sizeの次元
    NULL,              //NULLを指定すること
    globalWorkSize,    //全スレッド数
    localWorkSize,     //1グループのスレッド数
    0,                 //この関数が待機すべきeventの数
    NULL,              //この関数が待機すべき関数のリストへのポインタ
    NULL);             //この関数の返すevent

このwork size指定では、1つのwork itemを持つwork groupがwidth * heightだけ作られることになる。しかし、前回見たGPUの構造から、この指定は以下の理由で効率が悪いことがわかる。

・複数のwork itemをまとめて実行するSIMD-32等が効果を発揮できない
同じwork groupのwork itemしかまとめることはできない

・異なるwork groupだとどのsubsliceで実行されるわからない
メモリアクセスがばらばらで非効率

そこで、通常第0次元のwork itemを複数まとめて実行するよう指定する。このとき、一般的に32, 64など2の累乗がよいとされる。また、OpenCL 1.xでは、globalWorkSizeはlocalWorkSizeで割り切れる必要があることに注意して、localWorkSizeの倍数に切り上げる。

ここでは、work groupあたり32のwork itemを実行するよう指定する。

int ceil_int_div(int i, int div) {
    return (i + div - 1) / div;
}

//iをdivで切り上げ
int ceil_int(int i, int div) {
    return ceil_int_div(i, div) * div;
}

cl_uint ExecuteAddKernel(ocl_args_d_t *ocl, cl_uint width, cl_uint height) {
    cl_int err = CL_SUCCESS;

    size_t localWorkSize[2]  = { 32, 1 };
    size_t globalWorkSize[2] = {
        ceil_int(width,  localWorkSize[0]),
        ceil_int(height, localWorkSize[1])
    };

    (...略...)

切り上げを行うことで、work item数が実際の数を上回ることがあり得る。確保していないメモリ領域にアクセスしないよう、kernelコードを以下のように変更して、widthとheightは引数で与えるようにする。

__kernel void Add(__global int* pA, __global int* pB, __global int* pC, int width, int height) {
    const int x      = get_global_id(0);
    const int y      = get_global_id(1);

    if (x < width && y < height) {
        const int id = y * width + x;
        pC[id] = pA[id] + pB[id];
    }
}

---

kenrnel実行の所要時間を確認する。

work groupあたり1つのwork item = 33.11ms




work groupあたり32のwork item = 1.64 ms



というわけで劇的に高速化していることが確認できる。このように、work sizeの調整は重要であるが、これをOpenCL runtimeに自動的に調整させることも可能である。その際は、clEnqueueNDRangeKernel()のlocal_work_sizeにNULLを指定すればよい。

---

今回の変更はこちら。

---

続き>>

---

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)の使用

---

---

GPUのおおまかな構造はこんな感じ。i7 6700Kほか多くのCPUの内蔵GPUであるHD Graphics 530はGT2とよばれる規模のGPUで、24EUと言われる。この24EUのまとまりをsliceと呼び、これがリングバスを通してメモリコントローラにつながっている。



実際には1つのslice中は3つのsubsliceに分かれていて、1subsliceあたり8EUを搭載することで、24EUとなっている。



EUはいくつかの演算器の集合で、
・128bit SIMD演算器(ALU/FPU) x2
・Sendユニットx1 (メモリアクセスとかEUの外にアクセスする処理全般)
・Branchユニットx1
・256bitレジスタ x128 x 7スレッド分
を持つ。

2基の128bit SIMD演算器を合わせると、以下のような計算のいずれかを毎cycle発行可能である。
・16 x FP16
・8 x FP32
・2 x FP64 (片方のユニットしか実行できない)
・16 x INT16
・8 x INT32



EUでの計算は基本的にSIMD演算器を使うことになるので、命令もそれにそってSIMD型の命令となっている。コンパイラは、プログラムに応じて、8 / 16 / 32 のwork itemを1つのスレッドにまとめて実行する。(SIMD-8 / SIMD-16 / SIMD-32)

実際には、SIMD-32の演算器は存在しないので、例えばSIMD-8 x4などに分解されて実行される。

EUは最大7スレッドを(切り替えながら)同時に処理することができる。なるべく多くのスレッドをEUに突っ込むことで、演算器の使用率を高く保ち、演算効率を上げることができる。

8つのEUが集まってsubsliceを構成する。subsliceには8つのEUからのデータアクセスを処理するData Portが存在するが、読み書き双方向にそれぞれ64byte/cycleであり、EU8つの演算力と比べれば十分なデータ供給力があるとはいえず、なるべくEU内のレジスタを使用した演算を行うことが効率よく計算するためには必要であることがわかる。

ただし、Data Port以外にOpenCLのimageNdを使用した場合のみ使えるアクセス方法が存在し、これがいわゆるSamplerキャッシュ(テクスチャキャッシュ)である。これは読み込み専用であるが、64byte/cycleの帯域と4KBのL1キャッシュ + 24KBのL2キャッシュを持つ。



Data PortはEUからのメモリアクセスをまとめて効率化することが可能である。これをコアレスアクセスという。例えば、SIMD-16などであれば、そのスレッドは16work itemについて一度に処理するわけだが、その16work itemのメモリアクセスが1つのキャッシュライン(64byte以内)であれば、これをまとめてアクセスするすることで効率よくアクセスするというものである。

GPU演算でメモリアクセスの効率を上げることは非常に重要なため、連続するwork itemのメモリアクセスが必要とするキャッシュラインをなるべく少なくし、取り込んだデータに無駄がないようにする必要がある。基本的にメモリアクセスはキャッシュライン(64byteの塊)を単位として行われるので、これを意識すると効率が良くなる。

例1のように、連続するwork itemのアクセスが、ひとつのキャッシュラインに収まっていれば、1回のメモリアクセスですみ、効率が良い。

例1


例2のように、飛び飛びのアクセスであると、複数のキュッシュラインからデータを取ってくることになり、アクセス効率が低下する。また、読み込んでも使用していないデータが多く、無駄にメモリ帯域を使っていることになる。

例2


連続であればなんでも良いかというと、そういうわけではなく、先頭のアドレスがキャッシュラインとずれていると、キャッシュラインをまたぐことになよって、2回のメモリアクセスが必要になり、効率が悪化する。

例3



同じキャッシュラインへのアクセスであれば問題ないので、例4のように、逆順でも問題はない。

例4


このように、先頭のwork itemのアクセス先のアライメントをとり、work itemが連続したアクセスとすることで、なるべくキャシュラインをまたがないようにして、メモリアクセス効率を高めることが重要である。


最後に3つのsubsliceとL3キャッシュ(512KB)が1つのsliceを形成する。Data PortからのアクセスやSamplerからのアクセスはすべてこのL3キャッシュを通り、64byte/cycleでリングバスに接続される。



Shared Local Memoryはsubsliceあたり64KB(計192KB)まで利用可能な特殊なメモリ領域で、work group内で共有して利用することができる。また、通常のメモリアクセスよりも比較的多様なアクセスパターンを高速に処理できるという特徴を持つ。そのため、work group内のwork itemが、他のwork itemの計算結果を使用したい場合、このSLMを経由して渡してやると効率的にやり取りできる。また、work group内のwork itemが同じ値を何度も重複して利用する場合に、明示的にSLMにキャシュしてやると、高速化につながることがある。

ただ、このSLMがsubsliceの外にあるため、必ずDataPortを経由せざるを得ず、EUからのアクセス帯域に制限が生じていて、そこまで高速というわけではないので、注意が必要である。SLMはwork group内で共有すればよく、同じWork groupのwork itemは必ず同じsubsliceで実行されるのだとすれば、SLMはsubslice内にあればよいはずで、なぜSLMをsubsliceの外に実装しているのかが、Intelの実装で不思議な点である。(例えば、NVIDIA PascalのsharedメモリはSMの中にある)

ちなみに、sliceが1つなのはHD Graphics 530が24EUのGT2だからで、Skylakeにはさらに48EUのGT3、72EUのGT4があり、この場合はsliceの数が2、3と増えることになる。

---

というわけでざっと資料を見るとこんな感じ。

やはり、性能を出すにはwork groupの構成とか、work itemのメモリアクセスとか、そのあたりを注意する必要があるように思う。

---

続き>>

---

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)の使用

---

---

まずはVisual Studio 2015 Communityをインストールし、さらにこのへんからダウンロードしたIntel OpenCL SDKをインストールする。

インストールすると、VC++のプロジェクトテンプレートの中にOpenCLのテンプレートが登場する。



Code Builder Project... を選ぶと、設定が出る。

ここでは、"Use explicit Local Work Group Size"にチェックを入れた。



テンプレートは配列加算のサンプルだが、その入出力の配列のメモリに関する設定。デフォルトは"Use Host Memory"になっているが、これはZero Copyを使って、CPU-GPU間転送を省略するものであり、OpenCLの基本とはいえない。そこで、"Allocate & Copy Memory"に変更した。



Finishとすると、OpenCLのサンプルコードがもう動かす事ができる形で自動生成される。



OpenCLといえば、まず動かすまでに長すぎる呪文詠唱で嫌になるという印象が強かったが、最近では便利になったものだ。

---

とはいえ、どういう呪文を詠唱すればよいのか確認していく。見通しをよくするため、自動生成されたサンプルコードからエラー処理の一切を省き、ある程度簡略化してある。

まず"platform"の取得と選択 FindOpenCLPlatform()

//今回はIntelのGPUを探したいので、
//deviceType = CL_DEVICE_TYPE_GPU
//preferredPlatform = "Intel"
cl_platform_id FindOpenCLPlatform(
  const char* preferredPlatform, cl_device_type deviceType) {
    //platformの数を取得
    cl_uint numPlatforms = 0;
    clGetPlatformIDs(0, NULL, &numPlatforms);

    //実際にplatformのリストを取得
    std::vector platforms(numPlatforms);
    clGetPlatformIDs(numPlatforms, &platforms[0], NULL);

    for (cl_uint i = 0; i < numPlatforms; i++) {
        //platformの名前を取得する
        char buf[4096];
        clGetPlatformInfo(platform, CL_PLATFORM_NAME,
            sizeof(buf), buf, NULL);
        
        //ここではIntelのplatformが見つかったら
        if (strstr(buf, preferredPlatform) != 0) {
            //そのplatformがdeviceType(ここではGPU)を持つか、確認する
            cl_int err = clGetDeviceIDs(
                platforms[i], deviceType, 0, NULL, &numDevices);
            if (err == CL_SUCCESS && numDevices > 0) {
                return platforms[i];
            }
        }
    }
    return NULL;
}

つぎに取得したplatformを使って、opencl "context"を作成し、その"device"を取得する。さらに、そのdeviceに命令を投入するためのキューを作成する。

    //先程の関数でplatformを作成
    cl_platform_id platformId
        = FindOpenCLPlatform("Intel", deviceType);
    
    //コンテキスト情報
    cl_context_properties contextProperties[] = 
        { CL_CONTEXT_PLATFORM, (cl_context_properties)platformId, 0 };
        
    //コンテキストを作成
    ocl->context = clCreateContextFromType(
        contextProperties, deviceType, NULL, NULL, &err);

    //コンテキストのdeviceを取得
    err = clGetContextInfo(ocl->context, CL_CONTEXT_DEVICES,
        sizeof(cl_device_id), &ocl->device, NULL);

    //プロファイリングを有効にしてキューを作成
    cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE;
    ocl->commandQueue = clCreateCommandQueue(
        ocl->context, ocl->device, properties, &err);

次に、kernel(GPUで実行する関数)をビルドして実行可能にする必要がある。通常、kernelはソースコードを文字列として、clCreateProgramWithSource()関数に渡してビルドする。

Intelのサンプルコードのように、ソースファイルからそのまま文字列を読み込んでもいいが、面倒なので、リソースとして実行ファイルに埋め込むことを考える。

OpenCLTest.rcを新規作成し、ここに

CLDATA    KERNEL_DATA DISCARDABLE "Template.cl"

と書いてプロジェクトに登録し、このリソースを対象としてビルドを行い、関数名を指定してkernelを取得する。リソースを対象とする場合、#includeなどは使えないという制約が出るが、そこは…気にしないことにする。

//リソースからOpenCLのカーネルファイルを取得
HMODULE hmdl = GetModuleHandle(NULL);
HRSRC hrsrc = nullptr;
HGLOBAL hdata = nullptr;
const char *source = nullptr;
size_t src_size = 0;
if ( nullptr == (hrsrc = FindResource(hmdl, "CLDATA", "KERNEL_DATA"))
  || nullptr == (hdata = LoadResource(hmdl, hrsrc))
  || nullptr == (source = (const char *)LockResource(hdata))
  || 0       == (src_size = SizeofResource(hmdl, hrsrc))) {
	return CL_INVALID_VALUE;
}

//ソースから"program"を作成
ocl->program = clCreateProgramWithSource(ocl->context,
    1, (const char**)&source, &src_size, &err);

//ビルドを実行
clBuildProgram(ocl->program, 1, &ocl->device, "", NULL, NULL);

//ビルドした関数のうち、"Add"関数をkernelとして取得
ocl->kernel = clCreateKernel(ocl->program, "Add", &err);

さて、配列加算 C = A + B を行うが、そのサイズを width * height とする。この時、それぞれの配列について、CPU側の配列とは別に、GPU側にバッファを作成してやる必要がある。

ここで、自動生成されたsampleプログラムは、clEnqueueMapBufferを使った効率的な方法を用いていたが、ここではより基本的な、明示的にバッファへの転送を行うclEnqueueWriteBuffer / clEnqueueReadBuffer を用いた方法を見ていく。

int CreateBufferArguments(
    ocl_args_d_t *ocl,
    cl_int* inputA,  //CPU側の入力配列へのポインタ 
    cl_int* inputB,  //CPU側の入力配列へのポインタ 
    cl_int* outputC, //CPU側の出力配列へのポインタ
    cl_uint width, cl_uint height //配列のサイズ
) {
    cl_int err = CL_SUCCESS;
    ocl->srcA = clCreateBuffer(ocl->context,
        CL_MEM_READ_ONLY, //GPUからはこのメモリに対し、読み込みのみ行う
        sizeof(cl_uint) * width * height, //配列のサイズ
        nullptr,
        &err);  //エラー情報を受け取る

    ocl->srcB = clCreateBuffer(ocl->context,
        CL_MEM_READ_ONLY, //GPUからはこのメモリに対し、読み込みのみ行う
        sizeof(cl_uint) * width * height, //配列のサイズ
        nullptr,
        &err);  //エラー情報を受け取る

    ocl->dstMem = clCreateBuffer(ocl->context,
        CL_MEM_ALLOC_HOST_PTR, ////GPUからはこのメモリに対し、書き込みのみ行う
        sizeof(cl_uint) * width * height, //配列のサイズ
        nullptr,   //CPUから転送するものはない
        &err);

    //inputAから作成したバッファに転送
    err = clEnqueueWriteBuffer(ocl->commandQueue, //
        ocl->srcA, //転送先
        CL_FALSE,  //転送が終了するまで待機するか -> しない
        0,         //オフセット
        sizeof(cl_uint) * arrayWidth * arrayHeight, //転送サイズ
        inputA,   //転送元
        0,         //この関数が待機すべきeventの数
        nullptr,   //この関数が待機すべき関数のリストへのポインタ
        nullptr);  //この関数の返すevent

    //inputBから作成したバッファに転送
    err = clEnqueueWriteBuffer(ocl->commandQueue, //
        ocl->srcB, //転送先
        CL_FALSE,  //転送が終了するまで待機するか -> しない
        0,         //オフセット
        sizeof(cl_uint) * arrayWidth * arrayHeight, //転送サイズ
        inputB,    //転送元
        0,         //この関数が待機すべきeventの数
        nullptr,   //この関数が待機すべき関数のリストへのポインタ
        nullptr);  //この関数の返すevent

    return err;
}

ここからいよいよ関数を実行…の前に、関数の引数を一つずつ設定する必要がある。引数がたくさんあると、まあ、簡単に悲惨なコードが出来上がるわけで、正直、このあたりがOpenCLの、CUDAと比べてなかなかに面倒なところだ。

引数を積み終わった後、clEnqueueNDRangeKernelにより、指定のカーネルをGPUで実行するよう指示する。

//関数引数を設定
clSetKernelArg(ocl->kernel, 0, sizeof(cl_mem), (void *)&ocl->srcA);
clSetKernelArg(ocl->kernel, 1, sizeof(cl_mem), (void *)&ocl->srcB);
clSetKernelArg(ocl->kernel, 2, sizeof(cl_mem), (void *)&ocl->dstMem);

//work sizeの指定
//ここでは1要素に対して1 work item
//またグループあたり1 work item (実は効率的でない)
//width * heightの2次元でwork itemを作成
size_t globalWorkSize[2] = { width, height };
size_t localWorkSize[2] = { 1, 1 };

//タスクをキューに積む
clEnqueueNDRangeKernel(
    ocl->commandQueue, //タスクを投入するキュー
    ocl->kernel,       //実行するカーネル
    2,                 //work sizeの次元
    NULL,              //NULLを指定すること
    globalWorkSize,    //全スレッド数
    localWorkSize,     //1グループのスレッド数
    0,                 //この関数が待機すべきeventの数
    NULL,              //この関数が待機すべき関数のリストへのポインタ
    NULL);             //この関数の返すevent

work size(work itemの数)で、並列処理する数を指定し、3次元まで設定できる。

work itemは階層構造になっており、通常は複数のwork itemがwork groupを形成し、work groupが集まってwork item全体を構成する。OpenCL 1.2までは各次元のglobalWorkSizeは各次元のlocalWorkSizeで割りきれる必要があることに注意する。

この時、work group内の個々のwork itemのIDがlocal idで、groupによらない一意に決まるIDをglobal id という。下記の図はwork itemが1次元の場合である。



実際には、work itemは3次元まで設定可能になっていて、サンプルコードは2次元の例になっている。

同じwork groupのwork itemは、GPUで実行される際に必ず同じ演算ブロック(Intelでいうsubslice、NVIDIAでいうSM)で実行され、__localメモリ(IntelでいうShared Local Memory、NVIDIAでいうsharedメモリ)を共有する。また、連続する8, 16, または32のwork itemは同時に実行されるという特徴があり、こうした特徴はメモリアクセス効率を考える際に重要となる。

GPU上で実行するkernelコードはこのようになっていて、単なる配列の加算である。

__kernel void Add(__global int* pA, __global int* pB, __global int* pC) {
    const int x     = get_global_id(0); //0次元目についてglobal idを取得
    const int y     = get_global_id(1); //1次元目についてglobal idを取得
    const int width = get_global_size(0);

    const int id = y * width + x;

    pC[id] = pA[id] + pB[id];
}

width * height 分の各work itemがkernelを実行することになるから、2次元のwork itemの指定により、CPUでいうと以下のような2重ループを並列化したことに等しい。

for (int y = 0; y < height; y++) { //<- 1次元目の並列化
    for (int x = 0; x < width; x++) { //<- 0次元目の並列化
        const int id = y * width + x;
        pC[id] = pA[id] + pB[id];
    }
}

さて、clEnqueueNDRangeKernelにより、指定のカーネルをGPUで実行するよう指示したが、これはタスクをキューに積んだだけなので、GPUが実際にこのkernelを実行し、完了するには時間がかかる。キューのすべてのタスクが終了するまで待機するには、キューを指定してclFinish()を呼ぶ。

clFinish(ocl->commandQueue);

その後、GPUの計算結果を回収する。出力情報のあるGPUバッファをclEnqueueMapBuffer()によりmappingし、mappingされたCPU側の配列へのポインタが得られる。内部でGPUメモリ→CPUメモリへの転送が行われる。

    //inputBから作成したバッフからoutputCに転送
    err = clEnqueueReadBuffer(ocl->commandQueue, //
        ocl->dstMem, //転送元
        CL_FALSE,    //転送が終了するまで待機するか -> しない
        0,         //オフセット
        sizeof(cl_uint) * arrayWidth * arrayHeight, //転送サイズ
        outputC,   //転送先
        0,         //この関数が待機すべきeventの数
        nullptr,   //この関数が待機すべき関数のリストへのポインタ
        nullptr);  //この関数の返すevent
    
    //clEnqueueMapBufferの終了を待機
    clFinish(ocl->commandQueue);

     //outputCに計算結果が入っていて、CPUから操作できる

---

というわけで、ここまででやっと配列の加算を行うことができた。

まず、deviceを開き、OpenCL環境を初期化するところまでをまとめると、こんな感じ。



そして、OpenCLバッファの確保から、出力の取り出しまでをまとめると以下のようになる。



実際にはエラー処理などもあるので、結構長いコードになってしまう。CUDAと比べると、とにかく呪文(手続き)が長い、というのが汎用性・標準化を重視したOpenCLのつらいところだと思う。

---

今回のコードはこのへん。

---

続き>>

---

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)の使用

---