Microsoft PowerPoint - GTC2012-SofTek.pptx
|
|
|
- みりあ おいもり
- 7 years ago
- Views:
Transcription
1 GTC Japan 2012 PGI Accelerator Compiler 実践! PGI OpenACC ディレクティブを使用したポーティング 2012 年 7 月 加藤努株式会社ソフテック
2 本日の話 OpenACC によるポーティングの実際 OpenACC ディレクティブ概略説明 Accelerator Programming Model Fortran プログラムによるポーティング ステップ三つのディレクティブの利用性能チューニング PGI Accelerator Compiler 製品を使用 1
3 PGI OpenACC 対応コンパイラを使用 PGI Accelerator Compiler 製品 (x64+gpu) 内に実装 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 2012 年 7 月 OpenACC 正式版リリース PGI アクセラレータ コンパイラソフテック情報サイト 2
4 OpenACC Standard とは何か? 2011 年 11 月 NVIDIA, Cray, PGI, CAPS Accelerators 用のプログラミング API の標準仕様 Fortran, C/C++ 言語上で指定するコンパイラ ディレクティブ群 ユーザサイド開発者がアクセラレータで実行するコード部分をディレクティブで指定する ( コンパイラに対して ヒントを与える ) OpenACC コンパイラホスト側の処理をアクセラレータ (GPU) にオフロードするコード生成ホスト -- GPU 間のデータ転送コードの生成 2009 年リリース以来 実績を積んだ PGI Accelerator Compiler(directives) がベースとなっている 3
5 Accelerator Programming Model ホスト側 ハイブリッド構成 (CPU + Accelerator) GPU 側 CPU Main Memory Host_A(100) 重い計算部分の処理をオフロード 使用データを送る 結果データを戻す Overhead GPU Device Memory Device_A(100) Host GPU 間のメモリデータの転送が伴う データ転送のオーバーヘッド時間が伴う 4
6 OpenACC ディレクティブの主な構成 ホスト ( 処理 ) Accelerator 1 CPU 重い計算部分の処理をオフロード 3 GPGPU Main Memory 2 ( データ ) Device Memory 1 Accelerate Compute 構文 (offload 領域指示 ) 2 Data 構文 ( データ移動指示 ) 3 Loop 構文 (Mapping for parallel/vector, Tuning) 5
7 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 に行う 6
8 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 部分の指定 オフロードする並列対象領域の指 ( 一般に ループ部分 ) 7
9 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 部分の指定 オフロードする並列対象領域の指 ( 一般に ループ部分 ) 8
10 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 オフロードする並列対象領域の指 ( 一般に ループ部分 ) 9
11 三つの構文を使用してポーティング!$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 構文 オフロードする並列対象領域 ( 一般に ループ部分 ) 10
12 OpenACCを使用して GPU 上での実行を行うまでのプログラム ポーティングを行う (Fortran) 11
13 subroutine driver (u,f) * dx - grid spacing in x direction * dy - grid spacing in y direction ( 配列宣 等は省略 ) * Initialize data cpu0 = second() call initialize (n,m,alpha,dx,dy,u,f) * Solve Helmholtz equation ヤコビ反復プログラム call jacobi (n,m,dx,dy,alpha,relax,u,f,tol,mits) * Check error between exact solution call error_check (n,m,alpha,dx,dy,u,f) cpu1 = second() * Printout Elapsed time elapsed = (cpu1 -cpu0) * t_ac print '(/,1x,a,F10.3/)', & Elpased Time (Initialize + Jacobi solver + Check) : ',elapsed return end 三つのサブルーチン コール 手続間で配列の受渡し有り (u,f) 各ルーチン内で高速化を図る 時間の掛かっている場所は? 時間の掛かるループ内で 使用されている 配列 は何か? 手続間のデータの受渡しの状況を見る 12
14 Subroutine Jacobi の核心部分 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 Do while ループ内でステンシル計算 収束条件を満たしたら終了 内部は 2 次元の nested loop uold(i,j) 配列は 並列実行依存性なし u(i,j) 配列も依存性無し ストアのみ f(i,j) 配列も依存性無し 参照のみ error 変数は リダクション演算 並列依存性とは無し : u(i) = u(i) 有り : u(i) =u(i-1) 同じ 配列で定義 ~ 参照関係があるとき依存性の検討要 13
15 Jacobi ルーチンへの OpenMP directives 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 $ pgfortran -fastsse mp Minfo jacobi.f jacobi: 204, Parallel region activated 206, Parallel loop activated with static block schedule 207, Memory copy idiom, loop replaced by call to c_mcopy8 214, Barrier 215, Parallel loop activated with static block schedule 216, Generated 4 alternate versions of the loop Generated vector sse code for the loop Generated 4 prefetch instructions for the loop 223, Begin critical section End critical section Parallel region terminated 行番号 コンパイラメッセージ 実際に並列化とベクトル化を実装している 14
16 シリアル実行用コンパイルとその実行 OpenACC]$ pgfortran -O3 openmp.f -Minfo initialize: 139, Invariant if transformation 140, Invariant assignments hoisted out of loop jacobi: 207, Memory copy idiom, loop replaced by call to c_mcopy8 error_check: 262, Invariant if transformation 263, Invariant assignments hoisted out of loop OpenACC]$ a.out Input n,m - grid real*8 in x,y direction N= 5120 M= 5000 Input alpha - Helmholts constant Input relax - Successive over-relaxation parameter Input tol - error tolerance for iterative solver Input mits - Maximum iterations for solver Time measurement accuracy :.10000E-05 Total Number of Iterations 101 Residual E-011 Solution Error : E-004 コンパイル 実行 Elpased Time (Initialize + Jacobi solver + Check) :
17 シリアル実行用最適化コンパイルとその実行 OpenACC]$ pgfortran -fastsse openmp.f -Minfo initialize: 139, Invariant if transformation 140, Invariant assignments hoisted out of loop Loop not vectorized: mixed data types Unrolled inner loop 4 times jacobi: 207, Memory copy idiom, loop replaced by call to c_mcopy8 216, Generated 4 alternate versions of the loop Generated vector sse code for the loop Generated 4 prefetch instructions for the loop error_check: 262, Invariant if transformation 263, Invariant assignments hoisted out of loop Generated 2 alternate versions of the loop Generated vector sse code for the loop Generated a prefetch instruction for the loop [kato@photon29 OpenACC]$ a.out ( 省略 ) Residual E-011 Solution Error : E-004 SSE ベクトル化 コンパイル 実行 Elpased Time (Initialize + Jacobi solver + Check) :
18 OpenMP 並列実行用コンパイル OpenACC]$ pgf90 -fastsse -mp openmp.f -Minfo initialize: 138, Parallel region activated 139, Parallel loop activated with static block schedule 140, Loop not vectorized: mixed data types Unrolled inner loop 4 times 147, Parallel region terminated スレッド並列化 jacobi: 204, Parallel region activated 206, Parallel loop activated with static block schedule 207, Memory copy idiom, loop replaced by call to c_mcopy8 214, Barrier 215, Parallel loop activated with static block schedule 216, Generated 4 alternate versions of the loop Generated vector sse code for the loop Generated 4 prefetch instructions for the loop 226, Begin critical section SSE ベクトル化 End critical section Parallel region terminated error_check: 261, Parallel region activated 262, Parallel loop activated with static block schedule 263, Generated 2 alternate versions of the loop Generated vector sse code for the loop Generated a prefetch instruction for the loop 269, Begin critical section End critical section Parallel region terminated 17
19 OpenMP 並列実行 OpenACC]$ export OMP_NUM_THREADS=4 (4 スレッド実行 ) [kato@photon29 OpenACC]$ a.out Input n,m - grid real*8 in x,y direction N= 5120 M= 5000 Input alpha - Helmholts constant Input relax - Successive over-relaxation parameter Input tol - error tolerance for iterative solver Input mits - Maximum iterations for solver Time measurement accuracy :.10000E-05 Total Number of Iterations 101 Residual E-011 Solution Error : E-004 Elpased Time (Initialize + Jacobi solver + Check) :
20 シリアル OpenMP 並列実行性能 ( 倍精度演算 ) OpenMP と OpenACC 時間 ( 秒 ) 倍率 1 core スレッド (without SSE vector) -O core スレッド (with SSE vector) -fastsse x 2.0 OpenMP 4 core スレッド並列性能 -mp -fastsse 8.74 x 3.0 OpenMP 用オプション ベクトル最適化用オプション OpenMP 性能 OpenACC 性能 : Intel(R) Core(TM) i7 CPU 2.67GHz (Nehalem) 4core : (Host) 同上 : (GPU) NVIDIA GeForce GTX 580 PGI 12.5 を使用 19
21 プログラムのプロファイリング取得 コンパイル ( シリアル実 用 ) [kato@photon29]$ pgfortran -fastsse -Minfo=ccff jacobi.f -o jacobi コンパイル (OpenMP 並列実 用 ) [kato@photon29]$ pgfortran -fastsse -Minfo=ccff mp jacobi.f -o jacobi プロファイルデータの取得 (pgcollect utility でサンプリング ) [kato@photon29]$ pgcollect -time jacobi Input n,m - grid real*8 in x,y direction N= 5120 M= 5000 Input alpha - Helmholts constant Input relax - Successive over-relaxation parameter Input tol - error tolerance for iterative solver Input mits - Maximum iterations for solver Time measurement accuracy :.10000E-05 Total Number of Iterations 101 Residual E-011 Solution Error : E-004 実際の実行 Elpased Time (Initialize + Jacobi solver + Check) : target process has terminated, writing profile data プロファイリング ツール pgprof の実 [kato@photon29]$ pgprof -exe jacobi ( プロファイラ ツールの起動 ) 20
22 1 スレッド シリアル実行のプロファイル PGPROF PGI プロファイラ 計算コスト 99% Jacobi ルーチン 67% c_mcopy8 ルーチン 32% メモリアクセスのための PGI 組込ルーチン 21
23 4 スレッド並列実行のプロファイル バリア同期 計算コスト 97% 4 スレッドの各時間コスト Jacobi ルーチン 49% c_mcopy8 ルーチン 35% OpenMP バリア同期 13% 22
24 コンパイラ フィードバック情報 click! ループ このループは 8.91 秒 コンパイラフィードバック情報 Compute Intensity このループは 2.43 ベクトル化 最適化実施のメッセージ 23
25 ポーティングでの作業方針 1. Jacobi ルーチンの時間コストが 99% 占める 最初にこのルーチンの中の GPU 実行部分を特定して OpenACC ディレクティブを挿入 (targeting) 2. ホストと GPU 間のデータ移動を最小化する (GPU 上に計算に必要なデータを常駐化させる ) 3. NVIDIA GPU 用の Grid サイズ Block サイズのチューニングを行う 4. Jacobi ルーチン以外のルーチンに対しても OpenACC ディレクティブを適用する 5. プログラム全体にスコープ範囲を広げ ホストと GPU 間のデータ移動を最小化する 24
26 三つの構文を使って GPU 用に並列化する OpenACC ディレクティブを使用する!$acc data!$acc kernels!$acc loop do i = 1, n { 並列化可能なループ } end do!$acc end kernels!$acc end data Data 構文 Accelerate Compute 構文 Loop 構文 25
27 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 認識 ) 26
28 まず 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) 1 2 Accelerator 領域の開始 3 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 3 error = error + resid*resid end do!$acc end kernels Accelerator 領域の終了 * Error check k = k + 1 error = sqrt(error)/dble(n*m)! End iteration loop コンパイラは 以下のコードを自動生成 GPU 上のメモリに配列データエリアをアロケートホスト側のデータをGPU 側へコピーするホスト側から kernel プログラムを起動する GPU 上で計算した結果をホスト側に戻す GPU 上のデータをデアロケート 問題は? データ転送回数 27
29 PGI コンパイラ フィードバック情報 (-Minfo) pgfortran acc -fastsse Minfo=accel -o jacobi1.exe jacobi1.f jacobi: 行番号 204, Generating copyout(uold(1:n,1:m)) Generating copyin(u(:n,:m)) Generating copyout(u(2:n-1,2:m-1)) Generating copyin(f(2:n-1,2:m-1)) Generating compute capability 1.3 binary Generating compute capability 2.0 binary Accelerator kernel generated 213,!$acc loop gang, vector(8)! blockidx%y threadidx%y 214,!$acc loop gang, vector(8)! blockidx%x threadidx%x 使用レジスタ数使用 shared Mem 使用 const. Mem Occupancy per SM CC 1.3 : 32 registers; 640 shared, 28 constant; 50% occupancy CC 2.0 : 28 registers; 520 shared, 160 constant; 33% occupancy 222, Sum reduction generated for error 配列名 Host GPU 間配列データの転送命令生成 ネスト ループの並列分割の様子 (Grid/Block) 総和リダクション検出し リダクションコード生成 NVIDIA H/W Compute capability 使用特性 28
30 実行 コンパイル & 実行モジュール作成 make jacobi1.exe pgfortran -o jacobi1.exe jacobi1.f -fastsse -Minfo=accel acc jacobi1.exe と言うモジュールには Host 用コード +GPU 用コードが含まれる 実行 jacobi1.exe Input n,m - grid real*8 in x,y direction N= 5120 M= 5000 Input alpha - Helmholts constant Input relax - Successive over-relaxation parameter Input tol - error tolerance for iterative solver Input mits - Maximum iterations for solver Time measurement accuracy :.10000E-05 Total Number of Iterations 101 Residual E-011 Solution Error : E-004 Elpased Time (Initialize + Jacobi solver + Check) : FORTRAN STOP 29
31 PGI 環境変数 (Accelerator Profile) PGI_ACC_TIME $ export PGI_ACC_TIME=1 実行時に OpenACC 領域の実行プロファイル情報を出力する Accelerator Kernel Timing data プロファイル時間の単位 :μ 秒 jacobi( ルーチン名 ) 204: region entered 100 times time(us): total= init= region= Kernelの実 時間 (1.89 秒 ) kernels= data= w/o init: total= max= min= avg= データ転送時間 (15.52 秒 ) 番号 206: kernel launched 100 times grid: [640x625] block: [8x8] time(us): total= max=5204 min=5165 avg= : kernel launched 100 times grid: [640x625] block: [8x8] time(us): total= max=13200 min=13180 avg= : kernel launched 100 times grid: [1] block: [256] time(us): total=56786 max=570 min=565 avg= ループの Kernel の実 時間 (1.31 秒 ) Grid/Block 分割のサイズ 30
32 OpenACC 実行性能サマリー ( 倍精度演算 ) OpenMP 性能と OpenACC 性能 時間 ( 秒 ) 倍率 1 core スレッド (without SSE vector) -O core スレッド (with SSE vector) -fastsse OpenMP 4 core スレッド並列性能 -mp -fastsse 8.74 x 1.0 OpenACC ( 対象ループに kernels 構文のみ挿入 ) OpenMP 性能 OpenACC 性能 : Intel(R) Core(TM) i7 CPU 2.67GHz (Nehalem) 4core : (Host) 同上 : (GPU) NVIDIA GeForce GTX
33 ループ内にデータ転送があると転送の嵐 do while (k.le.maxit.and. error.gt. tol)!$acc kernels DO ループ群!$acc end kernels 収束するまで繰り返す外側ループ Host GPU データコピー場所 GPU 上の計算処理 GPU Host データバック場所 end do! End iteration loop Data 構文の利用 : 並列処理 と データ移動 の指示を分離する 32
34 ループの外でデータ転送を行う 明示的に データ構文 で転送指示!!$acc data copy(u) copyin(f)... do while (k.le.maxit.and. error.gt. tol)!$acc kernels DO ループ群 Host GPU データコピー場所 GPU 内のデータは常駐化 GPU 上の計算処理!$acc end kernels end do! End iteration loop!$acc end data GPU Host データバック場所 33
35 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 Acc データ領域 4 5 収束判定ループの外側で データ領域 を指定 GPU 上に使用データを 常駐 させる 収束ループが終了時に データをホストに戻す Host-GPU 間のデータ転送の削減 34
36 データ転送を 1 回だけにした場合のプロファイル Accelerator Kernel Timing data /home/kato/gpgpu/openmp/double/openacc/jacobi2.f jacobi 205: region entered 100 times 3つのkernel の存在 kernels 構 の領域の情報 time(us): total= init=3 region= kernels= data=0 w/o init: total= max=19442 min=19148 avg= : kernel launched 100 times grid: [640x625] block: [8x8] time(us): total= max=5197 min=5169 avg= : kernel launched 100 times grid: [640x625] block: [8x8] time(us): total= max=13186 min=13172 avg= : kernel launched 100 times grid: [1] block: [256] time(us): total=56761 max=569 min=566 avg=567 データ転送時間 (0 秒 ) 次は この時間をチューニングする /home/kato/gpgpu/openmp/double/openacc/jacobi2.f jacobi 199: region entered 1 time データ構 の領域のプロファイル情報 1 回のみ time(us): total= init=87485 region= data= データ転送時間 (0.11 秒 ) w/o init: total= max= min= avg=
37 OpenACC 実行性能サマリー ( 倍精度演算 ) OpenMP 性能と OpenACC 性能 時間 ( 秒 ) 倍率 1 core スレッド (without SSE vector) -O core スレッド (with SSE vector) -fastsse OpenMP 4 core スレッド並列性能 -mp -fastsse 8.74 x 1.0 OpenACC ( 対象ループに kernels 構文のみ挿入 ) OpenACC ( 繰返ループの外側に data 構文を挿入 ) 2.32 x 3.7 OpenMP 性能 OpenACC 性能 : Intel(R) Core(TM) i7 CPU 2.67GHz (Nehalem) 4core : (Host) 同上 : (GPU) NVIDIA GeForce GTX
38 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 コンパイラは 自動的に対象並列ループを CUDA の Thread-block/Grid に分割マッピングする ブロック分割等の mapping を明 的に変更することが可能 より良い性能を出すには gang, vector の並列スケジューリングを変えて試行錯誤が必要 37
39 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 38
40 実行プロファイル情報で性能評価 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=13189 min=13176 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=7365 min=7223 avg=7293 Device Name: GeForce GTX 580 ( 上記は倍精度計算 ) μ 秒 全体の実行時間 :2.32 秒 全体の実行時間 :1.32 秒 39
41 OpenACC 実行性能サマリー ( 倍精度演算 ) OpenMP 性能と OpenACC 性能 時間 ( 秒 ) 倍率 1 core スレッド (without SSE vector) -O core スレッド (with SSE vector) -fastsse OpenMP 4 core スレッド並列性能 -mp -fastsse 8.74 x 1.0 OpenACC ( 対象ループに kernels 構文のみ挿入 ) OpenACC ( 繰返ループの外側に data 構文を挿入 ) 2.32 x 3.7 OpenACC ( 対象ループを loop 節で並列 mapping 調整 ) 1.32 x 6.6 OpenMP 性能 OpenACC 性能 : Intel(R) Core(TM) i7 CPU 2.67GHz (Nehalem) 4core : (Host) 同上 : (GPU) NVIDIA GeForce GTX
42 プログラム全体にスコープ範囲を広げる subroutine driver (u,f) * dx - grid spacing in x direction * dy - grid spacing in y direction ( 配列宣 等は省略 ) * Initialize data cpu0 = second()!$acc data copy (u,f) call initialize (n,m,alpha,dx,dy,u,f) * Solve Helmholtz equation call jacobi (n,m,dx,dy,alpha,relax,u,f,tol,mits) * Check error between exact solution call error_check (n,m,alpha,dx,dy,u,f)!$acc end data cpu1 = second() * Printout Elapsed time elapsed = (cpu1 -cpu0) * t_ac print '(/,1x,a,F10.3/)', & Elpased Time (Initialize + Jacobi solver + Check) : ',elapsed return end u() と f() 配列が プログラム全体で使用される u() と f() 配列を Copyin to GPU 各手続上では u, f 配列に係わる計算処理を GPU kernel 化するだけ u() と f() 配列を Copyout to Host 41
43 コンパイラ フィードバック情報 ( データ構文に関するもののみ抽出 ) subroutine initialize (n,m,alpha,dx,dy,u,f) real*8 u(n,m),f(n,m),dx,dy,alpha!$acc kernels copyin(dx,dy,alpha) present(u,f) driver:!$acc loop gang private(xx,yy) 108, Generating copy(f(:,:)) do j = 1,m Generating copy(u(:,:))!$acc loop vector(256) ( 以下 省略 ) do i = 1,n initialize: xx = dx * real(i-1)! -1 < x < 1 152, Generating present(u(:,:)) yy = dy * real(j-1)! -1 < y < 1 Generating present(f(:,:)) u(i,j) = 0.0 ( 以下 省略 ) f(i,j) = -alpha *(1.0-xx*xx)*(1.0-yy*yy) jacobi: & - 2.0*(1.0-xx*xx)-2.0*(1.0-yy*yy) 215, Generating present_or_copyin(f(:,:)) Generating present_or_copy(u(:,:)) Generating local(resid)!$acc end kernels Generating local(uold(1:n,1:m)) ( 以下 省略 ) return error_check: 275, Generating present(u(:,:)) ( 以下 省略 ) present 節の意味 u() と f() 配列に関しては 既に GPU 上に存在していると言う意味 42
44 OpenACC 実行性能サマリー ( 倍精度演算 ) OpenMP 性能と OpenACC 性能 時間 ( 秒 ) 倍率 1 core スレッド (without SSE vector) -O core スレッド (with SSE vector) -fastsse OpenMP 4 core スレッド並列性能 -mp -fastsse 8.74 x 1.0 OpenACC ( 対象ループに kernels 構文のみ挿入 ) OpenACC ( 繰返ループの外側に data 構文を挿入 ) 2.32 x 3.7 OpenACC ( 対象ループを loop 節で並列 mapping 調整 ) 1.32 x 6.6 OpenACC (mainプログラム上に data 構文 & present 節使用 ) 1.23 x 7.1 OpenMP 性能 OpenACC 性能 : Intel(R) Core(TM) i7 CPU 2.67GHz (Nehalem) 4core : (Host) 同上 : (GPU) NVIDIA GeForce GTX
45 ポーティング 開発時における便利なツール等 NVIDIA Visual Profiler ストリーム イベント等挙動の視覚的な把握詳細なカーネル特性の把握 PGI 環境変数 ACC_NOTIFY kernel 動作の実行時履歴の出力 PGI_ACC_DEBUG 実行時の CUDA システムコールのイベントログ出力 44
46 NVIDIA Visual Profiler を使う make jacobi4.exe コンパイル pgfortran -o jacobi4.exe jacobi4.f -fastsse -acc -ta=nvidia:cuda4.1 which nvvp /usr/local/cuda/bin/nvvp CUDA toolkit 4.1 を使用するように指示する nvvp (NVIDIA Visual Profiler の起動 ) 起動 実行モジュール jacobi4.exe の指定 NVIDIA Visual Profiler 4.1 を使用する 45
47 NVIDIA CUDA Visual Profiler(1) データコピーが頻繁! GPU 特性 全体性能特性 この stream 全体の挙動が色別で分かるデータ転送 ( カーキー色 ) が卓越 46
48 NVIDIA CUDA Visual Profiler(2) Compute kernels の実行が主体 この stream 全体の挙動が色別で分かるカーネル実行 ( ピーコックブルー色 ) が卓越 Kernel の実行特性 47
49 NVIDIA CUDA Visual Profiler(3) Timeline の詳細 Kernel の実行特性 個々のイベント特性の詳細 48
50 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) で実行されているか 確認できる 49
51 PGI 環境変数 ( 実行時のイベントログ ) PGI_ACC_DEBUG (PGI 2013 以降 ) $export PGI_ACC_DEBUG=1 (disable したい場合は 0) 実行時のPGIのCUDAシステムコールのイベントログを出力 [kato@photon29]$ export PGI_ACC_DEBUG=1 [kato@photon29]$ jacobi4.exe pgi_cu_init() found 2 devices pgi_cu_init( file=acc_init.c, function=acc_init, line=41, startline=1, endline=-1 ) pgi_cu_init() will use device 0 (V2.0) pgi_cu_init() compute context created initialize nvidia pgi_cu_init( file=/home/kato/gpgpu/openmp/double/openacc/jacobi4.f, function=driver, line=107, startline=69, endline=129 ) pgi_acc_dataon(devptr=0x1,hostptr=0x7ff535c48230,offset=0,0,stride=1,5120,size=5120x5000, extent=5120x5000,eltsize=8,lineno=107,name=f,flags=0xf00=create+present+copyin+copyout) NO map for host:0x7ff535c48230 pgi_cu_alloc(size= ,lineno=107,name=f) pgi_cu_alloc( ) returns 0x map dev:0x host:0x7ff535c48230 size: offset:0 data[dev:0x7ff535c48230 host:0x size: ] (line:107 name:f) pgi_cu_launch_a(func=0xaf6f40, params=0x7fff8d2e1190, bytes=72, sharedbytes=0) First arguments are:
52 終わり 51
53 本ドキュメントに記述された各製品名は 各社の商標または登録商標です Copyright 2012 SofTek Systems Inc. All Rights Reserved. 52
Microsoft PowerPoint - GDEP-GPG_softek_May24-2.pptx
G-DEP 第 3 回セミナー PGI OpenACC Compiler PGIコンパイラ使用の実際 新しい OpenACC によるプログラミング 2012 年 5 月 加藤努株式会社ソフテック OpenACC によるプログラミング GPU / Accelerator Computing Model のデファクト スタンダードへ OpenACC Standard 概略説明 Accelerator Programming
GPU チュートリアル :OpenACC 篇 Himeno benchmark を例題として 高エネルギー加速器研究機構 (KEK) 松古栄夫 (Hideo Matsufuru) 1 December 2018 HPC-Phys 理化学研究所 共通コードプロジェクト
GPU チュートリアル :OpenACC 篇 Himeno benchmark を例題として 高エネルギー加速器研究機構 (KEK) 松古栄夫 (Hideo Matsufuru) 1 December 2018 HPC-Phys 勉強会 @ 理化学研究所 共通コードプロジェクト Contents Hands On 環境について Introduction to GPU computing Introduction
OpenACCによる並列化
実習 OpenACC による ICCG ソルバーの並列化 1 ログイン Reedbush へのログイン $ ssh reedbush.cc.u-tokyo.ac.jp l txxxxx Module のロード $ module load pgi/17.3 cuda ログインするたびに必要です! ワークディレクトリに移動 $ cdw ターゲットプログラム /srcx OpenACC 用のディレクトリの作成
TSUBAME2.0 における GPU の 活用方法 東京工業大学学術国際情報センター丸山直也第 10 回 GPU コンピューティング講習会 2011 年 9 月 28 日
TSUBAME2.0 における GPU の 活用方法 東京工業大学学術国際情報センター丸山直也第 10 回 GPU コンピューティング講習会 2011 年 9 月 28 日 目次 1. TSUBAMEのGPU 環境 2. プログラム作成 3. プログラム実行 4. 性能解析 デバッグ サンプルコードは /work0/gsic/seminars/gpu- 2011-09- 28 からコピー可能です 1.
CUDA 連携とライブラリの活用 2
1 09:30-10:00 受付 10:00-12:00 Reedbush-H ログイン GPU 入門 13:30-15:00 OpenACC 入門 15:15-16:45 OpenACC 最適化入門と演習 17:00-18:00 OpenACC の活用 (CUDA 連携とライブラリの活用 ) CUDA 連携とライブラリの活用 2 3 OpenACC 簡単にGPUプログラムが作成できる それなりの性能が得られる
Slides: TimeGraph: GPU Scheduling for Real-Time Multi-Tasking Environments
計算機アーキテクチャ第 11 回 マルチプロセッサ 本資料は授業用です 無断で転載することを禁じます 名古屋大学 大学院情報科学研究科 准教授加藤真平 デスクトップ ジョブレベル並列性 スーパーコンピュータ 並列処理プログラム プログラムの並列化 for (i = 0; i < N; i++) { x[i] = a[i] + b[i]; } プログラムの並列化 x[0] = a[0] + b[0];
演習1: 演習準備
演習 1: 演習準備 2013 年 8 月 6 日神戸大学大学院システム情報学研究科森下浩二 1 演習 1 の内容 神戸大 X10(π-omputer) について システム概要 ログイン方法 コンパイルとジョブ実行方法 OpenMP の演習 ( 入門編 ) 1. parallel 構文 実行時ライブラリ関数 2. ループ構文 3. shared 節 private 節 4. reduction 節
Copyright 2009, SofTek Systems, Inc. All rights reserved.
PGI Visual Fortran Release 9.0 SofTek Copyright 2009, SofTek Systems, Inc. All rights reserved. \\\ \\ \\\ \\ \\ SofTek Systems, Inc \ SofTek Systems, Inc SofTek Systems, Inc SofTek Systems, Inc SofTek
概要 OpenACC とは OpenACC について OpenMP, CUDA との違い OpenACC の指示文 並列化領域指定指示文 (kernels/parallel) データ移動指示文 ループ指示文 OpenACC の実用例 実習 コンパイラメッセージの見方 OpenACC プログラムの実装
第 74 回お試しアカウント付き 並列プログラミング講習会 GPU プログラミング入門 in 名古屋 星野哲也 ( 助教 ) [email protected] 大島聡史 ( 助教 ) [email protected] 2016 年 3 月 14 日 ( 火 ) 東京大学情報基盤センター 概要 OpenACC とは OpenACC について OpenMP, CUDA との違い
1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境 Lin
Windows で始める CUDA 入門 GTC 2013 チュートリアル エヌビディアジャパン CUDA エンジニア森野慎也 1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境
Slide 1
CUDA プログラミングの基本 パート II - カーネル CUDA の基本の概要 パート I CUDAのソフトウェアスタックとコンパイル GPUのメモリ管理 パート II カーネルの起動 GPUコードの具体像 注 : 取り上げているのは基本事項のみです そのほか多数の API 関数についてはプログラミングガイドを ご覧ください GPU 上でのコードの実行 カーネルは C 関数 + 多少の制約 ホストメモリはアクセスできない戻り値型は
スライド 1
GPU クラスタによる格子 QCD 計算 広大理尾崎裕介 石川健一 1.1 Introduction Graphic Processing Units 1 チップに数百個の演算器 多数の演算器による並列計算 ~TFLOPS ( 単精度 ) CPU 数十 GFLOPS バンド幅 ~100GB/s コストパフォーマンス ~$400 GPU の開発環境 NVIDIA CUDA http://www.nvidia.co.jp/object/cuda_home_new_jp.html
インテル(R) Visual Fortran Composer XE 2013 Windows版 入門ガイド
Visual Fortran Composer XE 2013 Windows* エクセルソフト株式会社 www.xlsoft.com Rev. 1.1 (2012/12/10) Copyright 1998-2013 XLsoft Corporation. All Rights Reserved. 1 / 53 ... 3... 4... 4... 5 Visual Studio... 9...
OpenMP¤òÍѤ¤¤¿ÊÂÎó·×»»¡Ê£±¡Ë
2012 5 24 scalar Open MP Hello World Do (omp do) (omp workshare) (shared, private) π (reduction) PU PU PU 2 16 OpenMP FORTRAN/C/C++ MPI OpenMP 1997 FORTRAN Ver. 1.0 API 1998 C/C++ Ver. 1.0 API 2000 FORTRAN
OpenMP¤òÍѤ¤¤¿ÊÂÎó·×»»¡Ê£±¡Ë
2011 5 26 scalar Open MP Hello World Do (omp do) (omp workshare) (shared, private) π (reduction) scalar magny-cours, 48 scalar scalar 1 % scp. ssh / authorized keys 133. 30. 112. 246 2 48 % ssh 133.30.112.246
07-二村幸孝・出口大輔.indd
GPU Graphics Processing Units HPC High Performance Computing GPU GPGPU General-Purpose computation on GPU CPU GPU GPU *1 Intel Quad-Core Xeon E5472 3.0 GHz 2 6 MB L2 cache 1600 MHz FSB 80 GFlops 1 nvidia
概要 目的 CUDA Fortran の利用に関する基本的なノウハウを提供する 本チュートリアル受講後は Web 上で公開されている資料等を参照しながら独力で CUDA Fortran が利用できることが目標 対象 CUDA Fortran の利用に興味を抱いている方 前提とする知識 Fortran
CUDA Fortran チュートリアル 2010 年 9 月 29 日 NEC 概要 目的 CUDA Fortran の利用に関する基本的なノウハウを提供する 本チュートリアル受講後は Web 上で公開されている資料等を参照しながら独力で CUDA Fortran が利用できることが目標 対象 CUDA Fortran の利用に興味を抱いている方 前提とする知識 Fortran を用いた Linux
ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014
ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014 コンカレントな処理の実行 システム内部の複数の処理を 平行に実行する CPU GPU メモリ転送 カーネル実行 複数のカーネル間 ストリーム GPU 上の処理キュー カーネル実行 メモリ転送の並列性 実行順序 DEFAULT STREAM Stream : GPU
OpenACC
109 OpenMP/OpenACC, hoshino @ cc.u-tokyo.ac.jp nakajima @ cc.u-tokyo.ac.jp 1 n Reedbush n $ ssh -Y reedbush.cc.u-tokyo.ac.jp l txxxxx n module n $ module load pgi/18.7 # n n $ cdw n OpenACC_samples n $
01_OpenMP_osx.indd
OpenMP* / 1 1... 2 2... 3 3... 5 4... 7 5... 9 5.1... 9 5.2 OpenMP* API... 13 6... 17 7... 19 / 4 1 2 C/C++ OpenMP* 3 Fortran OpenMP* 4 PC 1 1 9.0 Linux* Windows* Xeon Itanium OS 1 2 2 WEB OS OS OS 1 OS
3次多項式パラメタ推定計算の CUDAを用いた実装 (CUDAプログラミングの練習として) Implementation of the Estimation of the parameters of 3rd-order-Polynomial with CUDA
3 次多項式パラメタ推定計算の CUDA を用いた実装 (CUDA プログラミングの練習として ) Estimating the Parameters of 3rd-order-Polynomial with CUDA ISS 09/11/12 問題の選択 目的 CUDA プログラミングを経験 ( 試行錯誤と習得 ) 実際に CPU のみの場合と比べて高速化されることを体験 問題 ( インプリメントする内容
CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン
CUDA 画像処理入門 エヌビディアジャパン CUDA エンジニア森野慎也 GTC Japan 2014 CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン RGB Y( 輝度 ) 変換 カラー画像から グレイスケールへの変換 Y = 0.299 R + 0.587
Microsoft PowerPoint - GPU_computing_2013_01.pptx
GPU コンピューティン No.1 導入 東京工業大学 学術国際情報センター 青木尊之 1 GPU とは 2 GPGPU (General-purpose computing on graphics processing units) GPU を画像処理以外の一般的計算に使う GPU の魅力 高性能 : ハイエンド GPU はピーク 4 TFLOPS 超 手軽さ : 普通の PC にも装着できる 低価格
コードのチューニング
OpenMP による並列化実装 八木学 ( 理化学研究所計算科学研究センター ) KOBE HPC Spring School 2019 2019 年 3 月 14 日 スレッド並列とプロセス並列 スレッド並列 OpenMP 自動並列化 プロセス並列 MPI プロセス プロセス プロセス スレッドスレッドスレッドスレッド メモリ メモリ プロセス間通信 Private Private Private
I I / 47
1 2013.07.18 1 I 2013 3 I 2013.07.18 1 / 47 A Flat MPI B 1 2 C: 2 I 2013.07.18 2 / 47 I 2013.07.18 3 / 47 #PJM -L "rscgrp=small" π-computer small: 12 large: 84 school: 24 84 16 = 1344 small school small
OpenMP¤òÍѤ¤¤¿ÊÂÎó·×»»¡Ê£²¡Ë
2013 5 30 (schedule) (omp sections) (omp single, omp master) (barrier, critical, atomic) program pi i m p l i c i t none integer, parameter : : SP = kind ( 1. 0 ) integer, parameter : : DP = selected real
GPGPU
GPGPU 2013 1008 2015 1 23 Abstract In recent years, with the advance of microscope technology, the alive cells have been able to observe. On the other hand, from the standpoint of image processing, the
GPU GPU CPU CPU CPU GPU GPU N N CPU ( ) 1 GPU CPU GPU 2D 3D CPU GPU GPU GPGPU GPGPU 2 nvidia GPU CUDA 3 GPU 3.1 GPU Core 1
GPU 4 2010 8 28 1 GPU CPU CPU CPU GPU GPU N N CPU ( ) 1 GPU CPU GPU 2D 3D CPU GPU GPU GPGPU GPGPU 2 nvidia GPU CUDA 3 GPU 3.1 GPU Core 1 Register & Shared Memory ( ) CPU CPU(Intel Core i7 965) GPU(Tesla
Microsoft PowerPoint - suda.pptx
GPU の HWアーキテクチャと高性能化手法 須田礼仁 ( 東京大学 ) 2011/03/22 GPU 高性能プログラミング GPU のハードウェアを理解する CUDA のソフトウェアを理解する CUDA でプログラムを書くのは難しくないが, CUDA で高速なプログラムを書くのは難しい どうすれば遅くなるかを理解する! 効果が大きいものから順に説明します 1 高性能プログラミングの手順 1. 現在のコードの,
( CUDA CUDA CUDA CUDA ( NVIDIA CUDA I
GPGPU (II) GPGPU CUDA 1 GPGPU CUDA(CUDA Unified Device Architecture) CUDA NVIDIA GPU *1 C/C++ (nvcc) CUDA NVIDIA GPU GPU CUDA CUDA 1 CUDA CUDA 2 CUDA NVIDIA GPU PC Windows Linux MaxOSX CUDA GPU CUDA NVIDIA
openmp1_Yaguchi_version_170530
並列計算とは /OpenMP の初歩 (1) 今 の内容 なぜ並列計算が必要か? スーパーコンピュータの性能動向 1ExaFLOPS 次世代スハ コン 京 1PFLOPS 性能 1TFLOPS 1GFLOPS スカラー機ベクトル機ベクトル並列機並列機 X-MP ncube2 CRAY-1 S-810 SR8000 VPP500 CM-5 ASCI-5 ASCI-4 S3800 T3E-900 SR2201
Microsoft PowerPoint - sales2.ppt
最適化とは何? CPU アーキテクチャに沿った形で最適な性能を抽出できるようにする技法 ( 性能向上技法 ) コンパイラによるプログラム最適化 コンパイラメーカの技量 経験量に依存 最適化ツールによるプログラム最適化 KAP (Kuck & Associates, Inc. ) 人によるプログラム最適化 アーキテクチャのボトルネックを知ること 3 使用コンパイラによる性能の違い MFLOPS 90
Vol.214-HPC-145 No /7/3 C #pragma acc directive-name [clause [[,] clause] ] new-line structured block Fortran!$acc directive-name [clause [[,] c
Vol.214-HPC-145 No.45 214/7/3 OpenACC 1 3,1,2 1,2 GPU CUDA OpenCL OpenACC OpenACC High-level OpenACC CPU Intex Xeon Phi K2X GPU Intel Xeon Phi 27% K2X GPU 24% 1. TSUBAME2.5 CPU GPU CUDA OpenCL CPU OpenMP
XcalableMP入門
XcalableMP 1 HPC-Phys@, 2018 8 22 XcalableMP XMP XMP Lattice QCD!2 XMP MPI MPI!3 XMP 1/2 PCXMP MPI Fortran CCoarray C++ MPIMPI XMP OpenMP http://xcalablemp.org!4 XMP 2/2 SPMD (Single Program Multiple Data)
GPGPUクラスタの性能評価
2008 年度理研 HPC シンポジウム第 3 世代 PC クラスタ GPGPU クラスタの性能評価 2009 年 3 月 12 日 富士通研究所成瀬彰 発表の概要 背景 GPGPU による高速化 CUDA の概要 GPU のメモリアクセス特性調査 姫野 BMT の高速化 GPGPU クラスタによる高速化 GPU Host 間のデータ転送 GPU-to-GPU の通信性能 GPGPU クラスタ上での姫野
Microsoft PowerPoint - 01_Vengineer.ppt
Software Driven Verification テストプログラムは C 言語で! SystemVerilog DPI-C を使えば こんなに便利に! 2011 年 9 月 30 日 コントローラ開発本部コントローラプラットフォーム第五開発部 宮下晴信 この資料で使用するシステム名 製品名等は一般にメーカーや 団体の登録商標などになっているものもあります なお この資料の中では トレードマーク
修士論文
AVX を用いた倍々精度疎行列ベクトル積の高速化 菱沼利彰 1 藤井昭宏 1 田中輝雄 1 長谷川秀彦 2 1 工学院大学 2 筑波大学 1 目次 1. 研究背景 目的 2. 実装, 実験環境 3. 実験 - 倍々精度ベクトル演算 - 4. 実験 - 倍々精度疎行列ベクトル積 - 5. まとめ 多倍長精度計算フォーラム 2 目次 1. 研究背景 目的 2. 実装, 実験環境 3. 実験 - 倍々精度ベクトル演算
Microsoft PowerPoint - GPGPU実践基礎工学(web).pptx
GPU のプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 GPU プログラミング環境 (CUDA) GPU プログラムの実行の流れ CUDA によるプログラムの記述 カーネル (GPU で処理する関数 ) の構造 記述方法とその理由 GPU 固有のパラメータの確認 405 GPU(Graphics Processing Unit) とは 画像処理専用のハードウェア 具体的には画像処理用のチップ
<4D F736F F F696E74202D D F95C097F D834F E F93FC96E5284D F96E291E85F8DE391E52E >
SX-ACE 並列プログラミング入門 (MPI) ( 演習補足資料 ) 大阪大学サイバーメディアセンター日本電気株式会社 演習問題の構成 ディレクトリ構成 MPI/ -- practice_1 演習問題 1 -- practice_2 演習問題 2 -- practice_3 演習問題 3 -- practice_4 演習問題 4 -- practice_5 演習問題 5 -- practice_6
熊本大学学術リポジトリ Kumamoto University Repositor Title GPGPU による高速演算について Author(s) 榎本, 昌一 Citation Issue date Type URL Presentation
熊本大学学術リポジトリ Kumamoto University Repositor Title GPGPU による高速演算について Author(s) 榎本, 昌一 Citation Issue date 2011-03-17 Type URL Presentation http://hdl.handle.net/2298/23539 Right GPGPU による高速演算について 榎本昌一 東京大学大学院工学系研究科システム創成学専攻
Microsoft PowerPoint - 演習1:並列化と評価.pptx
講義 2& 演習 1 プログラム並列化と性能評価 神戸大学大学院システム情報学研究科横川三津夫 [email protected] 2014/3/5 RIKEN AICS HPC Spring School 2014: プログラム並列化と性能評価 1 2014/3/5 RIKEN AICS HPC Spring School 2014: プログラム並列化と性能評価 2 2 次元温度分布の計算
研究背景 大規模な演算を行うためには 分散メモリ型システムの利用が必須 Message Passing Interface MPI 並列プログラムの大半はMPIを利用 様々な実装 OpenMPI, MPICH, MVAPICH, MPI.NET プログラミングコストが高いため 生産性が悪い 新しい並
XcalableMPによる NAS Parallel Benchmarksの実装と評価 中尾 昌広 李 珍泌 朴 泰祐 佐藤 三久 筑波大学 計算科学研究センター 筑波大学大学院 システム情報工学研究科 研究背景 大規模な演算を行うためには 分散メモリ型システムの利用が必須 Message Passing Interface MPI 並列プログラムの大半はMPIを利用 様々な実装 OpenMPI,
main.dvi
PC 1 1 [1][2] [3][4] ( ) GPU(Graphics Processing Unit) GPU PC GPU PC ( 2 GPU ) GPU Harris Corner Detector[5] CPU ( ) ( ) CPU GPU 2 3 GPU 4 5 6 7 1 [email protected] 45 2 ( ) CPU ( ) ( ) () 2.1
! 行行 CPUDSP PPESPECell/B.E. CPUGPU 行行 SIMD [SSE, AltiVec] 用 HPC CPUDSP PPESPE (Cell/B.E.) SPE CPUGPU GPU CPU DSP DSP PPE SPE SPE CPU DSP SPE 2
! OpenCL [Open Computing Language] 言 [OpenCL C 言 ] CPU, GPU, Cell/B.E.,DSP 言 行行 [OpenCL Runtime] OpenCL C 言 API Khronos OpenCL Working Group AMD Broadcom Blizzard Apple ARM Codeplay Electronic Arts Freescale
(Microsoft PowerPoint \215u\213`4\201i\221\272\210\344\201j.pptx)
AICS 村井均 RIKEN AICS HPC Summer School 2012 8/7/2012 1 背景 OpenMP とは OpenMP の基本 OpenMP プログラミングにおける注意点 やや高度な話題 2 共有メモリマルチプロセッサシステムの普及 共有メモリマルチプロセッサシステムのための並列化指示文を共通化する必要性 各社で仕様が異なり 移植性がない そして いまやマルチコア プロセッサが主流となり
Slide 1
GPUコンピューティング入門 2015.08.26 エヌビディア合同会社 CUDAエンジニア 村上真奈 エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90 分 ) CUDA 入門 (90 分 ) Agenda 2 エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90
2012年度HPCサマーセミナー_多田野.pptx
! CCS HPC! I " [email protected]" " 1 " " " " " " " 2 3 " " Ax = b" " " 4 Ax = b" A = a 11 a 12... a 1n a 21 a 22... a 2n...... a n1 a n2... a nn, x = x 1 x 2. x n, b = b 1 b 2. b n " " 5 Gauss LU
PGIコンパイラ導入手順
1 注意この資料は PGI compiler 18.10 が最新であるときに作成した資料を元にしています PGI compiler 19.4 がリリースされましたが インストール手順や利用手順は 18.10 と変わりません 資料中の 1810 を 194 に 18.10 を 19.4 に読み替えてください 2019 年 6 月版 2 大きく分けて以下の 3 つの方法が利用可能 1. 手元のウェブブラウザでダウンロードして
Intel® Compilers Professional Editions
2007 6 10.0 * 10.0 6 5 Software &Solutions group 10.0 (SV) C++ Fortran OpenMP* OpenMP API / : 200 C/C++ Fortran : OpenMP : : : $ cat -n main.cpp 1 #include 2 int foo(const char *); 3 int main()
MATLAB® における並列・分散コンピューティング ~ Parallel Computing Toolbox™ & MATLAB Distributed Computing Server™ ~
MATLAB における並列 分散コンピューティング ~ Parallel Computing Toolbox & MATLAB Distributed Computing Server ~ MathWorks Japan Application Engineering Group Takashi Yoshida 2016 The MathWorks, Inc. 1 System Configuration
第12回講義(2019年7月17日)
スパコンプログラミング (1)(Ⅰ) 1 OpenACC の紹介 Reedbush-H お試し 東京大学情報基盤センター准教授塙敏博 2019 年 7 月 17 日 ( 水 )10:25 12:10 2019/7/16 スパコンプログラミング (1) (Ⅰ) 講義日程 ( 工学部共通科目 ) 1. 4 月 9 日 : ガイダンス 2. 4 月 16 日 l 並列数値処理の基本演算 ( 座学 ) 3.
Microsoft Word ●IntelクアッドコアCPUでのベンチマーク_吉岡_ _更新__ doc
2.3. アプリ性能 2.3.1. Intel クアッドコア CPU でのベンチマーク 東京海洋大学吉岡諭 1. はじめにこの数年でマルチコア CPU の普及が進んできた x86 系の CPU でも Intel と AD がデュアルコア クアッドコアの CPU を次々と市場に送り出していて それらが PC クラスタの CPU として採用され HPC に活用されている ここでは Intel クアッドコア
スパコンに通じる並列プログラミングの基礎
2018.09.10 [email protected] ( ) 2018.09.10 1 / 59 [email protected] ( ) 2018.09.10 2 / 59 Windows, Mac Unix 0444-J [email protected] ( ) 2018.09.10 3 / 59 Part I Unix GUI CUI:
スパコンに通じる並列プログラミングの基礎
2018.06.04 2018.06.04 1 / 62 2018.06.04 2 / 62 Windows, Mac Unix 0444-J 2018.06.04 3 / 62 Part I Unix GUI CUI: Unix, Windows, Mac OS Part II 2018.06.04 4 / 62 0444-J ( : ) 6 4 ( ) 6 5 * 6 19 SX-ACE * 6
Slide 1
OpenACC CUDA による GPU コンピューティング Akira Naruse, 19 th Jul. 2018 成瀬彰 (Naruse, Akira) 自己紹介 2013 年 ~: NVIDIA シニア デベローパーテクノロジー エンジニア 1996~2013 年 : 富士通研究所 研究員など 専門 興味 : 並列処理 性能最適化 スパコン HPC GPU コンピューティング DeepLearning
OpenMP (1) 1, 12 1 UNIX (FUJITSU GP7000F model 900), 13 1 (COMPAQ GS320) FUJITSU VPP5000/64 1 (a) (b) 1: ( 1(a))
OpenMP (1) 1, 12 1 UNIX (FUJITSU GP7000F model 900), 13 1 (COMPAQ GS320) FUJITSU VPP5000/64 1 (a) (b) 1: ( 1(a)) E-mail: {nanri,amano}@cc.kyushu-u.ac.jp 1 ( ) 1. VPP Fortran[6] HPF[3] VPP Fortran 2. MPI[5]
Agenda GRAPE-MPの紹介と性能評価 GRAPE-MPの概要 OpenCLによる四倍精度演算 (preliminary) 4倍精度演算用SIM 加速ボード 6 processor elem with 128 bit logic Peak: 1.2Gflops
Agenda GRAPE-MPの紹介と性能評価 GRAPE-MPの概要 OpenCLによる四倍精度演算 (preliminary) 4倍精度演算用SIM 加速ボード 6 processor elem with 128 bit logic Peak: 1.2Gflops ボードの概要 Control processor (FPGA by Altera) GRAPE-MP chip[nextreme
OpenMPプログラミング
OpenMP 基礎 岩下武史 ( 学術情報メディアセンター ) 1 2013/9/13 並列処理とは 逐次処理 CPU1 並列処理 CPU1 CPU2 CPU3 CPU4 処理 1 処理 1 処理 2 処理 3 処理 4 処理 2 処理 3 処理 4 時間 2 2 種類の並列処理方法 プロセス並列 スレッド並列 並列プログラム 並列プログラム プロセス プロセス 0 プロセス 1 プロセス間通信 スレッド
PowerPoint Presentation
ヘテロジニアスな環境におけるソフトウェア開発 Agenda 今日の概要 ヘテロジニアスな環境の登場 ホモジニアスからヘテロジニアスへ ヘテロジニアスなアーキテクチャ GPU CUDA OpenACC, XeonPhi 自分のプログラムを理解するために デバッガ 共通の操作体験 TotalView 続きはブースで より速く ホモジーニアスな並列 HPC 銀河生成 金融のリスク計算 車の衝突解析 製薬
hotspot の特定と最適化
1 1? 1 1 2 1. hotspot : hotspot hotspot Parallel Amplifier 1? 2. hotspot : (1 ) Parallel Composer 1 Microsoft* Ticker Tape Smoke 1.0 PiSolver 66 / 64 / 2.76 ** 84 / 27% ** 75 / 17% ** 1.46 89% Microsoft*
1 4 1.1........................................... 4 1.2.................................. 4 1.3................................... 4 2 5 2.1 GPU.....
CPU GPU N Q07-065 2011 2 17 1 1 4 1.1........................................... 4 1.2.................................. 4 1.3................................... 4 2 5 2.1 GPU...........................................
23 Fig. 2: hwmodulev2 3. Reconfigurable HPC 3.1 hw/sw hw/sw hw/sw FPGA PC FPGA PC FPGA HPC FPGA FPGA hw/sw hw/sw hw- Module FPGA hwmodule hw/sw FPGA h
23 FPGA CUDA Performance Comparison of FPGA Array with CUDA on Poisson Equation ([email protected]), ([email protected]), ([email protected]), ([email protected]),
演習2
神戸市立工業高等専門学校電気工学科 / 電子工学科専門科目 数値解析 2017.6.2 演習 2 山浦剛 ([email protected]) 講義資料ページ h t t p://clim ate.aic s. riken. jp/m embers/yamaura/num erical_analysis. html 曲線の推定 N 次多項式ラグランジュ補間 y = p N x = σ N x x
スパコンに通じる並列プログラミングの基礎
2016.06.06 2016.06.06 1 / 60 2016.06.06 2 / 60 Windows, Mac Unix 0444-J 2016.06.06 3 / 60 Part I Unix GUI CUI: Unix, Windows, Mac OS Part II 0444-J 2016.06.06 4 / 60 ( : ) 6 6 ( ) 6 10 6 16 SX-ACE 6 17
目次 1 はじめに 本文書の概要 PVF ソフトウェアと VISUAL STUDIO PVF ソフトウェアの種類 MICROSOFT VISUAL STUDIO の日本語化について VISUAL STUDIO
PGI Visual Fortran のための Microsoft Visual Studio 導入ガイド 2016 年版 日本語環境の Visual Studio の構築について PGI インストール関係の日本語ドキュメントは 以下の URL に全てアーカイブしてありま す オンラインでご覧になりたい場合は 以下の URL にアクセスしてください http://www.softek.co.jp/spg/pgi/inst_document.html
4 倍精度基本線形代数ルーチン群 QPBLAS の紹介 [index] 1. Introduction 2. Double-double algorithm 3. QPBLAS 4. QPBLAS-GPU 5. Summary 佐々成正 1, 山田進 1, 町田昌彦 1, 今村俊幸 2, 奥田洋司
4 倍精度基本線形代数ルーチン群 QPBLAS の紹介 [index] 1. Introduction 2. Double-double algorithm 3. QPBLAS 4. QPBLAS-GPU 5. Summary 佐々成正 1, 山田進 1, 町田昌彦 1, 今村俊幸 2, 奥田洋司 3 1 1 日本原子力研究開発機構システム計算科学センター 2 理科学研究所計算科学研究機構 3 東京大学新領域創成科学研究科
CCS HPCサマーセミナー 並列数値計算アルゴリズム
大規模系での高速フーリエ変換 2 高橋大介 [email protected] 筑波大学計算科学研究センター 2016/6/2 計算科学技術特論 B 1 講義内容 並列三次元 FFT における自動チューニング 二次元分割を用いた並列三次元 FFT アルゴリズム GPU クラスタにおける並列三次元 FFT 2016/6/2 計算科学技術特論 B 2 並列三次元 FFT における 自動チューニング
