インテル® コンパイラーを使用した
OpenMP* による並列プログラミング
インテル® C++/Fortran コンパイラーのバージョン 19.1 を使用して GPU オフロードに備えましょう

セッション 6: oneAPI への第一歩
IA Software User Society (iSUS)
編集長 すがわら きよふみ

このセッションの目的
明示的な並列プログラミング手法として注目されてきた OpenMP* による並列プロ
グラミングに加え、インテル® コンパイラーがサポートする OpenMP* 4.0 と 4.5 の機能を使用したベクトル・プログラミングとオフロード・プログラミングの概要をリフレッシュし、インテル® コンパイラー V19.1 でサポートされる OpenMP* 5.0 の機能と実装を紹介します。さらに新たなアクセラレーター・デバイスへのオフロードについて考えますセッションの対象者
すでに OpenMP* でマルチスレッド・プログラミングを開発し、4.0 以降でサポートされる新たなベクトル化とオフロードを導入し、アプリケーションのパフォーマンス
向上を計画する開発者© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 2

セッションリスト
セッション

説明

はじめに
(13:30 – 14:30)

異なるバージョンのインテル® コンパイラーや異なるコンパイラー間で OpenMP* を使用する
注意点や制限について説明します

OpenMP* のタスク機能(14:30 – 15:30)

OpenMP* 3.1 で追加されたタスク機能が 4.0 から 4.5 でどのように進化したかを例を使用して説明し、最新の OpenMP* 5.0 で強化された新機能を紹介します

OpenMP* の SIMD 機能(13:40 – 14:30)

OpenMP* のスレッド化機能を使用してプログラマーがマルチスレッドの動作をプログラミン
グしたように、OpenMP* 4.0 からは omp simd を使用してプログラマーが明示的にベクトル化もできるようになりました。OpenMP* simd に関連する機能を 4.0 から 5.0 までの進化を追って紹介します

OpenMP* のオフロード機能(14:30 – 15:30)

OpenMP* 4.0 で追加されたオフロード機能を利用することで、これまで共有メモリー型並列処理に加え分散メモリー型の並列処理を表現できるようになりました。このセッションでは、注目されるヘテロジニアス・プログラミング環境での OpenMP* オフロード機能について説明
します

OpenMP* 5.0 の注目する機能 セッション2、3、4 でカバーされなかった OpenMP* 5.0 のそのほかの機能について紹介しま(13:30 – 14:15)
すインテル® C++/Fortran コン
パイラーのバージョン 19.1 を使用して GPU オフロードに備えましょう (14:15 – 15:30)

oneAPI 向けのデータ並列 C++ (DPC++) へ移行する前に、現行のインテル® C++/Fortran コ
ンパイラー V19.1 やインテル® oneAPI HPC ツールキットに含まれるベータ版インテル®
C++/Fortran コンパイラー 2021 を使用して簡単にインテル® グラフィックスへのオフロードを行うソフトウェアを開発および検証方法を紹介します© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 3

内容
はじめに (OpenMP* が必要とされる背景) と概要(OpenMP* とは、歴史、各バージョンの機能概要)◼ OpenMP* の各バージョンの機能(4.0、4.5 および 5.0 の注目される新機能)◼ 次世代インテル® コンパイラー (nextgen) の機能◼

OpenMP* 5.0 API シンタックス・クイック・リファレンス・カードの日本語訳を公開しました:
https://www.isus.jp/products/c-compilers/openmp-ref-5-0-0519-released/© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 4

プログラミングの
課題
⚫

⚫

⚫

⚫

さまざまなデータ構造を使用するハードウェ
アのセットが存在します

CPU

GPU

AI

FPGA

スカラー

ベクトル

行列

空間

残念ながら、それらに対する共通のプログラ
ミング言語や API はありません
また、それぞれのプラットフォームにおいてツールのサポートに一貫性もありません
プラットフォームごとに個別のソフトウェア
投資/習得が必要です

SVMS© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 5

OpenMP*
OpenCL*
OpenACC*…

プログラミングの挑戦#pragma omp target data device(0) map(alloc:tmp[:N])map(to:input[:N)) map(from:res){#pragma omp target device(0)#pragma omp parallel forfor (i=0; i<N; i++)tmp[i] = some_computation(input[i], i);

CPU

GPU

AI

FPGA

スカラー

ベクトル

行列

空間

update_input_array_on_the_host(input);#pragma omp target update device(0) to(input[:N])#pragma omp target device(0)#pragma omp parallel for reduction(+:res)for (i=0; i<N; i++)res += final_computation(input[i], tmp[i], i)}SVMS© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 6

準備するもの (ここでは Windows* を使用)
コンパイラー:◼ インテル® C++ または Fortran コンパイラー・バージョン 19.1 Update2
またはインテル® oneAPI ベース・ツールキット beta08 以降◼ インテル® oneAPI HPC ツールキット beta08 以降◼

グラフィックス・ドライバー:◼ インテル® oneAPI ベース・ツールキット用グラフィックス・ドライバー
https://software.intel.com/content/www/us/en/develop/articles/intel-oneapi-graphics-drivernow-available.html (英語)© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 7

次世代インテル® コンパイラーとは
◼

◼

◼

◼

ICC Classic と ICC NextGen (icx) は、各コンパイラー・ドライバー (icc、icpc、icl)
でも同じ呼び方をします
ICC NextGen は、コア C++ コンパイラーであるとともに、インテル® DPC++ コン
パイラーと新しい dpcpp ドライバーのベース・コンパイラーでもあります
ICC NextGen は、インテルの「次世代」コンパイラーです。ICC NextGen は ICC
Classic の単なる代替品ではないことに注意してください
ICC Classic から ICC NextGen への移行を可能な限りスムーズにするため、かなり工夫されていますが、 開発者が既存のアプリケーションを ICC Classic から ICC
NextGen へ移行してチューニングするには、ある程度の労力が必要になりますインテル® C++ コンパイラーの次世代コード・ジェネレーターのドキュメント:https://www.isus.jp/products/c-compilers/early-documentation-for-c-compiler-based-on-llvm/© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 8

この資料では Nextgen コンパイラーのOpenMP* 機能に注目します
ICC NextGen (icx) の OpenMP* オプション:/Qiopenmp / -fiopenmp は、OpenMP* parallel 構造と SIMD プラグマ/ディレク
ティブを認識してコンパイルし、 インテルの OpenMP* ランタイム・ライブラリーを使用します◼ OpenMP* 4.5/5.0 の target プラグマ/ディレクティブを使用する場合、/Qopenmp-targets=spir64 / -fopenmp-targets=spir64 を指定する必要があります◼ OpenMP* 4.5/5.0 の target ディレクティブは、インテル® oneAPI HPC ツール
キットに含まれる ICC NextGen コンパイラーによってのみ認識されます
上記の 2 つのコンパイラー・オプションを同時に指定します:• icc -qnextgen -fiopenmp -fopenmp-targets=spir64• icl /Qnextgen /Qiopenmp /Qopenmp-targets=spir64◼

© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 9

インテル® コンパイラー V19.1 のNextgen 機能はプレビュー機能です_OPENMP マクロが返す値が異なります© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 10

GPU オフロードについて (旧バージョンの機能)
bool Sobel::execute_offload(){
int w = COLOR_CHANNEL_NUM * image_width;
float *outp = this->output;
float *img = this->image;int iw = image_width;
int ih = image_height;#pragma omp target map(to: ih, iw, w) ¥map(tofrom: img[0:iw*ih*COLOR_CHANNEL_NUM], ¥outp[0:iw*ih*COLOR_CHANNEL_NUM])#pragma omp parallel for collapse(2)for (int i = 1; I < ih – 1; i++) {for (int k = COLOR_CHANNEL_NUM; k < (iw - 1) * COLOR_CHANNEL_NUM; k++) {float gx = 1 * img[k + (i - 1) * w -1 * 4]
+ 2 * img[k + (i - 1) * w +0 * 4]
+ 1 * img[k + (i - 1) * w +1 * 4]
- 1 * img[k + (i + 1) * w -1 * 4]
- 2 * img[k + (i + 1) * w +0 * 4]
- 1 * img[k + (i + 1) * w +1 * 4];
float gy = 1 * img[k + (i - 1) * w -1 * 4]
- 1 * img[k + (i - 1) * w +1 * 4]
+ 2 * img[k + (i + 0) * w -1 * 4]
- 2 * img[k + (i + 0) * w +1 * 4]
+ 1 * img[k + (i + 1) * w -1 * 4]
- 1 * img[k + (i + 1) * w +1 * 4];outp[i * w + k] = sqrtf(gx * gx + gy * gy) / 2.0;}}
return true;}© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

◼

◼

インテル® C++/Fortran
コンパイラーの V18.0 までは、
インテル® グラフィックスへのオフロードがサポートされていました/Qopenmp-offload:gfx(Windows*)、-qopenmp-offload=gfx(Linux*) オプションでターゲッ
トデバイスが指定できました| 11

target ディレクティブで depend 節を使用してGPU へ非同期オフロードする例// arr1 を初期化 - ターゲットへオフロード#pragma omp target map(from: arr1[0:SIZE]) depend(out:arr1) nowait#pragma omp parallel forfor (int i = 0; i < SIZE; i++) { arr1[i] += i; }// arr2 の初期化
#pragma omp task depend(out:arr2)#pragma omp parallel forfor (int i = 0; i < SIZE; i++) { arr2[i] += -i; }// ターゲット上で中間結果を計算#pragma omp target ¥map(to: arr1[0:SIZE], arr2[0:SIZE]) ¥map(from:arr3[0:SIZE]) ¥nowait depend(in:arr1, arr2)#pragma omp parallel forfor (int i = 0; i < SIZE; i++) { arr3[i] = arr1[i] + arr2[i]; }#pragma omp taskwait#pragma omp parallel forfor (int i = 0; i < SIZE; i++) { res[i] += arr3[i]; }© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

親タスク
(ホスト)
子タスク 1(GPU)
スポーン

子タスク 3
(ホスト)
依存関係

arr1

依存関係

arr2

子タスク 2
(GPU)

バリア

| 12

Classic コンパイラーで並列化#pragma omp parallel forfor (int i = 0; i < MAX; i++)for (int k = 0; k < MAX; k++)for (int j = 0; j < MAX; j++)C[i][j] += A[i][k] * B[k][j];© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 13

GPU オフロード向けにコードを変更#pragma omp target teams distribute parallel for ¥map(to: A, B) map(tofrom: C) thread_limit(128)for (int i = 0; i < MAX; i++)for (int k = 0; k < MAX; k++)for (int j = 0; j < MAX; j++)C[i][j] += A[i][k] * B[k][j];Classic コンパイラーは、このソースをコンパイルできません© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 14

Nextgen モードでコンパイル

これは、どこで実行されているのか?

© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 15

Nextgen モードで /Qiopenmp コンパイル

これは、どこで実行されているのか?

© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 16

/Qopenmp-targets:spir64 を追加

このパスは標準のビルド環境向けには設定されていません© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 17

ビルドには 2 つのオプションが必要/Qiopenmp (Windows*)、-qiopenmp (Linux*)◼ /Qopenmp-targets:spir64 (Windows*)、-qopenmp-targets=spir64 (Linux*)◼ omptarget.lib へのパス設定が必要◼

生成されるバイナリーは、CPU バイナリーと GPU バイナリーが結合された FAT バイナリーとなります© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 18

GPU で実行するには
OMP_TARGET_OFFLOAD={"MANDATORY" | "DISABLED" | "DEFAULT" }MANDATORY: GPU またはアクセラレーターでターゲット領域のコードを実行
DISABLED: CPU でターゲット領域のコードを実行
DEFAULT: デバイスが利用可能な場合に GPU で実行され、利用できない場合は CPU にターゲット領域のコードをフォールバックして実行

これは!?!?
© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 19

他のデバイスで実行するには (オプション)◼ LIBOMPTARGET_DEVICETYPE に cpu、gpu、accelerator を指定できますが、他のデバイスを使用するにはデバイスをサポートするライブラリーが必要になります◼ OMP_TARGET_OFFLOAD とも関連します© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 20

オフロードのプロファイル◼

オフロードされたかテストするには、LIBOMPTARGET_PROFILE 環境変数を設定
set LIBOMPTARGET_PROFILE=T◼

この出力から、GPU へのデータの送信時間 (Data-Write) と GPU からのデータの受信時間 (DataRead) が分かります。オフロード TARGET 領域 (EXEC) ― main() の ５６ 行目があり、このカーネルで費やされた時間が示されています© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 21

オフロードのデバッグ◼

◼

前述のプロファイルが表示されない場合、OpenMP* ランタイムが GPU ドライバーデバイスを検出できなかったことを意味します

さらに詳しくデバッグするには、環境変数 LIBOMPTARGET_DEBUG=5 を設定してアプリケーショ
ンを再度実行します
Libomptarget --> TargetOffloadPolicy = DEFAULT
Libomptarget --> Initialized OMPT
Libomptarget --> Loading RTLs...
Libomptarget --> Loading library 'omptarget.rtl.level0.dll'...
Target LEVEL0 RTL --> Target device type is set to GPU
Target LEVEL0 RTL --> omp_get_thread_limit() returned 2147483647
Target LEVEL0 RTL --> omp_get_max_teams() returned 0
Libomptarget --> Successfully loaded library 'omptarget.rtl.level0.dll'!
Libomptarget --> Optional interface: __tgt_rtl_data_submit_nowait
Libomptarget --> Optional interface: __tgt_rtl_data_retrieve_nowait
Libomptarget --> Optional interface: __tgt_rtl_manifest_data_for_region
Libomptarget --> Optional interface: __tgt_rtl_data_alloc_base
Libomptarget --> Optional interface: __tgt_rtl_create_buffer
Libomptarget --> Optional interface: __tgt_rtl_release_buffer
Libomptarget --> Optional interface: __tgt_rtl_run_target_team_nd_region
Libomptarget --> Optional interface: __tgt_rtl_run_target_region_nowait
Libomptarget --> Optional interface: __tgt_rtl_run_target_team_region_nowait
Libomptarget --> Optional interface: __tgt_rtl_run_target_team_nd_region_nowait
Libomptarget --> Optional interface: __tgt_rtl_create_offload_queue
Libomptarget --> Optional interface: __tgt_rtl_release_offload_queue© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 22

インテル® VTune™ プロファイラー 2020 で確認© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 23

インテル® VTune™ プロファイラー 2020 のGPU オフロード解析 (プレビュー)◼

ここで説明した手順は、
インテル® oneAPI HPC ツールキットに含まれる
ベータ版インテル® C++/Fortran コ
ンパイラー 2021 でも利用できます

ここでの手順は動作を確認したのみであり、実際にはより高速に動作させるため最適化が必要です© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 24

簡単な最適化 (1)#pragma omp data map(to: A, B) map(tofrom: C){…………
#pragma omp target teams distribute parallel for thread_limit(128)for (int i = 0; i < MAX; i++)for (int k = 0; k < MAX; k++)for (int j = 0; j < MAX; j++)C[i][j] += A[i][k] * B[k][j];}© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 25

簡単な最適化 (2)#pragma omp data map(to: A, B) map(tofrom: C){…………

#pragma omp target teams distribute thread_limit(128)for (int i = 0; i < MAX; i++)for (int k = 0; k < MAX; k++)#pragma omp parallel forfor (int j = 0; j < MAX; j++)C[i][j] += A[i][k] * B[k][j];}© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 26

DPC++ への移行

© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 27

oneAPI 環境への移行ステップ
インテル® oneAPI ベース・ツールキット
インテル® oneAPI ベース・
ツールキットを使用するアクセラレーター向けのプログラ
ミングでは、DPC++ コンパイ
ラーと新しい SYCL* ベースのコードを使用する必要があります◼ Fortran ユーザーは?

インテル® oneAPI ベース・ツールキット

ダイレクト・
プログラミング

API ベースのプログラミング

解析ツール

インテル® oneAPI DPC++
コンパイラーインテル® oneAPI DPC++ ライブラリー

インテル® VTune™ プロファイラー

インテル® oneAPI マス・カーネル・ライブラリーインテル® Advisor

インテル® DPC++
互換性ツール

Python* 向けインテル®
ディストリビューション
[オプション] oneAPI ベース・ツール
キット用インテル® FPGA アドオン

◼ OpenMP* ユーザーは?

インテル® oneAPI データ・
アナリティクス・ライブラリー

GDB*

インテル® oneAPI スレッディング・
ビルディング・ブロック
インテル® oneAPI ビデオ・
プロセシング・ライブラリー
インテル® oneAPI コレクティブ・
コミュニケーション・ライブラリー
インテル® oneAPI ディープ・ニューラル・
ネットワーク・ライブラリー
インテル® インテグレーテッド・
パフォーマンス・プリミティブ© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 28

最適化されたアプリケーション
インテル® oneAPI HPC ツールキット
ダイレクト・プログラミング
インテル® oneAPI DPC++
コンパイラーインテル® DPC++
互換性ツール
Python* 向けインテル®
ディストリビューション
[オプション] oneAPI ベース・
ツールキット用インテル®
FPGA アドオン

インテル® C++
コンパイラーと OpenMP*
インテル® Fortran
コンパイラーと OpenMP*

解析ツール

API ベースのプログラミング
インテル® oneAPI DPC++
ライブラリーインテル® MPI ライブラリー

インテル® oneAPI マス・
カーネル・ライブラリー
インテル® oneAPI データ・
アナリティクス・ライブラリー

インテル® VTune™ プロファイラー

インテル® Inspector

インテル® Advisor

インテル® Trace Analyzer &Collector

GDB*

インテル® Cluster Checker

インテル® oneAPI スレッディング・
ビルディング・ブロック
インテル® oneAPI ビデオ・
プロセシング・ライブラリー
インテル® oneAPI コレクティブ・
コミュニケーション・ライブラリー
インテル® oneAPI ディープ・ニュー
ラル・ネットワーク・ライブラリー
インテル® インテグレーテッド・
パフォーマンス・プリミティブ

インテル® oneAPI ベース・ツールキット© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 29

データ並列 C++ (DPC++) の例© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 30

データ並列 C++

標準ベースのクロスアーキテクチャー言語
さまざまな CPU とアクセラレーターに妥協のない並列プログラ
ミングの生産性とパフォーマンスを提供する言語•

同じデータ並列プログラミング・モデルですべての SVMS アーキ
テクチャーをサポートする機能と抽象化を提供

C++ ベース•

使い慣れた C/C++ 構文を使用できるため C++ の生産性が得られる•

Khronos Group の SYCL* を組み込むことでデータ並列処理とヘテロジニアス・プログラミングをサポート

コミュニティー・プロジェクトを通して推進される言語拡張•

データ並列プログラミングを容易にする拡張•

継続的な進化のためのオープンな共同開発

インテルのアーキテクチャーとコンパイラーの長年の経験に基づいて構築© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 31

DPC++ プログラムの例
単一ソース◼

ホストコードとヘテロジニアス・アクセラ
レーター・カーネルを同じソースファイルに混在して記述
ホスト
コード#include <CL/sycl.hpp>#include <iostream>
constexpr int num=16;
using namespace cl::sycl;
int main() {auto R = range<1>{ num };
buffer<int> A{ R };使い慣れた C++◼

queue{}.submit([&](handler& h) {auto out =A.get_access<access::mode::write>(h);h.parallel_for(R, [=](id<1> idx) {out[idx] = idx[0];});});auto result =A.get_access<access::mode::read>();for (int i=0; i<num; ++i)
std::cout << result[i] << "¥n";ライブラリー構造は、次の機能を提供:
アクセラレーター・コード
コンストラクト目的

queue

ワークを投入buffer

データ管理

parallel_for

並列処理

ホスト
コード

return 0;© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

}

| 32

OpenMP* から
データ並列 C++ へ#pragma omp target teams distribute parallel for ¥map(to: A, B) map(tofrom: C) thread_limit(128)for (int i = 0; i < MAX; i++)cl::sycl::gpu_selector device;
cl::sycl::queue deviceQueue(device);
cl::sycl::range<2> matrix_range(MAX, MAX);for (int k = 0; k < MAX; k++)for (int j = 0; j < MAX; j++)C[i][j] += A[i][k] * B[k][j];cl::sycl::buffer<TYPE, 2> bufferA((TYPE*) A, matrix_range);
cl::sycl::buffer<TYPE, 2> bufferB((TYPE*) B, matrix_range);
cl::sycl::buffer<TYPE, 2> bufferC((TYPE*) C, matrix_range);deviceQueue.submit([&](cl::sycl::handler& cgh){auto accessorA = bufferA.template get_access<sycl::access::mode::read>(cgh);auto accessorB = bufferB.template get_access<sycl::access::mode::read>(cgh);auto accessorC = bufferC.template get_access<sycl::access::mode::read_write>(cgh);
cgh.parallel_for(matrix_range,[=](cl::sycl::id<2> ind){
int k;
for(k=0; k<MAX; k++){accessorC[ind[0]][ind[1]] += accessorA[ind[0]][k] * accessorB[k][ind[1]];
}
});
});
© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 33

C / C++ コードの典型的な最適化を実装
cl::sycl::accessor<TYPE, 2, cl::sycl::access::mode::read_write, cl::sycl::access::target::local>
aTile(cl::sycl::range<2>(MATRIX_TILE_SIZE, MATRIX_TILE_SIZE), cgh);
cl::sycl::accessor<TYPE, 2, cl::sycl::access::mode::read_write, cl::sycl::access::target::local>
bTile(cl::sycl::range<2>(MATRIX_TILE_SIZE, MATRIX_TILE_SIZE), cgh);
cgh.parallel_for(cl::sycl::nd_range<2>(matrix_range, tile_range),[=](cl::sycl::nd_item<2> it){int k;
16 x 16 のタイルを定義してconst int numTiles = MAX/MATRIX_TILE_SIZE;
ブロック化を実装
const int row = it.get_local_id(0);const int col = it.get_local_id(1);const int globalRow = MATRIX_TILE_SIZE * it.get_group(0) + row;
const int globalCol = MATRIX_TILE_SIZE * it.get_group(1) + col;TYPE acc = 0.0;
for (int t=0; t<numTiles; t++){const int tiledRow = MATRIX_TILE_SIZE * t + row;
const int tiledCol = MATRIX_TILE_SIZE * t + col;aTile[row][col] = accessorA[globalRow][tiledCol];
bTile[row][col] = accessorB[tiledRow][globalCol];it.barrier(cl::sycl::access::fence_space::local_space);for (k = 0; k < MATRIX_TILE_SIZE; k++){acc += aTile[row][k] * bTile[k][col];}it.barrier(cl::sycl::access::fence_space::local_space);}
accessorC[globalRow][globalCol] = acc;});
© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 34

パフォーマンスは …
インテル® Core™ i7-9600K (インテル® UHD グラフィックス 630) で実行

インテル® Core™ i7-1065G (インテル® Iris® Plus グラフィックス) で実行

パフォーマンスは開発中の DPC++ およびライブラリーを使用して測定したものであり、今後変わる可能性があります© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 35

参考資料
◼

この資料の詳細は iSUS で公開中のベータ版インテル® C++ コンパイラーおよびベータ版インテル® Fortran コンパイ
ラー向けの GPU への OpenMP* オフロード導入
https://www.isus.jp/products/c-compilers/get-started-with-cpp-fortrancompiler-openmp/

を参照してください© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 36

まとめ
oneAPI 向けのデータ並列 C++ (DPC++) へ移行する前に、インテル® C++/Fortran コンパイラー V19.1 やインテル® oneAPI HPC ツールキットに含まれるベータ版インテル® C++/Fortran コンパイラー 2021 を使用して簡単にインテル® グラフィックスへのオフロードを行うソフトウェアを開発および検証できますIntel、インテル、Intel ロゴ、VTune は、アメリカ合衆国および / またはその他の国における Intel Corporation またはその子会社の商標です。
© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 37

インテル® oneAPI HPC ツールキット Beta09 の追記 (1)インテル® oneAPI HPC ツールキットに同梱されるインテル® コンパイラーで OpenMP* 5.0
が正式にサポートされました

クラシックモード

Nextgen モード
インテル® コンパイラーの Nextgen モードでは、クラシックモードでサポートされていた、いくつかのオプション
が未サポートとなっています。詳細は /Qnextgen-diag オプションで確認してください© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 38

インテル® oneAPI HPC ツールキット Beta09 の追記 (2)インテル® oneAPI HPC ツールキットに同梱されるインテル® コンパイラーで OpenMP* 5.0 が正式にサポートされたことにより、前回のセッションで紹介したいくつかの OpenMP* 5.0 の機能が実装されています#pragma omp parallel loop private(x) reduction(+:sum)for (i=0; i<num_steps; i++){x = (i + .5)*step;sum += 4.0/(1.+ x*x);}Classic© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

Nextgen| 39

OpenMP* 5.1 の補足と 5.0 対応の例題◼

OpenMP* 5.1 TR9 のドラフト仕様が公開されています:
https://www.openmp.org/wp-content/uploads/openmp-TR9.pdf◼

OpenMP* Application Programming Interface Examples
version 5.0.1 が公開されています:
https://www.openmp.org/openmp-updates/openmp-examples-updatedwith-5-0-features/© 2020 iSUS *その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

| 40

ソフトウェア・セミナー