1 広島大学 理学研究科 尾崎 裕介 石川 健一. 1. Graphic Processing Unit (GPU) とは？ 2. Nvidia CUDA programming model 3. GPU の高速化 4. QCD with CUDA 5. 結果 6. まとめ 2.

Presentation on theme: "1 広島大学 理学研究科 尾崎 裕介 石川 健一. 1. Graphic Processing Unit (GPU) とは？ 2. Nvidia CUDA programming model 3. GPU の高速化 4. QCD with CUDA 5. 結果 6. まとめ 2."— Presentation transcript:

1 1 広島大学 理学研究科 尾崎 裕介 石川 健一

2 1. Graphic Processing Unit (GPU) とは？ 2. Nvidia CUDA programming model 3. GPU の高速化 4. QCD with CUDA 5. 結果 6. まとめ 2

3 3 GPU を搭載した Graphic Card  価格： 5 ～ 10 万円  性能： 数百 GFLOPS ( 単精度 )  主に画像処理を行う PC パーツ  滑らかな描画  リアルタイム表示  100 ～ 200 基の processer による 超並列高速計算 O(a) 改良した Wilson-Dirac quark の solver を CUDA によって作成し、 GPU でどのくらい加速されたか見てみた Gyözö I. Egri, hep-lat/0611022 “Lattice QCD as a video game” 先行研究 ： → 単精度 本研究では倍精度 手軽に高性能

4 4 10 -15 10 -6 10 -12 CPU の倍精度 solver CPU の単精度 solver 10 -15 10 -6 10 -12 GPU の単精度 solver 単精度 solver を用いて倍精度 の結果を得る手法 反復改良法、 連立方程式 (Wilson-Dirac) Dx = b を倍精度で 解く 単精度で Dx=b を複数回解 く と倍精度の解が得られるように した方法 GPU ：単精度計算が非常に高 速 (300- 900GFlops) 単精度で解くところを GPU に 担当させると全体が スピード アップ！

5 L 次元ベクトルの和の計算例 (L=N×M) c = a + b //=== host から呼び出される GPU code ==== _global_ void vadd_kernel(float *a, float *b, float *c) { int idx = threadIdx.x+blockIdx.x*blockDim.x; c[idx] = a[idx] + b[idx]; } //==== host 側 code === void main() { …… // GPU 上にメモリ確保 cudaMalloc((void**)&a,….); ….. // c = a+b カーネルを GPU へ投げる // thread 数 /block ＝ N, block 数 =M で並列実行 vadd_kernel >>(a,b,c); } 高い並列度をうまく利用する必要がある 5  thread ： 最小の実行単位 (max 512/block)  thread block ： 同一の multiprocessor 上で 実行される thread の集まり (max 65535)  grid ： thread block の集まり 並列化されたカーネルの全体 thread 1 thread 2 thread 3 thread 4 ⋮ thread N block 1 block 2 block 3 block 4 ⋮ block M grid block

6 6 Nvidia CUDA Programming Guide より  できる限り並列化 → 1thread で 1 格子点の計算  できる限り高速なメモリアクセス → GPU 上の様々なメモリ領域の最適な使い方

7 7  Shared Memory  global Memory 高速 なメモリアクセス (4 clock cycles) read-write アクセス 同一 block 内の thread 間で共有 16KB/block device memory 上のメモリ 低速なメモリアクセス (400 ～ 600 clock cycles) read-write アクセス 全 thread 間で共有 Shared Memory の有効活用

8 8  1 格子点あたりのデータ量とロード回数 fermion ： 8 回 +(1 回 ) 3×4×2×4Byte=96Byte gauge link ： 2 回 3×(3-1)×2×4Byte×4=48Byte×4 SU(3) reconstruction method. clover 項 ： 1 回 21×2×2×4Byte=336Byte fermion を shared memory に乗せた 4×4×4×2×96Byte=12.3KB, (max 16KB/block) gauge link と clover は device memory からロード CUDA with QCD, programming strategy データの出入り： 1584 Byte 計算量： 1896 Flop Byte/Flop = 0.83 G80 バンド幅 : ～ 80GB/s 予想性能： 100 GFlops!! CUDA ブロックに 4 3 ×2 格子点をアサイン スレッド数 =128 スレッド

9 9 GPU ・・・ NVIDIA GeForce 8800 GTX CPU ・・・ Intel Core 2 @2.66GHz 354.6GFLOPS 21.3GFLOPS  O(a) 改良の Wilson-Dirac quark solver  Bi-CGStab 法  反復改良法  単精度部分を GPU が担当  even-odd preconditioning マシン構成 solver

10 10 GPU を用いた場合 さらに 1/7 に 格子サイズ 16 3 ×32 quench 0.15fm quark 質量 [MeV] 23 、 52 、 81 単精度 solver で 加速効果 10 -15 10 -12 10 -15 10 -6 10 -12 10 -15 10 -6 倍精度 単精度 23MeV 52MeV 81MeV GPU

11 11 CPU GPU quark 質量 23MeV 格子サイズ 4 3 ×8 8 3 ×16 16 3 ×32 最大性能 17GFLOPS 今回の結果 ただし、まだ速くなるはず → coalesced access

12 12 格子点 0 格子点 1 格子点 2 ⋮ 格子点 0 格子点 1 格子点 2 ⋮ 格子点 0 格子点 1 格子点 2 ⋮ 格子点 0 ⋮ thread 0 thread 1 thread 2 ⋮ thread 0 thread 1 thread 2 ⋮ 4,8,or 16Byte

13 13 Nvidia GeForce GTX 280 Core 2 Duo 3.0GHz (6MB) non coalesced access on shared memory 20GFLOPS 石川健一 solver coalesced access on texture cache 40 ～ 50GFLOPS hopping → 89GFLOPS clover → 100GFLOPS 倍精度 solver GPU solver 220 秒 ～ 10 秒 ×22

14 14  GPU を用いると気軽に高速計算が可能。 ← 格子 QCD でも  GPU は単精度計算が高速。  反復改良法を利用した GPU solver を作成した。 ← 倍精度の結果  作成した solver は O(a) の改良を行う clover 項を導入している。  GeForce 8800 GTX での結果  solver の計算性能は最大約 17GFOLPS 。  計算時間は Core 2 Duo 2.66GHz CPU の 1/7 。  GeForce GTX 280 での結果  coalesced access 導入後 40 ～ 50GFLOPS 。  Core 2 Duo 3.0GHz の 22 倍。  高速な計算には coalesced access が必要。  複数の GPU による計算。