Microsoft PowerPoint - 高速化WS_ver1.1.1

(1)

非静力学海洋モデルkinaco

のGPUによる高速化

山岸孝輝

1 _{, 松村義正}

2

1

_{高度情報科学技術研究機構}

2

_{東京大学大気海洋研究所}

平成28年度高速化ワークショップ

～「京」を中核とするHPCI、メニーコアを見据えて～

平成29年3月24日秋葉原UDXカンファレンス

(2)

発表の概要

GPUの基本

ハードの特徴

実行モデル・プログラミングモデル

性能を引き出すための基本

海洋モデルのGPUによる高速化事例

まとめ

2

（※本発表ではNVIDIA GPUを前提）

(3)

(4)

GPUの構造

単純な構造でリソースの大半は演算多数のコア

NVIDIA Kepler GPUのブロックダイアグラム

(5)

ハード性能（演算）

低クロックのコアを多数

バルクで高い演算性能、高並列計算

コア内パイプライン化あり

複雑な機構は無し

インオーダで処理

無し：プリフェッチ、分岐予測etc.

高い演算性能、高並列計算

単純な命令スケジューリング

(6)

ハード性能（メモリ･キャッシュ）

高いメモリバンド幅、高レイテンシ

GDDR5の性質

キャッシュ：CPUより弱い

（階層による）

演算にリソースを多く割り当てた結果

高いメモリバンド幅、高レイテンシ

弱いキャッシュ

6

(7)

CPU/GPU 実行モデル

各々のハードウェアの特徴によるもの

64 65 66 67 68 94 95 0 1 2 3 4 33 1 33 0 … … スレッド0 64 65 66 67 68 _… 94 95

CPU

GPU

1スレッドが連続データを処理

CPUコア数＝スレッド数

基本は1スレッドが1データを処理

GPUコア数<<スレッド数

逐次処理

並列処理

(8)

カーネルのCUDA化

__global__ void func(){

int i = threadidx.x; JST[i] = UTC[i] + 9;

GPU

void func(){ for(i=0;i<N;i++){ JST[i] = UTC[i] + 9; }

CPU

1スレッドが連続データを処理 1スレッドが1データを処理

関数利用時にス

レッドの総数・

形状を指定

組み込み命令でスレッド管理

書くだけなら結構簡単だが、スレッドとデータの対応付けは

完全に書き手に任されている

8

逐次処理

並列処理

ハードウェア・実行モデル・プログラミングモデルの理解の

もとでスレッドとデータの対応をつける必要あり

CUDA: NVIDIA GPU用の原語拡張

(9)

CPU 依存性のある命令

do i = 1, N

tmp = data( i )

out(i) = tmp + 2.0

end do

プリフェッチ、キャッシュ、ソフトウェアパ

イプラインなど様々な手段で対応

実行待ちデータ読み込み待ちが発生 1スレッドが逐次で処理時間

(10)

GPU スレッドの切り替え

0

1

2

3

0

1

2

3 ストールしたスレッド

→待機していた他のス

レッドに入れ替え

他のスレッドにもレジ

スタを割り当てて、準

備をさせておく

→外部メモリなどへの

待避無しにすぐに切

り替え可能

※本当は32スレッドごとに入れ替わる（ここでは省略）コアが常に稼働する状態 10 時間：実行：待ち(ストール) ：待機(active)

(11)

スレッドの切り替え（続き）

効率的にレイテンシを隠蔽するには

スレッドが多い方がよい

スレッド間の同期はなるべくとりたくない

スレッドを「自由に」

レイテンシを隠蔽出来れば、GPUの利点

「高い演算性能」と「広いメモリバンド幅」を

活用出来る

(12)

スレッドの階層とメモリの階層

各階層でどこまでデータを共有できるかを把握する

※レジスタとL1/シェアードメモリは複数のスレッド（最大で2048）でリソースを分け合うことに注意12 レジスタ 65K

スレッド

L1/ シェアードメモリ 48KB

ブロック

(例: 256スレッド) L2 1MB メモリ5GB

グリッド

(全スレッド)

……

NVIDIA K20c

(13)

コアレスアクセス

64 65 66 67 68 94 95 0 1 2 3 4 33 1 33 0 インデックスが連続したスレッド群アドレスが連続したメモリ領域 … … + 64 65 66 67 68 94 95 0 1 2 11 5 _{ex. ストライド2のアクセス} →メモリ帯域が半分無駄に!! … …

両方とも連続するようなアルゴリズムが望ましい

またはシェアードメモリの活用

(14)

スレッドレベルの並列性

thread 0

x = x + c

x = x + b

x = x + a

thread 1

y = y + c

y = y + b

y = y + a

thread 2

z = z + c

z = z + b

z = z + a

thread 3

w = w + c

w = w + b

w = w + a

4つの独立な命令

独立なスレッドを複数用意

スレッドを切り替えてレイテンシを隠蔽

14 Volkov (2010)より

(15)

命令レベルの並列性

thread 0

w = w + b

z = z + b

y = y + b

x = x + b

w = w + a

z = z + a

y = y + a

x = x + a

4つの独立な命令

独立な命令をスレッドの中で並列処理

(16)

スレッドレベルの並列性

vs 命令レベルの並列性

基本はデータ並列性

多数のスレッドを用意する

複数データ/1スレッドを試してみる

どこかに最適解がある

どちらが良いかは処理の内容次第

余計なデータロードの削減も期待できる

16

(17)

ここまでのまとめ

GPUのハード性能

Good: 高い演算性能、高バンド幅

Bad: 高レイテンシ、弱いキャッシュ

GPUの実行モデルとプログラミングモデル

スレッドの切り替えによるレイテンシの隠蔽

階層化されたスレッドとメモリ、両者の対応

その他性能を引き出すために重要なこと

コアレスアクセス

命令レベルの並列性

(18)

海洋モデルのGPUによる高速化事例

(19)

非静力学海洋モデル kinaco

ウェッデル海における南極低層

水形成の再現シミュレーション

- 3次元Navier-Stokes方程式、移流拡散方程式_{- 静水圧近似なし、3次元の流れを陽に計算} - 等方構造格子 - ポワソン/ヘルムホルツ方程式をマルチグリッド法を用いた前処理付きCG法で求解

- 詳細はMatsumura and Hasumi (2008)

GPU向けの最適化：

- Yamagishi and Matsumura (2016)

- SC15, SC16 Poster session

～1kmのスケールの運動

を詳細に表現

(20)

3種類のカーネル最適化事例を紹介

Case1: コアレスアクセス

Case2: 命令レベルの並列性

Case3: スレッドレベルの並列性

(21)

Case1：疎行列-ベクトル積

DO k=1, n3 DO j=1, n2

DO i=1, n1

out(i,j,k) = a(-3,i,j,k) * x(i, j, k-1) & + a(-2,i,j,k) * x(i, j-1,k ) & + a(-1,i,j,k) * x(i-1,j, k ) & + a( 0,i,j,k) * x(i, j, k ) & + a( 1,i,j,k) * x(i+1,j, k ) & + a( 2,i,j,k) * x(i, j+1,k ) & + a( 3,i,j,k) * x(i, j, k+1) END DO END DO END DO

-3 -2 -1 0

1

2

3 a(-3,i,j,k)～a( 3,i,j,k)

7点ステンシル計算に相当

係数行列係数の空間的配置 -3 3 1 -1 -2 2 0

キャッシュライン上に7

点をまとめる

CPU実装

(22)

GPUでコアレスアクセス

a(i,j,k,-3)

a(i+1,j,k,-3)

a(i+2,j,k,-3)

thread(id)

thread(id+1)

thread(id+2)

a(

-3:3

,i,j,k)

a(i,j,k,

-3:3

)

GPUの各スレッドがストライドアクセス（7間隔）

a(-3,i,j,k)

a(-3,i+1,j,k)

a(-3,i+2,j,k)

thread(id)

thread(id+1)

thread(id+2)

次元の入れ替えでコアレスアクセス

(23)

Case 2: 命令レベルの並列性

i = threadidx%x + blockdim%x * (blockidx%x-1) j = threadidx%y + blockdim%y * (blockidx%y-1) k = threadidx%z + blockdim%z * (blockidx%z-1) out(i,j,k) = a(i,j,k,-3) * x(i, j, k-1) &

+ a(i,j,k,-2) * x(i, j-1,k ) & + a(i,j,k,-1) * x(i-1,j, k ) & + a(i,j,k, 0) * x(i, j, k ) & + a(i,j,k, 1) * x(i+1,j, k ) & + a(i,j,k, 2) * x(i, j+1,k ) & + a(i,j,k, 3) * x(i, j, k+1)

できる限り多くのスレッドを設定

(i, j, k)

• 3次元スレッド(i, j, k) • 1スレッド for 1データ

(24)

命令レベルの並列性を確保

独立な命令をスレッド内で繰り返し

i = threadidx%x + blockdim%x * (blockidx%x-1) j = threadidx%y + blockdim%y * (blockidx%y-1) DO k=1, n3

out(i,j,k) = a(i,j,k,-3) * x(i, j, k-1) & + a(i,j,k,-2) * x(i, j-1,k ) & + a(i,j,k,-1) * x(i-1,j, k ) & + a(i,j,k, 0) * x(i, j, k ) & + a(i,j,k, 1) * x(i+1,j, k ) & + a(i,j,k, 2) * x(i, j+1,k ) & + a(i,j,k, 3) * x(i, j, k+1) END DO

命令を重ねてレイテンシ隠蔽

• 2次元スレッド(i, j) • 1スレッド for 1カラム

(i, j)

本カーネルでは2次元

スレッドの設定の方

が高速

₂₄

(25)

Case 3: スレッドレベルの並列性

物理過程の1カーネル処理の概要

※実際のカーネルはより複雑、要点だけ抽出

DO i,j,k loop

tx(i,j,k) = Fx( A(i,j,k) )

ty(i,j,k) = Fy( B(i,j,k) )

tz(i,j,k) = Fz( C(i,j,k) )

END DO

DO i,j,k loop

out(i,j,k) = Fout( tx(i,j,k), tx(i-1,j,k),

ty(i,j,k), ty(i,j-1,k),

tz(i,j,k), tz(i,j,k-1) )

END DO

k-1 i-1 j-1

(i,j,k)

(26)

CPU tuned

DO k loop

DO i,j loop

tz0(i,j) = tz1(i,j)

tx (i,j,k) = Fx( A(i,j,k) )

ty (i,j,k) = Fy( B(i,j,k) )

tz1(i,j) = Fz( C(i,j,k) )

END DO

DO i,j loop

out(i,j,k) = Fout( tx(i,j,k), tx(i-1,j,k),

ty(i,j,k), ty(i,j-1,k),

tz1(i,j) , tz0(i,j) )

END DO

2次元データでブロック化キャッシュに k軸のデータは使い回しする 26

(27)

GPU tuned

i=threadidx%x; j=threadidx%y

DO k loop

r_tz0 = r_tz1

r_tx

= Fx( A(i ,j ,k) )

r_ty

= Fy( B(i ,j ,k) )

r_tz1 = Fz( C(i ,j ,k) )

r_tx_im1 = Fx( A(i-1,j ,k) )

r_ty_jm1 = Fy( B(i ,j-1,k) )

out(i,j,k) = Fout( r_tx , r_tx_im1,

r_ty , r_ty_jm1,

r_tz1, r_tz0 )

END DO

2次元でスレッドを生成隣接格子上の計算を付加レジスタで確保 k-1 i-1 j-1

(i,j,k)

各スレッドが独立

同期が不要

1スレッド K軸ループ

(28)

Case3「次に」何をやるべきか

複数データ/1スレッド

レジスタ増えるがロードストア・演算が減少

シェアードメモリの利用

tx, ty, tzを各スレッドで計算させた後スレッド間

で共有

スレッド間で同期をとる必要あり

以上はプロセッサのリソースとのバランスが

重要

レジスタ、シェアードメモリ増加→割り当てス

レッド数減少→レイテンシ隠蔽が非効率化

28

(29)

海洋モデル高速化上記以外の施策

カーネルの分割・融合

シェアードメモリによるスレッド間での共有

CPU-GPU間転送の最小化

混合精度演算（連立方程式ソルバ前処理を単

精度化）

粒子追跡コードにてメモリアクセスの改善

テクスチャメモリの活用

参考：Yamagishi and Matsumura (2016),

SC（15, 16）Poster Session

(30)

性能評価事例実験設定

CPU (Fujitsu SPARC64VIIIfx) vs GPU

(NVIDIA K20c)

1 CPU vs 1 GPU

傾圧不安定を伴う対流混合実験

Visbeck et al. (1996)

外部強制: 温度forcing, コリオリ力

等方構造格子

size: (256, 256, 32)

時間ステップ/総時間

2min/5hours (150 steps)

256

32 3次元: 200万

2次元: 6万

スレッド数

(31)

性能比較

CPU

GPU

Speedup

all components

174.2

37.3

4.7 Poisson/Helmholtz

solver

36.8

10.5

3.5 others

137.4

26.8

5.1 経過時間[s]: CPU vs GPU

GPUはCPU比較で4.7倍高速

CPU GPU

Computational performance (GFLOPS) 7.7 42.3(dp)/3.8(sp)

GFLOPS/PEAK (%) 6.0 3.6(dp)/0.1(sp)

Memory throughput (GB/S) 22.2 114.1

演算, メモリ性能: CPU vs GPU

dp: double precision, _{sp: single precision}

(32)

まとめ

GPUの基本を説明

ハードの特長、実行モデル・プログラミング

モデル

GPUで性能を引き出すための基本

非静力学海洋モデルの高速化事例紹介

コアレスアクセス、命令レベルの並列性、レ

ジスタ活用によるスレッドレベルの並列性

紹介しきれなかった細かい工夫も多数

CPU比較で5倍弱の高速化

32

(33)

References

NVIDIA公式資料

CUDA C Best Practices Guide, CUDA C Programming Guide

Programming Massively Parallel Processors, Third Edition, David B. Kirk, Wen-mei W. Hwu, Morgan Kaufmann.

CUDA Cプロフェッショナルプログラミング, John Chengら, インプレス. V. Volkov. Better performance at lower occupancy. GPU Technology Conference 2010.

Dan Cyca. Essential CUDA Optimization Techniques. 2014. http://on-demand.gputechconf.com/gtc/2014/video/S4702-essential-cuda-optimization-techniques-acceleware-part-4.mp4 Matsumura, Y. and Hasumi, H., 2008. A non-hydrostatic ocean model with a scalable multigrid Poisson solver. Ocean Modelling. Yamagishi, T. and Matsumura, Y., 2016. GPU Acceleration of a Non-hydrostatic Ocean Model with a Multigrid Poisson/Helmholtz Solver. Procedia Computer Science 80, 1658-1669.

出川智啓, GPGPU実践プログラミング,

Microsoft PowerPoint - 高速化WS_ver1.1.1

非静力学海洋モデルkinaco

のGPUによる高速化

山岸孝輝

1

, 松村義正

2

高度情報科学技術研究機構

東京大学大気海洋研究所

平成28年度 高速化ワークショップ

～「京」を中核とするHPCI、メニーコアを見据えて～

平成29年3月24日 秋葉原UDXカンファレンス

発表の概要

GPUの基本

ハードの特徴

実行モデル・プログラミングモデル

性能を引き出すための基本

海洋モデルのGPUによる高速化事例

まとめ

（※本発表ではNVIDIA GPUを前提）

GPUの構造

単純な構造でリソースの大半は演算 多数のコア

ハード性能（演算）

低クロックのコアを多数

バルクで高い演算性能、高並列計算

コア内パイプライン化あり

複雑な機構は無し

インオーダで処理

無し：プリフェッチ、分岐予測etc.

高い演算性能、高並列計算

単純な命令スケジューリング

ハード性能（メモリ･キャッシュ）

高いメモリバンド幅、高レイテンシ

GDDR5の性質

キャッシュ：CPUより弱い

（階層による）

演算にリソースを多く割り当てた結果

高いメモリバンド幅、高レイテンシ

弱いキャッシュ

CPU/GPU 実行モデル

各々のハードウェアの特徴によるもの

CPU

GPU

1スレッドが連続データを処理

CPUコア数＝スレッド数

基本は1スレッドが1データを処理

GPUコア数<<スレッド数

逐次処理

並列処理

カーネルのCUDA化

GPU

CPU

関数利用時にス

レッドの総数・

形状を指定

書くだけなら結構簡単だが、スレッドとデータの対応付けは

完全に書き手に任されている

逐次処理

並列処理

ハードウェア・実行モデル・プログラミングモデルの理解の

もとでスレッドとデータの対応をつける必要あり

CPU 依存性のある命令

do i = 1, N

tmp = data( i )

out(i) = tmp + 2.0

end do

プリフェッチ、キャッシュ、ソフトウェアパ

イプライン など 様々な手段で対応

GPU スレッドの切り替え

0

1

2

3

0

0

1

2

3

ストールしたスレッド

→待機していた他のス

_{, 松村義正}

_{高度情報科学技術研究機構}

_{東京大学大気海洋研究所}

平成28年度高速化ワークショップ

平成29年3月24日秋葉原UDXカンファレンス

単純な構造でリソースの大半は演算多数のコア

イプラインなど様々な手段で対応

複数データ/1スレッドを試してみる