GPUコンピューティングの現状と未来成瀬 彰, HPC Developer Technology, NVIDIA

Summary
我々のゴールと方向性
ゴール実現に向けて進めている技術開発
Unified Memory, OpenACC
Libraries, GPU Direct

Keplerの機能紹介
Warp shuffle, Memory system
Hyper-Q, Dynamic Parallelism© NVIDIA Corporation 2013

Our Goals

電力効率© NVIDIA Corporation 2013

プログラミング簡易化
ポータビリティ

多数アプリをカバー様々な分野で使われるGPUコンピューティング© NVIDIA Corporation 2013

CUDA

パラレル・コンピューティング・プラットフォーム© NVIDIA Corporation 2013

CUDA

パラレル・コンピューティング・プラットフォーム© NVIDIA Corporation 2013

CUDA

パラレル・コンピューティング・プラットフォーム© NVIDIA Corporation 2013

GPUプログラミング言語© NVIDIA Corporation 2013

GPUプログラム実行環境© NVIDIA Corporation 2013

GPU ロードマップ
Now

32
16MaxwellDP GFLOPS per Watt

Unified Memory

8Kepler
Dynamic Parallelism

4Fermi

2FP64

1
0.5

Tesla
CUDA2008© NVIDIA Corporation 2013

2010

2012

2014

2016

Tesla K40
メモリ容量

より多くのアプリ

CPUクロック

電力状況により適切なクロックを選択

6GB

流体
解析

レンダリング
地震波解析

12GB© NVIDIA Corporation 2013

GPU BoostUnified Memory
現在© NVIDIA Corporation 2013

将来

Unified Memory
CPU code
void sortfile(FILE *fp, int N) {char *data = (char*)malloc(N);char *sorted = (char*)malloc(N);fread(data, 1, N, fp);

cpu_sort(sorted, data, N);

use_data(sorted);free(data); free(sorted)}© NVIDIA Corporation 2013

GPU code
void sortfile(FILE *fp, int N) {char *data = (char*)malloc(N);char *sorted = (char*)malloc(N);fread(data, 1, N, fp);
char *d_data, *d_sorted;cudaMalloc(&d_data, N);
cudaMalloc(&d_sorted, N);
cudaMemcpy(d_data, data, N, …);
gpu_sort<<<…>>>(d_sorted, d_data, N);
cudaMemcpy(sorted, d_sorted, N, …);
cudaFree(d_data); cudaFree(d_sorted);
use_data(sorted);free(data); free(sorted)}Unified Memory
GPU code (UVM)

CPU code
void sortfile(FILE *fp, int N) {char *data = (char*)malloc(N);char *sorted = (char*)malloc(N);fread(data, 1, N, fp);

void sortfile(FILE *fp, int N) {char *data = (char*)malloc(N);char *sorted = (char*)malloc(N);fread(data, 1, N, fp);

cpu_sort(sorted, data, N);

gpu_sort<<<…>>>(sorted, data, N);

use_data(sorted);free(data); free(sorted)

use_data(sorted);free(data); free(sorted)}© NVIDIA Corporation 2013}

OpenACC: ディレクティブCPUGPU

シンプル: ディレクティブ挿入
Program myscience... serial code ...
!$acc kernels
do k = 1,n1
do i = 1,n2... parallel code ...
enddo
enddo
!$acc end kernels...
End Program myscience

オリジナル
Fotrran/C コード© NVIDIA Corporation 2013

OpenACCCompilerHint

パワフル: 少ない労力、コンパイラが並列化

オープン: 多数のベンダのアクセラ

レータをサポート
NVIDIA, AMD, (soon) Intel

OpenACC の特徴
親しみやすいプログラミングモデル

多分野への応用

ヘテロジニアスアーキテクチャオープンスタンダード

X86 and ARM
AMD, Intel, NVIDIA© NVIDIA Corporation 2013

プログラマは並列化に注力(アーキテクチャ向け最適化はコンパイラが実施)
OpenACCによるアプリ高速化事例事例 (ORNL and Tokyo Tech)(dual-CPU nods vs. CPU+GPU)

S3D

Combustion• Tuned top 3 kernels for GPUs (90% of runtime)•• End result: 2.2X faster with K20X vs. dual AMD node•• Kepler is 6X faster than Fermi•• Improved performance of CPU-only version by 50%© NVIDIA Corporation 2013

NICAM
Weather/Climate

Tuned top kernels using CUDA, then OpenACC
CUDA result: 3.1x faster on GPU vs. CPU node
OpenACC result (preliminary ) = 69-77% of CUDA
• More portable, more maintainable• Full OpenACC port in progress
Results from Cray/ORNL and Tokyo Tech

OpenACC対応状況
Geology

Weather/Climate/
Ocean

Plasma &Combustion

Fluid Dynamics /Cosmology

Quantum
Chemistry

AWP-ODCCAM-SE

Cloverleaf

CHIMERACASTEPEMGS ELANCOSMO Physics

GENEPMH bv DELPASSGAMESS CCSD(T)*Seismic CPML*FIM

GTC

DNS

GAUSSIAN

SPECFM3D

GEOS-5

LULESHMiniGHOST

MiniMD

Harmonie

S3D

RAMSESONETEPHBM

UPACS

Quantum Espresso

ICON

X-ECHONICAMNEMO GYRE
NIM
PALM-GPU

ROMS
WRF© NVIDIA Corporation 2013

CUBLAS: 逆行列計算
LAPACK準拠API
cublas<t>getrfBatched() … LU分解
cublas<t>getriBatched() … 逆行列計算
多数の小サイズ行列用
12Speed-up

1086420(*) 行列サイズ:64, 行列数:1000© NVIDIA Corporation 2013

CPU-1core(2.8GHz,MKL)

GPU(K20,naïve)GPU(K20,cublas)

NVIDIA GPUDirect™
データ移動を最適化する技術ファミリーGPUDirect™ Shared GPU and System memory
ノード内のメモリコピー負荷を削減

GPUDirect™ Peer-to-Peer
ノード内の別GPUのメモリを直接アクセス
ノード内のGPU-to-GPUメモリ転送を加速

GPUDirect™ RDMA
ノード間でGPU-to-GPU RDMA通信© NVIDIA Corporation 2013

GPUDirect™ Shared GPU and System Memory
GPU writes to pinned main memory 1
CPU copies main memory 1 to main memory 2
Network driver reads main memory 2

GPU writes to pinned main memory
Network driver reads from main memory

NetworkCPU

Chip
set

GPU

GPU
Memory© NVIDIA Corporation 2013

Main
Mem

NetworkCPU

1With GPUDirect™

12

Without GPUDirect™

Chip
setGPUGPUMemory

Main
Mem

GPUDirect™ (Peer-to-Peer)

Fermi以上

ホストメモリを仲介せずにデータ移動 (ノード内)© NVIDIA Corporation 2013

GPUDirect™ RDMA

Kepler以上

ホストメモリを仲介せずにデータ移動 (ノード間)
SystemMemoryCPUServer 1© NVIDIA Corporation 2013

GDDR5
MemoryGDDR5MemoryGDDR5MemoryGDDR5Memory

GPU1

GPU2

GPU2

GPU1

PCI-e

SystemMemoryCPU

PCI-e

Network
Card

Network

Network
Card

Server 2

GPU-aware MPI
MPI関数だけでGPU-to-GPU通信を可能に
MPI_Send(), MPI_Recv()にデバイスメモリを指定を可能に

通信処理の最適化からプログラマを解放
パイプライン転送 (DeviceHost, HostHost, HostDevice)

Without GPU-aware MPI
送信:
cudaMemcpy( s_buf, s_device, size, … );MPI_Send( s_buf, size, … );受信:

MPI_Recv( r_buf, size, … );cudaMemcpy( r_device, r_buf, size, … );© NVIDIA Corporation 2013

With GPU-aware MPI
送信:
MPI_Send( s_device, size, …);
受信:

MPI_Recv( r_device, size, …);

GPU-aware MPI Libraries• GPUメモリからの送信・受信• 多くの集合通信に対応• 利用可能な最も良い転送方式を選択
MVAPICH

Open MPI

IBM Platform
ComputingComputing

IBM Platform MPI© NVIDIA Corporation 2013

Versions:• MVAPICH2 1.9• OpenMPI 1.7.2• IBM Platform MPI V9.1.2 (Free Community Edition)
Reference• NVIDIA GPUDirect Technology Overview

Kepler
3216Maxwell

DP GFLOPS per Watt

Unified Memory8Kepler
Dynamic Parallelism

4Fermi

2FP6410.5

Tesla
CUDA2008© NVIDIA Corporation 201320102012

2014

2016

性能と電力: Fermi  Kepler
Fermi(M2090)Kepler(K20X)

ピーク演算性能(DP)0.665 TFLOPS

1.31 TFLOPS

x2

ピーク演算性能(SP)1.33 TFLOPS

3.95 TFLOPS

x3

最大メモリバンド幅

177 GB/s

250 GB/s

x1.4

TDP

225 Watt

235 Watt

x1© NVIDIA Corporation 201327

FermiからKeplerへ
KeplerFermi

最大３倍の電力あたり性能© NVIDIA Corporation 2013

28

SM(Fermi)  SMX(Kepler)© NVIDIA Corporation 2013

Fermi
(M2090)Kepler(K20X)CUDAコア

32192

コア周波数

1.3GHz最大スレッド数1536

2048

x1.3

最大スレッドブロック数

816x2

32ビットレジスタ数

32 K

64 Kx2L2容量0.75 MB

1.5 MBx2x6

0.73GHz x0.6

29

Keplerで強化された機能
Warp Shuffle
Memory System
Atomics Operations
Read-only Cache

Hyper-Q
Concurrency
Overlapping

Dynamic Parallelism© NVIDIA Corporation 201330

Warp Shuffle
他スレッドのレジスタの読み出しを可能に
対象: 同一ワープ内のスレッド (32スレッド)

共有メモリ不要のスレッド間データ交換__syncthreads()も不要に

Kepler世代(CC 3.0以上)から利用可能© NVIDIA Corporation 201331

4種類の関数idx, up, down, xor
a b c d e f g h__shfl()__shfl_up()__shfl_down()__shfl_xor()

h d f e a c c ba b a b c d e f

c d e f g h g hc d a b g h e fIndexed
any-to-any

Shift right/up tonth neighbour

Shift left/down
to nth neighbour

Butterfly (XOR)
exchange© NVIDIA Corporation 2013

Shuffleの効果(scan)
スレッドブロック内scan(prefix sum)
入力

0

1

2

34

出力

01

3

711

SMEM (Shared Memory)for (ofst = 1; ofst < BLOCK_SIZE; ofst *= 2) {if (idx >= ofst)
smem[idx] += smem[idx - ofst];__syncthreads();}© NVIDIA Corporation 2013

SHFL (Shuffle)for (ofst = 1; ofst < WARP_SIZE; ofst *= 2) {if (idx >= ofst)val += __shfl_up(val,ofst,WARP_SIZE);}
if (idx % WARP_SIZE == WARP_SIZE - 1)
smem[idx/WARP_SIZE] = val;__syncthreads();
if (idx < NUM_WARP) {sum = smem[idx];for (ofst = 1; ofst < NUM_WARP; ofst *= 2) {if (idx >= ofst)sum += __shfl_up(sum,ofst,NUM_WARP);}
smem[idx] = sum;}__syncthreads();
if (idx/WARP_SIZE > 0)
val += smem[idx/WARP_SIZE - 1];

33

Parallel scan
01234567891011121314150

12345678101214161820220

12346810121518212428

32

360

1246912162025

30

36

4249

56

64

0

136101521

28364555

66

7891105120© NVIDIA Corporation 2013

Parallel scan
012341367864915

9102281727

2261112131415

Warp scan

Warp scan

Warp scan

Warp scan
053812253954385466

120

Warp scan
2860© NVIDIA Corporation 2013

136101521

28364555667891105120

Shuffleの効果事例(scan)
Scan (fp32)
73倍の性能UP

6Time (ms)

5x3.0

43210
SMEM

4,096B© NVIDIA Corporation 2013

SHFL
smem per block

Tesla K20
グリッド形状 (26, 1, 1)
ブロック形状 (1024, 1, 1)
1000回実行

128B
36

Shuffleの効果(reduction)
スレッドブロック内reduction
reductionコード例 (ワープ内)
SMEM (Shared Memory)idx = threadIdx.x;for (mask = WARP_SIZE/2 ; mask > 0 ; mask >>= 1) {if (idx < mask)smem[idx] += smem[idx ^ mask];__syncthreads();}

SHFL (Shuffle)for (mask = WARP_SIZE/2 ; mask > 0 ; mask >>= 1) {var = __shfl_xor( var, mask, WARP_SIZE );}© NVIDIA Corporation 201337

Shuffleの効果事例(reduction)
Reduction within TB (fp32)
52.4倍の性能UP

Time (ms)

4x2.4

3210
SMEM

4,096B© NVIDIA Corporation 2013

SHFL
smem per block

Tesla K20
ブロック形状 (1024, 1, 1)
グリッド形状 (26, 1, 1)
1000回実行

128B
38

Atomic Operations
サポートタイプ・データ型の拡張
グローバルメモリ上のAtomic操作を高速化
複数カーネルに分離していた処理を単一カーネルで効果確認
16M要素reduction
データ型はfloat© NVIDIA Corporation 2013smem[idx] = input[g_idx];for (mask = BLOCK_SIZE/2; mask > 0; mask /= 2) {if (idx < mask) {smem[idx] += smem[idx ^ mask];}__syncthreads();}
if (idx == 0) {atomicAdd( output, smem[idx] );}39

Atomic Operations効果事例
Reduction (Sum, SP, 16M elements)
6

FermiからKepler
で2.1倍の性能UP

Time (ms)

5x2.1

4

Shuffle命令併用
で3.7倍の性能UP

x3.7

321

ブロック形状 (1024, 1, 1)

0Fermi(C2075)© NVIDIA Corporation 2013

Kepler(K20)

Kepler(K20)with SHFLECC off40

Read-Only(RO) CacheSMSMEMTEXThreads

L1L2 cache

ReadTEXonly

Texture API
CUDA Arrays

一般的なRead-Onlyキャッシュとして使用可能
Kepler以降
コンパイラに指示

DRAM© NVIDIA Corporation 201341

２つの使い方組み込み関数: __ldg()__global__ kernel( int* output, int* input ){...
output[idx] = ... + input[idx__ldg( &input[idx+ delta]++delta]...; ) + ...;...
}型修飾子: const __restrict____global__ kernel( int* output, int*
constinputint* )__restrict__ input ){...
output[idx] = ... + input[idx + delta] + ...;...
}
© NVIDIA Corporation 2013

42

RO Cacheの効果
Himeno BMT
19ポイント・ステンシル

テストコード共有メモリを使用せずにCUDA化
jacobi_kernel( ..., float* p, ... );jacobi_kernel( ..., const float* __strict__ p, ... );© NVIDIA Corporation 201343

GFLOPSRO Cacheの効果事例(Himeno BMT)
100908070
60
50
4030
20
100

GFLOPS (Himeno BMT)

25%性能UP
25%

Without RO cache Without RO cache(Fermi: C2075)(Kepler: K20)
© NVIDIA Corporation 2013

With RO cache(Kepler: K20)

Himeno BMT問題サイズ: L
ブロック形状 (128, 2, 1)

ECC off44Hyper-Q
Stream Queue Mgmt

より多くのカーネルを同時実行可能に

CR

Z

B

QYA

PX

Stream Queue Mgmt
CRZ

B

Q

YA

PX

SM© NVIDIA Corporation 2013

Grid Management Unit
Pending & Suspended Grids

CUDA
Generated
Work

1000s of pending grids

Work Distributor

Work Distributor

16 active grids

32 active grids

SM

SMFermi

SM

SMX

SMX

SMX

KeplerSMX

45

Without Hyper-Q (Fermi)
Stream 1

Kernel A, B, C
Stream 2

A—B—C P—Q—R X—Y—Z
Single Hardware Work Queue

Kernel P, Q, R
Stream 3

Kernel X, Y, Z

最多16同時実行
制限: 同時実行できるのはストリーム端のカーネル© NVIDIA Corporation 201346

With Hyper-Q (Kepler)Stream 1

A—B—C

Kernel A, B, C
Stream 2

P—Q—R

Kernel P, Q, R
Stream 3

X—Y—Z

Kernel X, Y, Z

Multiple Hardware Work Queue

最多32同時実行(偽の)ストリーム依存性から開放© NVIDIA Corporation 201347

小カーネル同時実行テストコード
cudaStream_t stream[nStreams];for (i = 0 ; i < nStreams ; ++i) {// ストリーム生成
cudaCreateStream( &stream[i] );}dim3 gdim( 1, 1, 1 );dim3 bdim( 1024, 1, 1 );for (i = 0 ; i < nStreams ; ++i) {// カーネル１を投入
kernel_1<<<gdim, bdim, 0, stream[i]>>>( ... );// カーネル２を投入
kernel_2<<<gdim, bdim, 0, stream[i]>>>( ... );// カーネル３を投入
kernel_3<<<gdim, bdim, 0, stream[i]>>>( ... );}© NVIDIA Corporation 201348

小カーネル同時実行テスト(Fermi)

Tesla C2075

部分的に同時実行
シングルハードウェアキューの制約© NVIDIA Corporation 201349

小カーネル同時実行テスト(Kepler)

Tesla K20

全カーネル(ストリーム)を同時実行
これまでより簡単に同時実行が可能に© NVIDIA Corporation 201350

データ転送とカーネル実行のオーバーラップ
3つの処理をオーバーラップ可能
データ転送 (Host to Device)
カーネル実行
データ転送 (Device to Host)
cudaMemcpy( a_dev, a_host, all, cudaMemcpyHostToDevice );kernel_1<<<gdim, bdim>>>( c_dev, a_dev, all );cudaMemcpy( b_dev, b_host, all, cudaMemcpyHostToDevice );kernel_2<<<gdim, bdim>>>( c_dev, b_dev, all );cudaMemcpy( c_host, c_dev, all, cudaMemcpyDeviceToHost );© NVIDIA Corporation 201351

データ転送とカーネル実行のオーバーラップ
パイプライン化
cudaStream_t stream[nStreams];for (s = 0 ; s < nStreams ; ++s) {cudaCreateStream( &stream[s] );}s = 0;for (p = 0 ; p < nPipeline; ++p) {cudaMemcpyAsync( a_dev[p], a_host[p], part, cudaMemcpyHostToDevice, stream[s] );kernel_1<<<gdim, bdim, 0, stream[s]>>>( c_dev[p], a_dev[p], part );cudaMemcpyAsync( b_dev[p], b_host[p], part, cudaMemcpyHostToDevice, stream[s] );kernel_2<<<gdim, bdim, 0, stream[s]>>>( c_dev[p], b_dev[p], part );cudaMemcpyAsync( c_host[p], c_dev[p], part, cudaMemcpyDeviceToHost, stream[s] );s = (s+1) % nStreams;}

パイプラインコードもHyper-Qで効率化© NVIDIA Corporation 201352オーバーラップ実行テスト(Fermi)

パイプライン化前

Tesla C2075

データ転送とカーネル実行のオーバーラップ無し© NVIDIA Corporation 201353

オーバーラップ実行テスト(Fermi)

パイプライン後

Tesla C2075

データ転送とカーネル実行、相応の時間でオーバーラップ
カーネル実行の間に隙間© NVIDIA Corporation 201354

オーバーラップ実行テスト(Kepler)

完全にオーバーラップ

Tesla K20

カーネル実行の間に空き無し
Hyper-Qの効果© NVIDIA Corporation 201355

Dynamic Parallelismとは？
GPUからカーネルを起動する仕組み
Dynamically
Simultaneously
Independently

CPUGPUFermi© NVIDIA Corporation 2013

実行時のデータ値に基づくカーネル起動
複数スレッドから同時に起動
スレッド毎に独自グリッドで起動

CPU

GPU

KeplerDynamic Parallelismの動作イメージ
CPU

GPU

CPUがきめ細かく制御© NVIDIA Corporation 2013

CPU

GPU

GPUが自律的にに動作Dynamic Parallelism コードサンプル
void rec_func( ... ){...

rec_func( ... );rec_func( ... );...

}© NVIDIA Corporation 2013__global__ void rec_func( ... ){...
if ( blockIdx.x == 0 ) {cudaStreamCreate( &st0 );cudaStreamCreate( &st1 );rec_func<<< ..., st0 >>>( ... );rec_func<<< ..., st1 >>>( ... );cudaDeviceSynchronize();cudaStreamDestroy( st0 );cudaStreamDestroy( st1 );}__syncthreads();
...
}

まとめ
我々のゴール
電力効率, プログラミング簡易化, 多数アプリ

その達成のために開発している技術
Unified Memory, OpenACC
Libraries, GPU Direct, GPU-aware MPI

Keplerの機能
Warp shuffle, Memory system
Hyper-Q, Dynamic Parallelism© NVIDIA Corporation 2013