Microsoft PowerPoint - GDEP-GPG_softek_May24-2.pptx

Transcription

1 G-DEP 第 3 回セミナー PGI OpenACC Compiler PGIコンパイラ使用の実際 新しい OpenACC によるプログラミング 2012 年 5 月 加藤努株式会社ソフテック OpenACC によるプログラミング GPU / Accelerator Computing Model のデファクト スタンダードへ OpenACC Standard 概略説明 Accelerator Programming Model OpenMP と OpenACC 何が違うのか? Accelerator 構成のボトルネックは? OpenACC directives 並列領域の定義データ転送の操作ループ分割の指示 Directives 構文の説明 三つの主要 directives を覚える データの移動指示を行うためのテクニック OpenACC 使用の実際例 1

2 アクセラレータと OpenACC Standard 2 アクセラレータとは Many Cores Multi threading ホスト 本日の前提 アクセラレータ NVIDIA GPU Intel MIC Others 3

3 Accelerator Programming Model ホスト ハイブリッド構成 (CPU+GPU) Accelerator CPU 重い計算部分の処理をオフロード GPGPU Main Memory Device Memory 今日の話 Accelerator のためのコード開発を支援するディレクティブ群 OpenACC directives 4 OpenACC Standard とは何か? 2011 年 11 月 NVIDIA, Cray, PGI,CAPS Accelerators 用のプログラミング API の標準仕様 Fortran, C/C++ 言語上で指定するコンパイラ ディレクティブ群 ユーザサイド開発者がアクセラレータで実行するコード部分をディレクティブで指定する ( コンパイラに対して ヒントを与える ) OpenACC コンパイラホスト側の処理をアクセラレータ (GPU) にオフロードするコード生成ホスト -- GPU 間のデータ転送コードの生成 2009 年リリース以来 実績を積んだ PGI Accelerator Compiler(directives) がベースとなっている 5

4 OpenACC の利点 特長 可搬性 ( ポータビリティ ) ホスト OS, ホスト CPUs, アクセラレータ コンパイラ間の portability 異なるプラットフォームに変わっても 性能 の可搬性を提供 ハイレベル / 高級言語 CUDA, OpenCL or 低レベル GPU programming の知識必要なしソースコードの変更が必要ない 高生産性クイックな開発ステップと段階的に行える開発 OpenMP-like ヘテロジニアスなプログラムシングル ソース ベース (CPU 用 GPU 用と分ける必要なし ) Easy, Fast, Portable! 6 OpenACC は どんな機能を提供するのか? プログラム内に ディレクティブ でヒントを与えると コンパイラは 次のような機能を組込んだホスト側コードとアクセラレータ側コードを自動生成する アクセラレータ (GPU) をイニシャライズ ホスト (CPU) とアクセラレータ間の データ と プログラム の転送の管理 ホスト (CPU) とアクセラレータ間の データ を監視 ホスト (CPU) とアクセラレータ間の 仕事 を管理 仕事 を並列に分割して GPU ハードウェアの並列構造にマッピング さらに 性能最適化を試みる ( キャッシュ等 ) 7

5 OpenMP の取組 2012/3/27 アナウンス OpenACC 知見妥当性確認ユーザの声反映 取り込み OpenMP 4.0 新バージョン (2012 年 ) 8 OpenMP と OpenACC プログラミング モデル 9

6 OpenMP Programming Model マルチコア 共有メモリ上のスレッド並列 (CPU) ホスト マルチコア CPU CPU 処理の分割のみ Main Memory Main Memory データは全て共有メモリ内に 処理の分割のみを考えれば良い データは均一のアドレス空間上でアクセス可能 10 Accelerator Programming Model ホスト側 ハイブリッド構成 (CPU+GPU) GPU 側 CPU Main Memory Host_A(100) 重い計算部分の処理をオフロード 使用データを送る 結果データを戻す Overhead GPGPU Device Memory Device_A(100) Host GPU 間のメモリデータの転送が伴う データ転送のオーバーヘッド時間が伴う 11

7 並列化で必要なプログラム記述 共有メモリ構成 (OpenMP on Host) アクセラレータ構成 (Host+GPU) 処理系 処理の並列化 処理の並列化 (Kernel 記述 ) データ系 なし データの移動 12 Accelerator Memory model の重要な視点 アクセラレータ モデルでは 何を考える必要があるか? ホスト側のデータ空間とGPU 側のデータ空間の二つの空間が必要 ホストとGPU 間のデータ転送が必要 新たなoverhead 時間 データ移動に係わる指示 (directive) が必要 データ移動の最小化を行うために最適化するための要素 ホスト側メモリとGPU 側メモリ間の転送帯域 (PCI 帯域 ) に制約がある 性能加速性があるコード領域かどうかの判別が必要 ループ内の 計算密度 のレベルで加速性が変化 GPU 側のメモリ容量に制約がある ( 現在 < 6GB) 非常に大きな配列を伴う コード領域 のオフロード化に制約 必要な場合 複数のGPUデバイス上のMPI 化が必要 13

8 CPU+Accelerator 構成のボトルネックは? ホスト側 ハイブリッド構成 (CPU+GPU) GPU 側 CPU Main Memory 処理の offload 化 GPGPU Device Memory 2 データの再利用 resident 1 データ転送回数を少なくする工夫 Host GPU 間のデータ転送帯域 (PCI バス性能 ) の遅さ 14 $ pgaccelinfo PGI command pgaccelinfo -- CUDA Driver Version: 4010 (CUDA driver version) NVRM version: NVIDIA UNIX x86_64 Kernel Module Device Number: 0 Device Name: Tesla C2075 Device Revision Number: 2.0 Global Memory Size: Number of Multiprocessors: 14 Number of Cores: 448 Concurrent Copy and Execution: Yes Total Constant Memory: Total Shared Memory per Block: Registers per Block: Warp Size: 32 Maximum Threads per Block: 1024 Maximum Block Dimensions: 1024, 1024, 64 Maximum Grid Dimensions: x x Maximum Memory Pitch: Texture Alignment: Clock Rate: Execution Timeout: Integrated Device: Can Map Host Memory: B 512B 1147 MHz Yes No Yes GPU device の特性の出力コマンド CUDA driver のチェック GPU ボードのチェック 性能要素のチェック CUDA 機能チェック PCI バス性能チェック 15

9 PGI command pgaccelinfo ( 続 ) Compute Mode: default Concurrent Kernels: Yes ECC Enabled: No Memory Clock Rate: 1566 MHz Memory Bus Width: 384 bits L2 Cache Size: bytes Max Threads Per SMP: 1536 Async Engines: 2 Unified Addressing: Yes Initialization time: 7325 microseconds Current free memory: Upload time (4MB): 821 microseconds ( 651 ms pinned) Download time: 948 microseconds ( 649 ms pinned) Upload bandwidth: 5108 MB/sec (6442 MB/sec pinned) Download bandwidth: 4424 MB/sec (6462 MB/sec pinned) CPU memory GPU memory 間の転送帯域 = PCI Express x16 高々 5GB/sec しかない! ちなみに 現在の CPU のメモリ帯域は 20~50GB/sec 16 OpenACC ディレクティブ 17

10 program main integer :: n! size of the vector real,dimension(:),allocatable :: a! the vector real,dimension(:),allocatable :: r! the results integer :: i n = allocate(a(n)) allocate(r(n)) do i = 1,n a(i) = i*2.0!$acc kernels do i = 1,n r(i) = sin(a(i)) ** 2 + cos(a(i)) ** 2!$acc end kernels print *, r(1000) end 2 行のディレクティブ挿入でコード生成 $ pgfortran -acc -Minfo test.f90 main: 12, Generating copyout(r(1:100000)) Generating copyin(a(1:100000)) Generating compute capability 1.0 binary Generating compute capability 2.0 binary 13, Loop is parallelizable Accelerator kernel generated 13,!$acc loop gang, vector(256)! blockidx%x threadidx%x オフロードする並列対象領域の指 ( 一般に ループ部分 ) GPU 側へのデータコピー GPU 用の並列化 Host 側へデータバック 自動的 かつ Implicit に行う 18 Kernel と言う言葉!$acc kernels do i = 1,n r(i) = sin(a(i)) ** 2 + cos(a(i)) ** 2!$acc end kernels アクセラレータへオフロードする並列対象 ( 一般に ループ構成部分 ) Accelerator(GPU) 用の実行サブ ルーチンとして切り出す ひとつの kernel ルーチン という 19

11 OpenACC ディレクティブの主な構成 ホスト CPU 3 ( 処理 ) 1 重い計算部分の処理をオフロード Accelerator GPGPU Main Memory 2 ( データ ) Device Memory 1 Accelerate Compute 構文 (offload 領域指示 ) 2 Data 構文 ( データ移動指示 ) 3 Loop 構文 (Mapping for parallel/vector, Tuning) 20 1 Accelerate Compute 構文 program main integer :: n! size of the vector real,dimension(:),allocatable :: a! the vector real,dimension(:),allocatable :: r! the results integer :: i n = allocate(a(n)) allocate(r(n)) do i = 1,n a(i) = i*2.0!$acc data copyin(a(1:n)),copyout(r)!$acc kernels!$acc loop gang(32),vector(64) do i = 1,n r(i) = sin(a(i)) ** 2 + cos(a(i)) ** 2!$acc end kernels!$acc end data -- Fortran -- 主要な三つのディレクティブ 1 並列実行 kernel 部分の指定 オフロードする並列対象領域の指 ( 一般に ループ部分 ) 21

12 2 Data 構文 program main integer :: n! size of the vector real,dimension(:),allocatable :: a! the vector real,dimension(:),allocatable :: r! the results integer :: i n = allocate(a(n)) allocate(r(n)) do i = 1,n a(i) = i*2.0!$acc data copyin(a(1:n)),copyout(r)!$acc kernels!$acc loop gang(32),vector(64) do i = 1,n r(i) = sin(a(i)) ** 2 + cos(a(i)) ** 2!$acc end kernels!$acc end data -- Fortran -- 主要な三つのディレクティブ 2データ移動指示 1 並列実行 kernel 部分の指定 オフロードする並列対象領域の指 ( 一般に ループ部分 ) 22 3 Loop 構文 program main integer :: n! size of the vector real,dimension(:),allocatable :: a! the vector real,dimension(:),allocatable :: r! the results integer :: i n = allocate(a(n)) allocate(r(n)) do i = 1,n a(i) = i*2.0!$acc data copyin(a(1:n)),copyout(r)!$acc kernels!$acc loop gang(32),vector(64) do i = 1,n r(i) = sin(a(i)) ** 2 + cos(a(i)) ** 2!$acc end kernels!$acc end data -- Fortran -- 主要な三つのディレクティブ 2データ移動指示 1 並列実行 kernel 部分の指定 3mapping for para/vector オフロードする並列対象領域の指 ( 一般に ループ部分 ) 23

13 int main( int argc, char* argv[] ) { int n; /* size of the vector */ float *restrict a; /* the vector */ float *restrict r; /* the results */ float s, c; int i; n = ; a = (float*)malloc(n*sizeof(float)); r = (float*)malloc(n*sizeof(float)); for( i = 0; i < n; ++i ) a[i] = (float)(i+1) * 2.0f; #pragma data copy (a[0:n]), copyout(r) #pragma acc kernels { #pragma acc loop gang, vector(128) } for( i = 0; i < n; ++i ){ s = sinf(a[i]); c = cosf(a[i]); r[i] = s*s + c*c; } C プログラムの場合 オフロードする並列対象領域の指 ( 一般に ループ部分 ) -- C-- 主要な三つのディレクティブ 2 データ移動指示 1 並列実行 kernel 部分の指定 3mapping for para/vector 24 三つの構文を覚える!$acc data copyin(a(1:n)),copyout(r)!$acc kernels!$acc loop gang(32),vector(64) do i = 1,n r(i) = sin(a(i)) ** 2 + cos(a(i)) ** 2!$acc end kernels!$acc end data Data 構文 Accelerate Compute 構文 Loop 構文 オフロードする並列対象領域 ( 一般に ループ部分 ) 25

14 OpenACC ディレクティブの表記法 Fortran!$acc directive-name [clause [,] clause] ] ( 構造化コード ブロック )!$acc end directive-name Clause= 節 ] ( 他の表記法 ) c$acc directive-name [clause [,] clause] ] *$acc directive-name [clause [,] clause] ] C #pragma acc directive-name [clause [,] clause] ] { 構造化コード ブロック } PGI Accelerator directive と同じ 26 Accelerate Compute 構文 アクセラレータ上で並列実行を行う領域を指定するための directive その領域内で使用する データ を転送するコードを生成 (Implicit xfer) 2 種類のディレクティブ Parallel 構文 ( 指定領域を並列に実行する ) C #pragma acc parallel { 並列実行領域 } Fortran!$acc parallel { 並列実行領域 }!$acc end parallel Kernels 構文 ( 複数の並列カーネルループを実行 ) C #pragma acc kernels { 並列ループ領域 } Fortran!$acc kernels { 並列ループ領域 }!$acc end kernels 開発初期の段階は まず kernels を使用する 27

15 Parallel 構文 OpenMP の!$omp parallel と同じコンセプト CPU スレッド 数を生成する 個々のworkerが冗長に実行 Any work-sharing loop gangs of workers 数を解放!$acc parallel Gang!$acc end parallel Gang Gang Multithreaded Workers gangs of workers 並列領域指定された!$acc parallel [clause ] if( condition ) private( list ) async( expression ) firstprivate( list ) num_gangs( expression ) reduction( operator:list ) num_workers( expression ) Any data clause vector_length( expression ) CPU スレッド 28 Parallel 構文の clauses( 節 ) num_gangs ( expression ) num_workers ( expression ) 生成する gang の数 (CUDA griddim) 各 gang 内で生成する worker の数 (CUDA blockdim) vector_length ( expression ) 各 worker 内の vector 長 (SIMD 実行 ) private ( list ) firstprivate ( list ) reduction ( operator:list ) 各 gang 上でプライベート変数の生成 ホスト側からの初期値をプライベート変数にセットする gang 間でプライベート コピーを使ってリダクション処理を行う 29

16 順番に実行Kernels 構文 CUDAの kernelの単位 と同じコンセプト (nested) ループを対象各ループをkernel 化する CPU スレッド コンパイラは 指定領域内のループ構造を独立のカーネルコードとして生成する!$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) Kernel 1 Gangs do i=1,n s=s+ a(i,j.k) Kernel 2!$acc kernels [clause ] If( condition ) async( expression ) Any data clause!$acc end kernels CPU スレッド Workers 30 OpenACC Parallelism 3 階層 Gang Worker Vector PEs : プロセッサ要素の基本集合体 each PE : マルチスレッディング処理 Each thread of the PE : ベクトル命令処理 あまり深刻に考えなくてよい = H/W とのマッピングが難しい worker thread PE Multithreaded Vector 命令可能 Gang Worker NVIDIA case PEs => Streaming Multiprocessors thread Multithreading => warps within SM (PE 内の Max. Threads per Thread Block) Vector_length => warp 内スレッド数 32 31

17 Gang 分割 + Worker or Vector 例えば 2 次元配列空間 ( 計算空間 )A(n,n) J 列で Gang 分割 j アクセラレータ H/W の並列実行体にマッピング i SM Worker の形 ( 分割法で異なる ) Vector length NVIDIA の場合は gang/vector の 2 階層でよい 同一 gang 内で worker or vector 特性によりマルチスレッディング処理 32 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) 33

18 parallel 構文内の Loop clauses( 節 ) collapse ( n ) gang worker vector seq Independent private ( list ) reduction ( operator:list ) 次の n の段数の nested ループに collapse を適用 並列領域の gangs 数によって イテレーションを分割引数はなし Parallel 構文の節で指定する gang 内の workers 数によって イテレーションを分割引数はなし Parallel 構文の節で指定する SIMD モードでイテレーションを実行 gang or worker のどちらかと共に使用する ( 但し 実装依存あり ) GPU 内でループをシーケンシャルに実行 次のループは依存性がないことを指示する ループの各イテレーションに対して list の変数のコピーを生成し プライベート変数を生成 イテレーション間でプライベート コピーを使ってリダクション処理を行う PGI の場合は 多くのリダクション演算を 動認識可能 34 kernels 構文内の Loop clauses( 節 ) collapse ( n ) gang ( expression ) worker ( expression ) vector ( length ) seq Independent private ( list ) reduction ( operator:list ) 次の n の段数の nested ループに collapse を適用 並列領域の gangs によって イテレーションを分割 gang 内の workers によって イテレーションを分割 SIMD モードでイテレーションを実行 length で strip-mining( 分割 ) する GPU 内で対象ループをシーケンシャルに実行 次のループは依存性がないことを指示する ループの各イテレーションに対して list の変数のコピーを生成し プライベート変数を生成 イテレーション間でプライベート コピーを使ってリダクション処理を行う PGI の場合は 多くのリダクション演算を 動認識可能 35

19 Combined Directives (parallel + loop) の複合ディレクティブ ( 指定した直下のループを対象とする ) parallel loop C #pragma acc parallel loop Fortran!$acc parallel loop do loop [!$acc end parallel loop ] kernels loop C #pragma acc kernels loop Fortran!$acc kernels loop do loop [!$acc end kernels loop ] 36 役目 Data 構文 Acceleratorのデバイスメモリ上に必要な配列 スカラ変数エリアを割付 Region 入口で ホストからデバイスへデータをコピー Region 出口で デバイスからホストへデータをコピーバック 明示的な指示のために使用 データのコピー動作発生 end data 文指定の場所で Host 側へのコピーバック発生 region entry data 文指定の場所で!$acc data デ!$acc kernels ータ領!$acc kernels 域region exit CPU スレッド!$acc end kernels!$acc end kernels!$acc end data CPU スレッド A(10) A(10) A(10) CPU 側 Memory Device 側 Memory CPU 側 Memory 37

20 Data 構文の clause ( 節 ) コピーの動作方法 データの存置状況を指示するための clause の意味 Host Memory 空間 Accelerator Device memory Memory 空間!$acc data ータ領域!$acc end data Host CPUスレッドデCPU スレッド present create A(10) A(10) A(10) copyin copyout copy 38 Data 構文の clause ( 節 ) Fortran C!$acc data [clause ]...!$acc end data #pragma acc data [clause ] Clause( 節 ) copy(list) copyin(list) copyout(list) create (list) present (list) present_or_copy(list) present_or_copyin(list) present_or_copyout(list) present_or_create(list) deviceptr(list) copyin + copyout 入口で hostからaccelerator 側にコピーするだけ出口で acceleratorからhostにデータをコピーバック accelerator 側のみで使用するローカル変数の割付指示既にaccelerator 側に存在していることを指示 accelerator 側にデータの存在のテストを行う (pcopy) accelerator 側にデータの存在のテストを行う (pcopyin) accelerator 側にデータの存在のテストを行う (pcopyout) accelerator 側にデータの存在のテストを行う (pcreate) list がデバイス側のポインタであることを宣言 39

21 データ転送を行う場所を意識する データ転送の最小化 ループの繰り返し毎に 以下の 1 と 2 のデータ転送が繰り返される 1Host > GPUへデータを転送するポイント Implicit Data xfer Accelerator(GPU) 上で Kernel code 実 2GPU -> Host へデータを転送するポイント 転送時間が大きくなり 性能が出ない! 40 Data 構文を使用する意味とは? データ転送 と 並列処理 を分ける Explicit Data xfer 1Host > GPUへデータを転送するポイント loop の繰り返し中 A,B,C 配列 ( 変数 ) は GPU 上のメモリに常駐する このタイミングではデータ転送は行わないコードが作成される loop の繰り返し終了後 A,B,C 配列 ( 変数 ) をホスト側へ戻す 2GPU -> Host へデータを転送するポイント 例えば A, B, C 配列の場合 41

22 プロシジャー間 Accelerator 側データ常駐化 時間積分ループの前で GPU 側へデータを転送する ( 配列を割付する ) データ転送が 1 回のみ (x, y, z 配列 )= データを常駐化させる Call されたルーチン側で アクセラレータ処理 x,y,z 配列データは GPU 上のデータを使用する ホスト側へデータ バック転送 Data 構文 present 節の利用 42 Data 構文 present 節の意味 Host 側メモリ GPU 側メモリ A(100) copyin copyout resident A(100) present と言う状態 Sub routine (GPU 上で計算 ) Main routine (host 側 ) A!$acc data copy (a(100))!$acc data present(a(100)) A A A call call call プログラムの流れ time A!$acc end data デバイス側のメモリ内の既存データを使用する 43

23 データの同期 update directive( 実行文 ) Host 側メモリ GPU 側メモリ A(100) copyin copyout resident A(100) present と言う状態 Sub routine (GPU 上で計算 )!$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)) Subarray 指定も可!$acc end data 44 Data 構文の活用 (C プロシジャー間 ) present 節の活用この時点で a をGPU 側で割付 かつ データコピー #pragma acc data copy(a[0:n]) { init( a, n );... 1 Main process( a, n ); のルーチンでも a を使用 } 既に a はGPU 上に存在して void init( float* a, int n ) いることをcompilerに伝える Subroutine { データ転送は行わない #pragma acc kernels loop present(a[0:n]) for( int i = 0; i < n; ++i ) a[i] = sinf((float)i); } Accelerator(GPU) 上のメモリにデータを常駐化させる 45

24 Data 構文 Present 節 (C 言語 ) typedef float *restrict *restrict MAT; typedef float *restrict VEC; main() { MAT aa; VEC bb[100]; ( 略 ) ( 109) ( 110) { ( 111) smooth( aa, bb, w0, w1, w2, n, m, 100 ); ( 112) } ( 略 ) Callee 側 Caller 側 引数 プロシジャー間のデータの引き渡し方法 Caller 側で Accelerator 側の Memory 割付を行い 必要なデータをコピーしておく ( 15) void smooth( MAT a, VEC b[100], float w0, float w1, float w2, int n, int m, int niters ) ( 16) { ( 21) for( iter = 1; iter < niters; ++iter ) { ( 22)... ( 並列化対象ループ )a[i][j], b[i][j] を使用 } ( 39) } 46 Data 構文 Present 節 (C 言語 ) 続き typedef float *restrict *restrict MAT; typedef float *restrict VEC; main() { MAT aa; VEC bb[100]; ( 略 ) ( 109) #pragma acc data copyout (aa[1:n-2][0:m]), copy(bb[0:n][0:n]) ( 110) { ( 111) smooth( aa, bb, w0, w1, w2, n, m, 100 ); ( 112) } ( 略 ) Callee 側 引数 Compiler message main: 109. Generating copyout(aa[1:n-2][0:m]) Generating copy(bb[0:n][0:n]) Visible device copy smooth: 20. Generating present(b[0:n][0:n]) Generating present(a[1:n-2][0:m]) ( 15) void smooth( MAT a, VEC b[100], float w0, float w1, float w2, int n, int m, int niters ) ( 16) { ( 20) #pragma acc data present ( a[1:n-2][0:m], b[0:n][0:n] ) ( 21) #pragma acc kernels { ( 22) for( iter = 1; iter < niters; ++iter ) { 既にGPU 上に割付られたa[],b[] を使用する... } } ( 39) } 47

25 Data 構文 Present 節 (C 言語 ) 続き void smooth( MAT a, VEC b[100], float w0, float w1, float w2, int n, int m, int niters ) { int i, j, iter; } #pragma acc data present ( a[1:n-2][0:m], b[0:n][0:n] ) #pragma acc kernels { for( iter = 1; iter < niters; ++iter ) { #pragma acc loop gang(8) vector(8) for( i = 1; i < n-1; ++i ) { for( j = 1; j < m-1; ++j ) { a[i][j] = w0 * b[i][j] + w1*(b[i-1][j] + b[i+1][j] + b[i][j-1] + b[i][j+1]) + w2*(b[i-1][j-1] + b[i-1][j+1] + b[i+1][j-1] + b[i+1][j+1]); } } for( i = 1; i < n-1; ++i ) { for( j = 1; j < m-1; ++j ) { b[i][j] = a[i][j]; } } } } smooth: 20, Generating present (b[0:n][0:n]) Generating present (a[1:n-2][0:m]) 25, Loop is parallelizable 26, Loop is parallelizable Accelerator kernel generated 25, #pragma acc loop gang(8), vector(8) /* blockidx.y threadidx.y */ 26, #pragma acc loop gang, vector(8) /* blockidx.x threadidx.x */ 32, Loop is parallelizable 33, Loop is parallelizable Accelerator kernel generated 32, #pragma acc loop gang, vector(8) /* blockidx.y threadidx.y */ 33, #pragma acc loop gang, vector(8) /* blockidx.x threadidx.x */ 48 Data 構文 Present 節 (C 言語 ) 続き 実行プロファイル結果 Accelerator Kernel Timing data smooth 21: region entered 1 time time(us): total=3559 kernels=1700 data=0 26: kernel launched 99 times grid: [13x8] block: [8x8] time(us): total=898 max=11 min=9 avg=9 33: kernel launched 99 times grid: [13x13] block: [8x8] time(us): total=802 max=14 min=8 avg=8 main 109: region entered 1 time Kernel 実 時間 データ転送時間を意味する time(us): total= init= region=5151 data=1149 w/o init: total=5151 max=5151 min=5151 avg=5151 Smooth 側では データ転送は行われていない Host-GPU 間のデータ転送は main 側で行っている 49

26 C/C++ Extensions for Array Notation OpenACC プラグマ上での部分配列の指定方法 (C の subarrays) Array [ 開始 index : 長さ ( 個数 ) ] PGI Accelerator notation ( 開始点 0からn 個の表記方法 ) x[0:n-1] OpenACC notation x[0:n] ( 開始点 21から100 個の表記方法 ) x[21:120] x[21:100] OpenACC は Intel s Array Notation for C を採用 データの並びが 連続 であることを保証すること Fortran の場合は Fortran 準拠 notation 例 (A (1:n,1:m)) をそのまま使う 50 部分配列の転送指定は 連続 であること OpenACC の data 構文の約束事 部分配列の転送指定では メモリのアドレスが連続並びになるように指定する Fortran 配列 a(100,100) として配列宣言 内部の98x98を転送したい copyin(a(2:99,2:99)) : PGI Accelerator ModelではOK copyin(a(1:100,2:99)): OpenACC 仕様での指定方法 (100x98) full index 指定 Fortran:Column-major Format ( 列方向に並ぶ順序 ) C : Row-major Format ( 方向並び ) C 配列 a[0:99][0:99] として配列宣言 内部の 98x98 を転送したい copyin(a[1:98][1:98]) : PGI Accelerator ModelではOK copyin(a[1:98][0:100]) : OpenACC 仕様での指定方法 (98x100) OpenACC notation 51

27 Parallel construct の非同期 (async) 実行 use openacc integer,parameter :: n=512 real(4) :: A(n,n), B(n,n), C(n,n), D(n,n), E(n,n) integer(4) :: handle handle = 2!$acc data copyin(a,b), copyout(e), create(c,d)! Task1 : C = A + B!$acc kernels async(handle) do i=1, n do j=1, n C(i,j) = A(i,j) + B(i,j) Task1!$acc end kernels!! Task 2 : D = D + A * B multiply!$acc kernels async(handle) do i = 1, n do j = 1, n D(i,j) = 0.0 do i = 1, n Task2 do j = 1, n do k = 1, n D(i,j) = D(i,j) + A(i,k)*B(k,j)!$acc end kernels! つのカーネル実 ブロックの同期を取る!$acc wait(handle)!!!!call acc_async_wait (handle) あるいは API 使用! Final Task : E = C + D!$acc kernels!$acc loop independent do i=1, n!$acc loop independent do j=1, n E(i,j) = C(i,j) + D(i,j)!$acc end kernels!$acc end data (E の処理 ) end 整数値!$acc parallel async(handle) or!$acc kernels async(handle)...!$acc wait(handle) 同期後 52 その他の OpenACC 構文 API 等 Data Declare Directives (data 構文を宣言文形式 ) device_resident clause Cache directives ( 最内側ループの変数をキャッシュ化 ) Executable Directives update directives ( 任意に Host GPU 間のデータの同期 ) wait directive Runtime Library Routine Environment Variables export ACC_DEVICE_TYPE=NVIDIA ( 使用デバイス名指定 ) export ACC_DEVICE_NUM=0 ( 使用デバイス番号指定 ) 53

28 OpenACC プログラミング例 54 三つの構文を使って並列化して見る!$acc data!$acc kernels!$acc loop do i = 1, n { 並列化可能なループ } end do!$acc end kernels!$acc end data Data 構文 Accelerate Compute 構文 Loop 構文 55

29 PGI コンパイラ OpenACC 用オプション OpenACC directive を認識 Fortran $ pgfortran acc Minfo fast {source}.f90 あるいは $ pgfortran ta=nvidia Minfo fast {source}.f90 (PGI Accelerator directives あるいは OpenACC directives を認識 ) C (C99) 現在 C++ には 実装していない $ pgcc acc Minfo fast {source}.c あるいは $ pgcc ta=nvidia Minfo fast {source}.c (PGI Accelerator directives あるいは OpenACC directives 認識 ) 56 PGI コンパイラ オプション -Minfo=accel Compiler Feedback jacobi: 215, Generating present_or_copyin(f(:,:)) Generating present_or_copy(u(:,:)) Generating local(resid) Generating local(uold(1:n,1:m)) 223, Generating compute capability 1.3 binary Generating compute capability 2.0 binary 225, Loop is parallelizable 227, Loop is parallelizable Accelerator kernel generated 225,!$acc loop gang! blockidx%x 227,!$acc loop vector(256)! threadidx%x データの転送に関して CC 1.3 : 14 registers; 76 shared, 12 constant, 0 local memory bytes; 100% occupancy CC 2.0 : 18 registers; 4 shared, 92 constant, 0 local memory bytes; 100% occupancy 233, Loop is parallelizable 235, Loop is parallelizable Accelerator kernel generated 233,!$acc loop gang(16), vector(16)! blockidx%y threadidx%y 235,!$acc loop gang, vector(16)! blockidx%x threadidx%x 並列化 block 分割に関して NVIDIA H/W 特性 CC 1.3 : 26 registers; 2176 shared, 36 constant, 0 local memory bytes; 50% occupancy CC 2.0 : 26 registers; 2056 shared, 144 constant, 0 local memory bytes; 66% occupancy 57

30 PGI 環境変数 PGI_ACC_TIME $export PGI_ACC_TIME=1 実行時に OpenACC 領域の実行プロファイル情報を出力する jacobi 223: region entered 100 times 単位 :μ 秒 time(us): total= init=10 region= kernels= data=0 w/o init: total= max=10787 min=10544 avg= : kernel launched 100 times grid: [5000] block: [256] time(us): total= max=2515 min=2508 avg= : kernel launched 100 times grid: [320x16] block: [16x16] time(us): total= max=7665 min=7628 avg= : kernel launched 100 times grid: [1] block: [256] time(us): total=1447 max=24 min=13 avg=14 driver 108: region entered 1 time time(us): total= init=94492 region= data= w/o init: total= max= min= avg= PGI 環境変数 ACC_NOTIFY $export ACC_NOTIFY=1 実行中 アクセラレータ上のkernel 動作実行履歴を出力する launch kernel file=/home/kato/jacobi4.f function=initialize line=154 device=0 grid=20 block=256 launch kernel file=/home/kato/jacobi4.f function=jacobi line=227 device=0 grid=5000 block=256 launch kernel file=/home/kato/jacobi4.f function=jacobi line=235 device=0 grid=320x16 block=16x16 launch kernel file=/home/kato/jacobi4.f function=jacobi line=240 device=0 grid=1 block=256 Kernel 実行が行われているか Kernel はどのような並列分割 (grid, thread block) で実行されているか 確認できる 59

31 ヤコビ反復プログラム例 error = 10.0 * tol k = 1 do while (k.le.maxit.and. error.gt. tol) error = 0.0!$omp parallel default(shared)!$omp do do j=1,m do i=1,n uold(i,j) = u(i,j)!$omp do private(resid) reduction(+:error) do j = 2,m-1 do i = 2,n-1 resid = (ax*(uold(i-1,j) + uold(i+1,j)) & & + ay*(uold(i,j-1) + uold(i,j+1)) + b * uold(i,j) - f(i,j))/b u(i,j) = uold(i,j) - omega * resid error = error + resid*resid end do!$omp nowait!$omp end parallel * Error check k = k + 1 error = sqrt(error)/dble(n*m)! End iteration loop 60 ヤコビ反復 (OpenMP) プログラム例 error = 10.0 * tol k = 1 do while (k.le.maxit.and. error.gt. tol) error = 0.0!$omp parallel default(shared)!$omp do do j=1,m do i=1,n uold(i,j) = u(i,j)!$omp do private(resid) reduction(+:error) do j = 2,m-1 do i = 2,n-1 resid = (ax*(uold(i-1,j) + uold(i+1,j)) & & + ay*(uold(i,j-1) + uold(i,j+1)) + b * uold(i,j) - f(i,j))/b u(i,j) = uold(i,j) - omega * resid error = error + resid*resid end do!$omp nowait!$omp end parallel * Error check k = k + 1 error = sqrt(error)/dble(n*m)! End iteration loop Work-sharing Work-sharing 61

32 まず Kernels directive を挿入してみる error = 10.0 * tol k = 1 収束判定ループ do while (k.le.maxit.and. error.gt. tol) error = 0.0!$acc kernels do j=1,m do i=1,n uold(i,j) = u(i,j) & & do j = 2,m-1 do i = 2,n-1 resid = (ax*(uold(i-1,j) + uold(i+1,j)) + ay*(uold(i,j-1) + uold(i,j+1)) + b * uold(i,j) - f(i,j))/b u(i,j) = uold(i,j) - omega * resid error = error + resid*resid end do!$acc end kernels 1 2 Accelerator 領域の開始 3 3 Accelerator 領域の終了 * Error check k = k + 1 error = sqrt(error)/dble(n*m)! End iteration loop コンパイラは 以下のコードを 動生成 GPU 上のメモリに配列データエリアをアロケートホスト側のデータをGPU 側へコピーするホスト側から kernel プログラムを起動する GPU 上で計算した結果をホスト側に戻す GPU 上のデータをデアロケート 問題は? データ転送回数 62 Data Directive を使用する error = 10.0 * tol k = 1!$acc data copy(u)!$acc+ copyin(f) create(uold) do while (k.le.maxit.and. error.gt. tol) error = 0.0 * Copy new solution into old!$acc kernels kernels 並列領域の開始 do j=1,m do i=1,n uold(i,j) = u(i,j) & & do j = 2,m-1 do i = 2,n-1 resid = (ax*(uold(i-1,j) + uold(i+1,j)) + ay*(uold(i,j-1) + uold(i,j+1)) + b * uold(i,j) - f(i,j))/b u(i,j) = uold(i,j) - omega * resid error = error + resid*resid end do!$acc end kernels 1 2 kernels 並列領域の終了 3 3 * Error check k = k + 1 error = sqrt(error)/dble(n*m)! End iteration loop!$acc end data region Accelerator データ領域 4 5 収束判定ループの外側で データ領域 を指定 GPU 上に使用データを 常駐 させる 収束ループが終了時に データをホストに戻す Host-GPU 間のデータ転送の削減 63

33 Loop Directives で並列動作を調整 error = 10.0 * tol k = 1!$acc data copy(u(1:n,1:m))!$acc+ copyin(f(1:n,1:m)) create(uold(1:n,1:m)) do while (k.le.maxit.and. error.gt. tol) error = 0.0 * Copy new solution into old!$acc kernels Accelerator 並列領域の開始 do j=1,m do i=1,n uold(i,j) = u(i,j)!$acc loop gang, vector(8) do j = 2,m-1!$acc loop gang, vector(8) do i = 2,n-1 resid = (ax*(uold(i-1,j) + uold(i+1,j)) & & + ay*(uold(i,j-1) + uold(i,j+1)) + b * uold(i,j) - f(i,j))/b u(i,j) = uold(i,j) - omega * resid error = error + resid*resid end do!$acc end kernels Accelerator 並列領域の終了 * Error check k = k + 1 error = sqrt(error)/dble(n*m)! End iteration loop!$acc end data region コンパイラは 動的に対象並列ループを CUDA の Thread-block/Grid に分割マッピングする ブロック分割等の mapping を変更することが可能 より良い性能を出すには gang, vector の並列スケジューリングを変えて試行錯誤が必要 64 コンパイラによるアクセラレータ翻訳情報 $ pgfortran -fast -Minfo=accel -acc jacobi.f -Minfo=accel : アクセラレータ翻訳情報 -acc : OpenACC 用コンパイル jacobi: 217, Generating local(uold(:,:)) Generating local(resid) Generating copyin(f(:,:)) Generating copy(u(:,:)) 235, Loop is parallelizable 237, Loop is parallelizable Accelerator kernel generated 235,!$acc loop gang vector(8)! blockidx%y threadidx%y 237,!$acc loop gang vector(8)! blockidx%x threadidx%x!$acc loop ( 235) do j = 2,m-1!$acc loop ( 237) do i= 2,n-1 ( 238) resid = (ax*(uold(i-1,j) + uold(i+1,j)) ( 239) & + ay*(uold(i,j-1) + uold(i,j+1)) ( 240) & + b * uold(i,j) - f(i,j))*b1b ( 241) u(i,j) = uold(i,j) - omega * resid ( 242) error = error + resid*resid ( 243) end do ( 244)!$acc end region CC 1.3 : 26 registers; 640 shared, 32 constant, 0 local memory bytes; 50% occupancy CC 2.0 : 26 registers; 520 shared, 144 constant, 0 local memory bytes; 33% occupancy 242, Sum reduction generated for error 65

34 Accelerator ループ マッピングを変更する 例えば Grid size (16 x16) Block size (16 x16) jacobi: 217, Generating local(uold(:,:)) Generating local(resid) Generating copyin(f(:n,:m)) Generating copy(u(:n,:m)) 235, Loop is parallelizable 237, Loop is parallelizable Accelerator kernel generated 235,!$acc loop gang(16), vector(16)! blockidx%y threadidx%y 237,!$acc loop gang(16), vector(16)! blockidx%y threadidx%y loop scheduling 節を変更!$acc loop gang(16) vector(16) ( 235) do j = 2,m-1!$acc loop gang(16) vector(16) ( 237) do i= 2,n-1 ( 238) resid = (ax*(uold(i-1,j) + uold(i+1,j)) ( 239) & + ay*(uold(i,j-1) + uold(i,j+1)) ( 240) & + b * uold(i,j) - f(i,j))*b1b ( 241) u(i,j) = uold(i,j) - omega * resid ( 242) error = error + resid*resid ( 243) end do ( 244)!$acc end region CC 1.3 : 26 registers; 2176 shared, 36 constant, 0 local memory bytes; 50% occupancy CC 2.0 : 26 registers; 2056 shared, 144 constant, 0 local memory bytes; 66% occupancy 242, Sum reduction generated for error 66 実行プロファイル情報で性能評価 loop scheduling(grid/block size) の変更で性能が変わる 235,!$acc loop gang, vector(8)! blockidx%y threadidx%y 237,!$acc loop gang, vector(8)! blockidx%x threadidx%x 237: kernel launched 100 times grid: [640x625] block: [8x8] time(us): total= max=12910 min=12872 avg= ,!$acc loop gang(16), vector(16)! blockidx%y threadidx%y 237,!$acc loop gang(16), vector(16)! blockidx%x threadidx%x 237: kernel launched 100 times grid: [16x16] block: [16x16] time(us): total= max=7501 min=7362 avg=7427 Device Name: GeForce GTX 580 ( 上記は倍精度計算 ) 全体の実行時間 :1.95 秒 全体の実行時間 :1.32 秒 67

35 OpenACC 実行性能サマリー ( 倍精度演算 ) OpenMP と OpenACC 時間 ( 秒 ) 倍率 OpenMP 1 core スレッド (without SSE vector) -O OpenMP 1 core スレッド (with SSE vector) -fastsse 8.18 OpenMP 4 core スレッド並列性能 7.75 x 1.0 OpenACC ( 対象ループに kernels 構文のみ挿入 ) OpenACC ( 繰返ループの外側に data 構文を挿入 ) 2.38 x 3.3 OpenACC ( 対象ループを loop 節で並列 mapping 調整 ) 1.46 OpenACC (mainプログラム上にdata 構文 & present 節使用 ) 1.32 x 5.9 OpenMP 性能 OpenACC 性能 : Intel(R) Core(TM) i GHz (Sandy-bridge) 4core : (Host) Intel(R) Core(TM) i7 CPU 2.67GHz (Nehalem) : (GPU) NVIDIA GeForce GTX 580 PGI 12.3 を使用 68 PGI OpenACC コンパイラのリリース PGI Accelerator Compiler 製品 (x64+gpu) 内に実装 現在 PGI 12.3/12.4/12.5 for Linux/Windows ベータ バージョン使用可能 PGI 12.6 にて 正式製品版リリース (Fortran/C99)(6 月下旬 ) PGI アクセラレータコンパイラ製品 (PGI Accelerator Fortran/C/C++) 1. OpenACC コンパイラ (Fortran, C99) 2. PGI Accelerator Programming Model (directiveベース) 3. PGI CUDA Fortran 4. PGI CUDA-x86 for C/C++ compatible& superset PGI アクセラレータ コンパイラソフテック情報サイト 69

36 終わり 70 付録 71

37 OpenACC directives TIPS for Maximum Performance 72 並列化を う対象は ループ (Nestedループ) である ループの iteration 計算において データ依存性が存在しないこと ループ カウントは GPU-CPU 間のデータ転送オーバーヘッドを無視できる程 十分な大きさが必要である ループ内の Compute Intensity(F.P. 演算数 / メモリアクセス数 ) が大きい程 性能加速性が高い 多次元配列を使用する場合は 素直な添字形式が良い 昔のベクトル機用最適化手法である 添字の1 次元化表現は 災い多し 多次元配列内の要素並びは 連続 であること Stride-1のアクセスができる演算形式は高いベクトル性能を得ることが可能 ループ内のポインタ演算は避けること ループ構造は rectangular 構成が望ましい Triangular Loop の場合は Host<->GPU 間のデータ転送に注意する While loop より for /do loop( 構造化ループ ) へ ループ内のローカル配列により依存性がある場合は 配列のプライベート化を図る Accelerator 領域では サブルーチン コールがある場合 並列化できない この場合は 当該ルーチン 関数のインライン展開を うこと NVIDIA GPUの場合 CUDA cores を常にビジーにするようなマルチスレッディングの状態を作り出せるような並列性を確保すること ( スレッド ブロックのサイズ グリッドの構成を調整 ) 73

38 配列インデックスの一次元化の使用 (1) ( 1) module arrays ( 2) implicit none ( 3) parameter M=2048, N=1024 ( 4) real(8),dimension(m*n):: a ( 5) real(8),dimension(m,n):: b ( 6) end module arrays ( 7) ( 8) program obstacle ( 9) use arrays ( 10) integer(4) :: i,j,idx ( 11) ( 12)!$acc kernels ( 13) do i = 1, M ( 14) do j = 1, N ( 15) idx = ((i-1)*m)+j ( 16) a(idx) = b(i,j) ( 17) ( 18) ( 19)!$acc end kernels ( 20) end program $ pgf90 -acc -Minfo obstacles.f90 obstacle: 12, Generating copyout(a(:)) Generating copyin(b(:,:)) 13, Parallelization would require privatization of array 'a(:)' 14, Parallelization would require privatization of array 'a(:)' Accelerator kernel generated 13,!$acc loop seq 14,!$acc loop seq シーケンシャル実行 too slow! ソースの引用は OpenAcc-standard.org から許可取得済 74 配列インデックスの一次元化の使用 (2) ( 1) module arrays ( 2) implicit none ( 3) parameter M=2048, N=1024 ( 4) real(8),dimension(m*n):: a ( 5) real(8),dimension(m,n):: b ( 6) end module arrays ( 7) ( 8) program obstacle ( 9) use arrays ( 10) integer(4) :: i,j,idx ( 20)!$acc kernels ( 21)!$acc loop independent ( 22) do i = 1, M ( 23) do j = 1, N ( 24) idx = ((i-1)*m)+j ( 25) a(idx) = b(i,j) ( 26) ( 27) ( 28)!$acc end kernels ( 29) end program $ pgf90 -acc -Minfo obstacles.f90 obstacle: 20, Generating copyout(a(:)) Generating copyin(b(:,:)) 強制的に並列化 22, Loop is parallelizable 23, Loop carried reuse of 'a' prevents parallelization Inner sequential loop scheduled on accelerator Accelerator kernel generated 22,!$acc loop gang, vector(256)! blockidx%x threadidx%x 23,!$acc loop seq ソースの引用は OpenAcc-standard.org から許可取得済 外側で並列化 しかし 効率が悪い! 75

39 配列インデックスの一次元化の使用 (3) ( 1) module arrays ( 2) implicit none ( 3) parameter M=2048, N=1024 ( 4) real(8),dimension(m*n):: a ( 5) real(8),dimension(m,n):: b ( 6) end module arrays ( 7) ( 8) program obstacle ( 9) use arrays ( 10) integer(4) :: i,j,idx ( 20)!$acc kernels ( 21) do i = 1, M ( 22)!$acc loop independent ( 23) do j = 1, N ( 24) idx = ((i-1)*m)+j ( 25) a(idx) = b(i,j) ( 26) ( 27) ( 28)!$acc end kernels ( 29) end program $ pgf90 -acc -Minfo obstacles.f90 obstacle: 20, Generating copyout(a(:)) Generating copyin(b(:,:)) 強制的に並列化 21, Parallelization would require privatization of array 'a(:)' 23, Loop is parallelizable Accelerator kernel generated 21,!$acc loop seq 23,!$acc loop gang, vector(256)! blockidx%x threadidx%x 内側で並列化 しかし 効率が悪い! ソースの引用は OpenAcc-standard.org から許可取得済 76 配列インデックスの一次元化の使用 (4) ( 1) module arrays ( 2) implicit none ( 3) parameter M=2048, N=1024 ( 4) real(8),dimension(m,n):: a ( 5) real(8),dimension(m,n):: b ( 6) end module arrays ( 7) ( 8) program obstacle ( 9) use arrays ( 10) integer(4) :: i,j,idx ( 11) ( 12)!$acc kernels ( 13) do j = 1, N ( 14) do i = 1, M ( 15) a(i,j) = b(i,j) ( 16) ( 17) ( 18)!$acc end kernels ( 19) end program 素直な添字に $ pgf90 -acc -Minfo obstacles2.f90 obstacle: 12, Generating copyout(a(:,:)) Generating copyin(b(:,:)) 13, Loop is parallelizable 14, Loop is parallelizable Accelerator kernel generated 13,!$acc loop gang, vector(16)! blockidx%y threadidx%y 14,!$acc loop gang, vector(16)! blockidx%x threadidx%x 外側 内側で並列化多数のスレッドを起動可能となる ソースの引用は OpenAcc-standard.org から許可取得済 77

40 ( 12)!$acc kernels ( 13) i = 0 ( 14) do, while (.not.found) ( 15) i = i + 1 ( 16) if (A(i).eq. 102) then ( 17) found = i ( 18) endif ( 19) ( 20)!$acc end kernels While Loops の使用 $ pgf90 -acc -Minfo while.f90 while: 12, Generating copyin(a(:)) 14, Loop carried scalar dependence for 'found' at line 19 Accelerator kernel generated 14,!$acc loop seq シーケンシャル実 too slow! Rectangular loop 形態に書き換える ( 12)!$acc kernels $ pgf90 -acc -Minfo while2.f90 ( 13) do i = 1, N while: ( 14) if (A(i).eq. 102) then 12, Generating copyin(a(:)) ( 15) found(i) = i Generating copyout(found(:)) ( 16) else 13, Loop is parallelizable ( 17) found(i) = 0 Accelerator kernel generated ( 18) endif 13,!$acc loop gang, vector(256) ( 19)! blockidx%x threadidx%x ( 20)!$acc end kernels 21, maxval reduction inlined ( 21) print *, 'Found at ', maxval(found) ソースの引用は OpenAcc-standard.org から許可取得済 78 Triangular Loops ( 11)!$acc kernels loop ( 12) do I = 1, M ( 13) do J = I, N Triangular loop ( 14) a(i,j) = i + j ( 15) ( 16) ( 17)!$acc end kernels ( 18) ( 19)!$acc kernels loop copy(a) ( 20) do I = 1, M ( 21) do J = I, N ( 22) a(i,j) = i + j ( 23) ( 24) ( 25)!$acc end kernels ( 26) end program $ pgf90 -acc -Minfo triangular.f90 tri: a 配列 :device ->host 転送 11, Generating copyout(a(:,:)) 12, Loop is parallelizable Accelerator kernel generated 12,!$acc loop gang, vector(36)! blockidx%x threadidx%x 13, Loop is parallelizable a 配列 : host-> device ->host 転送 19, Generating copy(a(:,:)) 20, Loop is parallelizable Accelerator kernel generated 20,!$acc loop gang, vector(36)! blockidx%x threadidx%x 21, Loop is parallelizable Triangular loop コンパイラは必要とするデータの転送しか わないホスト側のA 配列のLower triangle 明 的なデータ転送を指 する必要あり部分に ごみデータ が転送される (Visible device copy) ソースの引用は OpenAcc-standard.org から許可取得済 79

41 ( 11)!$acc kernels ( 12) do i = 1, M ( 13) do j = 1, N ( 14) idx = i+j ( 15) A(i,j) = idx ( 16) ( 17) ( 18)!$acc end kernels ( 19) print *, idx, A(1,1), A(M,N) スカラ idx をプライベート化内部的にローカル配列にプロモートする ( 11)!$acc kernels ( 12) do i = 1, M ( 13)!$acc loop private(idx) ( 14) do j = 1, N ( 15) idx = i+j ( 16) A(i,j) = idx ( 17) ( 18) ( 19)!$acc end kernels ( 20) print *, idx, A(1,1), A(M,N) ソースの引用は OpenAcc-standard.org から許可取得済 live-out スカラ変数 並列領域外でその値を参照する $ pgf90 -acc -Minfo live-out.f90 liveout: 11, Generating copyout(a(:,:)) 12, Loop is parallelizable 13, Inner sequential loop scheduled on accelerator Accelerator kernel generated 12,!$acc loop gang, vector(32)! blockidx%x threadidx%x 13,!$acc loop seq 14, Accelerator restriction: induction variable live-out from loop: idx 15, Accelerator restriction: induction variable live-out from loop: idx $ pgf90 -acc -Minfo live-out2.f90 liveout: 12, Loop is parallelizable 14, Loop is parallelizable Accelerator kernel generated 12,!$acc loop gang, vector(32)! blockidx%x threadidx%x 14,!$acc loop gang, vector(8)! blockidx%y threadidx%y 80 ローカル配列のプライベート化 (1) ( 12)!$acc kernels ( 13) do i = 1, M ( 14) do j = 1, N ( 15) do jj = 1, 10 ( 16) tmp(jj) = jj ( 17) end do ( 18) A(i,j) = sum(tmp) ( 19) ( 20) ( 21)!$acc end kernels tmp() 配列は i, j index に沿わない配列 $ pgf90 -acc -Minfo private.f90 private: 12, Generating copyout(a(:,:)) Generating copyout(tmp(:)) 13, Parallelization would require privatization of array 'tmp(:)' 14, Parallelization would require privatization of array 'tmp(:)' Accelerator kernel generated 13,!$acc loop seq IループとJループで並列化されない 14,!$acc loop seq 15, Loop is parallelizable 18, sum reduction inlined Loop is parallelizable ソースの引用は OpenAcc-standard.org から許可取得済 並列領域外でその値を参照する 81

42 ローカル配列のプライベート化 (2) ( 12)!$acc kernels ( 13) do i = 1, M ( 14)!$acc loop private(tmp) ( 15) do j = 1, N ( 16) do jj = 1, 10 ( 17) tmp(jj) = jj ( 18) end do ( 19) A(i,j) = sum(tmp) ( 20) ( 21) ( 22)!$acc end kernels tmp() 配列を j index を有する 2 次元配列にプロモート $ pgf90 -acc -Minfo private2.f90 private: 12, Generating copyout(a(:,:)) 13, Loop is parallelizable 15, Loop is parallelizable Accelerator kernel generated 13,!$acc loop gang, vector(16)! blockidx%x threadidx%x 15,!$acc loop gang, vector(16)! blockidx%y threadidx%y 16, Loop is parallelizable 19, sum reduction inlined Loop is parallelizable ソースの引用は OpenAcc-standard.org から許可取得済 82 関数 サブルーチンの呼び出し 並列化対象ループ内に 関数 サブルーチン が存在すると並列化できない これは GPUデバイス側の制約である device 用の linker が 存在してないこととハードウェアのサポートが無いことに因る ループ内に関数等をインライン展開する必要がある 手動でソース上でインライン展開 コンパイラオプション Minline or Mipa=inline で自動インライン function domul (x,y) result(mul) real(4) :: x,y real(4) :: mul mul = x*y end function!$acc kernels do i = 1,100 a(i) = float(i) * do i = 1,100 b(i) = domul(a(i),a(i))!$acc end kernels 並列化不可 83

43 ループ内の Function Call ( 1) function domul (x,y) result(mul) ( 2) real(4) :: x,y ( 3) real(4) :: mul ( 4) mul = x*y ( 5) end function ( 7) program main ( 8) real(4) :: a(100), b(100) ( 9)!$acc kernels ( 10) do i = 1,100 ( 11) a(i) = float(i) * ( 12) ( 14) do i = 1,100 ( 15) b(i) = domul(a(i),a(i)) ( 16) ( 17)!$acc end kernels ( 18) print *, b ( 19) end 並列化できず $ pgf90 -fast -acc -Minfo function_call.f90 PGF90-W-0155-Accelerator region ignored; see -Minfo messages (function_call.f90: 9) main: 9, Accelerator region ignored 14, Accelerator restriction: function/procedure calls are not supported 15, Accelerator restriction: unsupported call to 'domul' 84 Function Call ( 自動インライン化 -Minline) ( 7) program main ( 8) real(4) :: a(100), b(100) ( 9)!$acc kernels ( 10) do i = 1,100 ( 11) a(i) = float(i) * ( 12) ( 14) do i = 1,100 ( 15) b(i) = domul(a(i),a(i)) ( 16) ( 17)!$acc end kernels 自動インライン $ pgf90 -fast -acc -Minfo function_call.f90 -Minline main: 9, Generating copyout(b(:)) Generating copyout(a(:)) 10, Loop is parallelizable Accelerator kernel generated 10,!$acc loop gang, vector(96)! blockidx%x threadidx%x 14, Loop is parallelizable Accelerator kernel generated 並列化されたインライン化 14,!$acc loop gang, vector(96)! blockidx%x threadidx%x 15, domul inlined, size=2, file function_call.f90 (1) 85

44 Function Call (-Mipa=inline 使用 ) Filename: function_call_sub2.f90 ( 1) subroutine domul (x,y,mul) ( 2) real(4) :: x(100),y(100) ( 3) real(4) :: mul(100) ( 4) do i=1, 100 ( 5) mul(i) = x(i)*y(i) ( 6) end do ( 7) end subroutine -Mipa=inline を使ってインライン化 動インラインファイル名が異なる つのファイル間を Filename: function_call2.f90 ( 1) program main ( 2) real(4) :: a(100), b(100) ( 3)!$acc kernels ( 4) do i = 1,100 ( 5) a(i) = float(i) * ( 6) ( 7) ( 8) call domul (a,a,b) ( 9)!$acc end kernels ( 10) print *, b(1),b(100) ( 11) end 自$ pgf90 -fast -acc -Mipa=inline -Minfo function_call_sub2.f90 -c $ pgf90 -fast -acc -Mipa=inline -Minfo function_call_sub2.o function_call2.f90 IPA: no IPA optimizations for 1 source files IPA: Recompiling function_call2.o: stale object file main: 3, Generating copyout(a(:)) Generating copyout(b(:)) 4, Loop is parallelizable Accelerator kernel generated 4,!$acc loop gang, vector(96)! blockidx%x threadidx%x 8, domul inlined, size=5 (IPA) file function_call_sub2.f90 (1) 8, Loop is parallelizable Accelerator kernel generated 8,!$acc loop gang, vector(96)! blockidx%x threadidx%x 手続間最適化 IPA 86 ポインタを利用した演算を避ける (C 言語 ) Accelerate Compute 領域内のポインタを利用した演算ループは 並列化できない void memcpy(float *restrict A, *restrict B, int count) { float* ptra =A; float* ptrb =B; while (count --) { *ptra++ = *ptrb++; } return; } ソースの引用は OpenAcc-standard.org から許可取得済 void memcpy(float *restrict A, *restrict B, int count) { #pragma acc kernels { for ( int i=0 ; i <count; ++i ) { A[i] = B[i] ; } } return; 配列は添字表現へ 87

45 restrict 修飾子の利用 (C 言語 ) void vec_add_gpu( float *restrict r, float *a, float *b, int n ) { #pragma acc kernels for copyin(a[0:n],b[0:n]) copyout(r[0:n]) for( int i = 0; i < n; ++i ) r[i] = a[i] + b[i]; } float *restrict r : C99 restrict 修飾子の利用 C 言語は ポインタをかなり自由に使える言語であり 別名なども自由使用できる 各ポインタ間のオーバーラップもあり得る こうしたことが コンパイラの最適化を妨げる原因となっている *restrict を指定して 演算における データの依存性 競合 がないことをコンパイラに指示するために使用する もちろん 依存性がある場合は 並列化はできないため *restrict を使用してはならない 88 ランタイム エラー (Device error) call to cumemcpydtoh returned error 700: Launch failed!$acc region do i = 1, M do j = 1, N A(i,j) = B(i,j+1) << 配列境界外!$acc end region GPU 上で実行時 配列境界を越えたアクセスにより ランタイムがエラーを返した -Mbounds オプションで実 時チェック call to cumemcpy2d returned error 1: Invalid value parameter(n=1024,m=512) real :: A(M,N), B(M,N)...!$acc kernels loop copyout(a), copyin(b(0:n,1:m+1)) <<< Bad bounds for the copyin do i = 1, M do j = 1, N A(i,j) = B(i,j+1)!$acc end region 添字の誤り GPU デバイスへデータ転送時のディレクティブの記述ミス 89

46 PGI コンパイラ製品情報 URL ソフテック PGI コンパイラ製品ホームページ PGI アクセラレータ製品概要 PGI コンパイラ技術コラム PGI Accelerator Programming Model を使う CUBLAS CUFFT を CUDA Fortran から使う MPI 並列の話 その他 90 本ドキュメントに記述された各製品名は 各社の商標または登録商標です Copyright 2012 SofTek Systems Inc. All Rights Reserved. 91