Microsoft PowerPoint - GTC2012-SofTek.pptx

Size: px
Start display at page:

Download "Microsoft PowerPoint - GTC2012-SofTek.pptx"

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

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

More information

GPU チュートリアル :OpenACC 篇 Himeno benchmark を例題として 高エネルギー加速器研究機構 (KEK) 松古栄夫 (Hideo Matsufuru) 1 December 2018 HPC-Phys 理化学研究所 共通コードプロジェクト

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

More information

スライド 1

スライド 1 GTC Japan 2013 PGI Accelerator Compiler 新 OpenACC 2.0 の機能と PGI アクセラレータコンパイラ 2013 年 7 月 加藤努株式会社ソフテック 本日の話 OpenACC ディレクティブで出来ることを改めて知ろう! OpenACC 1.0 の復習 ディレクティブ操作で出来ることを再確認 OpenACC 2.0 の新機能 プログラミングの自由度の向上へ

More information

OpenACCによる並列化

OpenACCによる並列化 実習 OpenACC による ICCG ソルバーの並列化 1 ログイン Reedbush へのログイン $ ssh reedbush.cc.u-tokyo.ac.jp l txxxxx Module のロード $ module load pgi/17.3 cuda ログインするたびに必要です! ワークディレクトリに移動 $ cdw ターゲットプログラム /srcx OpenACC 用のディレクトリの作成

More information

TSUBAME2.0 における GPU の 活用方法 東京工業大学学術国際情報センター丸山直也第 10 回 GPU コンピューティング講習会 2011 年 9 月 28 日

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.

More information

CUDA 連携とライブラリの活用 2

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プログラムが作成できる それなりの性能が得られる

More information

Microsoft PowerPoint - GDEP-GPG_softek_May24-1.pptx

Microsoft PowerPoint - GDEP-GPG_softek_May24-1.pptx G-DEP 第 3 回セミナー ツールで始める GPGPU なぜ ディレクティブベースのプログラミングが有望なのか? 失敗しない並列プログラミングの始め 2012 年 5 月 加藤努株式会社ソフテック なぜ GPU プログラミング? GPU computing 時代の本流に! 今 何が起きているのか? HPC アーキテクチャはどこへ向かうのか? 確実に移行中 Why? CPU + GPU ハイブリッド型へ

More information

Slides: TimeGraph: GPU Scheduling for Real-Time Multi-Tasking Environments

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];

More information

演習1: 演習準備

演習1: 演習準備 演習 1: 演習準備 2013 年 8 月 6 日神戸大学大学院システム情報学研究科森下浩二 1 演習 1 の内容 神戸大 X10(π-omputer) について システム概要 ログイン方法 コンパイルとジョブ実行方法 OpenMP の演習 ( 入門編 ) 1. parallel 構文 実行時ライブラリ関数 2. ループ構文 3. shared 節 private 節 4. reduction 節

More information

Copyright 2009, SofTek Systems, Inc. All rights reserved.

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

More information

概要 OpenACC とは OpenACC について OpenMP, CUDA との違い OpenACC の指示文 並列化領域指定指示文 (kernels/parallel) データ移動指示文 ループ指示文 OpenACC の実用例 実習 コンパイラメッセージの見方 OpenACC プログラムの実装

概要 OpenACC とは OpenACC について OpenMP, CUDA との違い OpenACC の指示文 並列化領域指定指示文 (kernels/parallel) データ移動指示文 ループ指示文 OpenACC の実用例 実習 コンパイラメッセージの見方 OpenACC プログラムの実装 第 74 回お試しアカウント付き 並列プログラミング講習会 GPU プログラミング入門 in 名古屋 星野哲也 ( 助教 ) hoshino@cc.u-tokyo.ac.jp 大島聡史 ( 助教 ) ohshima@cc.u-tokyo.ac.jp 2016 年 3 月 14 日 ( 火 ) 東京大学情報基盤センター 概要 OpenACC とは OpenACC について OpenMP, CUDA との違い

More information

1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境 Lin

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 コンピューティング環境

More information

HPC143

HPC143 研究背景 GPUクラスタ 高性能 高いエネルギー効率 低価格 様々なHPCアプリケーションで用いられている TCA (Tightly Coupled Accelerators) 密結合並列演算加速機構 筑波大学HA-PACSクラスタ アクセラレータ GPU 間の直接通信 低レイテンシ 今後のHPCアプリは強スケーリングも重要 TCAとアクセラレータを搭載したシステムに おけるプログラミングモデル 例

More information

Slide 1

Slide 1 CUDA プログラミングの基本 パート II - カーネル CUDA の基本の概要 パート I CUDAのソフトウェアスタックとコンパイル GPUのメモリ管理 パート II カーネルの起動 GPUコードの具体像 注 : 取り上げているのは基本事項のみです そのほか多数の API 関数についてはプログラミングガイドを ご覧ください GPU 上でのコードの実行 カーネルは C 関数 + 多少の制約 ホストメモリはアクセスできない戻り値型は

More information

スライド 1

スライド 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

More information

Microsoft PowerPoint - OpenMP入門.pptx

Microsoft PowerPoint - OpenMP入門.pptx OpenMP 入門 須田礼仁 2009/10/30 初版 OpenMP 共有メモリ並列処理の標準化 API http://openmp.org/ 最新版は 30 3.0 バージョンによる違いはあまり大きくない サポートしているバージョンはともかく csp で動きます gcc も対応しています やっぱり SPMD Single Program Multiple Data プログラム #pragma omp

More information

インテル(R) Visual Fortran Composer XE 2013 Windows版 入門ガイド

インテル(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...

More information

OpenMP¤òÍѤ¤¤¿ÊÂÎó·×»»¡Ê£±¡Ë

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

More information

OpenMP¤òÍѤ¤¤¿ÊÂÎó·×»»¡Ê£±¡Ë

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

More information

07-二村幸孝・出口大輔.indd

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

More information

Microsoft PowerPoint - 03_What is OpenMP 4.0 other_Jan18

Microsoft PowerPoint - 03_What is OpenMP 4.0 other_Jan18 OpenMP* 4.x における拡張 OpenMP 4.0 と 4.5 の機能拡張 内容 OpenMP* 3.1 から 4.0 への拡張 OpenMP* 4.0 から 4.5 への拡張 2 追加された機能 (3.1 -> 4.0) C/C++ 配列シンタックスの拡張 SIMD と SIMD 対応関数 デバイスオフロード task 構 の依存性 taskgroup 構 cancel 句と cancellation

More information

目次 1 はじめに 本文書の概要 WINDOWS 上での PVF ソフトウェアの実装 PVF コンパイラの利用方法 PVF コンパイラのコマンド オプションについて PVF コンパイラの起動 (Microsoft

目次 1 はじめに 本文書の概要 WINDOWS 上での PVF ソフトウェアの実装 PVF コンパイラの利用方法 PVF コンパイラのコマンド オプションについて PVF コンパイラの起動 (Microsoft PGI (Accelerator) Visual Fortran 2011 For Microsoft Visual Studio Windows 版 (Release 2011) - 入門ガイド - PGI インストール関係の日本語ドキュメントは 以下の URL に全てアーカイブしてあります オンラインでご覧になりたい場合は 以下の URL にアクセスしてください http://www.softek.co.jp/spg/pgi/inst_document.html

More information

概要 目的 CUDA Fortran の利用に関する基本的なノウハウを提供する 本チュートリアル受講後は Web 上で公開されている資料等を参照しながら独力で CUDA Fortran が利用できることが目標 対象 CUDA Fortran の利用に興味を抱いている方 前提とする知識 Fortran

概要 目的 CUDA Fortran の利用に関する基本的なノウハウを提供する 本チュートリアル受講後は Web 上で公開されている資料等を参照しながら独力で CUDA Fortran が利用できることが目標 対象 CUDA Fortran の利用に興味を抱いている方 前提とする知識 Fortran CUDA Fortran チュートリアル 2010 年 9 月 29 日 NEC 概要 目的 CUDA Fortran の利用に関する基本的なノウハウを提供する 本チュートリアル受講後は Web 上で公開されている資料等を参照しながら独力で CUDA Fortran が利用できることが目標 対象 CUDA Fortran の利用に興味を抱いている方 前提とする知識 Fortran を用いた Linux

More information

ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014

ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014 ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014 コンカレントな処理の実行 システム内部の複数の処理を 平行に実行する CPU GPU メモリ転送 カーネル実行 複数のカーネル間 ストリーム GPU 上の処理キュー カーネル実行 メモリ転送の並列性 実行順序 DEFAULT STREAM Stream : GPU

More information

OpenACC

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 $

More information

01_OpenMP_osx.indd

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

More information

Microsoft Word - openmp-txt.doc

Microsoft Word - openmp-txt.doc ( 付録 A) OpenMP チュートリアル OepnMP は 共有メモリマルチプロセッサ上のマルチスレッドプログラミングのための API です 本稿では OpenMP の簡単な解説とともにプログラム例をつかって説明します 詳しくは OpenMP の規約を決めている OpenMP ARB の http://www.openmp.org/ にある仕様書を参照してください 日本語訳は http://www.hpcc.jp/omni/spec.ja/

More information

3次多項式パラメタ推定計算の CUDAを用いた実装 (CUDAプログラミングの練習として) Implementation of the Estimation of the parameters of 3rd-order-Polynomial with CUDA

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 のみの場合と比べて高速化されることを体験 問題 ( インプリメントする内容

More information

Slide 1

Slide 1 CUDA プログラミングの基本 パート I - ソフトウェアスタックとメモリ管理 CUDA の基本の概要 パート I CUDAのソフトウェアスタックとコンパイル GPUのメモリ管理 パートII カーネルの起動 GPUコードの具体項目 注 : 取り上げているのは基本事項のみです そのほか多数の API 関数についてはプログラミングガイドを ご覧ください CUDA インストレーション CUDA インストレーションの構成

More information

CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン

CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン CUDA 画像処理入門 エヌビディアジャパン CUDA エンジニア森野慎也 GTC Japan 2014 CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン RGB Y( 輝度 ) 変換 カラー画像から グレイスケールへの変換 Y = 0.299 R + 0.587

More information

目次 1 はじめに 本文書の概要 WINDOWS 上での PVF ソフトウェアの実装 PVF コンパイラの利用方法 PVF コンパイラのコマンド オプションについて PVF コンパイラの起動 (Microsoft

目次 1 はじめに 本文書の概要 WINDOWS 上での PVF ソフトウェアの実装 PVF コンパイラの利用方法 PVF コンパイラのコマンド オプションについて PVF コンパイラの起動 (Microsoft PGI (Accelerator) Visual Fortran 2018 For Microsoft Visual Studio Windows 版 (Release 2018) - 入門ガイド - 2018 年 2 月版 (Rev. 18.1-A) 株式会社ソフテック HPC ソリューション部 (http://www.softek.co.jp/spg/) SofTek 目次 1 はじめに...

More information

Microsoft PowerPoint - GPU_computing_2013_01.pptx

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 にも装着できる 低価格

More information

コードのチューニング

コードのチューニング OpenMP による並列化実装 八木学 ( 理化学研究所計算科学研究センター ) KOBE HPC Spring School 2019 2019 年 3 月 14 日 スレッド並列とプロセス並列 スレッド並列 OpenMP 自動並列化 プロセス並列 MPI プロセス プロセス プロセス スレッドスレッドスレッドスレッド メモリ メモリ プロセス間通信 Private Private Private

More information

I I / 47

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

More information

OpenMP¤òÍѤ¤¤¿ÊÂÎó·×»»¡Ê£²¡Ë

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

More information

GPGPU

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

More information

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 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

More information

並列・高速化を実現するための 高速化サービスの概要と事例紹介

並列・高速化を実現するための 高速化サービスの概要と事例紹介 第 4 回 AVS 可視化フォーラム 2019 並列 高速化を実現するための 高速化サービスの概要と事例紹介 株式会社アーク情報システム営業部仮野亮ソリューション技術部佐々木竜一 2019.08.30 はじめに アーク情報システムの紹介 高速化サービスとは? 事例紹介 コンサルティングサービスについて アーク情報システムの紹介 設立 資本金 :1987 年 10 月 :3 億 600 万円 従業員数

More information

XACCの概要

XACCの概要 2 global void kernel(int a[max], int llimit, int ulimit) {... } : int main(int argc, char *argv[]){ MPI_Int(&argc, &argc); MPI_Comm_rank(MPI_COMM_WORLD, &rank); MPI_Comm_size(MPI_COMM_WORLD, &size); dx

More information

Microsoft PowerPoint - suda.pptx

Microsoft PowerPoint - suda.pptx GPU の HWアーキテクチャと高性能化手法 須田礼仁 ( 東京大学 ) 2011/03/22 GPU 高性能プログラミング GPU のハードウェアを理解する CUDA のソフトウェアを理解する CUDA でプログラムを書くのは難しくないが, CUDA で高速なプログラムを書くのは難しい どうすれば遅くなるかを理解する! 効果が大きいものから順に説明します 1 高性能プログラミングの手順 1. 現在のコードの,

More information

( CUDA CUDA CUDA CUDA ( NVIDIA CUDA I

(    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

More information

openmp1_Yaguchi_version_170530

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

More information

Microsoft PowerPoint - sales2.ppt

Microsoft PowerPoint - sales2.ppt 最適化とは何? CPU アーキテクチャに沿った形で最適な性能を抽出できるようにする技法 ( 性能向上技法 ) コンパイラによるプログラム最適化 コンパイラメーカの技量 経験量に依存 最適化ツールによるプログラム最適化 KAP (Kuck & Associates, Inc. ) 人によるプログラム最適化 アーキテクチャのボトルネックを知ること 3 使用コンパイラによる性能の違い MFLOPS 90

More information

PowerPoint プレゼンテーション

PowerPoint プレゼンテーション OpenMP 並列解説 1 人が共同作業を行うわけ 田植えの例 重いものを持ち上げる 田おこし 代かき 苗の準備 植付 共同作業する理由 1. 短時間で作業を行うため 2. 一人ではできない作業を行うため 3. 得意分野が異なる人が協力し合うため ポイント 1. 全員が最大限働く 2. タイミングよく 3. 作業順序に注意 4. オーバーヘッドをなくす 2 倍率 効率 並列化率と並列加速率 並列化効率の関係

More information

PowerPoint プレゼンテーション

PowerPoint プレゼンテーション 1 サンプルソースコードは ITO の /home/tmp/gpu 以下に置いてあります 実質的に 演習の答え となるものもあるので注意 PGI compiler 19.4 で実行した場合の出力に準拠 18.10 でも 19.4 でも基本的には同じであるが 18.10 では出力されていた Accelerator kernel generated ( 並列実行カーネルが作成できた旨 ) が出力されなくなったことを反映

More information

Microsoft PowerPoint - GPUシンポジウム _d公開版.ppt [互換モード]

Microsoft PowerPoint - GPUシンポジウム _d公開版.ppt [互換モード] 200/0/9 数値流体解析の並列効率とその GPU による高速化の試み 清水建設 ( 株 ) 技術研究所 PHAM VAN PHUC ( ファムバンフック ) 流体計算時間短縮と GPU の活用の試み 現 CPUとの比較によりGPU 活用の可能性 現 CPU の最大利用 ノード内の最大計算資源の利用 すべてCPUコアの利用 適切なアルゴリズムの利用 CPU コア性能の何倍? GPU の利用の試み

More information

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 /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

More information

1 OpenCL OpenCL 1 OpenCL GPU ( ) 1 OpenCL Compute Units Elements OpenCL OpenCL SPMD (Single-Program, Multiple-Data) SPMD OpenCL work-item work-group N

1 OpenCL OpenCL 1 OpenCL GPU ( ) 1 OpenCL Compute Units Elements OpenCL OpenCL SPMD (Single-Program, Multiple-Data) SPMD OpenCL work-item work-group N GPU 1 1 2 1, 3 2, 3 (Graphics Unit: GPU) GPU GPU GPU Evaluation of GPU Computing Based on An Automatic Program Generation Technology Makoto Sugawara, 1 Katsuto Sato, 1 Kazuhiko Komatsu, 2 Hiroyuki Takizawa

More information

CPU Levels in the memory hierarchy Level 1 Level 2... Increasing distance from the CPU in access time Level n Size of the memory at each level 1: 2.2

CPU Levels in the memory hierarchy Level 1 Level 2... Increasing distance from the CPU in access time Level n Size of the memory at each level 1: 2.2 FFT 1 Fourier fast Fourier transform FFT FFT FFT 1 FFT FFT 2 Fourier 2.1 Fourier FFT Fourier discrete Fourier transform DFT DFT n 1 y k = j=0 x j ω jk n, 0 k n 1 (1) x j y k ω n = e 2πi/n i = 1 (1) n DFT

More information

XcalableMP入門

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)

More information

NUMAの構成

NUMAの構成 GPU のプログラム 天野 アクセラレータとは? 特定の性質のプログラムを高速化するプロセッサ 典型的なアクセラレータ GPU(Graphic Processing Unit) Xeon Phi FPGA(Field Programmable Gate Array) 最近出て来た Deep Learning 用ニューロチップなど Domain Specific Architecture 1GPGPU:General

More information

GPGPUクラスタの性能評価

GPGPUクラスタの性能評価 2008 年度理研 HPC シンポジウム第 3 世代 PC クラスタ GPGPU クラスタの性能評価 2009 年 3 月 12 日 富士通研究所成瀬彰 発表の概要 背景 GPGPU による高速化 CUDA の概要 GPU のメモリアクセス特性調査 姫野 BMT の高速化 GPGPU クラスタによる高速化 GPU Host 間のデータ転送 GPU-to-GPU の通信性能 GPGPU クラスタ上での姫野

More information

NUMAの構成

NUMAの構成 共有メモリを使ったデータ交換と同期 慶應義塾大学理工学部 天野英晴 hunga@am.ics.keio.ac.jp 同期の必要性 あるプロセッサが共有メモリに書いても 別のプロセッサにはそのことが分からない 同時に同じ共有変数に書き込みすると 結果がどうなるか分からない そもそも共有メモリって結構危険な代物 多くのプロセッサが並列に動くには何かの制御機構が要る 不可分命令 同期用メモリ バリア同期機構

More information

Microsoft PowerPoint - 01_Vengineer.ppt

Microsoft PowerPoint - 01_Vengineer.ppt Software Driven Verification テストプログラムは C 言語で! SystemVerilog DPI-C を使えば こんなに便利に! 2011 年 9 月 30 日 コントローラ開発本部コントローラプラットフォーム第五開発部 宮下晴信 この資料で使用するシステム名 製品名等は一般にメーカーや 団体の登録商標などになっているものもあります なお この資料の中では トレードマーク

More information

修士論文

修士論文 AVX を用いた倍々精度疎行列ベクトル積の高速化 菱沼利彰 1 藤井昭宏 1 田中輝雄 1 長谷川秀彦 2 1 工学院大学 2 筑波大学 1 目次 1. 研究背景 目的 2. 実装, 実験環境 3. 実験 - 倍々精度ベクトル演算 - 4. 実験 - 倍々精度疎行列ベクトル積 - 5. まとめ 多倍長精度計算フォーラム 2 目次 1. 研究背景 目的 2. 実装, 実験環境 3. 実験 - 倍々精度ベクトル演算

More information

Microsoft PowerPoint - KHPCSS pptx

Microsoft PowerPoint - KHPCSS pptx KOBE HPC サマースクール 2018( 初級 ) 9. 1 対 1 通信関数, 集団通信関数 2018/8/8 KOBE HPC サマースクール 2018 1 2018/8/8 KOBE HPC サマースクール 2018 2 MPI プログラム (M-2):1 対 1 通信関数 問題 1 から 100 までの整数の和を 2 並列で求めなさい. プログラムの方針 プロセス0: 1から50までの和を求める.

More information

Javaセキュアコーディングセミナー東京 第3回 入出力(File, Stream)と例外時の動作 演習解説

Javaセキュアコーディングセミナー東京 第3回 入出力(File, Stream)と例外時の動作 演習解説 Java セキュアコーディングセミナー東京第 3 回入出力と例外時の動作 演習解説 2012 年 11 月 11 日 ( 日 ) JPCERT コーディネーションセンター脆弱性解析チーム戸田洋三 1 Hands-on Exercises コンパイルエラーに対処しよう ファイルからのデータ入力を実装しよう 2 Hands-on Exercise(1) サンプルコードの コンパイルエラーに対処しよう 3

More information

Microsoft PowerPoint - 高速化WS富山.pptx

Microsoft PowerPoint - 高速化WS富山.pptx 京 における 高速化ワークショップ 性能分析 チューニングの手順について 登録施設利用促進機関 一般財団法人高度情報科学技術研究機構富山栄治 一般財団法人高度情報科学技術研究機構 2 性能分析 チューニング手順 どの程度の並列数が実現可能か把握する インバランスの懸念があるか把握する タイムステップループ I/O 処理など注目すべき箇所を把握する 並列数 並列化率などの目標を設定し チューニング時の指針とする

More information

Microsoft PowerPoint - GPGPU実践基礎工学(web).pptx

Microsoft PowerPoint - GPGPU実践基礎工学(web).pptx GPU のプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 GPU プログラミング環境 (CUDA) GPU プログラムの実行の流れ CUDA によるプログラムの記述 カーネル (GPU で処理する関数 ) の構造 記述方法とその理由 GPU 固有のパラメータの確認 405 GPU(Graphics Processing Unit) とは 画像処理専用のハードウェア 具体的には画像処理用のチップ

More information

<4D F736F F F696E74202D D F95C097F D834F E F93FC96E5284D F96E291E85F8DE391E52E >

<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

More information

熊本大学学術リポジトリ Kumamoto University Repositor Title GPGPU による高速演算について Author(s) 榎本, 昌一 Citation Issue date Type URL Presentation

熊本大学学術リポジトリ 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 による高速演算について 榎本昌一 東京大学大学院工学系研究科システム創成学専攻

More information

Microsoft PowerPoint - 演習1:並列化と評価.pptx

Microsoft PowerPoint - 演習1:並列化と評価.pptx 講義 2& 演習 1 プログラム並列化と性能評価 神戸大学大学院システム情報学研究科横川三津夫 yokokawa@port.kobe-u.ac.jp 2014/3/5 RIKEN AICS HPC Spring School 2014: プログラム並列化と性能評価 1 2014/3/5 RIKEN AICS HPC Spring School 2014: プログラム並列化と性能評価 2 2 次元温度分布の計算

More information

研究背景 大規模な演算を行うためには 分散メモリ型システムの利用が必須 Message Passing Interface MPI 並列プログラムの大半はMPIを利用 様々な実装 OpenMPI, MPICH, MVAPICH, MPI.NET プログラミングコストが高いため 生産性が悪い 新しい並

研究背景 大規模な演算を行うためには 分散メモリ型システムの利用が必須 Message Passing Interface MPI 並列プログラムの大半はMPIを利用 様々な実装 OpenMPI, MPICH, MVAPICH, MPI.NET プログラミングコストが高いため 生産性が悪い 新しい並 XcalableMPによる NAS Parallel Benchmarksの実装と評価 中尾 昌広 李 珍泌 朴 泰祐 佐藤 三久 筑波大学 計算科学研究センター 筑波大学大学院 システム情報工学研究科 研究背景 大規模な演算を行うためには 分散メモリ型システムの利用が必須 Message Passing Interface MPI 並列プログラムの大半はMPIを利用 様々な実装 OpenMPI,

More information

main.dvi

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 toyohiro@isc.kyutech.ac.jp 45 2 ( ) CPU ( ) ( ) () 2.1

More information

ex04_2012.ppt

ex04_2012.ppt 2012 年度計算機システム演習第 4 回 2012.05.07 第 2 回課題の補足 } TSUBAMEへのログイン } TSUBAMEは学内からのログインはパスワードで可能 } } } } しかし 演習室ではパスワードでログインできない設定 } 公開鍵認証でログイン 公開鍵, 秘密鍵の生成 } ターミナルを開く } $ ssh-keygen } Enter file in which to save

More information

! 行行 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

! 行行 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

More information

(Microsoft PowerPoint \215u\213`4\201i\221\272\210\344\201j.pptx)

(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 共有メモリマルチプロセッサシステムの普及 共有メモリマルチプロセッサシステムのための並列化指示文を共通化する必要性 各社で仕様が異なり 移植性がない そして いまやマルチコア プロセッサが主流となり

More information

Slide 1

Slide 1 GPUコンピューティング入門 2015.08.26 エヌビディア合同会社 CUDAエンジニア 村上真奈 エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90 分 ) CUDA 入門 (90 分 ) Agenda 2 エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90

More information

2012年度HPCサマーセミナー_多田野.pptx

2012年度HPCサマーセミナー_多田野.pptx ! CCS HPC! I " tadano@cs.tsukuba.ac.jp" " 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

More information

PGIコンパイラ導入手順

PGIコンパイラ導入手順 1 注意この資料は PGI compiler 18.10 が最新であるときに作成した資料を元にしています PGI compiler 19.4 がリリースされましたが インストール手順や利用手順は 18.10 と変わりません 資料中の 1810 を 194 に 18.10 を 19.4 に読み替えてください 2019 年 6 月版 2 大きく分けて以下の 3 つの方法が利用可能 1. 手元のウェブブラウザでダウンロードして

More information

02_C-C++_osx.indd

02_C-C++_osx.indd C/C++ OpenMP* / 2 C/C++ OpenMP* OpenMP* 9.0 1... 2 2... 3 3OpenMP*... 5 3.1... 5 3.2 OpenMP*... 6 3.3 OpenMP*... 8 4OpenMP*... 9 4.1... 9 4.2 OpenMP*... 9 4.3 OpenMP*... 10 4.4... 10 5OpenMP*... 11 5.1

More information

Intel® Compilers Professional Editions

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()

More information

演習1

演習1 神戸市立工業高等専門学校電気工学科 / 電子工学科専門科目 数値解析 2019.5.10 演習 1 山浦剛 (tyamaura@riken.jp) 講義資料ページ http://r-ccs-climate.riken.jp/members/yamaura/numerical_analysis.html Fortran とは? Fortran(= FORmula TRANslation ) は 1950

More information

MATLAB® における並列・分散コンピューティング ~ Parallel Computing Toolbox™ & MATLAB Distributed Computing Server™ ~

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

More information

第12回講義(2019年7月17日)

第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.

More information

Microsoft Word ●IntelクアッドコアCPUでのベンチマーク_吉岡_ _更新__ doc

Microsoft Word ●IntelクアッドコアCPUでのベンチマーク_吉岡_ _更新__ doc 2.3. アプリ性能 2.3.1. Intel クアッドコア CPU でのベンチマーク 東京海洋大学吉岡諭 1. はじめにこの数年でマルチコア CPU の普及が進んできた x86 系の CPU でも Intel と AD がデュアルコア クアッドコアの CPU を次々と市場に送り出していて それらが PC クラスタの CPU として採用され HPC に活用されている ここでは Intel クアッドコア

More information

スパコンに通じる並列プログラミングの基礎

スパコンに通じる並列プログラミングの基礎 2018.09.10 furihata@cmc.osaka-u.ac.jp ( ) 2018.09.10 1 / 59 furihata@cmc.osaka-u.ac.jp ( ) 2018.09.10 2 / 59 Windows, Mac Unix 0444-J furihata@cmc.osaka-u.ac.jp ( ) 2018.09.10 3 / 59 Part I Unix GUI CUI:

More information

スパコンに通じる並列プログラミングの基礎

スパコンに通じる並列プログラミングの基礎 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

More information

この時お使いの端末の.ssh ディレクトリ配下にある known_hosts ファイルから fx.cc.nagoya-u.ac.jp に関する行を削除して再度ログインを行って下さい

この時お使いの端末の.ssh ディレクトリ配下にある known_hosts ファイルから fx.cc.nagoya-u.ac.jp に関する行を削除して再度ログインを行って下さい 20150901 FX10 システムから FX100 システムへの変更点について 共通... 1 Fortran の変更点... 2 C/C++ の変更点... 4 C の変更点... 5 C++ の変更点... 7 共通 1. プログラミング支援ツールの更新 -FX システムについて旧バージョンのプログラミング支援ツールは利用できません 下記からダウンロードの上新規インストールが必要です https://fx.cc.nagoya-u.ac.jp/fsdtfx100/install/index.html

More information

Slide 1

Slide 1 OpenACC CUDA による GPU コンピューティング Akira Naruse, 19 th Jul. 2018 成瀬彰 (Naruse, Akira) 自己紹介 2013 年 ~: NVIDIA シニア デベローパーテクノロジー エンジニア 1996~2013 年 : 富士通研究所 研究員など 専門 興味 : 並列処理 性能最適化 スパコン HPC GPU コンピューティング DeepLearning

More information

AICS 村井均 RIKEN AICS HPC Summer School /6/2013 1

AICS 村井均 RIKEN AICS HPC Summer School /6/2013 1 AICS 村井均 RIKEN AICS HPC Summer School 2013 8/6/2013 1 背景 OpenMP とは OpenMP の基本 OpenMP プログラミングにおける注意点 やや高度な話題 2 共有メモリマルチプロセッサシステムの普及 共有メモリマルチプロセッサシステムのための並列化指示文を共通化する必要性 各社で仕様が異なり 移植性がない そして いまやマルチコア プロセッサが主流となり

More information

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)) 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]

More information

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 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

More information

OpenMPプログラミング

OpenMPプログラミング OpenMP 基礎 岩下武史 ( 学術情報メディアセンター ) 1 2013/9/13 並列処理とは 逐次処理 CPU1 並列処理 CPU1 CPU2 CPU3 CPU4 処理 1 処理 1 処理 2 処理 3 処理 4 処理 2 処理 3 処理 4 時間 2 2 種類の並列処理方法 プロセス並列 スレッド並列 並列プログラム 並列プログラム プロセス プロセス 0 プロセス 1 プロセス間通信 スレッド

More information

PowerPoint Presentation

PowerPoint Presentation ヘテロジニアスな環境におけるソフトウェア開発 Agenda 今日の概要 ヘテロジニアスな環境の登場 ホモジニアスからヘテロジニアスへ ヘテロジニアスなアーキテクチャ GPU CUDA OpenACC, XeonPhi 自分のプログラムを理解するために デバッガ 共通の操作体験 TotalView 続きはブースで より速く ホモジーニアスな並列 HPC 銀河生成 金融のリスク計算 車の衝突解析 製薬

More information

XACC講習会

XACC講習会 www.xcalablemp.org 1 4, int array[max]; #pragma xmp nodes p(*) #pragma xmp template t(0:max-1) #pragma xmp distribute t(block) onto p #pragma xmp align array[i] with t(i) int array[max]; main(int argc,

More information

hotspot の特定と最適化

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*

More information

1 4 1.1........................................... 4 1.2.................................. 4 1.3................................... 4 2 5 2.1 GPU.....

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...........................................

More information

<4D F736F F D20332E322E332E819C97AC91CC89F090CD82A982E78CA982E9466F E393082CC8D5C91A291CC90AB945C955D89BF5F8D8296D85F F8D F5F E646F63>

<4D F736F F D20332E322E332E819C97AC91CC89F090CD82A982E78CA982E9466F E393082CC8D5C91A291CC90AB945C955D89BF5F8D8296D85F F8D F5F E646F63> 3.2.3. 流体解析から見る Fortran90 の構造体性能評価 宇宙航空研究開発機構 高木亮治 1. はじめに Fortran90 では 構造体 動的配列 ポインターなど様々な便利な機能が追加され ユーザーがプログラムを作成する際に選択の幅が広がりより便利になった 一方で 実際のアプリケーションプログラムを開発する際には 解析対象となる物理現象を記述する数学モデルやそれらを解析するための計算手法が内包する階層構造を反映したプログラムを作成できるかどうかは一つの重要な観点であると考えられる

More information

デジタル表現論・第4回

デジタル表現論・第4回 デジタル表現論 第 4 回 劉雪峰 ( リュウシュウフォン ) 2016 年 5 月 2 日 劉 雪峰 ( リュウシュウフォン ) デジタル表現論 第 4 回 2016 年 5 月 2 日 1 / 14 本日の目標 Java プログラミングの基礎 出力の復習 メソッドの定義と使用 劉 雪峰 ( リュウシュウフォン ) デジタル表現論 第 4 回 2016 年 5 月 2 日 2 / 14 出力 Systemoutprint()

More information

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 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 (lijiang@sekine-lab.ei.tuat.ac.jp), (kazuki@sekine-lab.ei.tuat.ac.jp), (takahashi@sekine-lab.ei.tuat.ac.jp), (tamukoh@cc.tuat.ac.jp),

More information

演習2

演習2 神戸市立工業高等専門学校電気工学科 / 電子工学科専門科目 数値解析 2017.6.2 演習 2 山浦剛 (tyamaura@riken.jp) 講義資料ページ 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

More information

スパコンに通じる並列プログラミングの基礎

スパコンに通じる並列プログラミングの基礎 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

More information

a b GE(General Erectrics) 9 4 irr (JAPLA 2009/12) Example1 120 P = C r + C 2 (1 + r) C t 1 (1 + r) t 1 + C t + F (1 + r) t 10

a b GE(General Erectrics) 9 4 irr (JAPLA 2009/12) Example1 120 P = C r + C 2 (1 + r) C t 1 (1 + r) t 1 + C t + F (1 + r) t 10 1 SHIMURA Masato 2010 9 27 1 1 2 CF 6 3 10 *1 irr irr irr(inner rate of return)function is able to written only few lines,and it is very powerful and useful for simulate unprofitable business model. 1

More information

,4) 1 P% P%P=2.5 5%!%! (1) = (2) l l Figure 1 A compilation flow of the proposing sampling based architecture simulation

,4) 1 P% P%P=2.5 5%!%! (1) = (2) l l Figure 1 A compilation flow of the proposing sampling based architecture simulation 1 1 1 1 SPEC CPU 2000 EQUAKE 1.6 50 500 A Parallelizing Compiler Cooperative Multicore Architecture Simulator with Changeover Mechanism of Simulation Modes GAKUHO TAGUCHI 1 YOUICHI ABE 1 KEIJI KIMURA 1

More information

目次 1 はじめに 本文書の概要 PVF ソフトウェアと VISUAL STUDIO PVF ソフトウェアの種類 MICROSOFT VISUAL STUDIO の日本語化について VISUAL STUDIO

目次 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

More information

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, 奥田洋司 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 東京大学新領域創成科学研究科

More information

CCS HPCサマーセミナー 並列数値計算アルゴリズム

CCS HPCサマーセミナー 並列数値計算アルゴリズム 大規模系での高速フーリエ変換 2 高橋大介 daisuke@cs.tsukuba.ac.jp 筑波大学計算科学研究センター 2016/6/2 計算科学技術特論 B 1 講義内容 並列三次元 FFT における自動チューニング 二次元分割を用いた並列三次元 FFT アルゴリズム GPU クラスタにおける並列三次元 FFT 2016/6/2 計算科学技術特論 B 2 並列三次元 FFT における 自動チューニング

More information

GPU CUDA CUDA 2010/06/28 1

GPU CUDA CUDA 2010/06/28 1 GPU CUDA CUDA 2010/06/28 1 GPU NVIDIA Mark Harris, Optimizing Parallel Reduction in CUDA http://developer.download.nvidia.com/ compute/cuda/1_1/website/data- Parallel_Algorithms.html#reduction CUDA SDK

More information