TadaoYamaokaの日記

山岡忠夫Homeで公開しているプログラムの開発ネタを中心に書いていきます。

OpenCLのサンプルコード

ほぼ個人用メモ。

「改訂新版 OpenCL入門 1.2対応 マルチコアCPU・GPUのための並列プログラミング」を買って、OpenCLをお勉強中。


OpenCL 1.2とバージョンが古いが、KataGoもOpenCL 1.2で実装されているので問題ない。
そもそもNVIDIAOpenCLに対してやる気がなく、1.2しかサポートしていないようだ。
OpenCL 2.x support plans? - CUDA Programming and Performance - NVIDIA Developer Forums


以下、サンプルコードのWindowsでのビルド手順。

OpenCLSDK

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';
}
)%%"

まとめ

サンプルコードをWindowsでビルドして実行した。
OpenCLは、OpenCL C言語ソースコードを文字列としてランタイムに渡して動的にコンパイルされる仕組みになっていることを理解した。