2013年7月
加藤 努
株式会社 ソフテック
PGI Accelerator Compiler
GTC Japan 2013
新OpenACC 2.0の機能と
PGIアクセラレータコンパイラ
本日の話
OpenACC 1.0の復習
ディレクティブ操作で 出来ることを再確認OpenACC
2.0の新機能
プログラミングの自由度の 向上へPGI Accelerator Compiler 2013
• PGI Accelerator Programming Model 準拠
• OpenACC 1.0 準拠
• OpenACC 2.0 の機能は、毎月のリビジョンで順次追加
• プラスPGI Extension（先取り機能）
OpenACCディレクティブで出来ることを
改めて知ろう！
PGI Accelerator / OpenACC Release Map
2007 2010 2013
NVIDIA CUDA release
2008 2009 2011 2012
PGI Accelerator Programming Model 公開 (OpenACCの原型モデル）
PGI Accelerator Compiler リリース (PGI 2010)
OpenACC API 1.0 release
OpenACC 1.0準拠 PGI Compiler リリース(PGI 2012) OpenACC 2.0
順次取込 (PGI 2013)
OpenMP 4.0 with extensions for Accelerators
OpenACC API 2.0 final
PGI 提供開始
PGIは、2009年より提供開始
OpenACC Programming Model
CPU
Main
Memory
Device
Device
Memory
② データの移動管理
ホスト側
Accelerator側
オフロードモデル
重い計算部分の 処理をオフロード 「重い処理部分」をアクセラレータ側で処理するモデル① 処理の並列化(分割指示）
並列処理による
スループット改善
OpenACCディレクティブの主な構成
CPU
Main
Memory
GPGPU
Device
Memory
ホスト
Accelerator
重い計算部分の 処理をオフロード① Accelerate Compute 構文 (offload 領域指示)
② Data 構文 (データ移動指示)
③ Loop 構文 (Mapping for parallel/vector, Tuning)
①
②
③
(処理)
OpenACC 1.0 ディレクティブの復習
まだ、ハードルが高いあなたに！
① どこをオフロードの対象とするか？
コンパイラに対する指示、ユーザが行う最初のステップ② データ・マネージメントの方法
データの移動指示 ・コンパイラが自動的に行う ・ユーザによる明示的な指示 ホスト上、デバイス上にコピー③ 並列処理の分割方法の指示
Mapping (並列チューニング）のdirectives後は、コンパイラに任せる
…
① 並列化(オフロード）領域の対象は？
A. 並列化可能な「ループ」、「ループの集合体」
for( i = 0; i < nrows; ++i ){ double val = 0.0;
int nstart = rowindex[i]; int nend = rowindex[i+1];
for( n = nstart; n < nend; ++n )
val += m[n] * v[colndx[n]];
r[i] = val;
}
for( i = 0; i < n; ++i )
a[i] = b[i] + c[i];
for (i = 0; i < n; i++){ for (j = 0; j < n; j++){ for (k = 0; k < n; k++){ C[i*n+j] = A[i*n+k] * B[k*n+j]; } } }
Non-tightly Nested Loops
tightly Nested Loops Single Loops
for( i = 0; i < n; ++i )
a[i] = b[i] + cdd[i];
for( i = 1; i < n-1; ++i )
c[i] = 0.5 * a[i-1] + a[i+1];
Ajacent(隣接) Loops
Q. どこをオフロードの対象とするのか？
(どこにdirectiveを入れるか）
プログラム・シーケンス上でのオフロード対象
for (cgit = 1; cgit <= cgitmax; cgit++) {rho0 = rho; d = 0.0;rho = 0.0;
for (j = 1; j <= lastrow-firstrow+1; j++) {
sum = 0.0;
for (k = rowstr[j]; k < rowstr[j+1]; k++) {
sum = sum + a[k]*p[colidx[k]];
} w[j] = sum; q[j] = sum; } ……(snip)…… for (j = 1; j <= lastcol-firstcol+1; j++) { d = d + p[j]*q[j]; } alpha = rho0 / d; (ホスト上でスカラ演算） for (j = 1; j <= lastcol-firstcol+1; j++) {
double new_rj = r[j] - alpha*q[j]; z[j] = z[j] + alpha*p[j];
r[j] = new_rj;
rho = rho + new_rj*new_rj; }
……(snip)……
beta = rho / rho0; (ホスト上でスカラ演算） ……(snip)…… } 外側繰り返しループ 外側繰り返しループ Nested loop (並列化可能） Single loop (並列化可能） Single loop (並列化可能） OpenACC directive OpenACC directive OpenACC directive
並列化(オフロード）できないループ
do k = 1, d3 do jj = 0, d2 - fftblock, fftblock do j = 1, fftblock do i = 1, d1 y1(j,i) = x(i,j+jj,k) enddo enddocall cfftz (is, logd1, d1, y1, y2) do j = 1, fftblock do i = 1, d1 xout(i,j+jj,k) = y1(j,i) enddo enddo enddo enddo Nested loop (並列化可能） ループ内に Procedure Callがあると 並列化出来ない Kループはオフロード化できない Nested loop (並列化可能） ループ内に Function call • OpenACC API 1.0 では並列化できない
• NVIDIA CUDA 5.0からGPUカーネルのリンカー機能が追加
されたことのより、
OpenACC 2.0で可能
となる並列化を行うには、インライン展開 するしかなかった！
② OpenACCにおけるデータ移動の基本
「ホストとデバイス」間のデータ交換のタイミングポイントは？
!$acc end kernels
do k=1,n do j=1,n do i=1,n a(i,j.k)=b(i,j,k)*c(i,j,k) enddo enddo enddo !$acc kernels
Gang Gang Gang
!$acc parallel
!$acc end parallel
コンパイラのデフォルト動作
OpenACC 並列化構文（acc kernels や acc parallel) の始点と終点で その領域内で使われている配列、スカラ変数を解析し、「ホスト」と
「デバイス」間のデータコピーを行うコードを暗黙に自動生成
• ユーザが明示的にコピー配列名を指示(copy clause)することができる
• live スカラ変数は、デフォルトでcopyin/copyout されている (スカラ変数の移動管理は明示的に必要なし）
データ・マネージメントの基本的考え方
「ホスト」と「デバイス」間のデータの移動を最小化すること
デバイス上で可能な限りのデータ常駐化 ホスト側から一度、デバイスへコピー(copyin)された配列データは、デバイス 側の計算で役目が終わらない限り、ホスト側へデータ(copyout)を戻さないプログラミング上の鉄則
Fortran !$acc data [clause …]
...
!$acc end data
C #pragma acc data [clause …] { ...}
Data 構文 directive の活用
ユーザが明示的にデータ移動を管理するための構文
デバイス内にデータを常駐化する
!$acc data copy (a(100))
A(100) A(100) Host側メモリ デバイス側メモリ Main routine (host側) A A A A A デバイス側 の処理
!$acc end data
copyin copyout resident time “present” と言う状態
call call call
!$acc data present(a(100))
「Data 構文」で挟まれたプログラム領域に適用できる
プログラムの流れ①
①
②
②
データの同期 update directive(実行文)
!$acc data copy (a(100))
A(100) A(100)
Host側メモリ デバイス側メモリ
A
A A A
A
!$acc end data
copyin copyout
resident
“present” と言う状態
call call call
A
!$acc data present(a(100))
!$acc update host (a(10:20))
!$acc update device (a(10:20))
Device to host Host to device Main routine (host側) デバイス側 の処理 細かな操作が可能
実際のプログラムでdirectiveの適用
int main(int argc, char **argv) { read_input();
domain(); setcoeff();
#pragma acc data create(u,rsd,frct,flux,a,b,c,d)
copyin(ce) copyout(u) { setbv(); setiv(); erhs(); ssor();
} /* end acc data */ error();
｝
Nested loop (並列化可能） データの移動管理のdirective
OpenACC data directive
OpenACC directive static void ssor(void) {
for (istep = 1; istep <= itmax; istep++) { /* perform SSOR iteration */
#pragma acc kernels present(rsd)
for (i = ist; i <= iend; i++) { for (j = jst; j <= jend; j++) { for (k = 1; k <= nz - 2; k++) { for (m = 0; m < 5; m++) { rsd[i][j][k][m] = dt * rsd[i][j][k][m]; } } } } (以下、略） Accelerator Data 領域管理下 function call 並列領域指定のdirective Mainルーチン＋Subroutineタイプ
③ 並列処理の対象と方法を指示する構文
アクセラレータ上で
並列実行を行う領域を指定
するためのdirective
その領域内で使用する「データ」を転送するコードを生成(Implicit xfer)
Parallel 構文
(１つのカーネル構成でコード生成)
C #pragma acc parallel {並列実行領域}
Fortran !$acc parallel {並列実行領域}
!$acc end parallel
Kernels 構文
(tightly loopsを対象、複数のカーネルで構成)
C #pragma acc kernels {並列ループ領域}
Fortran !$acc kernels
{並列ループ領域} !$acc end kernels
並列処理分割の指示を行う Loop構文
Fortran !$acc loop [clause …]
C #pragma acc loop [clause …]
Clause(節) collapse (n) gang [ ( expression ) ] worker [ ( expression ) ] vector [ ( expression ) ] seq independent private( list )
reduction ( operator : list)
• 直下の「ループ」の並列の分割方法(gang, worker, vector)指示
• Acceleratorハードウェアの並列構造にマッピングを行うためのclause
• ループ内プライベートとなる変数の宣言
• リダクション処理の指定
OpenACC Parallelism ３階層
worker
Gang
PEs => Streaming Multiprocessors
Gang
Worker
Vector
PEs ：プロセッサ要素の基本集合体
each PE ：マルチスレッディング処理
Each thread of the PE : ベクトル命令処理
Multithreaded Worker thread thread あまり深刻に考えなくてよい = H/W とのマッピングが難しい Vector命令（長さ） PE
Multithreading => warps within SM
Vector_length => warp内スレッド数32
(PE内の Max. Threads per Thread Block)
並列処理分割の指示 loop構文
#pragma acc kernels
#pragma acc loop independent
for (i = 1; i < n+1; i++) {
#pragma acc loop independent
for (j = 1; j < n+1; j++) {
#pragma acc loop independent
for (k = 1; k < n+1; k++) { a1[i*sz*sz+j*sz+k] = ( a0[i*sz*sz+(j-1)*sz+k] + a0[i*sz*sz+(j+1)*sz+k] + …省略… a0[i*sz*sz+j*sz+(k-1)] + a0[i*sz*sz+j*sz+(k+1)] ) * fac; } } }
Accelerate Compute
構文
Loop
構文
161, Loop is parallelizableAccelerator kernel generated
161, #pragma acc loop gang /* blockIdx.x */
165, #pragma acc loop vector(128) /* threadIdx.x */
(並列対象指示）
(並列分割指示）
並列分割の指示 loop構文
#pragma acc kernels
#pragma acc loop independent gang(32)
for (i = 1; i < n+1; i++) {
#pragma acc loop independent
for (j = 1; j < n+1; j++) {
#pragma acc loop independent vector(32)
for (k = 1; k < n+1; k++) { a1[i*sz*sz+j*sz+k] = ( a0[i*sz*sz+(j-1)*sz+k] + a0[i*sz*sz+(j+1)*sz+k] + …省略… a0[i*sz*sz+j*sz+(k-1)] + a0[i*sz*sz+j*sz+(k+1)] ) * fac; } } }
Accelerate Compute
構文
Loop
構文
161, Loop is parallelizableAccelerator kernel generated
161, #pragma acc loop gang(32) /* blockIdx.x */ 165, #pragma acc loop vector(32) /* threadIdx.x */
(並列対象指示）
(並列分割指示）
OpenACC API 2.0
OpenACC API 2.0 New features
•
「手続きの呼び出し」 と 「分離コンパイル」
(効果) プログラム移行性向上、ライブラリとしてモジュール化可能•
入れ子（Nested)並列
(効果) parallel/kernels構文を有するルーチン・モジュールを構成可能•
デバイス特有なチューニング, 複数のデバイス使用
(効果) デバイス種に依存しない「性能互換性」の向上•
グローバル・データへの対応
(効果) ルーチン間での「広域データ」の使用の単純化•
データ・マネージメント自由度向上
(効果) unstructured data lifetimes（自由にデータ領域を構築）
•
新しい Data API ルーチン
(
より自由にデータ移動が可能）
•
Loop 構文のparallelism 3階層の明確化と auto、tile節の追加
(効果) 並列性3階層の整理、ユーザ・チューニングの簡易化•
非同期操作の向上
現在、手続きの呼び出しがあると？
Accelerator Compute領域の 中の call 文を含むループは、 並列化できない。OpenACC 1.0では、関数、手続きをソースレベルで
インライン展開をするしかなかった。
OpenACC 2.0 では可能になる。
#pragma acc parallel loop num_gangs(200)…for ( int i = 0; i < n; ++i ) { v[i] += rhs[i];
matvec( v, x, a, i, n );
// デバイスへの function call
OpenACC 2.0 Function Call support
#pragma acc routine worker
extern void matvec(float* v, float* x, … ); …
#pragma acc parallel loop num_gangs(200)…
for ( int i = 0; i < n; ++i ){ v[i] += rhs[i];
matvec( v, x, a, i, n );
// デバイスへの function call
} omp end parallel
① 「手続きのプロトタイプ」に指示 worker レベルの並列処理を 指示する例 Function 内でworkerレベルの並列処理を指示する場合は、その上位の parallel/kernels 構文内の上位loopの並列レベルは、gangレベルである 必要がある。同じ workerレベルは駄目。
OpenACC 2.0 “routine” directive
#pragma acc routine worker
extern void matvec(float* v, float* x, … ); …
#pragma acc parallel loop num_gangs(200)…
for ( int i = 0; i < n; ++i ){ v[i] += rhs[i];
matvec( v, x, a, i, n );
// デバイスへの function call
} omp end parallel
#pragma acc routine worker
void matvec( float* v, float* x, float* a, int i, int n ) { float xx = 0;
#pragma acc loop reduction(+:xx)
for ( int j = 0; j < n; ++j ) xx += a[i*n+j]*v[j]; x[i] = xx; } ① 「手続きのプロトタイプ」に指示 worker レベルの並列処理を 指示する例 ② 「手続き」への”routine”指示
“orphaned” loop directive 分離コンパイルも可能
OpenACC 2.0 Nested Parallelism (1)
#pragma acc routine
extern void matvec(float* v,float* x, … ); …
#pragma acc parallel loop …
for ( int i = 0; i < n; ++i )
matvec( v, x, i, n );end parallel
#pragma acc routine void matvec( … ) {
#pragma acc parallel loop
for ( int i = 0; i < n; ++i ) { … }
parallel or kernels 領域内の ループ内から call function function側でも、 parallel or kernels loop を構成できる 実際は、 “i “ 毎にデバイスの中の シングル・スレッドで次々に開始する
matvec内で、指定された parallelism を launch する ①
OpenACC 2.0 Nested Parallelism (2)
#pragma acc routine
extern void matvec(float* v,float* x, … ); …
#pragma acc parallel loop num_gangs(1) … {
matvec( v0, x0, i, n );e
matvec( v1, x1, i, n );end
matvec( v2, x2, i, n );e
}
#pragma acc routine void matvec( … ) {
#pragma acc parallel loop
for ( int i = 0; i < n; ++i ) { … }
一つのデバイス・スレッドで 次々に matvec を起動する
①
②
異なるデバイスへの指示 device_type
#pragma acc parallel loop num_gangs(200)…for ( int i = 0; i < n; ++i ) { v[i] += rhs[i];
matvec( v, x, a, i, n ); } omp end parallel
#pragma acc parallel loop \
device_type(nvidia) num_gangs(200) \ device_type(radeon) num_gangs(400) …
for ( int i = 0; i < n; ++i ) { v[i] += rhs[i]; matvec( v, x, a, i, n ); } omp デバイス毎にチューニング 複数の異なるtarget 用に clause を指示する device_type(target) … 例えば、NVIDIA用と AMD Radeon用に mapping を変える
Global dataのハンドリング
float a[1000000];
#pragma acc routine worker extern void matvec(…);
…
#pragma acc parallel loop for( i = 0; i < m; ++i ) { matvec( v, x, i, n ); }
extern float a[];
#pragma acc routine worker
void matvec ( float* v, float* x,int i, int n ) { #pragma acc loop worker
for( int j = 0; j < n; ++j ) x[i] += a[i*n+j]*v[j]; }
分離コンパイル機能の追加により
ルーチン間に跨る「グローバル変数」の データ管理をサポートする必要がある。declare文を使う４つの方法
① Global Static dataのハンドリング declare文
float a[1000000];
#pragma acc declare create(a) #pragma acc routine worker extern void matvec(…);
…
#pragma acc parallel loop for( i = 0; i < m; ++i ) { matvec( v, x, i, n );
}
extern float a[];
#pragma acc declare create(a) #pragma acc routine worker
void matvec ( float* v, float* x,int i, int n ) { #pragma acc loop worker
for( int j = 0; j < n; ++j ) x[i] += a[i*n+j]*v[j]; } 広域変数を宣言した後にdirectiveを指定 declare create(a) ・static dataをホストとデバイス両方に生成 ・ホストとデバイス間の明示的なデータ同期は update directive を使用して維持する declare文は routineの「宣言セクション」で使用 プログラム、ルーチン内のデータ領域等 で implicit に visible copy を生成
② Global Pointer のハンドリング declare文
float *a;
#pragma acc declare create(a) #pragma acc routine worker extern void matvec(…);
…
#pragma acc parallel loop \ copyin(a[0:n*m])
for( i = 0; i < m; ++i ) { matvec( v, x, i, n );
}
extern float *a;
#pragma acc declare create(a) #pragma acc routine worker
void matvec ( float* v, float* x,int i, int n ) { #pragma acc loop worker
for( int j = 0; j < n; ++j ) x[i] += a[i*n+j]*v[j]; } ポインタを宣言した後にdirectiveを指定 declare create(a) ・staticにポインタをホストとデバイス両方に生成 ・copyはユーザの責任で行う ・data clause中にcopy等を明示するとデバイス にデータ割付とコピーを行うコードを生成 ホスト側の割付、データ管理は ユーザの責任で行う
③ Global dataのハンドリング declare link節
float a[1000000];
#pragma acc declare link(a) #pragma acc routine worker extern void matvec(…);
…
#pragma acc parallel loop \ copyin(a[0:n*m])
for( i = 0; i < m; ++i ) { matvec( v, x, i, n );
}
extern float a[];
#pragma acc declare link(a) #pragma acc routine worker
void matvec ( float* v, float* x,int i, int n ) { #pragma acc loop worker
for( int j = 0; j < n; ++j ) x[i] += a[i*n+j]*v[j]; } link ()の引数は必ず、global 変数 declare link(a) ・ホスト上にデータ生成 ・デバイス上のデータへのリンク生成 ・「data clause」中にcopy等を明示する時だけ デバイスにデータ割付とコピーを行うコードを 生成
④ device_resident in declare文
float a[1000000];
#pragma acc declare device_resident(a) #pragma acc routine worker
extern void matvec(…); …
#pragma acc parallel loop for( i = 0; i < m; ++i ) { matvec( v, x, i, n );
}
extern float a[];
#pragma acc declare device_resident(a) #pragma acc routine worker nohost
void matvec ( float* v, float* x,int i, int n ) { #pragma acc loop worker
for( int j = 0; j < n; ++j ) x[i] += a[i*n+j]*v[j]; } declare device_resident(a) ・static dataをデバイス上のみに生成 ・ホスト側には a[ ] は生成されない 他に、 declare link( … ) 変数を宣言した後にdirectiveを指定 device上のみで動作、“nohost”が必要
“enter data” と “exit data” 構文
#pragma acc data copyin(a[0:n]) \ create(b[0:n]) {
… 構造化ブロック ….. }
#pragma acc enter data copyin(a[0:n]) \ create(b[0:n])
…
#pragma acc exit data delete(a[0:n]) …
#pragma acc exit data copyout(b[0:n]) データ領域の開始 ”enter” データ領域の終了 ”exit”
Data lifetime
を自由に設定する
OpenACC 1.0 では、「構造化ブロック」の前後 にデータ領域を指定するのみを任意の場所に
Unstructured data lifetimes
#pragma acc data copyin(a[0:n]) \ create(b[0:n]) { … 構造化ブロック ….. } void init (…) { …
#pragma acc enter data copyin(a[0:n]) \ create(b[0:n]) … } … void fini (…) { …
#pragma acc exit data copyout(b[0:n]) … } init ルーチンの中で ”enter”
Data lifetime
を自由に設定する
任意の場所に! fini ルーチンの中で ”exit”新しい data API ルーチンの提供
• acc_copyin ( ptr, bytes ) • acc_create ( ptr, bytes ) • acc_copyout ( ptr, bytes ) • acc_delete ( ptr, bytes ) • acc_is_present ( ptr, bytes ) • acc_update_device ( ptr, bytes ) • acc_update_local ( ptr, bytes ) • acc_deviceptr ( ptr ) • acc_hostptr ( devptr )• acc_map_data ( devptr, hostptr, bytes ) • acc_unmap_data ( hostptr )
data directive 構文の copy 節等の機能が、
APIルーチンとして提供
loop構文
OpenACC並列3階層をより厳密に整理
gangループ ： gangループを含んではならない ： 一番外側ループへの適用とする workerループ ： gang、workerループを含んではならない vectorループ ： gang、worker、vectorループを含んではならない ： 一番内側ループへの適用とする autoの新設 ： コンパイラが選択するgang / worker / vector の実行モードの用語を明確に定義
gang redundant vs. gang partitioned modeworker single vs. worker partitioned mode vector single vs. vector partitioned mode
Nested loops における tiling
!$acc parallel !$acc loop gang do i = 1, n
!$acc loop vector do j = 1, m
a(j,i) = (b(j-1,i)+b(j+1,i)+ & b(j,i-1)+b(j,i+1))*0.25 enddo
enddo
!$acc end parallel arallel
!$acc parallel
!$acc loop tile(64, 4) gang vector do i = 1, n
do j = 1, m
a(j,i) = (b(j-1,i)+b(j+1,i)+ b(j,i-1)+b(j,i+1))*0.25 enddo
enddo
!$acc end parallel
ループのタイル化
64 x 4 タイルのworker構成を gang 間で並列に実行 tile clause (64 x 4) 内側ループ (j) 外側ループ (i) • 2次元における分割化指示 • cacheの有効利用OpenACC API 機能向上
•
様々な OS 、様々なホスト、様々なデバイス上で使用可能とする
•
出来得る限り、実行部分を「デバイス」側で可能にする
•
プログラミングの自由度を高めてゆくこと
The OpenACC™ Application Programming Interface (OpenACC API) for offloading programs written in C, C++ and Fortran programs
from a host CPU to an attached accelerator device
PGI OpenACC 機能ロードマップ(予定）
2014
2012 7月 2013 7月
CUDA 4.2 CUDA 5.0 CUDA 5.ｘ
Kepler AMD Radeon NVIDIA
Maxell Xeon Phi Unstructured Data lifetimes OpenACC 2.0 AMD Radeon/APU Acc Func Call Nested Acc Compute OpenACC 1.0 Xeon Phi OpenACC 2.0 NVIDIA OpenACC AMD Radeon ACC Fortran texture Multi-device improve C++ OpenACC OpenACC 2.0 仕様決定 Code Generation Tuning
2.0 取込 対応 プラットフォーム ※予告なしに変更することがあります AMD APU 機能
