OpenCLのはじめかたなど

# はじめに  
Windows環境でgitとMicrosoft Visual Studio Community 2022をインストールしているものとします。  
次の四点について記述します。  
* OpenCLのインポートライブラリとヘッダファイルの入手方法  
* 乱数を生成してみる  
* 構造体とメンバ変数のアライメント  
* Intel製iGPU用OpenCLを使用する際の注意  

# OpenCLのインポートライブラリとヘッダファイルの入手方法  
[Getting started with OpenCL on Microsoft Windows](https://github.com/KhronosGroup/OpenCL-Guide/blob/main/chapters/getting_started_windows.md) にあるように、  
* KhronosGroupのOpenCL-SDKリポジトリからビルドする方法  
* vcpkg経由でビルドする方法  

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

* [GPUOpen-LibrariesAndSDKs/OCL-SDK](https://github.com/GPUOpen-LibrariesAndSDKs/OCL-SDK/releases)からダウンロードする  
* 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して、それをもとに画像データを生成し吐き出すものを作ってみます。  

```OpenCL: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](https://www.digitalpad.co.jp/~takechin/betasue.html)に[Susie32 PbmPlus Plug-in](https://cetus.sakura.ne.jp/softlab/software/ifpnm.html)を加えて利用するとよいかもしれません。  

```c++: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を呼ぶとよいでしょう。  

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