サンプルコード

はじめに

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

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

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

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

Orochi

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

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

Orochiのコード例(Orochiのサンプルより)

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;
}

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

たった一つのデータのために手間がかかる…

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;

ホストメモリとデバイスメモリは別モノのため仕方ないのですが、各プラットフォームでは通常、相互によりアクセスしやすくなるような仕組みがあります。

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を初期化する。

int driver, result;
hipewInit(&driver, &result, HIPEW_INIT_HIPRTC | HIPEW_INIT_HIPDRIVER);

HIPを初期化する。

hipInit(0);

(オプション)HIPデバイスの総数を取得。

int countDevice;
hipGetDeviceCount(&countDevice);

HIPデバイスの取得

hipDevice_t hipDevice;
hipDeviceGet(&hipDevice, 0);

カーネルコード

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]);
	}
}

(カーネル)プログラムを作成する。

hiprtcProgram prog = nullptr;
const char* funcName = "testKernel";

hiprtcCreateProgram(&prog, code, funcName, 0, 0, 0);

カーネルコードをコンパイルする。

std::vector<const char*> opts; //コンパイルオプションリスト
opts.push_back("-I ../");

hiprtcCompileProgram(prog, opts.size(), opts.data());

コンパイル済みバイナリを(デバイス)モジュールに読み込む。

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マネージドメモリを確保して初期化。

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)など各種パラメータを指定してカーネルを起動する。

const void* args[] = { &managedFloat };  //カーネルに渡す引数
hipModuleLaunchKernel(function, 1, 1, 1, 4, 1, 1, 0, NULL, (void**)args, 0);

デバイス上のすべてのコマンドが終了するのを待機。

hipDeviceSynchronize(void);

計算結果にアクセス

for (int i = 0; i < 4; i++)
{
	std::cout << "managedFloat[" << i << "]: " << managedFloat[i] << std::endl;
}

managedFloat[0]: 0.268941
managedFloat[1]: 0.437824
managedFloat[2]: 0.468791
managedFloat[3]: 0.48438

確保したデバイスメモリの解放

hipFree((hipDeviceptr_t)managedFloat);

初期化と簡単なカーネルの実行ですが、一応動作していることが確認できるはずです。

(試しているときに気が付いた)制限

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

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