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

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

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/ “Lattice QCD as a video game” 先行研究 ： → 単精度 本研究では倍精度 手軽に高性能

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

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

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  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 GPU ・・・ NVIDIA GeForce 8800 GTX CPU ・・・ Intel Core 354.6GFLOPS 21.3GFLOPS  O(a) 改良の Wilson-Dirac quark solver  Bi-CGStab 法  反復改良法  単精度部分を GPU が担当  even-odd preconditioning マシン構成 solver

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

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

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 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  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 による計算。