Microsoft PowerPoint - GDEP-GPG_softek_May24-1.pptx

Size: px

Start display at page:

Download "Microsoft PowerPoint - GDEP-GPG_softek_May24-1.pptx"

つかさつちた
5 years ago
Views:

1 G-DEP 第 3 回セミナーツールで始める GPGPU なぜディレクティブベースのプログラミングが有望なのか? 失敗しない並列プログラミングの始め 2012 年 5 月加藤努株式会社ソフテックなぜ GPU プログラミング? GPU computing 時代の本流に! 今何が起きているのか? HPC アーキテクチャはどこへ向かうのか? 確実に移行中 Why? CPU + GPU ハイブリッド型へ現状の CPU を知るマルチコア CPU の限界性能劣化の原因 GPU を知る性能能力の違い GPU 活用のための現実的な勘所プログラミングを知る並列化が必須一般の並列化手法はディレクティブベース適用実現性の判断納得確信 OpenMP OpenACC 失敗しない並列プログラミングの始め = 背景をきちんと理解した上でうこと 1

2 現状の CPU を知る現在の CPU の限界 2 今のマルチコア CPU の性能現在の最先端プロセッサ (Intel / AMD) 高クロック化 SSE/AVXによるベクトル化マルチコア化 (2,4,6,8,12,16) 理論値年プロセッサ名周波数 Core 数 GFLOPS 電力 2000 Pentium 4 2.0GHz W 2011 Sandy Bridge 3.3GHz W HPC の用途では高機能演算機構複数のコアがあってもその能力を十分に活かしきれない 3

3 今のプロセッサ性能の現実現在の最先端プロセッサ (Intel Sandy Bridge) 姫野ベンチマーク ( 単精度 ) 性能 (GFLOPS) マルチスレッド並列 (OpenMP) With SSE/AVX (-fastsse mp) Without SSE (-O2 mp) シングルスレッド実行 1 ベクトル機構未使用だと並列効果あり With SSE/AVX ( fastsse) 5.96 Without SSE (-O2) スレッドでできるだけ速く実行 ( ベクトル機構 ) 宝の持ち腐れベクトル演算機構 (SSE~AVX) マルチコア並列性能の律速要因 < メモリ帯域使い切り性能飽和メモリ帯域 (DDR3) PGI 12.4 / Intel(R) Core(TM) i GHz 4 メモリ帯域が性能を支配姫野ベンチの単精度性能 vs. 倍精度性能 (GFLOPS) マルチスレッド並列 (OpenMP) 単精度 with SSE 倍精度 with SSE シングルスレッド実行 1 単精度 with SSE 5.96 倍精度 with SSE 3.21 倍精度になると性能が約半分に低下メモリ帯域を使い尽くしている ( 余裕がない ) ベクトル演算機構 (SSE~AVX) を使い切ればマルチコア並列効果が薄れる現マルチコア CPU の限界ベクトル演算機構を使わないとマルチコア並列効果が現れる 5

4 最新プロセッサ技術の理想と現実プロセッサメーカー曰く SSE 命令からAVX 命令へベクトル性能が2 倍に上がるよ! マルチコアを増やす並列効果が上がるよ! 本当かよ! 無理無理これらの innovation を否定する訳ではないが HPC 用途の一般ユーザコードにとってはあまり意味がない ( 高価なプロセッサを買っている!) メモリ帯域の制約多くのユーザにとっての懸念 6 ユーザの要求数値モデルの大型化高解像解析を要求ユーザプログラムレガシーなプログラム資産が存在 1 コア / シングルスレッドしか使っていない他の CPU コアは死んでいるマルチコア & マルチスレッド化たとえマルチコア上の並列化を行っても理想性能を享受できず! 計算が終わらない! 7

5 現在の HPC システムの状況プロセッサシングルスレッド性能の高速化を競ってきた歴史高性能 MPU 開発の限界プロセッサ & システム熱 = 電消費高性能 MPU とマルチコアだけでは今後の性能 Scalability を満たせない特に HPC 用途ではアーキテクチャを変革中現在の Programming Model だけでは対応できない 8 今後の HPC システム構成高い電力効率と高い絶対パフォーマンス本当に使えるのかな? CPU Accelerator 4~8cores 程度 Multi-core PCI Express Many-core Memory Device Memory NVIDIA GPU ハイブリッド型アーキテクチャ Intel MIC 9

6 今後の HPC システムの方向付け今後の H/W 方向性 ( ほとんどのベンダーの共通認識 ) Many-Cores CPU + Accelerator ハイブリッド NVIDIA GPU Intel MIC ユーザが直面するレガシーソフトウェアへの対応できるだけ簡単に! Re-compile and Run で性能が出せるツールは存在しない並列化を行いプログラムの再構築が必須並列性の抽出必須 10 GPU を知る CPU と GPU のハードウェア的能力の違いマルチコア CPU 並列挙動時のしがらみ NVIDIA GPU スレッドの自由奔放さ 11

7 CPU と GPU の能力の比較 ( ポンチ絵 ) C.I. = CPU 4~16 cores 21~ 50GB/s << 演算コアメモリ帯域 3.8~8 倍 GPU 500~1500 cores ~170GB/s 90W~135W 消費電力 200W 前後 12 マルチコア並列での性能劣化要因 CPU GPU 演算コアスレッドが並列に実行共有キャッシュメモリ Coherency 隠蔽技術 21~ 50GB/s メモリアクセスの遅さ (latency=200~800cycle) ~170GB/s 大量のレジスタ群 & cache 13

マルチコア並列での性能劣化要因 14 GPU CPU 21~ 50GB/s メモリアクセスの遅さ (latency=200~800cycle) メモリ演算コア自由に実行できず共有キャッシュキャッシュ内容の同期問題 Coherency しがらみ大量のレジスタ群 & cache ~170GB/s 隠蔽技術 16core 限界だよ!

8 マルチコア並列での性能劣化要因 14 GPU CPU 21~ 50GB/s メモリアクセスの遅さ (latency=200~800cycle) メモリ演算コア自由に実行できず共有キャッシュキャッシュ内容の同期問題 Coherency しがらみ大量のレジスタ群 & cache ~170GB/s 隠蔽技術 16core 限界だよ! キャッシュコヒーレンシの維持 False sharing NVIDIA GPU スレッドの自由奔放さ 15 GPU CPU 21~ 50GB/s メモリ演算コア共有キャッシュ Coherency 基本は自由に動くしがらみ大量のレジスタ群 & cache ~170GB/s 遅延隠蔽技術依存性無しの処理大量のマルチスレッド ready-to-run ループ内演算仕事の固まりメモリ遅延時間スレッド

9 行列積計算の性能 (CPU vs. GPU) 性能の違いはどの程度? 単精度倍精度使用プログラミング手法 CPU (GFLOPS) Size : 4096 x 4096 GPU (GFLOPS) 性能比 PGI ACC. Directives PGI CUDA Fortran PGI Fortran + CUBLAS PGI ACC. Directives PGI CUDA Fortran PGI Fortran + CUBLAS CPU : Intel(R) Core i GHz / 4cores 4 スレッド並列性能 GPU : NVIDIA(R) GeForce GTX 580@1544 MHz / 512cores PGI 12.4 Compiler 16 この価値を認めるかどうか?GPU の性能加速性を得る条件この逆は性能が出ない条件 17

10 NVIDIA GPU 上の並列性能に係わる要素並列化可能演算であること Data Parallel 型 GPU 並列化の対象はループ内並列時のデータ依存がないこと多数のスレッドを発生させること 1 次元より 2 次元 2 次元より 3 次元ループ長が長いことメモリ負荷が小性能加速性は大きい計算密度が大きいメモリデータの再利用 (locality) 18 多数のマルチスレッドを発生させること生成するスレッド数に関係する do i = 1, n r(i) = a(i) * 2.0 1スレッド計算の固まり enddo ( ロジック ) データ依存性無しの保証 N 個 GPU 大量のマルチスレッド ready-to-run do i = 2, n-1 do j = 2, m-1 enddo enddo N x M 個 a(i,j) = w0 * b(i,j) + & w1 * (b(i-1,j) + b(i,j-1) + b(i+1,j) + b(i,j+1)) + & w2 * (b(i-1,j-1) + b(i-1,j+1) + b(i+1,j-1) + b(i+1,j+1)) 1 スレッド計算の固まり Device Memory 大量のレジスタ群と cache Important 小さなループカウントでは性能効果は望めない 19

11 GPU/CPU の演算時間比較 do i = 1, n r(i) = sin(a(i)) ** 2 + cos(a(i)) ** 2 enddo 生成スレッド数少ない do i = 2, n-1 do j = 2, m-1 a(i,j) = w0 * b(i,j) + & w1 * (b(i-1,j) + b(i,j-1) + b(i+1,j) + b(i,j+1)) + & w2 * (b(i-1,j-1) + b(i-1,j+1) + b(i+1,j-1) + b(i+1,j+1)) enddo enddo ループ長マイクロ秒 N GPU CPU N / M GPU CPU Core i7 GTX580 Important GPU 上での処理には CPU~GPU 間のデータ転送時間も含まれる (Overhead) 20 GPU 上での性能加速性 GPU 500~1500 cores 自由に動けるスレッド ( コア ) が十二分にある演算コアメモリ ~160GB/s Latency 200~600cycle メモリアクセスの負荷が小さいと 8 コア程度のマルチコアのレベルとは違う一方メモリバウンドな特性の場合は Important CPU のメモリ帯域対 GPU メモリ帯域比で性能加速性が支配される性能の勘所 3.8~8 倍 21

12 演算密度 (Computational Intensity) ループ内の演算数とメモリのロードストア数との比率を表し演算とメモリ参照のバランスを見るための指標 C.I. = 浮動小数点演算数メモリアクセス数 do i = 1, n a(i) = b(i) + c(i) end do メモリアクセス 3 回演算 1 回 C.I. = C.I. が小さい程メモリバウンドな特性を帯びる実効メモリ帯域が性能を支配する演算器が遊ぶ 22 姫野ベンチマークの場合 Computation Intensity for GPU ( 292) do loop=1,nn ( 293) gosa= 0.0 ( 294) do k=2,kmax-1 ( 295) do j=2,jmax-1 ( 296) do i=2,imax-1 ( 297) s0=a(i,j,k,1)*p(i+1,j,k) & ( 298) +a(i,j,k,2)*p(i,j+1,k) & ( 299) +a(i,j,k,3)*p(i,j,k+1) & ( 300) +b(i,j,k,1)*(p(i+1,j+1,k)-p(i+1,j-1,k) & ( 301) -p(i-1,j+1,k)+p(i-1,j-1,k)) & ( 302) +b(i,j,k,2)*(p(i,j+1,k+1)-p(i,j-1,k+1) & ( 303) -p(i,j+1,k-1)+p(i,j-1,k-1)) & ( 304) +b(i,j,k,3)*(p(i+1,j,k+1)-p(i-1,j,k+1) & ( 305) -p(i+1,j,k-1)+p(i-1,j,k-1)) & ( 306) +c(i,j,k,1)*p(i-1,j,k) & ( 307) +c(i,j,k,2)*p(i,j-1,k) & ( 308) +c(i,j,k,3)*p(i,j,k-1)+wrk1(i,j,k) ( 309) ss=(s0*a(i,j,k,4)-p(i,j,k))*bnd(i,j,k) ( 310) GOSA=GOSA+SS*SS ( 311) wrk2(i,j,k)=p(i,j,k)+omega *SS ( 312) enddo ( 313) enddo ( 314) enddo ~/Himeno> pgf90 -fast -Minfo=intensity himenobmtxp.f90 jacobi: 292, Intensity = [symbolic], and not printable, try the -Mpfi -Mpfo options 294, Intensity = [symbolic], and not printable, try the -Mpfi -Mpfo options 295, Intensity = [symbolic], and not printable, try the -Mpfi -Mpfo options 296, Intensity = 1.06 ~/Himeno> pgf90 -fast -Minfo=intensity himenobmtxp.f90 -Mpfo jacobi: 292, Intensity = , Intensity = , Intensity = , Intensity = 1.06 C.I. > 2.0 大きい程良い C.I. < 1.0??? Important GPU の場合は C.I. が大きい程性能加速率は高くなる 23

13 今までのまとめと今後のユーザの選択肢 HPC H/W アーキテクチャ変革中 1 マルチコアCPUの性能限界 HPC アプリケーション用 CPU Accelerator Multi-core Many-core 4~8cores 程度 PCIe Memory Device Memory HPC 業界への最大の貢献 2 NVIDIA CUDA 環境実効性 ( 性能 ) を実証 Throughput computing 現在のマルチコア CPU の延長シングルスレッド性能に頼る Many-Cores CPU + Accelerator ハイブリッドどうなると思う? プログラムの並列化が必須 24 レガシーなアプリケーションプログラムマルチスレッド並列化に投資していないプログラムシングルスレッド or 並列スケールしないプログラム現在のマルチコアの能力すら享受していない H/W の進化 = Many-cores 今後のメニーコアの能力も使えない Important 何もしなければ性能的に取り残されてゆく恐れ生産性の停滞研究の遅延競争力の低下 25

14 プログラミング方法を知る並列プログラミングの話 26 今後のプログラミングの方向性どうであれプログラムの並列化がどうしても必要 Re-compile and Run で性能が出せる generic なツールは存在しない明示的に並列化を行いプログラムの再構築が必須データパラレル ( ループレベル並列性 ) OpenMP/OpenACC, CUDA タスクパラレル ( 大きな仕事の固まりを分配実行 ) MPI 27

今まで行ってきた並列化メソッド過去のプログラム資産 (Legacy Codes) FORTRAN77 Fortran90/95/03 C/C++ Directives 挿入 Single thread 性能 Auto-Vector Multi-Core

Many-cores NVIDIA, Intel MIC 28 ( 復習 ) タスク並列とデータ並列 Program タスク並列データ並列 ( ループレベル並列 ) Single Program Multi Data ( 同じプログラムを異なるデータで )

方向とも 2 分割 do j = 2,m-1 do i = 2,n-1 resid = (ax*(uold(i-1,j) + uold(i+1,j)) & & enddo + ay*(uold(i,j-1) + uold(i,j+1)) + b *

15 今まで行ってきた並列化メソッド過去のプログラム資産 (Legacy Codes) FORTRAN77 Fortran90/95/03 C/C++ Directives 挿入 Single thread 性能 Auto-Vector Multi-Core 性能 OpenMP CPU の高速化 SSE/AVX マルチコア化並列化必須 Multi-Node 性能 MPI 今までの延長線上で可能か? Many-cores NVIDIA, Intel MIC 28 ( 復習 ) タスク並列とデータ並列 Program タスク並列データ並列 ( ループレベル並列 ) Single Program Multi Data ( 同じプログラムを異なるデータで ) 複数の解析領域 ( タスク ) に分割 do j=1,m do i=1,n uold(i,j) = u(i,j) enddo enddo GPGPU ループ分割並列処理 k j k j i 1 process i 4 process 全体領域 i,j 方向とも 2 分割 do j = 2,m-1 do i = 2,n-1 resid = (ax*(uold(i-1,j) + uold(i+1,j)) & & enddo + 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 k j k j Important ループ分割並列処理 i 6 process i 9 process i, 方向 2 分割 j 方向 3 分割 i, 方向 3 分割 j 方向 3 分割ループ単位が並列対象 29

16 並列化を行うまでの手間 -MPI - FORTRAN77(Legacy) Include 文で変数配列の取り込み配列サイズの指定 COMMON 文共通配列の指定 ---- param.inc file parameter n=1000,m=500 common /array/ a(n,m),b(n,m),c(n,m) program main include param.inc.. do j=,1,n do i=1,n (... ) end do dnd do 1 Fortran90/95 1. 配列サイズを自由に変更できるように Dynamic Allocation 型の配列へ変更 2. Moduleを活用し COMMON 化 2 1. 解析空間を領域分割するロジック追加 2. 各領域用の配列演算を相対アドレスで実行可能になるように変更 3. DOループインデックスを相対アドレス添字に変更 4. 並列実行時にデータ依存性が存在する部分を回避する 5. 袖との通信のためにMPIルーチン追加分割対象の配列を使用するルーチンの全てを変更する必要がある段階的なポーティングができず 30 並列化を行うまでの手間 OpenMP( ディレクティブ型 ) FORTRAN77(Legacy) Include 文で変数配列の取り込み配列サイズの指定 COMMON 文共通配列の指定 ---- param.inc file parameter n=1000,m=500 common /array/ a(n,m),b(n,m),c(n,m) program main include param.inc #pragma omp parallel for do j=,1,n do i=1,n (... ) end do dnd do 1 1 Fortran90/95 1. 配列サイズを自由に変更できるように Dynamic Allocation 型の配列へ変更 2. Moduleを活用し COMMON 化 2 1. スレッド並列の対象となるループを探すホットスポットを調べる ( プロファイル ) 2. Do ループ内に並列計算上の依存性がないかを調べる 3. 並列実行したいループに並列化ヒント ( ディレクティブ ) を挿入する 4. テスト & 修正 & 並列化範囲を広げるソースに大きな変更がほとんどない段階的なポーティングが可能 31

17 パラレルデータパラレル現在主流の並列プログラミング方法 MPI プロセス並列ソースを人が修正 ( 領域分割型へ ) 明示的プロセス間通信関数の追加関連する全ルーチン修正必要タOpenMP OpenACC スレッド並列ディレクティブ挿入コンパイラが並列コード生成段階的に部分ポーティング CUDA OpenCL 性能最適言語 FFT/BLAS 並列ライブラリ置換スクGPU 並列プログラミンSource Portability 要 kernel コード作成 +API 追加 x86 系 CPU では動作しない (PGI では可能 ) それなりの作業工数と知識が必要グGPGPU Computing 利用層の拡大ライブラリルーチンとの I/F を記述汎用性あるライブラリが必要部分的なソース修正 32 生産性が高い開発環境が求められる! Many Legacy Codes CAE Research/Expert スパコンセンター大学 General Scientists/Engineers life ( 大学企業 ) Oil 気象物理化学 CUDA C/Open CL CUDA Fortran OpenACC easy productivity portability CFD 時間 2007 年 2009 年 2011 年 2012 年現在 33

18 Scientist, Engineer にとっての並列化手法計算機科学の専門家ではない既存のレガシーコードが多い! CUDA C/Fortran は開発コストが掛かる! 1. 手間暇をあまり掛けずにそれなりの性能を! 2. 新しい H/W 技術革新があっても以下のことが大事プログラムソース資産の可搬性 ( ポータビリティ ) 性能の可搬性この 20 年間で最も成功したプログラミングモデル MPI 明示的プロセス並列 OpenMP ディレクティブベース ( 生産性 ) ( ポータブル性 ) ( 生産性 ) ( ポータブル性 ) 34 GPU Accelerator 用にもディレクティブベース CUDA GPGPU 対応 PGI アクセラレータコンパイラ製品 CAPS HMPP Workbench 2010 年リリース 2 年以上の実装実績 PGI Accelerator Programming Model (directiveベース) 1. 手間暇をあまり掛けずにそれなりの性能を! CPU/GPU メモリ帯域性能比の 3~8 倍程度が妥当 2. プログラムソース資産の可搬性ディレクティブはコメント行扱いソースは変化せず OpenACC 標準化で性能の可搬性あり 3. 段階的ポーティング可能 GPU プログラムをポータブルかつ高生産性に! OpenACC Standard 35

Accelerator Directives をベースとして策定された OpenACC コンパイラ PGI ベータリリース中 36

19 OpenACC Standard とは 2012/3/27 ディレクティブベース ( ソース上に指示行を挿入する形態 ) (2011 年 11 月 ) GPU 用 Programming for Fortran/C/C++ PGI Accelerator Directives をベースとして策定された OpenACC コンパイラ PGI ベータリリース中 36 OpenMP の取組 2012/3/27 アナウンス OpenACC 知見妥当性確認反映取り込み OpenMP 4.0 新バージョン (2012 年 ) 37

$acc kernels loop gang, vector(128) do j = 1,1000 do i = 1,1000 c(i,j) = a(i,j) + b(i,j) end do end do!

20 GPU 上の本来のプログラミングホスト側ハイブリッド構成 (CPU+GPU) GPU 側 CPU Main Memory 重い計算部分の処理をオフロード使用データを送る結果データを戻す GPU Device Memory Host_A(100) Device_A(100) Host GPU 間のメモリデータの転送が伴うデータ転送のオーバーヘッド時間が伴う 38 directive OpenACC コンパイラは何を行うのか? (1) a 配列 b 配列を GPU へコピー!$acc kernels loop gang, vector(128) do j = 1,1000 do i = 1,1000 c(i,j) = a(i,j) + b(i,j) end do end do!$acc end kernels (2) 処理を分割し GPU の並列実行単位にマッピング j=1 ベクトル処理 j=2 j=3 ベクトル処理ベクトル処理 j=4 ベクトル処理スレッドブロックに分割 Runtime 監視 j=1000 ベクトル処理 j index 毎に 1x1000 のグリッドで構成 CUDA cores GPU (3) c 配列を Host へコピー 39

21 二種類の並列機構 NVIDIA GPU の並列処理機構 A1 + B1 = C1 A2 + B2 = C2 A3 + B3 = C3 A4 + B4 = C4 NVIDIA Warp ( ベクトル長 ) A1 B1 A2 B2 + = A3 B3 A4 B4 C1 C2 C3 C4 ベクトル長 SIMD 幅とも言う ( この例は 4 ) 各ベクトル要素は同じ処理を行う =32 固定キューイングされ独に実されるベクトル処理を行うスレッドブロック do j = 1,1000 do i = 1,1000 c(i,j) = a(i,j) + b(i,j) end do end do j=1 ベクトル処理 j=2 j=3 j=4 ベクトル処理ベクトル処理ベクトル処理 j index 毎に独立にスレッドで並列処理 1x1000 のグリッドで構成 j=1000 ベクトル処理 CUDA cores 40 プログラムの並列化を行うとは? 41

22 ループレベルで並列化すると言うこと (1) Q. 並列化可能とは? 計算の順番に結果が依存しないこと i 添字に関する順序で並列化並列化 for( i = 0; i < n; i++ ) { A[i] = A[i] + B[i]; } ループ内配列に依存性がなければ for( i = 0; i < n; i++ ) { A[i] = A[i-1] + B[i]; } 回帰参照という同じ配列で定義 ~ 参照関係があるとき依存性の検討要 MPI 並列化 OpenMP 並列化ベクトル化が可能であると言うことになる 42 ループレベルで並列化すると言うこと (2) Q. 並列化を行うとは? ループ内の依存性を排除するように変更 Q. ディレクティブベースで並列化を行うとは? ループの前にヒント ( ディレクティブ ) を挿入する OpenMP #pragma omp parallel for for( i = 0; i < n; i++ ) { A[i] = A[i] + B[i]; } OpenACC #pragma acc kernels loop for( i = 0; i < n; i++ ) { A[i] = A[i] + B[i]; } コンパイラが並列コードを生成するマルチコア用並列コード Accelerator(GPU) 用並列コード 43

23 ディレクティブベースでないと辛い場面!$acc kernels loop gang, vector(128) do k = 2,nz-1 km = k-1 k0 = k kp = k+1 do j=2,ny-1 jp=j+1 j0=j jm=j-1 do i=2,nx-1 ip=i+1 i0=i im=i-1 gxxc = lgxx(i0,j0,k0) gxyc = lgxy(i0,j0,k0) gxzc = lgxz(i0,j0,k0) gyyc = lgyy(i0,j0,k0) gyzc = lgyz(i0,j0,k0) gzzc = lgzz(i0,j0,k0) & & & & & & kxxc = 1.5D0*ADM_kxx_stag_p(i,j,k) -0.5D0*ADM_kxx_stag_p_p(i,j,k) kxyc = 1.5D0*ADM_kxy_stag_p(i,j,k) -0.5D0*ADM_kxy_stag_p_p(i,j,k) kxzc = 1.5D0*ADM_kxz_stag_p(i,j,k) -0.5D0*ADM_kxz_stag_p_p(i,j,k) kyyc = 1.5D0*ADM_kyy_stag_p(i,j,k) -0.5D0*ADM_kyy_stag_p_p(i,j,k) kyzc = 1.5D0*ADM_kyz_stag_p(i,j,k) -0.5D0*ADM_kyz_stag_p_p(i,j,k) kzzc = 1.5D0*ADM_kzz_stag_p(i,j,k) -0.5D0*ADM_kzz_stag_p_p(i,j,k) 例えば 600ステップ位のdoループを CUDA Fortran/C 言語で書き換えることは大変な労力が必要 & detg_tempxx = gyyc*gzzc-gyzc*gyzc detg_tempxy = gxzc*gyzc-gxyc*gzzc detg_tempxz = -gxzc*gyyc+gxyc*gyzc detg_tempyy = gxxc*gzzc-gxzc*gxzc detg_tempyz = gxyc*gxzc-gxxc*gyzc detg_tempzz = gxxc*gyyc-gxyc*gxyc detg_detcg = (detg_tempxx*gxxc+ detg_tempxy*gxyc+detg_tempxz*gxzc) detg_detg = detg_psi4**3*detg_detcg uppermet_fdet = 1d0/(detg_psi4*detg_detcg) uppermet_uxx = detg_tempxx*uppermet_fdet uppermet_uxy = detg_tempxy*uppermet_fdet uppermet_uxz = detg_tempxz*uppermet_fdet uppermet_uyy = detg_tempyy*uppermet_fdet uppermet_uyz = detg_tempyz*uppermet_fdet uppermet_uzz = detg_tempzz*uppermet_fdet fac = idx2*dxdg_psi4 delgb111 = i2dx*(lgxx(ip,j0,k0)-lgxx(im,j0,k0)) さらに 500 行位続く ( 省略 ) dda_dzzda = idz2*(lalp(i0,j0,kp)-2d0*lalp(i0,j0,k0)+ & lalp(i0,j0,km)) cdcda_cdzzda = (dda_dzzda-gamma133*da_dxda-gamma233* & da_dyda-gamma333*da_dzda) dkdt_dkzzdt = lalp(i0,j0,k0)*(ricci_r33-2*kk33+kzzc & *trk_trk)-cdcda_cdzzda ADM_kzz_stag(i,j,k) = ADM_kzz_stag_p(i,j,k)+ & dkdt_dkzzdt*dt end do end do end do 44 5x in 5 Hours: Porting a 3D Elastic Wave Simulator to GPUs Using PGI Accelerator SEISMIC_CPML by Mathew Colgrove, PGI Applications Engineer Version MPI Processes OpenMP Threads GPUs Execution Time (sec) Programming Time (min) Original Host ACC Step x1.00 ACC Step ACC Step ACC Step Problem Size: 101x641x128 System Information: 4 Core Intel Core-i7 920 Running at 2.67Ghz with 2 Tesla C2070 GPU Compiler: PGI 2012 version 12.3 x

マルチコア CPU 今日のまとめ確実に移行アーキテクチャの変化 CPU + Accelerator ハイブリッドアーキテクチャ Single-thread 歴史が証明否が応でも並列化必須プログラミングモデル ( 手法 ) とは多くのユーザが使えるものでなければならない Multi-threads アセンブリ専用言語現実解ディレクティブベースによる並列化 1.

24 マルチコア CPU 今日のまとめ確実に移行アーキテクチャの変化 CPU + Accelerator ハイブリッドアーキテクチャ Single-thread 歴史が証明否が応でも並列化必須プログラミングモデル ( 手法 ) とは多くのユーザが使えるものでなければならない Multi-threads アセンブリ専用言語現実解ディレクティブベースによる並列化 1. 高級言語かつコードの可搬性が良いこと様々な platform/device に対しても同じソースで対応可能 2. 生産 ( 容易 ) 性に優れていること基本的に directive(hint) を与えるだけ Step by Step 3. 現実的かつ妥当な性能向上が得られること CPU/GPU メモリ帯域性能比の 3~8 倍程度 OpenMP / OpenACC 46 そして No 今後のIntel / AMD CPUの行方これ以上に速くなるのか? 性能はメモリ帯域に依存 YES 本当に GPGPU / Intel MIC の many-core solution に向かうのか? 何もしなければ YES 並列化プログラムへ移行する価値はあるのか? YES 自分のプログラムは高速化から取り残されていく並列化プログラミングが必須の時代に突入! 47

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 クアッドコア