Intel内蔵GPUをoneAPI Level Zeroで直接叩いてみる（環境構築編）

はじめに

手のひらサイズのミニPCでよく使われている Intel N100 などの CPU は、基本的に Intel UHD Graphics などの内蔵グラフィックス機能を持っています。

この Intel CPU 内蔵グラフィックス機能を GPGPU として活用できないだろうかと考えて調べたところ、oneAPI の一部として提供されている Level Zero API (oneAPI/L0) が使えそうだと分かりました。

ただ Level Zero API は英語の仕様書¹は公開されているものの、実際の使用例やサンプルを探してもなかなか見つからず、導入するのに苦労しました。せっかくなので、Level Zero API の導入方法について記事として残したいと思います。

セットアップ環境は Ubuntu 22.04 LTS を使用します。

oneAPI Level Zero

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

oneAPI はインテルが提唱するマルチアーキテクチャプログラミングモデルであり、SYCL/DPC++ などの言語を用いて、単一のソースコードで CPU, GPU, FPGA が混在する環境に対応可能といった特徴を持ちます。 ただし、ここでは忘れてください。

本記事では、oneAPI の一部である Level Zero API を単独で使います。
Level Zero API を単独で使う場合、OpenCL のようにカーネルプログラムとホストプログラムをそれぞれ用意する形となります。

1. カーネルプログラム
   GPU 上で動作させるプログラム。インテルが提供するツール (ocloc) を用いて、OpenCL C 形式で書かれたコードをカーネルプログラム用バイナリ形式 (SPIR-V) に変換することができます。
2. ホストプログラム
   GPU デバイスの初期化や共有メモリの確保、カーネルプログラムのロードや実行などの制御を行うプログラム。一般的な C/C++ を使ってコードを書きます。

OpenCL と何が違うの？

OpenCL と比較して、Level Zero API を使用するメリット・デメリットを示します。

メリット

多機能

例えば Level Zero API では、以下のような機能が用意されています。

• コマンドの即時実行 (Immediate Command Lists)
• メモリアドレス範囲を指定したバリア (Memory Barriers)
• 推奨されるグループサイズの自動算出 (zeKernelSuggestGroupSize)

共有メモリについても、明示的な同期処理なしでアクセスできるため便利です。
（ OpenCL でも Fine-grained buffer がありますが、今回使用した環境では非対応）

解析ツール (Intel VTune Profiler) のサポート

Intel VTune Profiler と相性が良く、GPU 処理の解析がやりやすいと感じました。
（ OpenCL でも解析できるようですが、Linux カーネルの再構築などが必要で複雑）

デメリット

対応ハードウェアが少ない

FPGA や AI 用アクセラレータなどを除いた一般向け CPU/GPU に限定して考えると、
Level Zero は Intel GPU のみに対応しており、現状 NVIDIA, AMD GPU は非対応のようです。

また Intel GPU であっても、古い世代の Atom 系 CPU（Apollo Lake など）内蔵グラフィックスの場合は対応していません。

環境構築

ここでは、以下の PC を使用します。

項目	内容
CPU	Intel N100 (Alder Lake-N)
GPU	Intel UHD Graphics (Intel N100 CPU 内蔵グラフィックス)
OS	Ubuntu 22.04 LTS
Kernel	HWEカーネル 5.19

Ubuntu 22.04 LTS 初期リリースのカーネル 5.15 は Intel N100 に非対応のようです。uname -r でカーネルバージョンを確認できます。

Intel GPU ドライバ導入 (Ubuntu 22.04 LTS)

パッケージレポジトリの追加

wget -qO - https://repositories.intel.com/graphics/intel-graphics.key | sudo gpg --dearmor --output /usr/share/keyrings/intel-graphics.gpg
echo 'deb [arch=amd64,i386 signed-by=/usr/share/keyrings/intel-graphics.gpg] https://repositories.intel.com/graphics/ubuntu jammy arc' | sudo tee  /etc/apt/sources.list.d/intel.gpu.jammy.list
sudo apt-get -y update

必要なパッケージのインストール

sudo apt-get -y install intel-opencl-icd intel-level-zero-gpu \
  intel-media-va-driver-non-free libmfx1 libmfxgen1 libvpl2 \
  libegl-mesa0 libegl1-mesa libegl1-mesa-dev libgbm1 libgl1-mesa-dev libgl1-mesa-dri \
  libglapi-mesa libgles2-mesa-dev libglx-mesa0 libigdgmm12 libxatracker2 mesa-va-drivers \
  mesa-vdpau-drivers mesa-vulkan-drivers va-driver-all

sudo apt-get -y install libigc-dev intel-igc-cm libigdfcl-dev libigfxcmrt-dev

ここでPCを再起動します。再起動後に、デバイスファイル /dev/dri/renderD128 が存在することを確認してください。

デバイスファイルのグループを確認

stat -c "%G" /dev/dri/render*    # "render" が返るはず

ユーザの所属するグループを確認

groups ${USER}    # ユーザが所属するグループ一覧が返る

ユーザの所属するグループに render が含まれていない場合、render グループに追加

sudo gpasswd -a ${USER} render
newgrp render

GPU ドライバが正しくインストールされていることを確認

sudo apt-get install clinfo
clinfo

ここで、以下のようなメッセージが表示されることを確認します。

Number of platforms                               1
  Platform Name                                   Intel(R) OpenCL Graphics
  Platform Vendor                                 Intel(R) Corporation
  Platform Version                                OpenCL 3.0
  Platform Profile                                FULL_PROFILE
  (省略)

Number of devices                                 1
  Device Name                                     Intel(R) UHD Graphics
  Device Vendor                                   Intel(R) Corporation
  Device Vendor ID                                0x8086
  Device Version                                  OpenCL 3.0 NEO
  Driver Version                                  23.17.26241.33
  (省略)

Intel GPU ドライバのアップデート

パッケージレポジトリからインストールした 23.17.26241.33 ですが、Intel VTune Profiler のセルフチェックで failed が発生しました。どうやら少し古いため Level Zero の一部機能 (zeCommandListHostSynchronize など) に非対応のようです。

ここでは GitHub から 23.35.27191.9 のパッケージを取得してアップデートします。
※ 最新ではなく 23.35.27191.9 を使用した理由は、Ubuntu 20.04 や ROS noetic 対応を考慮したためです。（詳細は後述）

mkdir ~/neo && cd ~/neo
wget https://github.com/intel/intel-graphics-compiler/releases/download/igc-1.0.15136.4/intel-igc-core_1.0.15136.4_amd64.deb
wget https://github.com/intel/intel-graphics-compiler/releases/download/igc-1.0.15136.4/intel-igc-opencl_1.0.15136.4_amd64.deb
wget https://github.com/intel/compute-runtime/releases/download/23.35.27191.9/intel-level-zero-gpu_1.3.27191.9_amd64.deb
wget https://github.com/intel/compute-runtime/releases/download/23.35.27191.9/intel-opencl-icd_23.35.27191.9_amd64.deb
wget https://github.com/intel/compute-runtime/releases/download/23.35.27191.9/libigdgmm12_22.3.11.ci17747749_amd64.deb
sudo dpkg -i *.deb

clinfo でドライバのバージョンが 23.35.27191.9 に更新されていることを確認します。

Level Zero Loader

パッケージレポジトリから Level Zero Loader をインストールすることもできますが、バージョン v1.11 と少し古いです。

ドライバ 23.35.27191.9 が対応する Level Zero Loader v1.14 を使いたいので、GitHub からソースコードを取得して自前でビルドします。

ビルド用パッケージのインストール

sudo apt-get -y install build-essential cmake

Level Zero v1.14 のダウンロード、ビルド、インストール

mkdir ~/levelzero && cd ~/levelzero
wget https://github.com/oneapi-src/level-zero/archive/refs/tags/v1.14.0.tar.gz
tar zxvf v1.14.0.tar.gz
cd level-zero-1.14.0/
mkdir build && cd build
cmake ..
cmake --build . --config Release --target package
sudo dpkg -i *.deb

Level Zero Loader に付属するサンプルを利用して動作確認します。
上記のビルド時に build/bin フォルダ内に生成される zello_world を実行します。

cd ~/levelzero/level-zero-1.14.0/build/bin
./zello_world

ここで、以下のようなメッセージが表示されることを確認します。

Driver initialized.
zelLoaderGetVersions number of components found: 1
Version 0
Name: loader
Major: 1
Minor: 14
(省略)

Congratulations, the device completed execution!

サンプルコードのビルド、実行

GPU を用いた演算を、実際に動かしてみましょう。
今回使ったソースコードを 付録 に載せています。

ファイル構造

lzsample
├── vectadd.cl
└── lzsample.cpp

カーネルプログラム

vectadd.cl は、OpenCL C 形式で書かれたベクトル加算のカーネルプログラムです。
下記のコマンドにより、SPIR-V 形式ファイル vectadd.spv に変換します。

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

ここで "adl-n" は CPU の種別 (Alder Lake-N) を示します。
ocloc compile --help を確認して、使用する CPU に合わせてください。

<device_type> can be: bdw, skl, kbl, cfl, apl, bxt, glk, whl, aml, cml, icllp, lkf, ehl, jsl, tgllp, rkl, adl-s, adl-p, adl-n, dg1, acm-g10, ats-m150, dg2-g10, acm-g11, ats-m75, dg2-g11, acm-g12, dg2-g12, pvc-sdv, pvc, mtl-m, mtl-s, mtl-p, gen11, gen12lp, gen8, gen9, xe, xe-hpc, xe-hpg, xe-lpg, ip version  or hexadecimal value with 0x prefix

ホストプログラム

lzsample.cpp は、カーネルプログラムのロードや実行、共有メモリの確保などを行います。
ここでは、引数で指定された回数（デフォルト：1000000）だけ、ベクトル加算のカーネルプログラムを実行します。

ビルド、実行

gcc lzsample.cpp -o lzsample -lze_loader -lstdc++
./lzsample 1000000

共有ライブラリに関するエラーが出る場合、LD_LIBRARY_PATH を確認してください。

(例) LD_LIBRARY_PATH の追加

export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/lib/

ここで、以下のメッセージが表示されることを確認します。

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

GPU 使用率の確認

GPU を用いて演算されていることを確認するために、プログラム実行中の GPU 使用率を見てみましょう。

Intel GPU ツールのインストール、実行

sudo apt-get -y install intel-gpu-tools
sudo intel_gpu_top -s 100    # 100ms毎に更新

別のターミナルを開き、lzsample を実行します。ここで ./lzsample 10000000 のように回数を増やすと分かりやすくなります。

このとき、「Render/3D」の BUSY % が上昇することを確認します。

docker 環境で動かす場合

詳細は省きますが、docker を用いて Level Zero API の開発環境を構築することもできます。
docker を使用する場合、以下のようになります。

• ホスト環境に Intel GPU ドライバを導入
• コンテナ環境に Intel GPU ドライバ、Level Zero Loader、およびビルド用ツールを導入

コンテナ起動時のオプション
例：docker run -it --device=/dev/dri --name u2004tmp ubuntu:20.04 /bin/bash

• Intel GPU を使用するために --device=/dev/dri が必要
• --gpus は不要。これは NVIDIA 用のため付けるとエラーとなる

ホスト環境とコンテナ環境の Intel GPU ドライバは、バージョンを揃える必要があります。

OS のバージョンも基本的には揃えますが、以下のようなバージョンが異なる環境でも Level Zero API が動作することを確認できました。

• ホスト環境：Ubuntu 22.04 LTS
• コンテナ環境：Ubuntu 20.04 LTS あるいは ROS noetic

Intel GPU ドライバのバージョンは、ホスト環境とコンテナ環境の両方にインストール可能な 23.35.27191.9 を選択しました。

備考：Intel GPU ドライバ導入 (Ubuntu 20.04 LTS)

docker コンテナ内の root ユーザでの操作を想定しています。

Ubuntu 20.04 用のパッケージレポジトリ追加

apt-get -y install gpg-agent wget

wget -qO - https://repositories.intel.com/graphics/intel-graphics.key | gpg --dearmor --output /usr/share/keyrings/intel-graphics.gpg
echo 'deb [arch=amd64 signed-by=/usr/share/keyrings/intel-graphics.gpg] https://repositories.intel.com/graphics/ubuntu focal-legacy main' | tee  /etc/apt/sources.list.d/intel.gpu.focal-legacy.list
apt-get -y update

必要なパッケージのインストール

apt-get -y install intel-opencl-icd intel-level-zero-gpu \
  intel-media-va-driver-non-free libmfx1 libmfxgen1 libvpl2 \
  libegl-mesa0 libegl1-mesa libegl1-mesa-dev libgbm1 libgl1-mesa-dev libgl1-mesa-dri \
  libglapi-mesa libgles2-mesa-dev libglx-mesa0 libigdgmm11 libxatracker2 mesa-va-drivers \
  mesa-vdpau-drivers mesa-vulkan-drivers va-driver-all

apt-get -y install libigc-dev intel-igc-cm libigdfcl-dev libigfxcmrt-dev

バージョン (22.43.24595.35) が古いため、23.35.27191.9 にアップデートします。

mkdir ~/neo && cd ~/neo
wget https://github.com/intel/intel-graphics-compiler/releases/download/igc-1.0.15136.4/intel-igc-core_1.0.15136.4_amd64.deb
wget https://github.com/intel/intel-graphics-compiler/releases/download/igc-1.0.15136.4/intel-igc-opencl_1.0.15136.4_amd64.deb
wget https://github.com/intel/compute-runtime/releases/download/23.35.27191.9/intel-level-zero-gpu_1.3.27191.9_amd64.deb
wget https://github.com/intel/compute-runtime/releases/download/23.35.27191.9/intel-opencl-icd_23.35.27191.9_amd64.deb
wget https://github.com/intel/compute-runtime/releases/download/23.35.27191.9/libigdgmm12_22.3.11.ci17747749_amd64.deb
dpkg -i *.deb

おわりに

本記事では Level Zero API の導入方法について紹介しました。
Intel N100 CPU 内蔵グラフィックスを用いて、Ubuntu 22.04 LTS 環境で GPGPU を動かすことができました。また docker を用いて Ubuntu 20.04 LTS や ROS noetic での動作も実現できました。

次回は Level Zero API の具体的な使用方法、プログラムの書き方について説明していきたいと考えています。

参考文献

付録

vectadd.cl

vectadd.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];
}

lzsample.cpp

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

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

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

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

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

  // 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);

  // module 作成
  ze_module_handle_t hModule = createModule(hContext, hDevice, "vectadd.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);

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

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

  // 実行
  executeCmdList(hCmdQueue, hCmdList);

  // カーネル関数 vectadd からの出力 c の値を確認
  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(hCmdList);
  zeCommandQueueDestroy(hCmdQueue);
  zeContextDestroy(hContext);

  return 0;
}

1. https://spec.oneapi.io/level-zero/latest/index.html ↩