GPU を用いた lattice 計算

GPUを用いたlattice計算 広島大学　理学研究科　尾崎　裕介石川　健一

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

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

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

CUDA Programming Model 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<<<M,N>>>(a,b,c); } 高い並列度をうまく利用する必要がある block grid • thread ：最小の実行単位 • (max 512/block) • thread block ： • 同一のmultiprocessor上で • 実行されるthreadの集まり • (max 65535) • grid ：thread blockの集まり • 　　　　　並列化されたカーネルの全体

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

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

CUDA with QCD, programming strategy • 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 CUDA ブロックに 43×2 格子点をアサイン　　　　　　　　　　　　スレッド数=128 スレッドデータの出入り：1584 Byte 計算量：1896 Flop Byte/Flop = 0.83 G80バンド幅: ～80GB/s 予想性能：　100 GFlops!! fermion を shared memory に乗せた 4×4×4×2×96Byte=12.3KB, (max 16KB/block) gauge link と clover は device memory からロード

マシン構成及びsolverの詳細 • マシン構成 • solver 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

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

Performance (Volume) • quark 質量 23MeV • 格子サイズ • 43×8 • 83×16 • 163×32 GPU CPU 今回の結果最大性能 17GFLOPS ただし、まだ速くなるはず →　coalesced access

Coalesced Access 4,8,or 16Byte

最新の結果 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 ×22 ～10秒 220秒

まとめ • 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による計算。

GPU を用いた lattice 計算