スライド 1

Transcription

1 GTC Japan 2013 PGI Accelerator Compiler 新 OpenACC 2.0 の機能と PGI アクセラレータコンパイラ 2013 年 7 月 加藤努株式会社ソフテック

2 本日の話 OpenACC ディレクティブで出来ることを改めて知ろう! OpenACC 1.0 の復習 ディレクティブ操作で出来ることを再確認 OpenACC 2.0 の新機能 プログラミングの自由度の向上へ PGI Accelerator Compiler 2013 PGI Accelerator Programming Model 準拠 OpenACC 1.0 準拠 OpenACC 2.0 の機能は 毎月のリビジョンで順次追加 プラス PGI Extension( 先取り機能 ) 1

3 PGI Accelerator / OpenACC Release Map PGI Accelerator Programming Model 公開 (OpenACC の原型モデル ) PGI Accelerator Compiler リリース (PGI 2010) OpenACC 1.0 準拠 PGI Compiler リリース (PGI 2012) PGI 提供開始 OpenACC 2.0 順次取込 (PGI 2013) NVIDIA CUDA release OpenACC API 1.0 release OpenACC API 2.0 final PGI は 2009 年より提供開始 OpenMP 4.0 with extensions for Accelerators 2

4 OpenACC Programming Model オフロードモデル 重い処理部分 をアクセラレータ側で処理するモデル ホスト側 CPU 重い計算部分の処理をオフロード Accelerator 側 Device Main Memory Device Memory 並列処理によるスループット改善 1 処理の並列化 ( 分割指示 ) 2 データの移動管理 3

5 OpenACC ディレクティブの主な構成 ホスト ( 処理 ) Accelerator 1 CPU 重い計算部分の処理をオフロード 3 GPGPU Main Memory 2 ( データ ) Device Memory 1 Accelerate Compute 構文 (offload 領域指示 ) 2 Data 構文 ( データ移動指示 ) 3 Loop 構文 (Mapping for parallel/vector, Tuning) 4

6 OpenACC 1.0 ディレクティブの復習 まだ ハードルが高いあなたに! 1 どこをオフロードの対象とするか? コンパイラに対する指示 ユーザが行う最初のステップ 2 データ マネージメントの方法データの移動指示 コンパイラが自動的に行う ユーザによる明示的な指示ホスト上 デバイス上にコピー 3 並列処理の分割方法の指示 Mapping ( 並列チューニング ) の directives 後は コンパイラに任せる 5

7 1 並列化 ( オフロード ) 領域の対象は? Q. どこをオフロードの対象とするのか? ( どこに directive を入れるか ) A. 並列化可能な ループ ループの集合体 Single Loops for( i = 0; i < n; ++i ) a[i] = b[i] + c[i]; tightly Nested Loops OpenACC directive 挿入 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 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; Ajacent( 隣接 ) 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]; 6

8 プログラム シーケンス上でのオフロード対象 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) 外側繰り返しループ OpenACC directive Nested loop ( 並列化可能 ) OpenACC directive Single loop ( 並列化可能 ) OpenACC directive Single loop ( 並列化可能 ) 外側繰り返しループ 7

9 並列化 ( オフロード ) できないループ 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 enddo call 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 ループ内に Function call K ループはオフロード化できない Nested loop ( 並列化可能 ) ループ内に Procedure Call があると並列化出来ない Nested loop ( 並列化可能 ) 並列化を行うには インライン展開するしかなかった! OpenACC API 1.0 では並列化できない NVIDIA CUDA 5.0 から GPU カーネルのリンカー機能が追加されたことのより OpenACC 2.0 で可能となる 8

10 2 OpenACC におけるデータ移動の基本 ホストとデバイス 間のデータ交換のタイミングポイントは? コンパイラのデフォルト動作 OpenACC 並列化構文 (acc kernels や acc parallel) の始点と終点でその領域内で使われている配列 スカラ変数を解析し ホスト と デバイス 間のデータコピーを行うコードを暗黙に自動生成!$acc 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 end kernels!$acc parallel Gang Gang Gang!$acc end parallel ユーザが明示的にコピー配列名を指示 (copy clause) することができる live スカラ変数は デフォルトで copyin/copyout されている ( スカラ変数の移動管理は明示的に必要なし ) 9

11 データ マネージメントの基本的考え方 ホスト と デバイス 間のデータの移動を最小化すること プログラミング上の鉄則 デバイス上で可能な限りのデータ常駐化 ホスト側から一度 デバイスへコピー (copyin) された配列データは デバイス側の計算で役目が終わらない限り ホスト側へデータ (copyout) を戻さない ユーザが明示的にデータ移動を管理するための構文 Data 構文 directive の活用 Fortran!$acc data [clause ]...!$acc end data C #pragma acc data [clause ] {... Data 構文で指示された配列は data 構文領域の中で常駐化する 10

12 デバイス内にデータを常駐化する デバイス側の処理 Main routine (host 側 ) Host 側メモリ A(100) A!$acc data copy (a(100)) copyin copyout!$acc data present(a(100)) A A A call call call プログラムの流れ デバイス側メモリ resident A(100) time 2 A present と言う状態!$acc end data Data 構文 で挟まれたプログラム領域に適用できる 11

13 データの同期 update directive( 実行文 ) Host 側メモリ デバイス側メモリ A(100) copyin copyout resident A(100) present と言う状態!$acc data present(a(100)) デバイス側の処理 A A A Main routine (host 側 ) A call call call A A!$acc data copy (a(100)) 細かな操作が可能 Device to host!$acc update host (a(10:20)) Host to device!$acc update device (a(10:20))!$acc end data 12

14 実際のプログラムで 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(); function call Main ルーチン +Subroutine タイプ データの移動管理の directive OpenACC data directive Accelerator Data 領域管理下 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]; ( 以下 略 ) 並列領域指定の directive OpenACC directive Nested loop ( 並列化可能 ) 13

15 3 並列処理の対象と方法を指示する構文 アクセラレータ上で並列実行を行う領域を指定するための directive その領域内で使用する データ を転送するコードを生成 (Implicit xfer) 2 種類の Accelerator Compute 構文 Parallel 構文 (1 つのカーネル構成でコード生成 ) C #pragma acc parallel { 並列実行領域 Fortran!$acc parallel { 並列実行領域!$acc end parallel Kernels 構文 (tightly loops を対象 複数のカーネルで構成 ) C #pragma acc kernels { 並列ループ領域 Fortran!$acc kernels { 並列ループ領域!$acc end kernels 14

16 役目 並列処理分割の指示を行う Loop 構文 直下の ループ の並列の分割方法 (gang, worker, vector) 指示 Acceleratorハードウェアの並列構造にマッピングを行うためのclause ループ内プライベートとなる変数の宣言 リダクション処理の指定 Fortran C!$acc loop [clause ] #pragma acc loop [clause ] Clause( 節 ) collapse (n) gang [ ( expression ) ] worker [ ( expression ) ] vector [ ( expression ) ] seq independent private( list ) reduction ( operator : list) 15

17 OpenACC Parallelism 3 階層 Gang Worker Vector PEs : プロセッサ要素の基本集合体 each PE : マルチスレッディング処理 Each thread of the PE : ベクトル命令処理 あまり深刻に考えなくてよい = H/W とのマッピングが難しい PE worker thread Multithreaded Vector 命令 ( 長さ ) NVIDIA case Gang Worker PEs => Streaming Multiprocessors thread Multithreading => warps within SM (PE 内の Max. Threads per Thread Block) Vector_length => warp 内スレッド数 32 16

18 並列処理分割の指示 loop 構文 #pragma acc kernels #pragma acc loop independent ( 並列対象指示 ) for (i = 1; i < n+1; i++) { #pragma acc loop independent Loop 構文 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 構文 161, Loop is parallelizable Accelerator kernel generated 161, #pragma acc loop gang /* blockidx.x */ 165, #pragma acc loop vector(128) /* threadidx.x */ 17

19 並列分割の指示 loop 構文 #pragma acc kernels #pragma acc loop independent gang(32) ( 並列対象指示 ) for (i = 1; i < n+1; i++) { #pragma acc loop independent Loop 構文 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; 明示的に分割指示 161, Loop is parallelizable Accelerator kernel generated 161, #pragma acc loop gang(32) /* blockidx.x */ 165, #pragma acc loop vector(32) /* threadidx.x */ Accelerate Compute 構文 18

20 OpenACC API 2.0 in PGI Accelerator Compilers 19

21 OpenACC API 2.0 New features 手続きの呼び出し と 分離コンパイル ( 効果 ) プログラム移行性向上 ライブラリとしてモジュール化可能 入れ子 (Nested) 並列 ( 効果 ) parallel/kernels 構文を有するルーチン モジュールを構成可能 デバイス特有なチューニング, 複数のデバイス使用 ( 効果 ) デバイス種に依存しない 性能互換性 の向上 グローバル データへの対応 ( 効果 ) ルーチン間での 広域データ の使用の単純化 データ マネージメント自由度向上 ( 効果 ) unstructured data lifetimes( 自由にデータ領域を構築 ) 新しい Data API ルーチン ( より自由にデータ移動が可能 ) Loop 構文の parallelism 3 階層の明確化と auto tile 節の追加 ( 効果 ) 並列性 3 階層の整理 ユーザ チューニングの簡易化 非同期操作の向上 その他 20

22 現在 手続きの呼び出しがあると? #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 Accelerator Compute 領域の中の call 文を含むループは 並列化できない OpenACC 1.0 では 関数 手続きをソースレベルでインライン展開をするしかなかった OpenACC 2.0 では可能になる 21

23 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 1 手続きのプロトタイプ に指示 worker レベルの並列処理を指示する例 Function 内で worker レベルの並列処理を指示する場合は その上位の parallel/kernels 構文内の上位 loop の並列レベルは gang レベルである必要がある 同じ worker レベルは駄目 22

24 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 1 手続きのプロトタイプ に指示 worker レベルの並列処理を指示する例 2 手続き への routine 指示 分離コンパイルも可能 orphaned loop directive #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; 23

25 OpenACC 2.0 Nested Parallelism (1) #pragma acc routine extern void matvec(float* v,float* x, ); 1 #pragma acc parallel loop for ( int i = 0; i < n; ++i ) matvec( v, x, i, n );end parallel 実際は i 毎にデバイスの中のシングル スレッドで次々に開始する parallel or kernels 領域内のループ内から call function function 側でも parallel or kernels loop を構成できる #pragma acc routine void matvec( ) { 2 #pragma acc parallel loop for ( int i = 0; i < n; ++i ) { matvec 内で 指定された parallelism を launch する 24

26 OpenACC 2.0 Nested Parallelism (2) #pragma acc routine extern void matvec(float* v,float* x, ); 1 #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 一つのデバイス スレッドで次々に matvec を起動する #pragma acc routine void matvec( ) { 2 #pragma acc parallel loop for ( int i = 0; i < n; ++i ) { matvec 内で 指定された parallelism を launch する 25

27 異なるデバイスへの指示 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 複数の異なる target 用に clause を指示する device_type(target) デバイス毎にチューニング 例えば NVIDIA 用と AMD Radeon 用に mapping を変える #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 26

28 Global data のハンドリング float a[ ]; #pragma acc routine worker extern void matvec( ); #pragma acc parallel loop for( i = 0; i < m; ++i ) { matvec( v, x, i, n ); 分離コンパイル機能の追加により ルーチン間に跨る グローバル変数 のデータ管理をサポートする必要がある declare 文を使う 4 つの方法 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]; 27

29 1 Global Static data のハンドリング declare 文 float a[ ]; #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 ); declare 文は routine の 宣言セクション で使用 declareの機能プログラム ルーチン内のデータ領域等で implicit に visible copy を生成 広域変数を宣言した後に directive を指定 declare create(a) static data をホストとデバイス両方に生成 ホストとデバイス間の明示的なデータ同期は update directive を使用して維持する 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]; 28

30 2 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 ); ホスト側の割付 データ管理はユーザの責任で行う ポインタを宣言した後に directive を指定 declare create(a) static にポインタをホストとデバイス両方に生成 copy はユーザの責任で行う data clause 中に copy 等を明示するとデバイスにデータ割付とコピーを行うコードを生成 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]; 29

31 3 Global data のハンドリング declare link 節 float a[ ]; #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 ); link () の引数は必ず global 変数 declare link(a) ホスト上にデータ生成 デバイス上のデータへのリンク生成 data clause 中に copy 等を明示する時だけデバイスにデータ割付とコピーを行うコードを生成 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]; 30

32 4 device_resident in declare 文 float a[ ]; #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 ); 変数を宣言した後にdirectiveを指定 declare device_resident(a) static dataをデバイス上のみに生成 ホスト側には a[ ] は生成されない device 上のみで動作 nohost が必要 他に declare link( ) 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]; 31

33 enter data と exit data 構文 #pragma acc data copyin(a[0:n]) \ create(b[0:n]) { 構造化ブロック.. Data lifetime を自由に設定する OpenACC 1.0 では 構造化ブロック の前後にデータ領域を指定するのみ データ領域の開始 enter データ領域の終了 exit を任意の場所に #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]) 32

34 Unstructured data lifetimes #pragma acc data copyin(a[0:n]) \ create(b[0:n]) { 構造化ブロック.. init ルーチンの中で enter fini ルーチンの中で exit Data lifetime を自由に設定する 任意の場所に! void init ( ) { #pragma acc enter data copyin(a[0:n]) \ create(b[0:n]) void fini ( ) { #pragma acc exit data copyout(b[0:n]) 33

35 新しい data API ルーチンの提供 data directive 構文の copy 節等の機能が 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 ) 34

36 loop 構文 OpenACC 並列 3 階層をより厳密に整理 gang ループ : gang ループを含んではならない : 一番外側ループへの適用とする worker ループ : gang worker ループを含んではならない vector ループ : gang worker vector ループを含んではならない : 一番内側ループへの適用とする auto の新設 : コンパイラが選択する gang / worker / vector の実行モードの用語を明確に定義 gang redundant vs. gang partitioned mode worker single vs. worker partitioned mode vector single vs. vector partitioned mode an accelerator thread = a single vector lane of a single worker of a single gang 35

37 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 ループのタイル化 2 次元における分割化指示 cache の有効利用 64 x 4 タイルの worker 構成を gang 間で並列に実行 内側ループ (j) tile clause (64 x 4) 外側ループ (i)!$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 36

38 OpenACC API 機能向上 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 様々な OS 様々なホスト 様々なデバイス上で使用可能とする 出来得る限り 実行部分を デバイス 側で可能にする プログラミングの自由度を高めてゆくこと 今後も機能向上を続ける 37

39 PGI OpenACC 機能ロードマップ ( 予定 ) 予告なしに変更することがあります ACC Fortran texture 2.0 取込 OpenACC AMD Radeon OpenACC 2.0 AMD Radeon/APU 機能 Multi-device improve C++ OpenACC Code Generation Tuning Unstructured Acc Data Func Call lifetimes Nested Acc Compute OpenACC 1.0 Xeon Phi OpenACC 2.0 NVIDIA 対応プラットフォーム Kepler AMD Radeon OpenACC 2.0 仕様決定 AMD APU NVIDIA Maxell CUDA 4.2 CUDA 5.0 CUDA 5.x Xeon Phi 月 月

40 終わり 39

41 本ドキュメントに記述された各製品名は 各社の商標または登録商標です Copyright 2013 SofTek Systems Inc. All Rights Reserved. 40