ほぼ個人用メモ。
「改訂新版 OpenCL入門 1.2対応 マルチコアCPU・GPUのための並列プログラミング」を買って、OpenCLをお勉強中。
OpenCL 1.2とバージョンが古いが、KataGoもOpenCL 1.2で実装されているので問題ない。
そもそもNVIDIAはOpenCLに対してやる気がなく、1.2しかサポートしていないようだ。
OpenCL 2.x support plans? - CUDA Programming and Performance - NVIDIA Developer Forums
以下、サンプルコードのWindowsでのビルド手順。
OpenCLのSDK
CUDA Toolkitに含まれるため、追加でインストールはいらない。
※書籍では、GPU Computing SDKが必要とあるが、現在はCUDAに含まれる。
※ただし、OpenCLのサンプルプログラムは提供されなくなっている。
Where is gpu computing sdk ? - CUDA Setup and Installation - NVIDIA Developer Forums
プロジェクトの設定
OpenCLのインクルードファイルとライブラリは、CUDAをインストールすると以下の場所にある。
インクルード：C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.1\include
ライブラリ：C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.1\lib\x64
それぞれ、プロジェクトのインクルードとライブラリのディレクトリ設定に追加する。
リンクライブラリに、「OpenCL.lib」を追加する。
サンプルコード
書籍のサンプルコードでは、.clを動的にファイルから読むようになっていたが、ソースに埋め込むようにした。
KataGoの書き方を参考にして、
const char source_str[] = #include "hello.cl" ;
のようにして、文字列にした。
.clの先頭と末尾には、
R"%%( // OpenCLのコード )%%"
が必要になる。
修正したサンプルコード
hello.cpp
#include <iostream> #ifdef __APPLE__ #include <OpenCL/opencl.h> #else #include <CL/cl.h> #endif constexpr size_t MEM_SIZE = 128; int main() { cl_device_id device_id = nullptr; cl_context context = nullptr; cl_command_queue command_queue = nullptr; cl_mem memobj = nullptr; cl_program program = nullptr; cl_kernel kernel = nullptr; cl_platform_id platform_id = nullptr; cl_uint ret_num_devices; cl_uint ret_num_platforms; cl_int ret; char string[MEM_SIZE]; const char source_str[] = #include "hello.cl" ; size_t source_size = sizeof(source_str) - 1; const char* source_list[] = { source_str }; /* プラットフォーム・デバイスの情報の取得 */ ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms); ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices); /* OpenCLコンテキストの作成 */ context = clCreateContext(nullptr, 1, &device_id, nullptr, nullptr, &ret); /* コマンドキューの作成 */ command_queue = clCreateCommandQueue(context, device_id, 0, &ret); /* メモリバッファの作成 */ memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(char), nullptr, &ret); /* 読み込んだソースからカーネルプログラムを作成 */ program = clCreateProgramWithSource(context, 1, source_list, &source_size, &ret); /* カーネルプログラムをビルド */ ret = clBuildProgram(program, 1, &device_id, nullptr, nullptr, nullptr); /* OpenCLカーネルの作成 */ kernel = clCreateKernel(program, "hello", &ret); /* OpenCLカーネル引数の設定 */ ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&memobj); /* OpenCLカーネルを実行 */ ret = clEnqueueTask(command_queue, kernel, 0, nullptr, nullptr); /* メモリバッファから結果を取得 */ ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(char), string, 0, nullptr, nullptr); /* 結果の表示 */ std::cout << string << std::endl; /* 終了処理 */ ret = clFlush(command_queue); ret = clFinish(command_queue); ret = clReleaseKernel(kernel); ret = clReleaseProgram(program); ret = clReleaseMemObject(memobj); ret = clReleaseCommandQueue(command_queue); ret = clReleaseContext(context); return 0; }
hello.cl
R"%%(
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
__kernel void hello(__global char* string)
{
string[0] = 'H';
string[1] = 'e';
string[2] = 'l';
string[3] = 'l';
string[4] = 'o';
string[5] = ',';
string[6] = ' ';
string[7] = 'W';
string[8] = 'o';
string[9] = 'r';
string[10] = 'l';
string[11] = 'd';
string[12] = '!';
string[13] = '\0';
}
)%%"
