OpenCLのはじめかたなど

はじめに

Windows環境でgitとMicrosoft Visual Studio Community 2022をインストールしているものとします。
次の四点について記述します。

• OpenCLのインポートライブラリとヘッダファイルの入手方法
• 乱数を生成してみる
• 構造体とメンバ変数のアライメント
• Intel製iGPU用OpenCLを使用する際の注意

OpenCLのインポートライブラリとヘッダファイルの入手方法

Getting started with OpenCL on Microsoft Windows にあるように、

• KhronosGroupのOpenCL-SDKリポジトリからビルドする方法
• vcpkg経由でビルドする方法

その他、ヘッダファイルをOpenCL-HeadersとOpenCL-CLHPPからダウンロードし、インポートライブラリを次のように得るという方法もあります。

• GPUOpen-LibrariesAndSDKs/OCL-SDKからダウンロードする
• dumpbinとlibを使用してOpenCL.dllからインポートライブラリを生成する

KhronosGroupのOpenCL-SDKリポジトリからビルドする場合

まずはミ田キーを押して2022とでも入力し、x64 Native Tools Command Prompt for VS 2022を実行してコマンドプロンプトを開き、任意のディレクトリを作ってそこに行きます。
そして次のようにタイプします。
git clone --recursive https://github.com/KhronosGroup/OpenCL-SDK.git
cmake -G "Visual Studio 17 2022" -A x64 -T v143 -D CMAKE_INSTALL_PREFIX=./OpenCL-SDK/install -B ./OpenCL-SDK/build -S ./OpenCL-SDK
cmake --build OpenCL-SDK/build --config Release --target install -- /m /v:minimal

OpenCL-SDK\install 以下にインポートライブラリとヘッダファイルその他が生成されます。

vcpkg経由でビルドする場合

まずはコマンドプロンプトを開き任意のディレクトリを作ってそこに行き、次のようにしてvcpkgをインストールします。
git clone https://github.com/microsoft/vcpkg.git
cd vcpkg
.\bootstrap-vcpkg.bat

そしてopenclパッケージをインストールします。
.\vcpkg.exe --triplet=x64-windows install opencl

vcpkg\installed\x64-windows 以下にインポートライブラリとヘッダファイルその他が生成されます。

dumpbinとlibを使用してOpenCL.dllからインポートライブラリを生成する

まずはミ田キーを押して2022とでも入力し、x64 Native Tools Command Prompt for VS 2022を実行してコマンドプロンプトを開き、任意のディレクトリを作ってそこに行きます。
そして次のようにタイプします。
dumpbin /exports %SystemRoot%\System32\OpenCL.dll > temp.txt
メモ帳などを用い、OpenCL.defというテキストファイルを作成し、最初の部分を次のようにします。

LIBRARY OpenCL
EXPORTS

temp.txtに書き出した中にある関数名全てをOpenCL.defの1行にひとつずつ追記します。123個ほどあるかと思います。
関数名の追記とファイルの保存を終えたら次のようにタイプします。
lib /DEF:OpenCL.def /MACHINE:X64
32bit用のインポートライブラリを生成する場合はlib /DEF:OpenCL.def /MACHINE:X86でよいでしょう。

乱数を生成してみる

ふとした思い付きで1スレッドごとに乱数を得たいことがあるかもしれません。とりあえずホストで生成した乱数をもとにデバイスでxorshiftして、それをもとに画像データを生成し吐き出すものを作ってみます。

kernel.cl

// kernel.cl
uint GetRandomUint(uint *random_state);
float GetRandomUnitFloat(uint *random_state);

kernel void RandomTest(global float *dst, global uint *rnd, uint plane_size)
{
    size_t gid = get_global_id(0);
    uint random_state = rnd[gid];
	float f;

	f = GetRandomUnitFloat(&random_state);
	dst[gid] = f;

	f = GetRandomUnitFloat(&random_state);
	dst[plane_size + gid] = f;

	f = GetRandomUnitFloat(&random_state);
	dst[plane_size * 2 + gid] = f;

    rnd[gid] = random_state;
}

uint GetRandomUint(uint *random_state)
{
	uint x = *random_state;

	x ^= x << 13;
	x ^= x >> 17;
	x ^= x << 5;
	
	return *random_state = x;
}

float GetRandomUnitFloat(uint *random_state)
{
	uint x = GetRandomUint(random_state);

	return (float)(x & 0x7fff) / (0x7fff + 1.0f);
}

標準出力へ吐き出します。フォーマットはNetpbm(Pbmplus)です。
その際Windows PowerShellを使うとエンコーディング指定が面倒だったり処理が重くなるので、コマンド プロンプトを使いましょう。
画像を表示するにはSusie32にSusie32 PbmPlus Plug-inを加えて利用するとよいかもしれません。

main.cpp

// main.cpp
#include <algorithm>
#include <format>
#include <fstream>
#include <iostream>
#include <random>

#define CL_HPP_TARGET_OPENCL_VERSION 120
#define CL_HPP_MINIMUM_OPENCL_VERSION 120
#define CL_HPP_ENABLE_EXCEPTIONS
#define CL_HPP_CL_1_2_DEFAULT_BUILD
#include <CL/opencl.hpp>

const cl_uint kWidth = 1920;
const cl_uint kHeight = 1080;
const cl_uint kThreads = kWidth * kHeight;
const cl_uint kPlanePixels = kThreads;
const cl_uint kPlaneBytes = sizeof (float) * kPlanePixels;

void FillRandomUints(cl_uint *dst, size_t uints);
void OutputImage(cl_uint width, cl_uint height, const float *red, const float *green, const float *blue);

int main()
{
	std::vector<cl::Platform> platforms;
	cl::Platform::get(&platforms);

	std::vector<cl::Device> devices;
	platforms.at(0).getDevices(CL_DEVICE_TYPE_GPU, &devices);

	cl::Context ctx(devices.at(0));
	cl::CommandQueue command_queue(ctx);

	std::string program_string;
	{
		std::ifstream ifs("kernel.cl");
		program_string = std::string((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());
	}

	cl::Program program(ctx, program_string, true);
	cl::Kernel random_test_kernel(program, "RandomTest");

	cl::Buffer picture_buffer(ctx, CL_MEM_READ_WRITE, kPlaneBytes * 3);
	cl::Buffer random_buffer(ctx, CL_MEM_READ_WRITE| CL_MEM_HOST_WRITE_ONLY, sizeof (cl_uint) * kThreads);

	{
		void *p = command_queue.enqueueMapBuffer(random_buffer, CL_TRUE, CL_MAP_WRITE, 0, sizeof (cl_uint) * kThreads);
		FillRandomUints(static_cast<cl_uint *>(p), kThreads);
		command_queue.enqueueUnmapMemObject(random_buffer, p);
		command_queue.finish();
	}

	command_queue.enqueueFillBuffer(picture_buffer, cl_float(0), 0, kPlaneBytes * 3);
	command_queue.finish();

	random_test_kernel.setArg(0, picture_buffer);
	random_test_kernel.setArg(1, random_buffer);
	random_test_kernel.setArg(2, kPlanePixels);
	command_queue.enqueueNDRangeKernel(random_test_kernel, cl::NDRange(0), cl::NDRange(kThreads));
	command_queue.finish();

	{
		void *p = command_queue.enqueueMapBuffer(picture_buffer, CL_TRUE, CL_MAP_READ, 0, kPlaneBytes * 3);
		float *r = static_cast<float *>(p);
		float *g = r + kPlanePixels;
		float *b = g + kPlanePixels;

		OutputImage(kWidth, kHeight, r, g, b);
		command_queue.enqueueUnmapMemObject(picture_buffer, p);
	}
}

void FillRandomUints(cl_uint *dst, size_t uints)
{
	std::random_device seed_gen;
	std::mt19937 engine(seed_gen());

	while (uints--)
	{
		cl_uint v = engine();

		while (!v)
			v = engine();

		*dst++ = v;
	}
}

void OutputImage(cl_uint width, cl_uint height, const float *red, const float *green, const float *blue)
{
	std::cout << std::format("P3 {} {} 255\n", width, height);

	for (cl_uint y=0; y<height; ++y)
	{
		for (cl_uint x=0; x<width; ++x)
		{
			int r = std::clamp(static_cast<int>(255 * *red++), 0, 255);
			int g = std::clamp(static_cast<int>(255 * *green++), 0, 255);
			int b = std::clamp(static_cast<int>(255 * *blue++), 0, 255);

			std::cout << std::format("{} {} {}\n", r, g, b);
		}
	}

	std::cout.flush();
}

構造体とメンバ変数のアライメント

ホストプログラム - カーネル間で構造体をやり取りする際に留意すべきです。
例えばカーネル側ではfloat3を16バイトサイズ・16バイトアラインとしていますが、Visual C++とcl_platform.hの組合せでは特に何もしなければcl_float3を16バイトサイズ・4バイトアラインとしています。alignasキーワードを用いたりパディングするなりしてメンバ変数のアライメントを、また構造体の配列を扱う場合は構造体のアライメントも適切なものになるよう宣言しましょう。

Intel製iGPU用OpenCLを使用する際の注意

clFinishおよびclWaitForEventsでは、処理の完了を検出するのにビジーループでポーリングしているようです。
OpenCLをバックエンドとするSYCLでも同様です。
Direct3DとVulkanではこのような挙動はありません。
これを回避したいなら、clGetEventInfoでCL_EVENT_COMMAND_EXECUTION_STATUSを指定して状態を読み、適宜Sleepするとよいでしょう。
この際、Sleepの精度を上げるためにAvSetMmThreadCharacteristicsで"Pro Audio"を指定する、あるいはtimeBeginPeriodを呼ぶとよいでしょう。

Windows 11 23H2 (10.0.22631.3085)では、AvSetMmThreadCharacteristicsでSleepの精度は上がりませんでした。