HIP(ROCmバックエンド)がWindowsで(一応)使えます。

# サンプルコード

https://github.com/skyt301/hipewSample

# はじめに
**AMD**の提供している総合的な機械学習プラットフォームである**ROCm**は現状Windowsには対応していません。

https://www.amd.com/ja/graphics/servers-solutions-rocm

ただし、そのバックエンドである**HIP**(C++ Heterogeneous-Compute Interface for Portability)はいつの間にやらWindows上でひっそりと使えるようになっています。[^example]
せっかくだから使ってみましょう。


<details><summary>
2023/04/14時点で、もうすぐ対応するとの話が出てきたようです。</summary>

:::note warn
ただし、AMD時空で"もうすぐ"というのは、少なくとも半年以上、場合によっては1年以上先なのが通例です。
https://www.tomshardware.com/news/amd-rocm-comes-to-windows-on-consumer-gpus
:::
</details>

## Orochi
Windows上でHIPにアクセスする正式な方法としては**Orochi**と呼ばれるライブラリを使うことになっています。

https://github.com/GPUOpen-LibrariesAndSDKs/Orochi/

これ自体はHIP専用ではなく、**CUDA**とHIPを同一コードで扱えるようにするという趣旨のライブラリになっています。


<details><summary>Orochiのコード例(Orochiのサンプルより)</summary>

```C++:DeviceEnum/main.cpp
#include <Orochi/Orochi.h>
#include <Test/Common.h>

int main( int argc, char** argv )
{
	int a = oroInitialize( ( oroApi )( ORO_API_CUDA | ORO_API_HIP ), 0 );

	oroError e;
	e = oroInit( 0 );
	int nDevicesTotal;
	e = oroGetDeviceCount( &nDevicesTotal );
	ERROR_CHECK( e );
	int nAMDDevices;
	e = oroGetDeviceCount( &nAMDDevices, ORO_API_HIP );
	ERROR_CHECK( e );
	int nNVIDIADevices;
	e = oroGetDeviceCount( &nNVIDIADevices, ORO_API_CUDA );
	ERROR_CHECK( e );

	printf( "# of devices: %d\n", nDevicesTotal );
	printf( "# of AMD devices: %d\n", nAMDDevices );
	printf( "# of NV devices: %d\n\n", nNVIDIADevices );

	for( int i = 0; i < nDevicesTotal; i++ )
	{
		oroDevice device;
		e = oroDeviceGet( &device, i );
		ERROR_CHECK( e );

		char name[128];
		e = oroDeviceGetName( name, 128, device );
		ERROR_CHECK( e );

		oroDeviceProp props;
		e = oroGetDeviceProperties( &props, device );
		ERROR_CHECK( e );
		printf( "executing on %s (%s)\n", props.name, props.gcnArchName );

		oroCtx ctx;
		e = oroCtxCreate( &ctx, 0, device );
		ERROR_CHECK( e );

		//try kernel execution
		 oroFunction function;
		{
			const char* code = "extern \"C\" __global__ "
							   "void testKernel()"
							   "{ int a = threadIdx.x; printf(\"	thread %d running\\n\", a); }";
			const char* funcName = "testKernel";
			orortcProgram prog;
			orortcResult e;
			e = orortcCreateProgram( &prog, code, funcName, 0, 0, 0 );
			std::vector<const char*> opts;
			opts.push_back( "-I ../" );

			e = orortcCompileProgram( prog, opts.size(), opts.data() );
			if( e != ORORTC_SUCCESS )
			{
				size_t logSize;
				orortcGetProgramLogSize( prog, &logSize );
				if( logSize )
				{
					std::string log( logSize, '\0' );
					orortcGetProgramLog( prog, &log[0] );
					std::cout << log << '\n';
				};
			}
			size_t codeSize;
			e = orortcGetCodeSize( prog, &codeSize );

			std::vector<char> codec( codeSize );
			e = orortcGetCode( prog, codec.data() );
			e = orortcDestroyProgram( &prog );
			oroModule module;
			oroError ee = oroModuleLoadData( &module, codec.data() );
			ee = oroModuleGetFunction( &function, module, funcName );
		}

		void** args = {};
		oroError e = oroModuleLaunchKernel( function, 1, 1, 1, 32, 1, 1, 0, 0, args, 0 ); 
		oroDeviceSynchronize();

		oroApi api = oroGetCurAPI( 0 );
		printf( "executed on %s\n", api == ORO_API_HIP ? "AMD" : "NVIDIA" );
		e = oroCtxDestroy( ctx );
	}
	return 0;
}
```
</details>

ただこの趣旨のためか、HIPとCUDAの両方に存在するAPIしか使えるようになっておらず、特にメモリ周りの扱いは手間がかかってしまいます。
<details><summary>たった一つのデータのために手間がかかる…</summary>

```C++:
float a = 0.0f;
float b = 0.0f;
oroDeviceptr src, dest;
//デバイスメモリを計算元と計算先それぞれ確保
e = oroMalloc( &src, sizeof(float) );
e = oroMalloc( &dest, sizeof(float) );

//ホストメモリからデバイスメモリに内容をコピー
e = oroMemcpyHtoD( src, &a, sizeof(float));

//デバイスで計算
oroModuleLaunchKernel(...);
oroDeviceSynchronize(void);

//ホストで計算結果にアクセスするために、デバイスからホストにコピー
e = oroMemcpyDtoH(&b, dest, sizeof(float));

//ようやく結果にアクセスできます
std::cout << b << std::endl;
```
</details>
ホストメモリとデバイスメモリは別モノのため仕方ないのですが、各プラットフォームでは通常、相互によりアクセスしやすくなるような仕組みがあります。

## hipewってのがあるぞ…？
ところでOrochiのヘッダを眺めていると、どうやらOrochi側で実装していないだけでHIPのAPIはWindows上でもほぼすべて呼び出せるようになっているようです。
そのため`Orochi/contrib/hipew/`内のコードをそのまま使えばWindows上でHIPに直接アクセスできるようになります。
~~どうして…どうしてこっちを先に公開してくれないんですか…~~

# hipewからHIPにアクセスする。
完全なコードは https://github.com/skyt301/hipewSample にあります。
またHIPのAPIリファレンスは https://rocmdocs.amd.com/projects/HIP/en/develop/index.html にあります。
##### hipewを初期化する。
```C++
int driver, result;
hipewInit(&driver, &result, HIPEW_INIT_HIPRTC | HIPEW_INIT_HIPDRIVER);
```
##### HIPを初期化する。
```C++
hipInit(0);
```
##### (オプション)HIPデバイスの総数を取得。
```C++
int countDevice;
hipGetDeviceCount(&countDevice);
```
##### HIPデバイスの取得
```C++
hipDevice_t hipDevice;
hipDeviceGet(&hipDevice, 0);
```
##### カーネルコード
```Cuda:const char* code =
extern "C" {
    float sigmoid(float x)
	{
    	return 1 / (1 + expf(x));
	}
	__global__ void testKernel(float *a)
	{
		a[threadIdx.x] = sigmoid((float)1 / a[threadIdx.x]);
	}
}
```
##### (カーネル)プログラムを作成する。
```C++
hiprtcProgram prog = nullptr;
const char* funcName = "testKernel";

hiprtcCreateProgram(&prog, code, funcName, 0, 0, 0);
```
##### カーネルコードをコンパイルする。
```C++
std::vector<const char*> opts; //コンパイルオプションリスト
opts.push_back("-I ../");

hiprtcCompileProgram(prog, opts.size(), opts.data());
```
##### コンパイル済みバイナリを(デバイス)モジュールに読み込む。
```C++
size_t codeSize;
hiprtcGetCodeSize(prog, &codeSize);

std::vector<char> codec(codeSize);
//コンパイル済みバイナリを抽出
hiprtcGetCode(prog, codec.data());
//hiprtcProgramの破棄
hiprtcDestroyProgram(&prog);
hipModule_t module;
//コンパイル済みバイナリをデバイスにロード
hipError_t ee = hipModuleLoadData(&module, codec.data());
//実行可能デバイスカーネルを保持
hipModuleGetFunction(&function, module, funcName);
```
##### HIPマネージドメモリを確保して初期化。
```C++
float* managedFloat;
hipMallocManaged((hipDeviceptr_t*)&managedFloat, sizeof(float) * 4, hipMemAttachGlobal);
//マネージドメモリを初期化
managedFloat[0] = 1.0f;
managedFloat[1] = 4.0f;
managedFloat[2] = 8.0f;
managedFloat[3] = 16.0f;
```

##### gridDim(X,Y,Z), blockDim(X,Y,Z)など各種パラメータを指定してカーネルを起動する。
```C++
const void* args[] = { &managedFloat };  //カーネルに渡す引数
hipModuleLaunchKernel(function, 1, 1, 1, 4, 1, 1, 0, NULL, (void**)args, 0);
```
##### デバイス上のすべてのコマンドが終了するのを待機。
```C++
hipDeviceSynchronize(void);
```
##### 計算結果にアクセス
```C++
for (int i = 0; i < 4; i++)
{
	std::cout << "managedFloat[" << i << "]: " << managedFloat[i] << std::endl;
}
```
```bash
managedFloat[0]: 0.268941
managedFloat[1]: 0.437824
managedFloat[2]: 0.468791
managedFloat[3]: 0.48438
```

##### 確保したデバイスメモリの解放
```C++
hipFree((hipDeviceptr_t)managedFloat);
```
初期化と簡単なカーネルの実行ですが、一応動作していることが確認できるはずです。

# (試しているときに気が付いた)制限
Windows上ではカーネルコード内での`printf`や`assert`関数の使用や、`__host__`関数の定義は行えないようです。
おそらくWindows上にROCmが実装されたときに改めて利用可能になるのでしょう。

[^example]:Blender上でAMDGPU用プラグインのために導入されたようです。https://github.com/blender/blender/tree/main/extern/hipew