OpenMP* 4.x における拡張
オフロード© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

内容
• デバイス (アクセラレーター) 拡張• ⼊れ⼦の並列化制御• プロセッサーバインドとアフィニティーの制御

2© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

内容
• デバイス (アクセラレーター) 拡張–基本–データ移動–永続性–並⾏ (⾮同期) 実⾏– インテル® コンパイラーのオフロード向け⾔語拡張 (LEO)  OpenMP*3© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

内容
• デバイス (アクセラレーター) 拡張–基本–データ移動–永続性–並⾏ (⾮同期) 実⾏– インテル® コンパイラーのオフロード向け⾔語拡張 (LEO)  OpenMP*4© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

オフロード:他のプロセッサーで実⾏
ホスト上で動作するプログラムは、コードの特定のブロックを MIC で実⾏するため、ワークを “オフロード” する。 また、ホストはホストとデバイス間のデータ転送を指⽰するホスト上で実⾏されるアプリケーション

デバイス (別のプロセッサー) が割り当てられたワークを実⾏している間、ホストがアクティブ状態を保つのが理想的

x16
PCIe

ワーク
(ホスト)
5
© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

" ... ワークを実⾏し、指⽰された結果を転送 ... "

平⾏

ワーク
(mic)
s

u g a

t

e

c h

オフロードを制御

ワーク
(mic)

⾃動オフロードをさらに効率よく⾏うには、「修飾⼦」が必要(句、属性、指定⼦、キーワード)
６© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

割り当て/関連付け
MIC メモリー
データを MIC へ転送
実⾏
MIC 上の コード
MIC から
データを転送単⼀のオフロード宣⾔⽂

データ転送
割り当て/開放
平⾏性

x16
PCIe

ワーク 平⾏(ホスト)

MIC を検出開放
MIC メモリー

s

u g a

t

e

c h

基本オフロード

同期(デバイス選択、データ転送、およびデバイス上のストレージ）
データが構⽂の範囲内にある場合、変数はデバイスへ/から転送され、領域の入口/出口で割り当て/開放される
インテル LEO

a=b

領域# pragma offload target(mic:0){}7© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

OpenMP* 4.0 への移⾏
• 基本操作は同じ• シンタックスが異なる
インテル LEO# pragma offload target(mic:0){

OpenMP# pragma omp target device(0)
{

a=b
}

8
© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

a=b
}

s

u g a

t

e

c h

データ転送の⽅向• PCIe バスの帯域幅を抑える
インテル LEO#pragma offload target(mic:0) ¥in(b),out(a),inout(c){a=b;c=c*c}9© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

OpenMP* 4.0 への移⾏
• OpenMP の動作は同じ、シンタックスが異なるだけ

インテル LEO#pragma offload target(mic:0) ¥in(b),out(a),inout(c){OpenMP#pragma omp target device(0) ¥map(to:b),map(from:a),map(tofrom:c){

a=b;c=c*c}

a=b;c=c*c}
データ転送属性は省略可能

10© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

サンプルコード
Pi を求めるプログラムをインテル® コンパイラーのオフロード向け拡張と OpenMP 4.0 の機能を使⽤して記述
演習と⼿順© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

内容
• デバイス (アクセラレーター) 拡張–基本–データ移動–永続性–並⾏ (⾮同期) 実⾏– インテル® コンパイラーのオフロード向け⾔語拡張 (LEO)  OpenMP*12© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

データの永続性
割り当て/転送

割り当て/開放  永続性
データ転送

実⾏
MIC 上の コード

転送実⾏
MIC 上の コード

転送
転送/開放

13
© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

データの永続性 -- LEO と OpenMP* 4.0
インテル LEO

OpenMP
割り当て/転送

新しい永続性マッピングは、
オフロード宣⾔⽂を含むどこにでも作成できる

⾮構造化

{

実⾏
MIC 上の コード

新しい永続性マッピングは、この領域内で作成できない

転送
実⾏
MIC 上の コード

構造化

転送
転送/開放

14
© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

}*OpenMP の⾮構造化は、OpenMP 4.5 で利⽤可能
s

u g a

t

e

c h

LEO のデータ永続性
⾮構造化転送 :#pragma offload_transfer ¥in/out(vars:alloc_if(logical) free_if(logical))
真なら割り当て

真なら開放

オフロード宣⾔⽂でも指定可能 :#pragma offload ¥in/out(vars:alloc_if(logical) free_if(logical) )15© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

OpenMP* のデータ永続性
構造化転送 :
#progma target data device(0) ¥map([to|from|alloc]: vars){...
}

デフォルトのマッピングは、
変数 vars の "tofrom"16© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

OpenMP* 4.0 への移⾏
データの永続性
インテル LEO#pragma offload_transfer target(mic:0)&in(c:alloc_if(1) free_if(0))!データは、割り当てられ、コピーされ、開放されない...

#pragma offload target(mic:0) nocopy(c)a=b;c=c*c； ! a & b, 自動; c 永続...

OpenMP*#pragma omp target data device(0) &map(c)!デフォルトは tofrom!データは、割り当てられ、コピーされ、開放されない...
{#pragma omp target device(0)! a & b, 自動;
a=b;c=c*c； ! c 永続...

! Offload_transfer は、データをデバイスからコピーし開放#pragma offload_transfer target(mic:0)&out(c:alloc_if(0) free_if(1))

17© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

}!データ構造化ブロックの最後!デフォルトは from

s

u g a

t

e

c h

OpenMP* 4.5 における新しいデータの永続性
非構造化転送 :#progma omp target enter data device(0) ¥map(to|alloc: vars)
非構造化転送 :#progma omp target exit data device(0) ¥map(from|release|delete: vars)
delete = 完全に開放18© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

release = 参照カウントを減らす
s

u g a

t

e

c h

サンプルコード
プログラム内に複数のオフロード領域がある場合、デバイス上のデータをどのように継続利⽤、もしくは廃棄するか︖

オフロードのサンプルコード© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

内容
• デバイス (アクセラレーター) 拡張–基本–データ移動–永続性–並⾏ (⾮同期) 実⾏– インテル® コンパイラーのオフロード向け⾔語拡張 (LEO)  OpenMP*20© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

⾮同期オフロード
⾮同期実⾏ :

同期実⾏ :• ホストのスレッド/プ
ロセスは、オフロード
構⽂で完了を待機• ホストのスレッド/プロセス
は、オフロード構⽂を実⾏後、すぐに実⾏を継続し、waitで指⽰される場所で待機
ホスト
実行

オフロード宣言子

MIC
実行

MIC でのみ実行

時間

時間

ホスト
実行

オフロード (非同期)ホストと MIC は平行に
ここを実行

MIC
実行

ホストは wait で待機
オフロード完了
ホストはここから実行を続行
21© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

平⾏実⾏ -- 簡単
LEO• target … signal(id)• signal 句• ⾮同期に⾏う• 追加のスレッドは必要なし• signal の引数は待機する "プロセス" を特定wait(id) でプロセスを待機(単独もしくは target 構⽂)

22© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

OpenMP*• target … nowait
インテル® コンパイラー 16(OpenMP 4.5) でサポート• nowait 句• ⾮同期に⾏う• 追加のスレッドが必要• task 識別⼦なし
taskwait を使⽤
s

u g a

t

e

c h

同時実⾏
LEO (ホストと MIC -- 簡単)#pragma offload target(mic:0) signal(&isig)#pragma omp parallel num_threads(60)work(noff1, nend1)!mic で実行!T1

MIC

オフロードが完了するまでホストはポーズ

CPU

MIC への新たなワークも開始しない!T0#pragma omp parallel num_threads(16)work(noff2, nend2)!ホストで実行!T2#pragma offload_wait target(mic:0) wait(&isig)

インテルの⾮同期実装では、オフロード後即座にホストのほかの並列領域を実⾏することを許可
時間 (秒) は、ホスト上で計測
23

Total
⾮同期 off
⾮同期 on© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

T0
T1
T2
2.111171.12804
0.983141.08984
0.00002
1.01347
s

u g a

t

e

c h

同時実⾏
OpenMP* (ホストと MIC -- 簡単)#pragma omp target parallel num_threads(60) nowait
work(noff1, nend1)! mic で実行

#pragma omp parallel num_threads(16)work(noff2, nend2)!ホストで実行MIC

CPU#pragma omp taskwait• 複合構⽂に注意 (target と parallel)© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h 24

LEO のまとめ
操作
関数定義、グローバル

宣⾔⽂句もしくは修飾⼦
attributes/declspec

コード実⾏

offload

領域割り当て
データ転送
同時実⾏

offload_transfer

同期

offload_wait

alloc_if free_if
in, out, inout, length
signal、wait

データ転送
永続性
⾮同期実⾏
© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

su 25g atec hOpenMP* まとめ操作

宣⾔⽂句もしくは修飾⼦

関数定義、グローバル
コード実⾏

declare target

領域割り当てデータ転送
同時実⾏

target update

同期

taskwait

26© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

targetalloc, release, deletemap(to,from,tofrom) length
nowait、depend

データ転送
永続性
⾮同期実⾏
s

u g a

t

e

c h

OpenMP* 4.x への移⾏•

OpenMP 4.0 は、LEO (Language Extensions for Offload) と同様のデバイス Target 宣⾔を持っている•インテル® コンパイラー 15.0 は、OpenMP 4.0 に準拠している
Update 版は、4.5 のコンポーネントを組み込んでいる•

OpenMP 4.0 では、明⽰的な⾮同期句と⾮構造化データマッピン
グがサポートされていない。これらは、OpenMP 4.5 でサポートされる•

データの永続性はより簡単になるが、⾮構造化マッピングの制御はさらに必要となる。インテル® コンパイラー 16.0 は、⾮同期
句 (nowait) をサポートしている© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u 27g atec h

⾮同期オフロードの例
素数を求める計算をホスト(CPU)とデバイス(MIC) で⾮同期に同時実⾏する

オフロードのサンプルコード© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

内容
• デバイス (アクセラレーター) 拡張–基本–データ移動–永続性–並⾏ (⾮同期) 実⾏– インテル® コンパイラーのオフロード向け⾔語拡張 (LEO)  OpenMP*29© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

LEO  OpenMP 4.x (Fortran)
汎⽤

デバイス実⾏OFFLOADTARGET

デバイス指⽰target(MIC:#)

device(#)

デバイスメモリー割り当て、転送 IN( ) OUT( )例

データの永続性

INOUT( )map(TO:) map(FROM:) map(TOFROM:)

⾮構造化データ転送

OFFLOAD_TRANSFERUPDATE

⾮同期

SIGNAL(S)WAIT(S)NOWAIT*TASKWAIT!dir$ offload begin target(MIC:0) in(x) out(y) signal(sync)<mic work>!dir$ end offload<host work>!dir$ offload_wait target(mic:0) wait(sync)!$omp target device(0) map(to:x) map(from:y) nowait<device work>!$omp end target<host work>!$omp taskwait!dir$ attributes offload:mic :: x, y
!dir$ offload offload_transfer target(MIC:0) in(a: alloc/free...)out(b: alloc/free…) nocopy(c: alloc/free…)
更新:!dir$ offload offload_transfer target(MIC:0) in(a:…)out(b:…)!dir$ offload begin target(mic:0)
in(a:…)out(b:…)!$omp target declare (x, y)!$omp target data device(0) map(to: a)map(from: b) map(alloc: c)

グローバル
構造化*

!$omp target update device(0) to(a) from(b)!$omp target device(0) map(always,to:a)map(always,from: b)*nowait は、インテル® コンパイラー 16 の 4.5 の機能、*⾮構造化マッピング (target enter/exit data) は OpenMP 4.5 でサポート© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a
30

t

e

c h

LEO  OpenMP 4.x (C/C++)
汎⽤

デバイス実⾏OFFLOADTARGET

デバイス指⽰target(MIC:#)

device(#)

デバイスメモリー割り当て、転送 IN( ) OUT( )例

データの永続性

INOUT( )map(TO:) map(FROM:) map(TOFROM:)

⾮構造化データ転送

OFFLOAD_TRANSFERUPDATE

⾮同期

SIGNAL(S)WAIT(S)NOWAIT*TASKWAIT#pragma offload target(MIC:0) in(x) out(y) signal(sync){ <mic work> }#pragma omp target device(0) map(to:x) map(from:y) nowait{ <device work> }{ <host work> }#pragma offload_wait target(mic:0) wait(sync){ <host work> }#pragma omp taskwait#pragma offload_attribute( target(mic)) int x,y#pragma offload offload_transfer target(MIC:0) in(a: alloc/free...)out(b: alloc/free…) nocopy(c: alloc/free…)
更新:#pragma offload offload_transfer target(MIC:0) in(a:…)out(b:…)#pragma offload target(mic:0) in(a:…)out(b:…)#pragma omp target map(x, y)#pragma omp target data device(0) map(to: a)map(from: b) map(alloc: c)

グローバル
構造化*

#pragma omp target update device(0) to(a) from(b)#pragma omp target device(0) map(always,to: a)map(always,from: b)*nowait は、インテル® コンパイラー 16 の 4.5 の機能、*⾮構造化マッピング (target enter/exit data) は OpenMP 4.5 でサポート© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a
31

t

e

c h

内容
• デバイス (アクセラレーター) 拡張• ⼊れ⼦の並列化制御• プロセッサーバインドとアフィニティーの制御

32© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

OpenMP* 3.1 における並列領域の⼊れ⼦#pragma omp parallel#pragma omp parallel

OpenMP* 3.1 では、⼊れ⼦になった並列領
域の内側は、デフォルトでシングルスレッド
で実⾏されるOMP_NESTED 環境変数を true に設定すると、内側の領域もマルチスレッドで実⾏できるが、最⼤スレッド数は (外側のスレッド x 内側のスレッド) となり、オーバーサブスクライブとなる。スレッド数とアフィニティーの制御は困難
OpenMP* 4.0 では、数百スレッドを実⾏できるデバイスでの⼊れ⼦になった並列領域を制御するため、teams と distribute 句が追加された
サンプルコード: nest.c33© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

teams 構⽂• 複数レベルの並列デバイスをサポート• 構⽂ (C/C++):#pragma omp teams [節[[,]節],…]構造化ブロック

このプラグマの直後は、各チームのマスタースレ
ッドのみが実⾏し、ほかのチームメンバーは次の(⼊れ⼦構造の) 並列領域からのみ実⾏を開始します。そのため、実⾏中のスレッド数は
num_teams のみで、それぞれのスレッドは
omp_get_thread_num() == 0 になります• 構⽂ (Fortran):!$omp teams [節[[,]節],…]構造化ブロック• 節: num_teams(整数式)、thread_limit(整数式)、default(shared | none)、
private(リスト)、firstprivate(リスト)、shared(リスト)、
reduction(演算⼦ :リスト)© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

distribute 構⽂• ループ反復を複数のチームのマスタースレッドで分割• 構⽂ (C/C++):#pragma omp distribute [節[[,]節],…]構造化ブロック

このプラグマは、teams 構造内の緊密な⼊れ⼦構造の 1 つ以上のループに関連付けられます。
collapse を使⽤すると、omp for プラグマでcollapse 節を指定した場合と同様に、複数のループを 1 つの反復シーケンスに結合できます• 構⽂ (Fortran):!$omp distribute [節[[,]節],…]構造化ブロック• 節:

collapse(n)、 private(リスト)、firstprivate(リスト)、
dist_schedule(static [, chunk_size])© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

コプロセッサーへ SAXPY をオフロードする
int main(int argc, const char* argv[]) {float *x = (float*) malloc(n * sizeof(float));•floatSAXPY*y = (float*) malloc(n * sizeof(float));// Define scalars n, a, b & initialize x, y#pragma omp target data map(to:x[0:n]){#pragma omp target map(tofrom:y)#pragma omp teams num_teams(num_blocks) thread_limit(nthreads)
全てが同じことを⾏う
for (int i = 0; i < n; i += num_blocks){
for (int j = i; j < i + num_blocks; j++) {y[j] = a*x[j] + y[j];}}}}free(x); free(y); return 0;© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

コプロセッサーへ SAXPY をオフロードする
int main(int argc, const char* argv[]) {float *x = (float*) malloc(n * sizeof(float));float *y = (float*) malloc(n * sizeof(float));// Define scalars n, a, b & initialize x, y

• SAXPYコプロセッサー/アクセラレーター#pragma omp target data map(to:x[0:n]){#pragma omp target map(tofrom:y)#pragma omp teams num_teams(num_blocks) thread_limit(nthreads)
全てが同じことを⾏う#pragma omp distributefor (int i = 0; i < n; i += num_blocks){
ワークシェア (barrier なし)#pragma omp parallel forfor (int j = i; j < i + num_blocks; j++) {ワークシェア (barrier あり)
y[j] = a*x[j] + y[j];}}} free(x); free(y); return 0; }© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

コプロセッサーへ SAXPY をオフロードする• SAXPY複合構⽂
int main(int argc, const char* argv[]) {float *x = (float*) malloc(n * sizeof(float));float *y = (float*) malloc(n * sizeof(float));// Define scalars n, a, b & initialize x, y#pragma omp target map(to:x[0:n]) map(tofrom:y){#pragma omp teams distribute parallel for ¥num_teams(num_blocks) thread_limit(nthreads)for (int i = 0; i < n; ++i){y[i] = a*x[i] + y[i];}}}free(x); free(y); return 0;© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

⼊れ⼦並列とデバイスへの割り当て
オフロードのサンプルコード© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

内容
• デバイス (アクセラレーター) 拡張• ⼊れ⼦の並列化制御• プロセッサーバインドとアフィニティーの制御

40© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

スレッド・アフィニティー: プロセッサーのバインド• バインドの⽅針は、マシンとアプリケーションに依存する• スレッドを離して配置、例、異なるパッケージ• (おそらく) メモリー帯域幅を向上させる• (おそらく) 統合されたキャッシュサイズを改善• (おそらく) 同期構⽂のパフォーマンスを低下させる• スレッドを近づけて配置、例、キャッシュを共有する可能性がある 2 つのコアに隣接• (おそらく) 同期構⽂のパフォーマンスを向上させる• (おそらく) 利⽤可能なメモリー帯域幅とキャッシュサイズ (スレッドごとの) を
低下させる

41
© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

OpenMP* 4.0 におけるスレッド・アフィニティー
• OpenMP* 4.0 で、アフィニティーのコンセプトを導⼊ ...
• 1 つ以上のプロセッサー上で動作する⼀連のスレッド• ユーザーによって定義される• 事前定義された配置:• スレッド ハイパースレッドごとに 1 つの位置
• コア
物理コアごとに 1 つの位置
• ソケット プロセッサー・パッケージごとに 1 つの位置... そしてアフィニティーのポリシーは ....
• spread• close• master

OpenMP* スレッドをすべての位置に広く配置
OpenMP* スレッドをマスタースレッドの近辺にパック
OpenMP* スレッドをマスタースレッドを併置• ... そしてこれらの設定を制御する• 環境変数 OMP_PLACES と OMP_PROC_BIND• 並列領域向けに proc_bind 節

42© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

スレッド・アフィニティーの例• 例 (インテル® Xeon Phi™ コプロセッサー): 外部領域を分配し、内部領域を近く保つ
OMP_PLACES=cores(8)#pragma omp parallel proc_bind(spread)#pragma omp parallel proc_bind(close)

p0

p1

p2

p3

p4

p5

p6

p7

p0

p1

p2

p3

p4

p5

p6

p7

p0

p1

p2

p3

p4

p5

p6

p7

43
© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

アフィニティー制御の例
オフロードのサンプルコード

環境変数を使⽤してテスト
set MIC_ENV_PREFIX=MIC
set MIC_OMP_PROC_BIND=[master, close, spread]$> matmul_dist_para を実⾏

44© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

内容
• デバイス (アクセラレーター) 拡張• プロセッサーバインドとアフィニティーの制御• GFX コンパイラーと GFX へのオフロード

45© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

Gfx コンパイラー

インテル® グラフィックス・テクノロジー• プログラミング・モデル機能− 共有仮想メモリー− OpenMP* 4.0 の⼀部− ⾮同期プログラミング・サポートの改善• パフォーマンスの改善− 共有ローカルメモリー− 第 5 世代インテル® Core™ プロセッサー向けにチューニング− Gen ターゲット向けのベクトル化機能の改善• 利⽤法− Gfx_sys_check ツール− デバッグサポートの改善© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

46

Gfx コンパイラー

OpenMP* 4.0 offload サポートへの追加機能
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;}© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

利⽤⽅法︓− サブセットのみのサポート− “tofrom” と “to” を “pin” へマップ
− “-qopenmp-offload=gfx” を指定

s

u g a

t

e

c h

47

まとめ
OpenMP* 4.0 / 4.5 は、OpenMP における⼤きな⾶躍• 新しいレベルの並列性を導⼊• デバイス（MIC、GPU) への演算のオフロード• データの永続性を制御• ⾮同期実⾏を制御• デバイスによる異種システム構成をサポート© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

参考サイトインテル® ソフトウェア・フォーラム、ナレッジベース、記事、ツールのサポート(http://software.intel.com 参照、http://isus.jp 翻訳版)
記事の例:
- http://www.isus.jp/article/parallel-special/requirements-for-vectorizable-loops/（ループをベクトル化するための条件）
OpenMP* 4.5 Specification- http://www.openmp.org５/mp-documents/openmp-4.5.pdf
OpenMP* 3.1 仕様をカバーするオンライン・トレーニング⁃ http://www.isus.jp/online-training/© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

関連書籍
Structured Parallel Programming: Patterns for Efficient Computation
著者 Michael McCool, James Reinders ,Arch Robison 出版⽇: 2012 年 7 ⽉ 9 ⽇ | ISBN: 978-0-124159938『構造化並列プログラミング: 効率良い計算を⾏うためのパターン』
著者

マイケル・マックール／アーク・D・ロビソン／ジェームス・レインダース（共著）訳者

菅原 清⽂／エクセルソフト株式会社（共訳） | ISBN

978-4-87783-305-3

Intel® Xeon Phi™ Coprocessor High Performance Programming
著者 Jim Jeffers , James Reinders 出版⽇: 2013 年 3 ⽉ | ISBN: 978-0-124104143『インテル® Xeon Phi™ コプロセッサー ハイパフォーマンス・プログラミング』
著者

ジェームス・レインダース/アーク・D・ロビソン（共著）訳者

菅原 清⽂／エクセルソフト株式会社（共訳） | ISBN

978-4-87783-332-9

High Performance Parallelism Pearls
著者 Jim Jeffers , James Reinders

出版⽇: 2014 年 11 ⽉

簡単にインテル® Xeon Phi™ コプロセッサー・ファミリーの優れた並列性を利⽤してコードを実⾏できるため、最適化に集中し、ハイパフォーマンスを実現することが可能です。並列処理を細かくチューニングすることで、正しいアプリケーションを正しく効率良いアプリケーションにすることができます。インテル コーポレーションの並列プログラミング・エバンジェリストである
James Reinders とインテル コーポレーションのエンジニアである Jim Jeffers により執筆された最新の書籍は、69 ⼈の専⾨家の実際の経験を基に、インテルのマルチコアおよびメニーコア・プロセッサーを最⼤限に利⽤するための創意⼯夫を紹介しています。
© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

インテル® コンパイラーによるオフロード拡張の変更点 (1)インテル® C++ および Fortran コンパイラーのバージョン 16 で OpenMP*
4.0 のオフロード拡張をサポートするにあたり、LEO のデータ属性の扱いが変更されました#pragma offload target(mic) in(num_steps, step) inout(sum)#pragma omp parallel for simd reduction(+:sum) private(x)for (i=0;i< num_steps; i++){x = (i+0.5)*step;sum = sum + 4.0/(1.0+x*x);}

インテル® コンパイラーのバージョン 15.x ではオフロード時に明⽰的にデータの target 属性を記述する必要がありましたが、バージョン 16 以降
では省略できます52© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h

インテル® コンパイラーによるオフロード拡張の変更点 (2)in、out、inout のデータ属性を省略するとすべての変数に inout が適⽤される。

inout で転送されるのは、sum のみ戻る
53© 2016 iSUS. 無断での引⽤、転載を禁じます。* その他の社名、製品名などは、⼀般に各社の表⽰、商標または登録商標です。

s

u g a

t

e

c h