サンプルコード
はじめに
AMDの提供している総合的な機械学習プラットフォームであるROCmは現状Windowsには対応していません。
ただし、そのバックエンドであるHIP(C++ Heterogeneous-Compute Interface for Portability)はいつの間にやらWindows上でひっそりと使えるようになっています。1
せっかくだから使ってみましょう。
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 ↩