IntelコンパイラのOpenMPオフロード機能を使ってGPUで行列積を計算してみた

はじめに（経緯）

XLsoftが「OpenMP* プログラム GPU オフロード・チャレンジ・キャンペーン」なんてものをやっていて、WindowsでOpenMPのオフロード機能でIntelグラフィックスを叩けるということを知った。
普段からHPCな環境で並列計算をやっているHPCガチ勢の一人として、またGPUで汎用演算をすることについては色々とアレがソレでナニな民としてはとても気になるネタだったので、ちょっとだけ本気で試してみた。

実行環境

ThinkPad X1 Carbon 7th
Intel Core i7-10510U, Intel UHD グラフィックス内蔵
（電源設定で最高パフォーマンスにした状態、このあたりをいじるとどうなるのかまでは確認していない）

準備

oneAPIのBase ToolkitとHPC Toolkitをインストールする。
サーバが遅いのかなんなのか、Base ToolkitをOnline Installerでインストールをしようとしたら途中でエラー終了したので、両ToolkitのLocal Installerをインストールした。

後ほどVTuneを使うときにsampling driverが無効だというメッセージが出ていたので、これを参考に、cmd.exeを管理者権限で実行し、 C:\Program Files (x86)\Intel\oneAPI\vtune\2021.1-beta10\bin64>amplxe-sepreg.exe を適当に実行（-cオプションでインストール、-sオプションで確認）。
以下のように、なんとなくインストールされた感じ。

c:\Program Files (x86)\Intel\oneAPI\vtune\2021.1-beta10\bin64>amplxe-sepreg.exe -c
Checking platform...
Platform is genuine Intel: OK
Platform has SSE2: OK
Platform architecture: INTEL64
User has admin rights: OK
Drivers will be installed to C:\Windows\System32\Drivers\
Checking sepdrv5 driver path...OK
Checking sepdrv5 service...
Driver status: the sepdrv5 service is running
Checking sepdal driver path...OK
Checking sepdal service...
Driver status: the sepdal service is running
Checking socperf3 driver path...OK
Checking socperf3 service...
Driver status: the socperf3 service is running

c:\Program Files (x86)\Intel\oneAPI\vtune\2021.1-beta10\bin64>amplxe-sepreg.exe -s
Checking sepdrv5 driver path...OK
Checking sepdrv5 service...
Driver status: the sepdrv5 service is running
Checking sepdal driver path...OK
Checking sepdal service...
Driver status: the sepdal service is running
Checking socperf3 driver path...OK
Checking socperf3 service...
Driver status: the socperf3 service is running

インストールしたToolkitのバージョンはこんな感じ。

icc>icx -v
Intel(R) oneAPI DPC++/C++ Compiler Pro for applications running on Intel(R) 64, Version 2021.1 Beta Build 20201005
Copyright (C) 1985-2020 Intel Corporation. All rights reserved.

Intel(R) oneAPI DPC++ Compiler Pro 2021.1 (2020.8.0.1005)
Target: x86_64-pc-windows-msvc
Thread model: posix
InstalledDir: C:\PROGRA~2\Intel\oneAPI\compiler\latest\windows\bin

実験プログラム

ベタではあるが、単純な三重ループの行列積を用意し、最外ループを並列化。
ちなみに配列をmallocによる動的確保にしたらプログラムが正常終了しなかったり計算結果が正しく得られなくなったりした。map指示文に長さを指定してもダメだった。
このあたりの調査はまた気が向いたらということで……。

  float a[n][n], b[n][n], c[n][n];
  t_begin = omp_get_wtime();
#pragma omp target data map(to:a,b) map(tofrom:c)
#pragma omp target teams distribute parallel for private(j,k)
  for(i=0; i<n; i++){
    for(j=0; j<n; j++){
      for(k=0; k<n; k++){
        c[i][j] += a[i][k] * b[k][j];
      }
    }
  }
  t_end = omp_get_wtime();
  printf("time %f\n", (double)(t_end - t_begin));

コンパイル

スタートメニューから「Intel oneAPI command prompt for Intel 64 for Visual Studio 2017」を起動。ディレクトリを移動し、

icx matmul_ofl_static1.c /Qiopenmp /Qopenmp-targets:spir64 /nologo

でコンパイル。以下のようなwarningが出たが、解消方法は不明。

matmul_ofl_static1-b5aa07.obj : warning LNK4078: 複数の '__CLANG_OFFLOAD_BUNDLE__openmp-s' セクションが見つかりました。 これらは異なる属性 (40500040) を持っています。

実行時にset OMP_TARGET_OFFLOAD=DISABLEDと環境変数をセットしておくとCPUで実行、set OMP_TARGET_OFFLOAD=MANDATORYとしておくとGPUで実行、ついでにset LIBOMPTARGET_PROFILE=Tをセットしておくと簡単なプロファイル情報が得られるようだ。

問題サイズ（行列の行数と列数）を100にした状態で実行すると以下のような結果が得られた。
（現在のディレクトリの表示は省略。）

>set LIBOMPTARGET_PROFILE=T

>set OMP_TARGET_OFFLOAD=DISABLED

>matmul_ofl_static1
n = 100
time 0.001609

>set OMP_TARGET_OFFLOAD=MANDATORY

>matmul_ofl_static1.exe
n = 100
Target LEVEL0 RTL --> Warning: Invalid kernel timestamp bit width (0). Long-running kernels may report incorrect device time.
time 5.878747
LIBOMPTARGET_PROFILE for OMP DEVICE(0) Intel(R) Gen9, Thread 12788
-- Name: Host Time (msec), Device Time (msec)
-- DataAlloc: 0.166800, 0.166800
-- DataRead: 0.002000, 0.002000
-- DataWrite: 0.179900, 0.179900
-- Kernel#__omp_offloading_821ec21a_199401__Z4main_l35: 1.838900, 1.328498
-- ModuleBuild: 5786.827200, 5786.827200
-- Total: 5789.014800, 5788.504398

結果としては、GPUを使うと遅い。何度か実行してみたが、だいたい6秒くらい。CPUだと0.001秒のレベル。
ただ、

-- Kernel#__omp_offloading_821ec21a_199401__Z4main_l35: 1.838900, 1.328498

の部分がGPUカーネルっぽい気がするので、実際に計算している時間は1ミリ秒レベル、CPUとあまり変わらないのかもしれない。

そこで、問題サイズを1000に大きくして再測定してみた。

icc>set OMP_TARGET_OFFLOAD=DISABLED

icc>matmul_ofl_static1.exe
n = 1000
time 1.261088

icc>set OMP_TARGET_OFFLOAD=MANDATORY

icc>matmul_ofl_static1.exe
n = 1000
Target LEVEL0 RTL --> Warning: Invalid kernel timestamp bit width (0). Long-running kernels may report incorrect device time.
time 7.539387
LIBOMPTARGET_PROFILE for OMP DEVICE(0) Intel(R) Gen9, Thread 14048
-- Name: Host Time (msec), Device Time (msec)
-- DataAlloc: 0.441400, 0.441400
-- DataRead: 0.625600, 0.625600
-- DataWrite: 18.447000, 18.447000
-- Kernel#__omp_offloading_821ec21a_199401__Z4main_l35: 1320.353900, 1313.705408
-- ModuleBuild: 6106.610200, 6106.610200
-- Total: 7446.478100, 7439.829608```

やはり主要な計算部分だけ見ると1.2秒と1.3秒で同じくらいの性能が出ているようだ。
ここまで性能が似ていると、本当にGPUで動いているの？と逆に不安になるわけであるが、メッセージを見る限りは動いているよねえ……。

VTuneで見てみた

スタートメニューから「Intel VTune Profiler 2021.1」を起動。赤枠部分を適当に設定して実行。（引数に1000を与えているが、これはプログラム中では意味がなくて、実際には静的確保してある。）

DLLの読み込みエラーが出るのだが、設定する場所がわからないのでとりあえず放置。

いくつか警告は出るが、とりあえずプロファイリング自体は完了。（debugging informationが見つからない件、コンパイル時に/debugオプションを付けるなどしてみたが解決しなかった。）

全体サマリや

内訳を表示できた

が、これ結局どこを見ればいいんだろう……？
1秒くらいかかったと思われるGPUカーネルの場所も良くわからないという状況。例によってVTuneの測定箇所を限定する関数とかを入れないとダメかもね？

とりあえずの結論

なんとなくGPUが動いている気がするという状況にはなったが、それ以上ではないという感じ。行列積くらいパッとCPUより速くなってくれないと、わざわざ使おうとする人は……。
あと、やっぱりWindowsはコマンドプロンプトの使いにくさが最悪。

とりあえず、今回はここまで。次回があるかは不明。
もちろん、技術としては面白いネタなので、今後の発展に期待。