Intel内蔵GPUをoneAPI Level Zeroで直接叩いてみる（実践編）

はじめに

前回の記事で、Level Zero API の紹介および導入方法について説明しました。

今回は Level Zero API の具体的な使用方法、C++ を用いたホストプログラムの書き方について見ていきます。Level Zero Specification documentation v1.7.8 をもとに、重要だと感じた点や、興味を惹かれた点を主に解説したいと思います。

本記事では、基本的な C++ の知識は習得済みであることを前提とします。

カーネルプログラム

カーネルプログラムについては OpenCL C 形式で書かれたコードを Level Zero 向けに変換して使用することができます。 OpenCL C については、OpenCL C 仕様書 や市販の入門書などを参考にしてください。

ホストプログラム

前回の記事で動かしたサンプルコードの中身を見てみましょう。
サンプルコード lzsample.cpp では、以下の処理を順に行います。

• 初期化
  • LevelZero ドライバ初期化
  • driver ハンドラ取得
  • device ハンドラ取得
• context 作成
• command queue / command list 作成
• shared memory 割り当て
• module / kernel 作成、設定
• 実行
• クリーンアップ

各処理の説明

前置き

Level Zero API は、関数呼び出し時の戻り値としてエラーコードを返します。
処理成功時に ZE_RESULT_SUCCESS が返ることを確認するために、用意した CHECK_RESULT マクロを使用しています。

サンプルコードでは、分かりやすさを優先するためエラーハンドリングを簡略化しています。

CHECK_RESULT マクロ (抜粋)

#define CHECK_RESULT(CALL) validateCall(CALL, #CALL)

void validateCall(ze_result_t result, const char *msg) {
  if (result != ZE_RESULT_SUCCESS) {
    std::cout << "failed: " << msg << " code:" << result << std::endl;
    std::terminate();
  }
}

初期化

Level Zero API の仕様により、最初に必ず zeInit() を実行する必要があります。
その後 zeDriverGet() および zeDeviceGet() により driver ハンドラ および device ハンドラを取得します。

zeDriverGet() や zeDeviceGet() は、最初に要素数のみを取得してから、実際の要素を取得する流れとなります。環境によっては driver や device が複数存在する可能性がありますが、サンプルコードでは簡略化のため先頭の要素のみを使用しています。

初期化 (抜粋)

// LevelZero ドライバ初期化 (initLevelZero)
CHECK_RESULT(zeInit(ZE_INIT_FLAG_GPU_ONLY));

// driver ハンドラ取得 (getDriver)
uint32_t driverCount = 0;
CHECK_RESULT(zeDriverGet(&driverCount, nullptr));
std::vector<ze_driver_handle_t> allDrivers(driverCount);
CHECK_RESULT(zeDriverGet(&driverCount, allDrivers.data()));

// device ハンドラ取得 (getDevice)
uint32_t deviceCount = 0;
CHECK_RESULT(zeDeviceGet(hDriver, &deviceCount, nullptr));
std::vector<ze_device_handle_t> allDevices(deviceCount);
CHECK_RESULT(zeDeviceGet(hDriver, &deviceCount, allDevices.data()));

context 作成

context は、この後に出てくる command queue / command list や module あるいはメモリなどを管理するために使用される論理オブジェクトです。zeContextCreate() により context を作成します。

context 作成 (抜粋)

// context 作成 (createContext)
ze_context_handle_t hContext = nullptr;
ze_context_desc_t contextDescription = {ZE_STRUCTURE_TYPE_CONTEXT_DESC};
CHECK_RESULT(zeContextCreate(hDriver, &contextDescription, &hContext));

command queue / command list 作成

下記の図は Level Zero Specification documentation v1.7.8 より引用

ここでは command queue group, command queue, command list について説明します。

• command queue group は、物理的な入力ストリーム (実際のデバイス上のエンジン) を示します
• command queue は、論理的な入力ストリームを示します
• command list は、command queue で実行するための一連のコマンドを示します

Intel N100 CPU 内蔵グラフィックスでは command queue group の数は 1 ですが、複数存在する場合や、データコピー専用エンジンを持つ場合もあるようです。
まずは zeDeviceGetCommandQueueGroupProperties() によりカーネル実行可能な command queue group を探します。

command queue group 取得 (抜粋)

// command queue group 取得 (getCmdQueueGroup)

// command queue group の数を取得
uint32_t numQueueGroups = 0;
CHECK_RESULT(zeDeviceGetCommandQueueGroupProperties(hDevice, &numQueueGroups, nullptr));
std::vector<ze_command_queue_group_properties_t> allQueueGroups(numQueueGroups);
CHECK_RESULT(zeDeviceGetCommandQueueGroupProperties(hDevice, &numQueueGroups, allQueueGroups.data()));

// カーネル実行が可能な command queue group を探す
int cmdQueueGroupOrdinal = -1;
for (int i = 0; i < numQueueGroups; i++) {
  if (allQueueGroups[i].flags & ZE_COMMAND_QUEUE_GROUP_PROPERTY_FLAG_COMPUTE) {
    cmdQueueGroupOrdinal = i;
    break;
  }
}

zeCommandQueueCreate() により command queue を作成します。
ここで cmdQueueDesc.ordinal で command queue group と紐づける必要があります。
また、サンプルコードでは非同期で処理するため cmdQueueDesc.mode に ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS を指定しています。

zeCommandListCreate() により command list を作成します。
こちらも cmdListDesc.commandQueueGroupOrdinal で command queue group と紐づけが必要です。

command queue / command list 作成 (抜粋)

// command queue 作成 (createCmdQueue)
ze_command_queue_handle_t hCmdQueue = nullptr;
ze_command_queue_desc_t cmdQueueDesc = {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC};
cmdQueueDesc.ordinal = cmdQueueGroupOrdinal;
cmdQueueDesc.index = 0;
cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS;
CHECK_RESULT(zeCommandQueueCreate(hContext, hDevice, &cmdQueueDesc, &hCmdQueue));

// command list 作成 (createCmdList)
ze_command_list_handle_t hCmdList = nullptr;
ze_command_list_desc_t cmdListDesc = {ZE_STRUCTURE_TYPE_COMMAND_LIST_DESC};
cmdListDesc.commandQueueGroupOrdinal = cmdQueueGroupOrdinal;
CHECK_RESULT(zeCommandListCreate(hContext, hDevice, &cmdListDesc, &hCmdList));

shared memory 割り当て

GPU から読み書き可能なメモリを割り当てます。
ここでは zeMemAllocShared() により、CPU および GPU からそれぞれアクセス可能な共有メモリを使用します。

サンプルコードでは、以下のように sharedA, sharedB, sharedC としてそれぞれ shared memory の割り当てを行っています。

shared memory 割り当て (抜粋)

void * createSharedMem(ze_context_handle_t hContext, ze_device_handle_t hDevice, uint32_t bufSize)
{
  void *bufPtr;
  ze_device_mem_alloc_desc_t memAllocDesc = {ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC};
  memAllocDesc.ordinal = 0;
  ze_host_mem_alloc_desc_t hostDesc = {ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC};

  CHECK_RESULT(zeMemAllocShared(hContext, &memAllocDesc, &hostDesc, bufSize, 1, hDevice, &bufPtr));
  return bufPtr;
}

// shared memory 割り当て
uint32_t bufSize = sizeof(int32_t) * itemCount;
void *sharedA = createSharedMem(hContext, hDevice, bufSize);
void *sharedB = createSharedMem(hContext, hDevice, bufSize);
void *sharedC = createSharedMem(hContext, hDevice, bufSize);

module / kernel 作成、設定

module は、カーネルプログラムをまとめたもので、複数のカーネル関数を含むことができます。ここでは、あらかじめ用意したカーネルプログラム用バイナリ形式 (SPIR-V) ファイルをロードしてから、zeModuleCreate() により module を作成します。

module 作成 (抜粋)

// module 作成 (createModule)
ze_module_handle_t hModule = nullptr;
ze_module_build_log_handle_t buildLog = nullptr;
ze_module_desc_t moduleDesc = {ZE_STRUCTURE_TYPE_MODULE_DESC};
moduleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV;
moduleDesc.inputSize = length;
moduleDesc.pInputModule = reinterpret_cast<const uint8_t *>(spirvData.get());
moduleDesc.pBuildFlags = "";
CHECK_RESULT(zeModuleCreate(hContext, hDevice, &moduleDesc, &hModule, &buildLog));

kernel は、それぞれのカーネル関数を示します。 ここでは kernelDesc.pKernelName にカーネル関数名を指定して、zeKernelCreate() により kernel を作成します。

kernel 作成 (抜粋)

ze_kernel_handle_t createKernel(ze_module_handle_t hModule, const char *kernelName)
{
  ze_kernel_handle_t hKernel = nullptr;
  ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC};
  kernelDesc.pKernelName = kernelName;
  CHECK_RESULT(zeKernelCreate(hModule, &kernelDesc, &hKernel));
  return hKernel;
}

// kernel 作成
ze_kernel_handle_t hKernelVectAdd = createKernel(hModule, "vectadd");

作成した kernel に対して、カーネル関数の引数をそれぞれ設定する必要があります。

ここで使用するカーネル関数 vectadd は、以下のように3つの引数を持ちます。
__kernel void vectadd (__global int *a, __global int *b, __global int *c)
そのため、それぞれに対応する shared memory のポインタを zeKernelSetArgumentValue() により設定します。

kernel argument 設定 (抜粋)

void setKernelArgument(ze_kernel_handle_t hKernel, uint32_t index, uint32_t size, void *bufPtr)
{
  CHECK_RESULT(zeKernelSetArgumentValue(hKernel, index, size, &bufPtr));
}

// kernel argument 設定
setKernelArgument(hKernelVectAdd, 0, bufSize, sharedA);
setKernelArgument(hKernelVectAdd, 1, bufSize, sharedB);
setKernelArgument(hKernelVectAdd, 2, bufSize, sharedC);

また kernel に対して、ワークグループのサイズを設定する必要があります。
Level Zero API では、zeKernelSuggestGroupSize() を用いることで、ワークアイテム数に応じた推奨グループサイズを取得することができます。 ここでは、取得した推奨値をそのまま zeKernelSetGroupSize() により設定します。

kernel dispatch 設定 (抜粋)

void setKernelGroupSize(ze_kernel_handle_t hKernel, uint32_t itemsX, uint32_t itemsY, uint32_t itemsZ, ze_group_count_t *dispatch)
{
  // 推奨グループサイズを取得
  uint32_t groupSizeX;
  uint32_t groupSizeY;
  uint32_t groupSizeZ;
  CHECK_RESULT(zeKernelSuggestGroupSize(hKernel, itemsX, itemsY, itemsZ, &groupSizeX, &groupSizeY, &groupSizeZ));

  // グループサイズを設定
  CHECK_RESULT(zeKernelSetGroupSize(hKernel, groupSizeX, groupSizeY, groupSizeZ));

  // kernel dispatch 設定
  dispatch->groupCountX = itemsX / groupSizeX;
  dispatch->groupCountY = itemsY / groupSizeY;
  dispatch->groupCountZ = itemsZ / groupSizeZ;
}

// kernel dispatch 設定
ze_group_count_t dispatchVectAdd = {};
setKernelGroupSize(hKernelVectAdd, itemCount, 1, 1, &dispatchVectAdd);

実行

作成した kernel を実行するために、command list を利用します。
ここでは zeCommandListAppendLaunchKernel() により command list に対してカーネル起動コマンドを追加します。

command list の仕様により、コマンドの追加が完了したら zeCommandListClose() によりクローズする必要があります。

command list に対する kernel 起動コマンドの設定 (抜粋)

void prepareCmdList(ze_command_list_handle_t hCmdList, ze_kernel_handle_t hKernel, ze_group_count_t *dispatch)
{
  // command list リセット
  CHECK_RESULT(zeCommandListReset(hCmdList));
  // カーネル起動コマンドを追加
  CHECK_RESULT(zeCommandListAppendLaunchKernel(hCmdList, hKernel, dispatch, nullptr, 0, nullptr));
  // command list クローズ
  CHECK_RESULT(zeCommandListClose(hCmdList));
}

// command list に対して kernel 起動コマンドを設定
prepareCmdList(hCmdList, hKernelVectAdd, &dispatchVectAdd);

command list の準備ができたら、zeCommandQueueExecuteCommandLists() により command queue に対して実行を要求します。

サンプルコードでは、非同期で処理する command queue を使用しているため、zeCommandQueueSynchronize() により処理が完了するのを待ちます。

実行 (抜粋)

void executeCmdList(ze_command_queue_handle_t hCmdQueue, ze_command_list_handle_t hCmdList)
{
  // コマンド実行を要求
  CHECK_RESULT(zeCommandQueueExecuteCommandLists(hCmdQueue, 1, &hCmdList, nullptr));
  // コマンド実行が完了するのを待つ
  CHECK_RESULT(zeCommandQueueSynchronize(hCmdQueue, UINT64_MAX));
}

// 実行
executeCmdList(hCmdQueue, hCmdList);

zeCommandQueueSynchronize() の引数 timeout に UINT64_MAX を設定した場合、処理が完了するまでビジーループで待つことになります。一方、timeout に 0 を指定することで、下記のようにポーリングで確認することもできます。

ポーリングによる実行完了待ち (例)

while (true) {
  ze_result_t result = zeCommandQueueSynchronize(hCmdQueue, 0);
  if (result == ZE_RESULT_SUCCESS) {
    break;
  } else if (result == ZE_RESULT_NOT_READY) {
    usleep(1000);  /* 1ms sleep */
  } else {
    /* エラー処理 (省略) */
  }
}

クリーンアップ

最後に、作成した context などのクリーンアップを行います。

クリーンアップ (抜粋)

zeKernelDestroy(hKernelVectAdd);
zeModuleDestroy(hModule);
zeMemFree(hContext, sharedA);
zeMemFree(hContext, sharedB);
zeMemFree(hContext, sharedC);
zeCommandListDestroy(hCmdList);
zeCommandQueueDestroy(hCmdQueue);
zeContextDestroy(hContext);

GPU 演算

前回の記事で示したサンプルコードでは、以下のベクトル加算処理を行います。
sharedC[idx] = sharedA[idx] + sharedB[idx]; ( idx は 0 から 999999 まで)

ここで sharedA の値は全て固定値 (100) とし、sharedB の値は要素番号をそのまま代入します。sharedC の期待値は、以下のようになります。

idx	sharedA	sharedB	sharedC
0	100	0	100
1	100	1	101
...	...	...	...
999999	100	999999	1000099

サンプルコードでは、出力された sharedC の値が、上記の期待値と一致するかを確認しています。

command list を用いた複数コマンドの実行

command list を利用することで、複数のコマンドを一度にまとめて要求することができます。カーネル起動コマンド以外にも、メモリコピーやメモリアドレス範囲を指定したバリア設定などのコマンドも用意されており、うまく使うことで複雑な処理を GPU に任せることができます。

カーネル起動コマンドの逐次実行

実際に command list を使って、複数のカーネル起動コマンドを順に実行してみましょう。

ここでは、以下のカーネルプログラムを使用します。
vectadd はベクトル加算処理、vectincrement は各要素に 1 を加算するカーネル関数です。

vectadd2.cl

__kernel void vectadd (__global int *a, __global int *b, __global int *c)
{
  int idx = get_global_id(0);
  c[idx] = a[idx] + b[idx];
}
__kernel void vectincrement (__global int *src, __global int *dst)
{
  int idx = get_global_id(0);
  dst[idx] = src[idx] + 1;
}

ホストプログラム側は、vectadd からの出力を vectincrement の入力として順に実行させるように command list を用意します。

• shared memory 割り当て
  • sharedA, sharedB, sharedC に加えて、sharedD を追加
• kernel 作成
  • vectadd に加えて、vectincrement 用の kernel を作成
• kernel 設定
  • vectincrement 用の kernel argument 設定を追加
    • __kernel void vectincrement (__global int *src, __global int *dst)
      0番目引数 (src) に sharedC を設定する
      1番目引数 (dst) に sharedD を設定する
  • vectincrement 用の kernel dispatch 設定を追加
    • vectadd と同様に、推奨グループサイズを使用する
• command list に対する kernel 起動コマンドの設定
  • vectadd に加えて、vectincrement に対するカーネル起動コマンドを追加

以下のように zeCommandListAppendLaunchKernel() を使用して、vectadd と vectincrement をそれぞれ追加すれば良さそうですね。

command list に対する複数の kernel 起動コマンドの設定案

// command list リセット
CHECK_RESULT(zeCommandListReset(hCmdList));
// command list に対して kernel 起動コマンドを設定 (vectadd)
CHECK_RESULT(zeCommandListAppendLaunchKernel(hCmdList, hKernelVectAdd, &dispatchVectAdd, nullptr, 0, nullptr));
// command list に対して kernel 起動コマンドを設定 (vectincrement)
CHECK_RESULT(zeCommandListAppendLaunchKernel(hCmdList, hKernelVectInc, &dispatchVectInc, nullptr, 0, nullptr));
// command list クローズ
CHECK_RESULT(zeCommandListClose(hCmdList));

GPU 演算

sharedA の値は全て固定値 (100) とし、sharedB の値は要素番号をそのまま代入します。
sharedC および sharedD の期待値は、以下のようになります。

idx	sharedA	sharedB	sharedC	sharedD
0	100	0	100	101
1	100	1	101	102
...	...	...	...	...
999999	100	999999	1000099	1000100

出力された sharedC, sharedD の値が、上記の期待値と一致するかを確認します。

実行結果

まずは、要素数を 1000000 に設定して実行してみます。一見すると正しく動いているように見えますね。

実行結果 (うまくいくケース)

./lzsample2a 1000000
itemCount:1000000
compute finished. ok:1000000 ng:0

しかし、要素数を 1 に設定すると正しく動きませんでした。
また、要素数を 1000～10000 あたりに設定すると、タイミングによって正しく動いたり動かなかったりと怪しい挙動が見られました。

実行結果 (うまくいかないケース)

./lzsample2a 1
itemCount:1
failed. idx:0, sharedC:100, expectedC:100, sharedD:1, expectedD:101
compute finished. ok:0 ng:1

./lzsample2a 1000
itemCount:1000
failed. idx:64, sharedC:164, expectedC:164, sharedD:1, expectedD:165
（省略）
compute finished. ok:224 ng:776

この原因は、コマンドが実行される順番にあります。

command list で実行されるコマンドは、投入された順に処理が開始されます（※）が、完了のタイミングについては決まっていません。
そのため vectadd が完了する前に vectincrement が開始される可能性があり、このケースでは sharedD の値が意図しない結果となってしまいます。

※ オプション指定によっては、処理が開始される順番も入れ替わる可能性があるようです

バリア設定の追加

vectadd が完了した後に vectincrement を実行してほしい場合、どうすればよいでしょうか？ 様々な方法がありますが、ここではメモリアドレス範囲を指定したバリアを使ってみます。

以下のように zeCommandListAppendMemoryRangesBarrier() により、command list に対してバリアを追加します。ここでは vectadd からの出力である sharedC をバリア対象とします。

vectadd と vectincrement のカーネル起動コマンドの間に入れることで vectadd が完了するのを待つことができるとともに、キャッシュなどに残った古いデータが使われることを防止できます。

command list に対する複数の kernel 起動コマンドの設定 (バリア追加)

// command list リセット
CHECK_RESULT(zeCommandListReset(hCmdList));
// command list に対して kernel 起動コマンドを設定 (vectadd)
CHECK_RESULT(zeCommandListAppendLaunchKernel(hCmdList, hKernelVectAdd, &dispatchVectAdd, nullptr, 0, nullptr));
// sharedC に対するバリア設定
int barrierNum = 1;
std::vector<const void *> barrierPtrs = {sharedC};
std::vector<size_t> barrierSizes = {bufSize};
CHECK_RESULT(zeCommandListAppendMemoryRangesBarrier(hCmdList, barrierNum, barrierSizes.data(), barrierPtrs.data(), nullptr, 0, nullptr));
// command list に対して kernel 起動コマンドを設定 (vectincrement)
CHECK_RESULT(zeCommandListAppendLaunchKernel(hCmdList, hKernelVectInc, &dispatchVectInc, nullptr, 0, nullptr));
// command list クローズ
CHECK_RESULT(zeCommandListClose(hCmdList));

実行結果を見ると、要素数に関わらず正しく動いていることが分かります。

実行結果 (バリア設定を追加した場合)

./lzsample2b 1
itemCount:1
compute finished. ok:1 ng:0

./lzsample2b 1000
itemCount:1000
compute finished. ok:1000 ng:0

./lzsample2b 1000000
itemCount:1000000
compute finished. ok:1000000 ng:0

immediate command list

これまでは command list および command queue を利用して GPU に対して非同期で処理させる方法を紹介しましたが、immediate command list を使うことで命令をその場で実行することもできます。

immediate command list は、command list と command queue が合わさったような構造になっており、zeCommandListCreateImmediate() により作成します。同期的に処理を行うため cmdQueueDesc.mode に ZE_COMMAND_QUEUE_MODE_SYNCHRONOUS を指定します。

immediate command list 作成

// immediate command list 作成 (createCmdListImmediate)
ze_command_list_handle_t hCmdListImm = nullptr;
ze_command_queue_desc_t cmdQueueDesc = {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC};
cmdQueueDesc.ordinal = cmdQueueGroupOrdinal;
cmdQueueDesc.index = 0;
cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_SYNCHRONOUS;
CHECK_RESULT(zeCommandListCreateImmediate(hContext, hDevice, &cmdQueueDesc, &hCmdListImm));

通常の command list と同様に、zeCommandListAppendLaunchKernel() などを使ってコマンドを設定できます。ここで投入されたコマンドは即座に実行されます。

zeCommandQueueExecuteCommandLists() や zeCommandQueueSynchronize() は immediate command list では使用しません。

immediate command list に対する kernel 起動コマンドの設定

CHECK_RESULT(zeCommandListAppendLaunchKernel(hCmdListImm, hKernelVectAdd, &dispatchVectAdd, nullptr, 0, nullptr));

immediate command list を用いたカーネル実行

immediate command list を使用して vectadd を実行してみます。

• command queue / command list 作成
  • 通常の command list の代わりに immediate command list を作成
  • 通常の command queue は使用しない
• 実行
  • immediate command list に対するカーネル起動コマンドに変更

実行結果

immediate command list を使用して、正しく動作することを確認できました。

実行結果 (immediate command list を使用)

./lzsample2c 1000000
itemCount:1000000
compute finished. ok:1000000 ng:0

おわりに

本記事では Level Zero API の具体的な使用方法について紹介しました。

ホストプログラムの全体像および各処理の具体的な使用方法を示すとともに、コマンドの即時実行やメモリアドレス範囲を指定したバリア、推奨されるグループサイズの自動算出といった Level Zero API の特徴的な機能を実際に動かして確認することができました。

今回使ったソースコードを 付録 に載せています。

参考文献

付録

サンプルコード

ファイル構造

lzsample2
├── vectadd2.cl
└── lzsample2.cpp

ビルド、実行

カーネルプログラム（ -device adl-n は使用する CPU に合わせてください）

ocloc compile -file vectadd2.cl -device adl-n -spv_only -o vectadd2.spv

ホストプログラム（ -DSAMPLE_ADDINC, -DSAMPLE_ADDINC_BARRIER, -DSAMPLE_IMMEDIATE のいずれか1つを指定してください）

# command list を用いた複数コマンドの実行 (バリアなし)
gcc -DSAMPLE_ADDINC lzsample2.cpp -o lzsample2a -lze_loader -lstdc++
./lzsample2a 1000000

# command list を用いた複数コマンドの実行 (バリアあり)
gcc -DSAMPLE_ADDINC_BARRIER lzsample2.cpp -o lzsample2b -lze_loader -lstdc++
./lzsample2b 1000000

# immediate command list を用いたカーネル実行
gcc -DSAMPLE_IMMEDIATE lzsample2.cpp -o lzsample2c -lze_loader -lstdc++
./lzsample2c 1000000

vectadd2.cl

vectadd2.cl

__kernel void vectadd (__global int *a, __global int *b, __global int *c)
{
  int idx = get_global_id(0);
  c[idx] = a[idx] + b[idx];
}
__kernel void vectincrement (__global int *src, __global int *dst)
{
  int idx = get_global_id(0);
  dst[idx] = src[idx] + 1;
}

lzsample2.cpp

lzsample2.cpp

#include <level_zero/ze_api.h>
#include <iostream>
#include <fstream>
#include <vector>
#include <memory>

#define CHECK_RESULT(CALL) validateCall(CALL, #CALL)

void validateCall(ze_result_t result, const char *msg)
{
  if (result != ZE_RESULT_SUCCESS) {
    std::cout << "failed: " << msg << " code:" << result << std::endl;
    std::terminate();
  }
}

void initLevelZero(void)
{
  CHECK_RESULT(zeInit(ZE_INIT_FLAG_GPU_ONLY));
}

ze_driver_handle_t getDriver(void)
{
  uint32_t driverCount = 0;
  CHECK_RESULT(zeDriverGet(&driverCount, nullptr));
  if (driverCount == 0) {
    std::cout << "driver not found." << std::endl;
    std::terminate();
  }

  std::vector<ze_driver_handle_t> allDrivers(driverCount);
  CHECK_RESULT(zeDriverGet(&driverCount, allDrivers.data()));
  return allDrivers[0];
}

ze_device_handle_t getDevice(ze_driver_handle_t hDriver)
{
  uint32_t deviceCount = 0;
  CHECK_RESULT(zeDeviceGet(hDriver, &deviceCount, nullptr));
  if (deviceCount == 0) {
    std::cout << "device not found." << std::endl;
    std::terminate();
  }

  std::vector<ze_device_handle_t> allDevices(deviceCount);
  CHECK_RESULT(zeDeviceGet(hDriver, &deviceCount, allDevices.data()));
  return allDevices[0];
}

ze_context_handle_t createContext(ze_driver_handle_t hDriver)
{
  ze_context_handle_t hContext = nullptr;
  ze_context_desc_t contextDescription = {ZE_STRUCTURE_TYPE_CONTEXT_DESC};
  CHECK_RESULT(zeContextCreate(hDriver, &contextDescription, &hContext));

  return hContext;
}

uint32_t getCmdQueueGroup(ze_device_handle_t hDevice)
{
  // command queue group の数を取得
  uint32_t numQueueGroups = 0;
  CHECK_RESULT(zeDeviceGetCommandQueueGroupProperties(hDevice, &numQueueGroups, nullptr));
  if (numQueueGroups == 0) {
    std::cout << "command queue group not found." << std::endl;
    std::terminate();
  }

  std::vector<ze_command_queue_group_properties_t> allQueueGroups(numQueueGroups);
  CHECK_RESULT(zeDeviceGetCommandQueueGroupProperties(hDevice, &numQueueGroups, allQueueGroups.data()));

  // カーネル実行が可能な command queue group を探す
  int cmdQueueGroupOrdinal = -1;
  for (int i = 0; i < numQueueGroups; i++) {
    if (allQueueGroups[i].flags & ZE_COMMAND_QUEUE_GROUP_PROPERTY_FLAG_COMPUTE) {
      cmdQueueGroupOrdinal = i;
      break;
    }
  }
  if (cmdQueueGroupOrdinal < 0) {
    std::cout << "command queue group for compute not found." << std::endl;
    std::terminate();
  }

  return cmdQueueGroupOrdinal;
}

ze_command_queue_handle_t createCmdQueue(ze_context_handle_t hContext, ze_device_handle_t hDevice, uint32_t cmdQueueGroupOrdinal)
{
  ze_command_queue_handle_t hCmdQueue = nullptr;
  ze_command_queue_desc_t cmdQueueDesc = {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC};
  cmdQueueDesc.ordinal = cmdQueueGroupOrdinal;
  cmdQueueDesc.index = 0;
  cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS;
  CHECK_RESULT(zeCommandQueueCreate(hContext, hDevice, &cmdQueueDesc, &hCmdQueue));

  return hCmdQueue;
}

ze_command_list_handle_t createCmdList(ze_context_handle_t hContext, ze_device_handle_t hDevice, uint32_t cmdQueueGroupOrdinal)
{
  ze_command_list_handle_t hCmdList = nullptr;
  ze_command_list_desc_t cmdListDesc = {ZE_STRUCTURE_TYPE_COMMAND_LIST_DESC};
  cmdListDesc.commandQueueGroupOrdinal = cmdQueueGroupOrdinal;
  CHECK_RESULT(zeCommandListCreate(hContext, hDevice, &cmdListDesc, &hCmdList));

  return hCmdList;
}

ze_command_list_handle_t createCmdListImmediate(ze_context_handle_t hContext, ze_device_handle_t hDevice, uint32_t cmdQueueGroupOrdinal)
{
  ze_command_list_handle_t hCmdListImm = nullptr;
  ze_command_queue_desc_t cmdQueueDesc = {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC};
  cmdQueueDesc.ordinal = cmdQueueGroupOrdinal;
  cmdQueueDesc.index = 0;
  cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_SYNCHRONOUS;
  CHECK_RESULT(zeCommandListCreateImmediate(hContext, hDevice, &cmdQueueDesc, &hCmdListImm));

  return hCmdListImm;
}

void * createSharedMem(ze_context_handle_t hContext, ze_device_handle_t hDevice, uint32_t bufSize)
{
  void *bufPtr;

  ze_device_mem_alloc_desc_t memAllocDesc = {ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC};
  memAllocDesc.ordinal = 0;
  ze_host_mem_alloc_desc_t hostDesc = {ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC};

  CHECK_RESULT(zeMemAllocShared(hContext, &memAllocDesc, &hostDesc, bufSize, 1, hDevice, &bufPtr));
  return bufPtr;
}

ze_module_handle_t createModule(ze_context_handle_t hContext, ze_device_handle_t hDevice, const char *filePath)
{
  std::unique_ptr<char[]> spirvData = nullptr;
  size_t length;

  // SPIR-V ファイルを開き、カーネルプログラムのバイナリデータを取得
  std::ifstream file(filePath, std::ios::binary);
  if (file.is_open()) {

    file.seekg(0, file.end);
    length = file.tellg();
    file.seekg(0, file.beg);

    spirvData.reset(new char[length]);
    file.read(spirvData.get(), length);
    file.close();

  } else {
    std::cout << "spv file cannot open." << std::endl;
    std::terminate();
  }

  // module 作成
  ze_module_handle_t hModule = nullptr;
  ze_module_build_log_handle_t buildLog = nullptr;
  ze_module_desc_t moduleDesc = {ZE_STRUCTURE_TYPE_MODULE_DESC};
  moduleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV;
  moduleDesc.inputSize = length;
  moduleDesc.pInputModule = reinterpret_cast<const uint8_t *>(spirvData.get());
  moduleDesc.pBuildFlags = "";
  CHECK_RESULT(zeModuleCreate(hContext, hDevice, &moduleDesc, &hModule, &buildLog));

  // module ビルドログの削除
  zeModuleBuildLogDestroy(buildLog);

  return hModule;
}

ze_kernel_handle_t createKernel(ze_module_handle_t hModule, const char *kernelName)
{
  // kernel 作成
  ze_kernel_handle_t hKernel = nullptr;
  ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC};
  kernelDesc.pKernelName = kernelName;
  CHECK_RESULT(zeKernelCreate(hModule, &kernelDesc, &hKernel));

  return hKernel;
}

void setKernelArgument(ze_kernel_handle_t hKernel, uint32_t index, uint32_t size, void *bufPtr)
{
  CHECK_RESULT(zeKernelSetArgumentValue(hKernel, index, size, &bufPtr));
}

void setKernelGroupSize(ze_kernel_handle_t hKernel, uint32_t itemsX, uint32_t itemsY, uint32_t itemsZ, ze_group_count_t *dispatch)
{
  // 推奨グループサイズを取得
  uint32_t groupSizeX;
  uint32_t groupSizeY;
  uint32_t groupSizeZ;
  CHECK_RESULT(zeKernelSuggestGroupSize(hKernel, itemsX, itemsY, itemsZ, &groupSizeX, &groupSizeY, &groupSizeZ));

  // グループサイズを設定
  CHECK_RESULT(zeKernelSetGroupSize(hKernel, groupSizeX, groupSizeY, groupSizeZ));

  // kernel dispatch 設定
  dispatch->groupCountX = itemsX / groupSizeX;
  dispatch->groupCountY = itemsY / groupSizeY;
  dispatch->groupCountZ = itemsZ / groupSizeZ;
}

int main( int argc, char *argv[] )
{
  // 演算回数を指定
  uint32_t itemCount = 1000 * 1000;
  if (argc >= 2) {
    int val = std::stoi(argv[1]);
    if (val > 0) {
      itemCount = val;
    }
  }
  std::cout << "itemCount:" << itemCount << std::endl;

  // LevelZero ドライバ初期化
  initLevelZero();
  // driver ハンドラ取得
  ze_driver_handle_t hDriver = getDriver();
  // device ハンドラ取得
  ze_device_handle_t hDevice = getDevice(hDriver);

  // context 作成
  ze_context_handle_t hContext = createContext(hDriver);

#if defined(SAMPLE_ADDINC) || defined(SAMPLE_ADDINC_BARRIER)

  // command queue 作成
  uint32_t cmdQueueGroupOrdinal = getCmdQueueGroup(hDevice);
  ze_command_queue_handle_t hCmdQueue = createCmdQueue(hContext, hDevice, cmdQueueGroupOrdinal);
  // command list 作成
  ze_command_list_handle_t hCmdList = createCmdList(hContext, hDevice, cmdQueueGroupOrdinal);

  // shared memory 割り当て
  uint32_t bufSize = sizeof(int32_t) * itemCount;
  void *sharedA = createSharedMem(hContext, hDevice, bufSize);
  void *sharedB = createSharedMem(hContext, hDevice, bufSize);
  void *sharedC = createSharedMem(hContext, hDevice, bufSize);
  void *sharedD = createSharedMem(hContext, hDevice, bufSize);

  // module 作成
  ze_module_handle_t hModule = createModule(hContext, hDevice, "vectadd2.spv");
  // kernel 作成
  ze_kernel_handle_t hKernelVectAdd = createKernel(hModule, "vectadd");
  ze_kernel_handle_t hKernelVectInc = createKernel(hModule, "vectincrement");

  // kernel argument 設定
  setKernelArgument(hKernelVectAdd, 0, bufSize, sharedA);
  setKernelArgument(hKernelVectAdd, 1, bufSize, sharedB);
  setKernelArgument(hKernelVectAdd, 2, bufSize, sharedC);
  setKernelArgument(hKernelVectInc, 0, bufSize, sharedC);
  setKernelArgument(hKernelVectInc, 1, bufSize, sharedD);

  // kernel dispatch 設定
  uint32_t itemsX = itemCount;
  uint32_t itemsY = 1;
  uint32_t itemsZ = 1;
  ze_group_count_t dispatchVectAdd = {};
  setKernelGroupSize(hKernelVectAdd, itemsX, itemsY, itemsZ, &dispatchVectAdd);
  ze_group_count_t dispatchVectInc = {};
  setKernelGroupSize(hKernelVectInc, itemsX, itemsY, itemsZ, &dispatchVectInc);

  // command list リセット
  CHECK_RESULT(zeCommandListReset(hCmdList));
  // command list に対して kernel 起動コマンドを設定 (vectadd)
  CHECK_RESULT(zeCommandListAppendLaunchKernel(hCmdList, hKernelVectAdd, &dispatchVectAdd, nullptr, 0, nullptr));

#if defined(SAMPLE_ADDINC_BARRIER)
  // sharedC に対するバリア設定
  int barrierNum = 1;
  std::vector<const void *> barrierPtrs = {sharedC};
  std::vector<size_t> barrierSizes = {bufSize};
  CHECK_RESULT(zeCommandListAppendMemoryRangesBarrier(hCmdList, barrierNum, barrierSizes.data(), barrierPtrs.data(), nullptr, 0, nullptr));
#endif

  // command list に対して kernel 起動コマンドを設定 (vectincrement)
  CHECK_RESULT(zeCommandListAppendLaunchKernel(hCmdList, hKernelVectInc, &dispatchVectInc, nullptr, 0, nullptr));
  // command list クローズ
  CHECK_RESULT(zeCommandListClose(hCmdList));

  // カーネル関数 vectadd への入力 a, b の値を設定
  int32_t *sharedInA = reinterpret_cast<int32_t *>(sharedA);
  int32_t *sharedInB = reinterpret_cast<int32_t *>(sharedB);
  for (int i=0; i < itemCount; i++) {
    sharedInA[i] = 100;
    sharedInB[i] = i;
  }

  // コマンド実行を要求
  CHECK_RESULT(zeCommandQueueExecuteCommandLists(hCmdQueue, 1, &hCmdList, nullptr));
  // コマンド実行が完了するのを待つ
  CHECK_RESULT(zeCommandQueueSynchronize(hCmdQueue, UINT64_MAX));

  // カーネル関数 vectadd からの出力、および vectincrement からの出力の値を確認
  int32_t *sharedOutC = reinterpret_cast<int32_t *>(sharedC);
  int32_t *sharedOutD = reinterpret_cast<int32_t *>(sharedD);
  int okCount = 0;
  int ngCount = 0;
  for (int i=0; i < itemCount; i++) {
    int32_t expectedValC = 100 + i;
    int32_t expectedValD = 100 + i + 1;
    if ((sharedOutC[i] == expectedValC) && (sharedOutD[i] == expectedValD)) {
      okCount++;
    } else {
      ngCount++;
      std::cout << "failed. idx:" << i <<
        ", sharedC:" << sharedOutC[i] << ", expectedC:" << expectedValC <<
        ", sharedD:" << sharedOutD[i] << ", expectedD:" << expectedValD << std::endl;
    }
  }
  std::cout << "compute finished. ok:" << okCount << " ng:" << ngCount << std::endl;

  // クリーンアップ
  zeKernelDestroy(hKernelVectAdd);
  zeKernelDestroy(hKernelVectInc);
  zeModuleDestroy(hModule);
  zeMemFree(hContext, sharedA);
  zeMemFree(hContext, sharedB);
  zeMemFree(hContext, sharedC);
  zeMemFree(hContext, sharedD);
  zeCommandListDestroy(hCmdList);
  zeCommandQueueDestroy(hCmdQueue);

#elif defined(SAMPLE_IMMEDIATE)

  // immediate command list 作成
  uint32_t cmdQueueGroupOrdinal = getCmdQueueGroup(hDevice);
  ze_command_list_handle_t hCmdListImm = createCmdListImmediate(hContext, hDevice, cmdQueueGroupOrdinal);

  // shared memory 割り当て
  uint32_t bufSize = sizeof(int32_t) * itemCount;
  void *sharedA = createSharedMem(hContext, hDevice, bufSize);
  void *sharedB = createSharedMem(hContext, hDevice, bufSize);
  void *sharedC = createSharedMem(hContext, hDevice, bufSize);

  // module 作成
  ze_module_handle_t hModule = createModule(hContext, hDevice, "vectadd2.spv");
  // kernel 作成
  ze_kernel_handle_t hKernelVectAdd = createKernel(hModule, "vectadd");

  // kernel argument 設定
  setKernelArgument(hKernelVectAdd, 0, bufSize, sharedA);
  setKernelArgument(hKernelVectAdd, 1, bufSize, sharedB);
  setKernelArgument(hKernelVectAdd, 2, bufSize, sharedC);

  // kernel dispatch 設定
  uint32_t itemsX = itemCount;
  uint32_t itemsY = 1;
  uint32_t itemsZ = 1;
  ze_group_count_t dispatchVectAdd = {};
  setKernelGroupSize(hKernelVectAdd, itemsX, itemsY, itemsZ, &dispatchVectAdd);

  // カーネル関数 vectadd への入力 a, b の値を設定
  int32_t *sharedInA = reinterpret_cast<int32_t *>(sharedA);
  int32_t *sharedInB = reinterpret_cast<int32_t *>(sharedB);
  for (int i=0; i < itemCount; i++) {
    sharedInA[i] = 100;
    sharedInB[i] = i;
  }

  // 実行
  CHECK_RESULT(zeCommandListAppendLaunchKernel(hCmdListImm, hKernelVectAdd, &dispatchVectAdd, nullptr, 0, nullptr));

  // カーネル関数 vectadd からの出力 c の値を確認
  int32_t *sharedOutC = reinterpret_cast<int32_t *>(sharedC);
  int okCount = 0;
  int ngCount = 0;
  for (int i=0; i < itemCount; i++) {
    int32_t expectedVal = 100 + i;
    if (sharedOutC[i] == expectedVal) {
      okCount++;
    } else {
      ngCount++;
      std::cout << "failed. result:" << sharedOutC[i] << ", expect:" << expectedVal << std::endl;
    }
  }
  std::cout << "compute finished. ok:" << okCount << " ng:" << ngCount << std::endl;

  // クリーンアップ
  zeKernelDestroy(hKernelVectAdd);
  zeModuleDestroy(hModule);
  zeMemFree(hContext, sharedA);
  zeMemFree(hContext, sharedB);
  zeMemFree(hContext, sharedC);
  zeCommandListDestroy(hCmdListImm);

#else
#error "set -DSAMPLE_ADDINC or -DSAMPLE_ADDINC_BARRIER or -DSAMPLE_IMMEDIATE"
#endif

  zeContextDestroy(hContext);

  return 0;
}