MC計算の高速化について

神前 純一 (KEK)

LHC が切り拓く新しい物理@ICEPP
2009年4月3日

「きっかけ」と概要- 2008年1月頃
D. Rainwater氏から「MC計算をGPUを使って行うことにより、高速化することができないか」とのアイディアが出された。
このアイディアの検証を
K. Hagiwara、J. Kanzaki、N. Okamura（KEK）
T. Steltzer、D. Rainwater
で行っている。

- NVIDIA製の GPU board とホスト用の PC を導入- CUDAを使用し software 開発を行った- QED、QCDの基本過程の amplitude、そして断面積の計算を行い、CPU上で同様の計算を行う場合と比較し、結果、および計算速度についての検討を行った

GPUについて• GPU (Graphics Processing Unit): PCやworkstationで画像出力を担う
近年、Game 等の 3D Graphics の出力などのために高い能力が要求されるようになった。

• 主要なメーカーはNVIDIAとAMD/ATIだが、NVIDIAから提供されている開発環境の「CUDA」を利用することにより C/C++ を用いた科学技術計算への応用が容易になった。

- 既に物理・科学以外の幅広い分野で応用されている
特に image 処理、graphics 出力、金融関連の simulation、、- 物理・科学でも流体力学、天体物理学、医療分野での画像処理
KEKでも最近、CUDAに関連する講習会が行われた(2009年2月）

!"#$%&'(GPUの構造（Processor）• 1 )*!"#$%&'(+,– 8 -*./'0$%&'( (.1234$%&'() 567• Multi Processor– 1 )*!"#$%&'(5 1 )89*./'0:%'; (<=) >?@– A30BCDEFGHIJ./'0.KLM32NOから構成される- 8個の thread processor– 8192 - (G80) / 16384 (GT200) * 32 P'1/L.Q– 16 KB *RSTU2 (16 VN;567)
1つの multiprocessor
上で1つ以上の-NVIDIA
GT200 (T10)!"#thread block が実行される- 同じ命令が実行される• .WX3Y*Z[&'1– \]^_,`aZ[– 32 P'1bc,defcghi (IEEE754 jklmn)– FADD, FMUL o*hi+ IEEE754 pq,FMA (rs) tu

8 ./'0$%&'(SPSPSPSPSPSPSPSPSFU SFUDP– vi,wxy+zc{|*}– ($%) 240 &!'()*+,-).
RS– Special Function
Unit (SFU) EF~•zc{|,€•‚cJƒ>„multiprocessor
あたり一つの倍精度演算器• 最新のGPUではTU2– ($%) 1.44 GHz /012
3 1 TFLOPS/GPU, IEEE754
456…†,‡ˆ‰iŠ‹†•Œ•Ž5•G(jkl)
1 )*•kldefcghi– 1GT200+,!"#$%&'(tˆuE789+,-).:;<=!>56?@AB ($%)
86.4 GFLOPS/GPUprocessor により一つのGPUを構成• 複数の –multi
Š>‘’G,FMA (rs) tu
12Multi Processor– CDEFGHIJF' 512-bit K 800 MHz DDR B 102.4 GB/s

30 x 8 → 240 thread processor (GT200)'()*OP7QMRS

L,MN8CDE

!"#$%&'
GPUの構造（Thread）• Thread: プログラムの実行の最小単位• !"#$()*+,threadの集合。一つのmulti• Thread Block:– !"#$– -.(/012!"#$

blockDim.x• 1 ;<=2 ;<=>9? 3 ;<block の集合。CPUから一回の関数呼• Grid: Thread• 34#5@9A/B 512 !"#$び出しで実行することが出来る。
• C;<(/BDEF? (512, 512, 64)• GHI(34#5JK>LMN8O(!"#$P

から、自分の• 実行される各 threadQR(S'TU4V#DWX-.YL>NZ

blockDim.y– !"#$34#5– !"#$6789:(processor上で実行される!"#$34#5
g_#$

gridDim.x

gridDim.ythread ID（threadblock 内の自分の位置）、およ• [\]^%_`6\](ab)• __syncthreads() XQcNM12
び block ID （自分の所属するthread block の• BdY? blockDim.{x, y, z} Xef– g_#$ – !"#$34#56789:(出来る。
• 1 ;<>9? 2 ;<• BdY? gridDim.{x, y, z} Xef• Warp: 「同時に」同じ命令を実行する• /B (65535, 65535, 1)

grid 内での位置）を知ることがthread からなる物理的な単位（= 32 threads)\]^%_

GHI(!"#$34#5?=
GHI(S'TU4V#DWX-.YLMZ
21

GPUの構造（メモリ）!"#"$%• CPU（ホスト）とGPU（デバイス）のメモリ空間は異なり、お互いのメモリを直接アクセスす• !"#&'– CPU (()*) + GPU ($,-)) .!"#&'/る事はできない。
0123()*45$,-)!"#6789:;)→ 専用の <=23>=/?.@6ABC+/DE>FGH
API によりデータ転送を行う– ()*+$,-).'.$IJKL/MN.
Memory及びregisters•
API 6OPQ.RABC+SDE>TH)W^b_d;^s

…
YZ!"#

• $,-).!"#UV• Registers:– WX)J
同一の thread block• Shared– Memory:
YZ!"#• [\]^_!"#
内の threadによって共有される。register• `a.)W^bcd^:efTg)W^b'Dと同程度に高速。16KB/multiYZhi>TH• ,\:j\k#:*S1lim3WX)J+`
processor。Onchip
noepqDTH• 16 KB / r%]_d;^s• Global Memory:– tdI,%!"#• Oﬀ chip• [k]^_!"# (DRAM)• TuQ.)W^b'DYZhi>TH• 全ての thread 間で共有)W^b_d;^s

YZ!"#YZ!"#YZ!"#tdI,%!"#$,-)MN API D$IJKL()*!-\!"#!-\!"#vw^xyvw^xy

CPUj9

CPUj9

CPU
j9

CPU
j924

我々のGPU環境• 2008年5月 NVIDIA製GeForce 9800GTXを導入（2万8000円）• 2009年12月
GeForce
GTX280（5万円）• Host PC
Intel Core2 Duo
3GHz + 4GB
memoryFedora8 64bit

GTX2809800GTX

Multi Processor

30

16SP240

128

Global Memory

1GB500MB

Constant Memory

64KB

64KB

Shared Memory/block

16KB

16KB

Registers/block

16384

8192

Warp Size3232

Clock Rate

1.30GHz

1.67GHz

倍精度サポート有り比較のためのプログラム• CUDA環境にHELAS（Helicity Amplitude 計算ライブラリ）を移植（HeGET）• 乱数をCPUで生成 → GPUへ転送• GPUでPhase Space の生成＋PDF weightの計算＋Amplitude の計→ 断面積の計算を行い結果をCPUへ転送する。この段階の計算時間を測定する• CPUで動作する同じ構造の program を作成し、対応する部分の計算時間を測定し、比較する。

• 断面積計算の結果の検証のため MadGraph/MadEvent による断面積の計算を行った。
また FORTRAN version の HELAS を用い BASES により断面積を計算させる program を作成した。

比較のためのプログラム• MadGraph により amplitude を計算する FORTRAN code を作成し、それを GPU/CPU で使える C/C++ code へ変換する（いずれも ﬁnal state particle の数が多い場合は MadGraph 自身のmodiﬁcation が必要）
ー→ n-photons• QED: uuー

QCD: gg → n-gluons, uu → n-gluons, uu → uu + n-gluonsの各プロセスを用いて• 全断面積が正しく計算されているか• 計算時間の速度の比較を検証した。

• 計算速度は thread に割り当てられた register の数、また block あたりの thread の数に大きく依存する。

• pp collision @ 14 TeV+CTEQ6L1• 単精度での結果

QED Process• u uー → n-photons• n が大きくなると diagram の数が膨大になる# photons# diagrams =(# photons)!

total sigma(Heget/MG-ME)

221.000 ± 0.002

361.000 ± 0.0054240.990 ± 0.005

51200.998 ± 0.009

6720(1.025 ± 0.018)

75040(0.999 ± 0.035)• その様な場合でも amplitude---

program を分割することにより、compile 及び実行が可能

840320n=6, 7 は BASES による積分計算との比較• 簡単なパターンの process な
ので HeGET（HELAS）call とfor loop を用いた、その“permutation” で書くことが出来る• n≥6 の場合に amplitude
program がそのままではcompile 出来なかった→ CUDA compiler の問題？

QED Process• u uー → 4 photons
pTmax( !/012'!)/0 )*+,%.

+)!+

3$4$5

$%#&

/0678$95$%#":0;$;!<(=>94$%#$%$(+)!,

12324

$%$'

*+567284$%$&

9+:2:!;0<=83$%$"+)

!-)

*)

+))

+*)

,))

,*)

-))

-*)

.))

$!!

.*)
*))
('!"#$%&( !3456!"!#

$

#

"!
!

"!
!

!)/0 )*+,%-

'.

*)!*

7$8$9

$%#&

3:;<=$59$%#">:?$?!@(A458$%#$%$(*)!+12324$%$'*56782,4$%$&95:2:!;0<+,3*)$%$"!,

)

*)

+)

,)

-).)/)

0)

1)

2)
*))
('!"#$%&$!!

pTmin!"!#$#計算速度の比較(QED)

Process Time / Event [µ sec]

104
10

Event Process Time on GTX280
u u ! photons

3Permutation
102

CPU10Permutation

110

MadGraph (divided)

GPU-1

MadGraph
10-2

23456
Number of Photons

Preliminary78計算速度の比較(QED)180

CPU / GPU(GTX280)

Ratio of Process Time

160
140
MadGraph120

MadGraph (divided)

100
80

Permutation

60
40

Preliminary

20
0

2

3456
Number of Photons

78

QCD Processes• QCD coupling のみ• pTj>20GeV, |ηj|<2.5, pTjj>20GeV• Qren = Qfac = 20GeV# final jetsg g → gluonsーu u → gluonsu u → u u + gluons#diagram#color#diagram#color#diagram#color2663222345241861084510

120

159

2476

4057245720

1890120786240

計算速度の比較(QCD)

Process Time / Event [µ sec]

103

Event Process Time on GTX280
QCD Processes
uu
gg

102

uu

CPU
10uu
gg

1

10GPU-1

10-2uuPreliminary
234
Number of Jets in Final State5

計算速度の比較(QCD)
180

Ratio of Process TimeCPU / GPU(GTX280)

uu

160
140

uu

120

gg

100
80
60
40
200

Preliminary
234
Number of Jets in Final State5

結果とまとめ• Graphic Processor (GPU) の上でかなり広範囲の program を実行することが出来る• 計算結果についても問題ない。
• 同等の program を CPU と GPU で実行させ、計算時間を比較した。我々の環境で GPU を用いることで10倍から100倍(以上)の計算速度が得られた。
✴ 大きな規模の program の compile に問題がある可能性✴ 良い performance を得るには code に工夫(know-how)が必要- メモリアクセス- Loop 構造- 条件分岐- 実行時パラメタｰ✴ 倍精度計算のサポートが不足。

→ 何でも簡単に高速になる「夢の装置」ではないけれど、 いろいろな、そして大きな応用可能性がある。