openCLのコードを少し調べてみる（カーネルの実行）
openCLの最初の記事で動かしたコードについて見て行く。
動かしたコードのソース：OpenCLで行列の掛け算をしてみた : 試行錯誤な日々
そもそも何をするコードだったのか、、、
並列かしているのは、行列のかけ算、と書かれていたが行列を一次元配列で表現しているらしい。
それを踏まえた上でカーネルコードをまずは見てみる。
~ 113 /* Set source code of kernel*/ ~ 114 char *source_str = ~ 115 "__kernel void matrix_dot_matrix( \ ~ 116 __global const float* A, \ // 行列 ~ 117 __global const float* B, \ // 行列 ~ 118 __global float* Result, \ // かけ算結果いれる配列 ~ 119 const int wA, \ // Aのwidth (列数 ~ 120 const int wB) { \ // Bのwidth (列数 ~ 121 const int x = get_global_id(0); \ ~ 122 const int y = get_global_id(1); \ ~ 123 float value = 0; \ ~ 124 for (int i = 0; i < wA; ++i) { \ // 行列を一次元で書いているので少しわかりづらいが // y*wAはy行目を見ている、ということ ~ 125 int index_a = y * wA + i; \ // こっちはx列目を見ている、ということ。 ~ 126 int index_b = i * wB + x; \ // 下は各行列の要素取得して掛けているだけ ~ 127 float elementA = A[index_a]; \ ~ 128 float elementB = B[index_b]; \ ~ 129 value = value + elementA * elementB; \ ~ 130 } \ ~ 131 Result[wB * y + x] = value; \ ~ 132 }";
上のコードでopenCL specificな部分は大体以下のものかと思う。
~ 115 "__kernel void matrix_dot_matrix( \ __kernelはカーネルコードにつける。 ~ 116 __global const float* A, \ // 行列 ~ 117 __global const float* B, \ // 行列 ~ 118 __global float* Result, \ // かけ算結果いれる配列 __globalはデバイスのグローバルメモリにつける。このメモリは全てのwork groupから参照可能なデバイスメモリ。 ~ 121 const int x = get_global_id(0); \ ~ 122 const int y = get_global_id(1); \ get_global_idは、work itemのidを取得する函数。引数は次元。 CUDAでいうblockIdxとかthreadIdxを取ってきているイメージかな？
カーネル実行までの道程を読む
例のごとく、マニュアルは手元に、、、 前回（openCLのコードを少し調べてみる（デバイスの情報取得） - tkokamoの日記）の結果からopenCL 1.2であることが分かったので、1.2のmanualをみる。 OpenCL 1.2 Reference Pages
contextの作成
~ 135 /* Create OpenCL context */ ~ 136 context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
" Contexts are used by the OpenCL runtime for managing objects such as command-queues, memory, program and kernel objects and for executing kernels on one or more devices specified in the context."
Contextはよくあるものなので、特に説明は不要。
openClのruntimeがprogramなどを管理するために使う。
command queueの作成
~ 142 command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
コマンドキューは、デバイスに実行させる処理を入れるキュー。
今回のコードでは行列のかけ算をさせるカーネル一つを入れているが、
複数入れるとそれらは逐次dispatchされるんだろうか、、、
カーネル実行の準備
それぞれをいまいちちゃんと理解しないんだけど、実行前にいくつかやることがあるみたい。
~ 145 /* Create Kernel Program from the source */ ~ 146 program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t * )&source_size, &ret);
マニュアルを見ると、
- プログラムオブジェクトを作成してsource_strで指定されたプログラムをプログラムオブジェクトにロードする。
- プログラムオブジェクトはcontextを作成したデバイスに結びつけられる。
らしい。なんか必要そうなのはわかるけど、openCLのモデルを理解していないからふわっとしている。
~ 148 /* Build Kernel Program */ ~ 149 ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
プログラムをビルドする。まぁ複数デバイスサポートしているのこれは必要でしょう。
~ 159 /* Create OpenCL Kernel */ ~ 160 kernel = clCreateKernel(program, "matrix_dot_matrix", &ret);
カーネルオブジェクトの作成。ビルド後のプログラムのオブジェクト、と思えばよいのだろうか。
~ 167 /* Create Memory Buffer */ ~ 168 matrixAMemObj = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, matrixAMem Size, matrixA, &ret); ~ 169 matrixBMemObj = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, matrixBMem Size, matrixB, &ret); ~ 170 matrixRMemObj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, matrixRMemSize, NULL, &ret);
カーネルの引数となるデバイスメモリの獲得を行っている。
ここで指定しているのは、
- context
- 獲得方法を決めるフラグ
- CL_MEM_READ_ONLYはreadonlyのメモリ
- CL_MEM_COPY_HOST_PTRは獲得したデバイスメモリに第3引数のホストメモリの値をコピーする
- ホストメモリのポインタ
- メモリのサイズ
- エラーコードを格納する変数のアドレス
戻り値はbuffer object。フラグについては結構種類があるようだが、今後自分でプログラムを書くときに必要であれば調べてみる。
~ 176 /* Set OpenCL Kernel Parameters */ ~ 177 ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&matrixAMemObj); ~ 178 ret |= clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&matrixBMemObj); ~ 179 ret |= clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&matrixRMemObj); ~ 180 ret |= clSetKernelArg(kernel, 3, sizeof(int), (void *)&wA); ~ 181 ret |= clSetKernelArg(kernel, 4, sizeof(int), (void *)&wB);
カーネルに渡す引数の設定
カーネルの実行
~ 256 /* Execute OpenCL Kernel */ ~ 257 ret = clEnqueueNDRangeKernel(command_queue, kernel, workDim, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL);
command_queueにkernelをenqueueする。
この時指定するのは、（第3引数から）
- work itemが使う次元数。ここでは2次元
- global work offset。ここではNULLだが、指定されるとglobalIDの開始offsetが指定したIDからになる。
- global_work_size。各次元におけるwork sizeを指定する。全体的なwork sizeは各次元のサイズの積になる。この値については後述。
- local_work_size。work groupのサイズを指定する。この値については後述。openCLではworkは（こういう言い方でよいのか？）複数のwork groupから構成され、work groupは複数のwork itemから構成される。work group内ではデバイスのlocal memoryが使えるので適切にworkを分割すると計算性能の向上が図れるんだろう、、、多分
- num_events_in_wait_list。カーネルが実行される前に待ち合わせるイベントの数。同期に使われる。
- event_wait_list。待ち合わせるeventの配列。
- event。ここではNULLだが、ポインタが指定されるとここにenqueueしたカーネルのevent objectが入れられる。
local_work_sizeを決める為に、
まずはデバイスの最大work groupサイズと最大次元数を取得している。
~ 189 size_t maxWorkGroupSize;
~ 190 {
~ 191 ret = clGetDeviceInfo(device_id, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(maxWorkGroupSize), &maxWorkGroupSize, NULL);
~ 198 size_t* maxLocalSizes;
~ 199 {
~ 200 cl_uint maxWorkItemDims;
~ 201 ret = clGetDeviceInfo( device_id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(cl_uin t), &maxWorkItemDims, NULL);
次に得られた次元数分size_tの配列を獲得し、各次元の最大work itemサイズを取得している。
~ 206 maxLocalSizes = (size_t*)malloc(maxWorkItemDims * sizeof(size_t)); ~ 208 ret = clGetDeviceInfo(device_id, CL_DEVICE_MAX_WORK_ITEM_SIZES, maxWorkItemDims* sizeof(size_t), maxLocalSizes, NULL);
そして最終的にlocal work sizeを以下の処理で決めている。
~ 220 size_t localWR = wR, localHR = hR;
~ 221 while (maxLocalSizes[0] < localWR ||
~ 222 maxLocalSizes[1] < localHR ||
~ 223 localWR * localHR > maxWorkGroupSize) {
~ 224 if (maxLocalSizes[0] < localWR) {
~ 225 if (localWR % 2 == 0) {
~ 226 localWR /= 2;
~ 227 } else {
~ 228 localWR = getMaxCommonFactorOf2Pow(localWR);
~ 229 }
~ 230 } else if (maxLocalSizes[1] < localHR) {
~ 231 if (localWR % 2 == 0) {
~ 232 localHR /= 2;
~ 233 } else {
~ 234 localHR = getMaxCommonFactorOf2Pow(localHR);
~ 235 }
~ 236 } else if (localHR != 1) {
~ 237 if (localHR % 2 == 0) {
~ 238 localHR /= 2;
~ 239 } else {
~ 240 localHR = getMaxCommonFactorOf2Pow(localHR);
~ 241 }
~ 242 } else {
~ 243 if (localWR % 2 == 0) {
~ 244 localWR /= 2;
~ 245 } else {
~ 246 localWR = getMaxCommonFactorOf2Pow(localWR);
~ 247 }
~ 248 }
~ 249 }
結果のreadback
~ 264 ret = clEnqueueReadBuffer(command_queue, matrixRMemObj, CL_TRUE, 0, matrixRMemSize, matrixR , 0, NULL, NULL);
終了処理
~ 271 /* Finalization */ ~ 272 ret = clFlush(command_queue); ~ 273 ret = clFinish(command_queue); ~ 274 ret = clReleaseKernel(kernel); ~ 275 ret = clReleaseProgram(program); ~ 276 ret = clReleaseMemObject(matrixAMemObj); ~ 277 ret = clReleaseMemObject(matrixBMemObj); ~ 278 ret = clReleaseMemObject(matrixRMemObj); ~ 279 ret = clReleaseCommandQueue(command_queue); ~ 280 ret = clReleaseContext(context);
こんなところか、、、次は自前コードを動かそう。
