Transcript GPU上でのLattice計算

広島大学 理学研究科 尾崎 裕介
石川 健一
1
1.
2.
3.
4.
5.
6.
Graphic Processing Unit (GPU)とは？
Nvidia CUDA programming model
GPUの高速化
QCD with CUDA
結果
まとめ
2
主に画像処理を行うPCパーツ
滑らかな描画
リアルタイム表示
100～200基のprocesserによる
超並列高速計算
価格： 5～10万円
性能： 数百GFLOPS (単精度)
GPUを搭載した
Graphic Card
手軽に高性能
O(a)改良したWilson-Dirac quarkの
solverをCUDAによって作成し、
GPUでどのくらい加速されたか見てみた
先行研究 ：
Gyözö I. Egri, hep-lat/0611022
“Lattice QCD as a video game”
→ 単精度
本研究では倍精度
3
単精度solverを用いて倍精度
の結果を得る手法
反復改良法、
連立方程式(Wilson-Dirac)
Dx = b を倍精度で解く
単精度で Dx=b を複数回解く
と倍精度の解が得られるようにし
た方法
GPU：単精度計算が非常に高速
(300-900GFlops)
単精度で解くところをGPUに担当
させると全体がスピードアップ！
CPUの倍精度solver
10-15
CPUの単精度solver
10-6
10-12
GPUの単精度solver
10-15
10-6
10-15
10-12
4
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);
}
高い並列度をうまく利用する必要がある
thread 1
block 1
thread 2
block 2
thread 3
block 3
thread 4
block 4
⋮
⋮
thread N
block M
block
grid
thread ： 最小の実行単位
(max 512/block)
thread block ：
同一のmultiprocessor上で
実行されるthreadの集まり
(max 65535)
grid ：thread blockの集まり
並列化されたカーネルの全体
5
Nvidia CUDA Programming Guide より
できる限り並列化
→ 1thread で 1 格子点の計算
できる限り高速なメモリアクセス
→ GPU上の様々なメモリ領域の最適な使い方
6
Shared Memory
•高速なメモリアクセス
(4 clock cycles)
•read-write アクセス
•同一block内のthread間で共有
•16KB/block
global Memory
•device memory 上のメモリ
•低速なメモリアクセス
(400～600 clock cycles)
•read-write アクセス
•全thread間で共有
Shared Memory の有効活用
7
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 スレッド
fermion を shared memory に乗せた
4×4×4×2×96Byte=12.3KB,
(max 16KB/block)
gauge link と clover は device memory からロード
データの出入り：1584 Byte
計算量：1896 Flop
Byte/Flop = 0.83
G80バンド幅: ～80GB/s
予想性能： 100 GFlops!!
8
•マシン構成
GPU・・・NVIDIA GeForce 8800 GTX 354.6GFLOPS
CPU・・・Intel Core 2 @2.66GHz
21.3GFLOPS
•solver
O(a)改良のWilson-Dirac quark solver
Bi-CGStab 法
反復改良法
単精度部分をGPUが担当
even-odd preconditioning
9
•格子サイズ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
10
•quark 質量 23MeV
•格子サイズ
43×8
83×16
163×32
GPU
CPU
今回の結果
最大性能
17GFLOPS
ただし、まだ速くなるはず
→ coalesced access
11
4,8,or 16Byte
格子点 0
thread 0
格子点 0
thread 0
格子点 1
thread 1
格子点 0
thread 1
格子点 2
thread 2
格子点 0
thread 2
⋮
⋮
格子点 1
⋮
格子点 0
格子点 1
格子点 1
格子点 1
格子点 2
格子点 2
⋮
格子点 2
格子点 0
格子点 2
⋮
⋮
12
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
220秒
×22
GPU solver
～10秒
13
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による計算。
14