GPUコンピューティングの現状と未来

(1)

GPUコンピューティングの現状と未来

(2)

Summary

我々のゴールと方向性

ゴール実現に向けて進めている技術開発

Unified Memory, OpenACC

Libraries, GPU Direct

Keplerの機能紹介

Warp shuffle, Memory system

Hyper-Q, Dynamic Parallelism

(3)

Our Goals

プログラミング簡易化

ポータビリティ

多数アプリをカバー

(4)

(5)

CUDA

(6)

CUDA

(7)

CUDA

(8)

(9)

(10)

GPU ロードマップ

D P GFLOPS pe r W at t Kepler Tesla Fermi Maxwell Unified Memory Dynamic Parallelism FP64 CUDA 32 16 8 4 2 1 0.5 2008 2010 2012 2014 2016

Now

(11)

Tesla K40

メモリ容量

より多くのアプリ

CPUクロック

電力状況により

適切なクロックを選択

6GB

流体解析 地震波 解析 レンダ リング

12GB

GPU Boost

(12)

Unified Memory

(13)

Unified Memory

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) }

CPU 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) }

GPU code

(14)

Unified Memory

cpu_sort(sorted, data, N);

use_data(sorted);

free(data); free(sorted) }

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

use_data(sorted);

free(data); free(sorted) }

(15)

OpenACC: ディレクティブ

Program myscience ... serial code ... !$acc kernels do k = 1,n1 do i = 1,n2 ... parallel code ... enddo enddo

!$acc end kernels ...

End Program myscience

CPU GPU

オリジナル

Fotrran/C コード

OpenACC Compiler Hint

シンプル:

ディレクティブ挿入

パワフル:

少ない労力、

コンパイラが並列化

オープン:

多数のベンダのアクセラ

レータをサポート

(16)

OpenACC の特徴

オープンスタンダード

ヘテロジニアス

アーキテクチャ

X86 and ARM

AMD, Intel, NVIDIA

多分野への応用

親しみやすいプログラミングモデル

(17)

(アーキテクチャ向け最適化はコンパイラが実施)

プログラマは並列化に注力

OpenACCによるアプリ高速化事例事例 (ORNL and Tokyo Tech)

(dual-CPU nods vs. CPU+GPU)

S3D

Combustion

NICAM

Weather/Climate • 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%

• 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

(18)

OpenACC対応状況

Geology Weather/Climate/ Ocean Plasma & Combustion Fluid Dynamics / Cosmology Quantum Chemistry

AWP-ODC CAM-SE Cloverleaf CHIMERA CASTEP

EMGS ELAN COSMO Physics GENE PMH bv DELPASS GAMESS CCSD(T)

*Seismic CPML* FIM GTC DNS GAUSSIAN

SPECFM3D GEOS-5 LULESH MiniGHOST MiniMD

Harmonie S3D RAMSES ONETEP

HBM UPACS Quantum Espresso

ICON X-ECHO NICAM NEMO GYRE NIM PALM-GPU ROMS WRF

(19)

CUBLAS: 逆行列計算

LAPACK準拠API

cublas<t>getrfBatched() … LU分解

cublas<t>getriBatched() … 逆行列計算

多数の小サイズ行列用

0 2 4 6 8 10 12 CPU-1core (2.8GHz,MKL) GPU (K20,naïve) GPU (K20,cublas) Spee d -up (*) 行列サイズ:64, 行列数:1000

(20)

NVIDIA GPUDirect™

データ移動を最適化する技術ファミリー

GPUDirect™

Shared GPU and System memory

ノード内のメモリコピー負荷を削減

GPUDirect™

Peer-to-Peer

ノード内の別

GPUのメモリを直接アクセス

ノード内の

GPU-to-GPUメモリ転送を加速

GPUDirect™

RDMA

ノード間で

GPU-to-GPU RDMA通信

(21)

GPUDirect™ Shared GPU and System Memory

Without GPUDirect™

GPU writes to pinned main memory 1

CPU copies main memory 1 to main memory 2

Network driver reads main memory 2

With GPUDirect™

GPU writes to pinned main memory

Network driver reads from main memory

CPU GPU Chip set GPU Memory Network Main Mem

1

2

CPU GPU Chip set GPU Memory Network Main Mem

1

(22)

GPUDirect™ (Peer-to-Peer)

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

(23)

GPUDirect

™

_RDMA

Server 1

GPU1 GPU2 CPU GDDR5 Memory Memory GDDR5 Network Card System Memory PCI-e

Server 2

GPU1 GPU2 CPU GDDR5 Memory GDDR5 Memory Network Card System Memory PCI-e

Kepler以上

ホストメモリを仲介せずにデータ移動 (ノード間)

(24)

GPU-aware MPI

MPI関数だけでGPU-to-GPU通信を可能に

MPI_Send(), MPI_Recv()にデバイスメモリを指定を可能に

通信処理の最適化からプログラマを解放

パイプライン転送

(DeviceHost, HostHost, HostDevice)

送信:

cudaMemcpy( s_buf, s_device, size, … ); MPI_Send( s_buf, size, … );

受信:

MPI_Recv( r_buf, size, … );

cudaMemcpy( r_device, r_buf, size, … );

送信:

MPI_Send( s_device, size, …); 受信:

MPI_Recv( r_device, size, …);

(25)

GPU-aware MPI Libraries

• GPUメモリからの送信・受信

• 多くの集合通信に対応

• 利用可能な最も良い転送方式を選択

Versions:

• MVAPICH2 1.9

• OpenMPI 1.7.2

• IBM Platform MPI V9.1.2 (

Free Community Edition

)

Reference

• NVIDIA GPUDirect Technology Overview

MVAPICH

Open MPI IBM Platform

Computing

(26)

Kepler

D P GFLOPS pe r W at t Kepler Tesla Fermi Maxwell Unified Memory Dynamic Parallelism FP64 CUDA 32 16 8 4 2 1 0.5 2008 2010 2012 2014 2016

(27)

性能と電力: Fermi  Kepler

Fermi

(M2090)

Kepler

(K20X)

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

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

最大メモリバンド幅

_{177 GB/s}

_{250 GB/s}

TDP

225 Watt

235 Watt

27

x2

x3

x1.4

x1

(28)

FermiからKeplerへ

Fermi

Kepler

(29)

SM(Fermi)  SMX(Kepler)

Fermi

(M2090)

Kepler

(K20X)

CUDAコア

₃₂

₁₉₂

コア周波数

_1.3GHz

_0.73GHz

最大スレッド数

₁₅₃₆

₂₀₄₈

最大スレッドブロック数

₈

₁₆

32ビットレジスタ数

_{32 K}

_{64 K}

L2容量

_{0.75 MB}

_{1.5 MB}

29

x6

x0.6

x2

x1.3

(30)

Keplerで強化された機能

Warp Shuffle

Memory System

Atomics Operations

Read-only Cache

Hyper-Q

Concurrency

Overlapping

Dynamic Parallelism

(31)

Warp Shuffle

他スレッドのレジスタの読み出しを可能に

対象

: 同一ワープ内のスレッド (32スレッド)

共有メモリ不要のスレッド間データ交換

__syncthreads()も不要に

Kepler世代(CC 3.0以上)から利用可能

31

(32)

4種類の関数

idx, up, down, xor

a b c d e f g h a b a b c d e f Shift right/up to nth_neighbour c d e f g h g h Shift left/down to nth_neighbour c d a b g h e f Butterfly (XOR) exchange

__shfl()

__shfl_up()

__shfl_down() __shfl_xor()

h d f e a c c b Indexed

(33)

Shuffleの効果(scan)

スレッドブロック内scan(prefix sum)

33

for (ofst = 1; ofst < BLOCK_SIZE; ofst *= 2) { if (idx >= ofst)

smem[idx] += smem[idx - ofst]; __syncthreads();

}

SMEM (Shared Memory)

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];

SHFL (Shuffle)

0 入力 1 2 3 0 出力 1 3 7 4 11

(34)

Parallel scan

0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 0 1 2 3 4 5 6 7 8 10 12 14 16 18 20 22 0 1 2 3 4 6 8 10 12 15 18 21 24 28 32 36 0 1 2 4 6 9 12 16 20 25 30 36 42 49 56 64 0 1 3 6 10 15 21 28 36 45 55 66 78 91 105 120

(35)

6 28 66 120

0 1 3 6 4 9 15 22 8 17 27 38 12 25 39 54

Warp scan Warp scan Warp scan Warp scan

Warp scan

Parallel scan

0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15

6 22 38 54

(36)

Shuffleの効果事例(scan)

0 1 2 3 4 5 6 7 SMEM SHFL T ime ( ms )

Scan (fp32)

Tesla K20

グリッド形状 (26, 1, 1)

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

1000回実行

4,096B

smem per block

128B

3倍の性能UP

(37)

Shuffleの効果(reduction)

スレッドブロック内reduction

reductionコード例 (ワープ内)

SMEM (Shared Memory)

SHFL (Shuffle)

37

for (mask = WARP_SIZE/2 ; mask > 0 ; mask >>= 1) { var = __shfl_xor( var, mask, WARP_SIZE );

}

idx = threadIdx.x;

for (mask = WARP_SIZE/2 ; mask > 0 ; mask >>= 1) { if (idx < mask)

smem[idx] += smem[idx ^ mask]; __syncthreads();

(38)

Shuffleの効果事例(reduction)

0 1 2 3 4 5 SMEM SHFL T ime ( ms )

Reduction within TB (fp32)

Tesla K20

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

グリッド形状 (26, 1, 1)

1000回実行

4,096B

smem per block

128B

2.4倍の性能UP

(39)

Atomic Operations

サポートタイプ・データ型の拡張

グローバルメモリ上のAtomic操作を高速化

複数カーネルに分離していた処理を単一カーネルで

効果確認

16M要素reduction

データ型は

float

39 smem[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] ); }

(40)

Atomic Operations効果事例

0 1 2 3 4 5 6

Fermi(C2075) Kepler(K20) Kepler(K20) with SHFL

T

ime (

ms

)

Reduction (Sum, SP, 16M elements)

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

FermiからKepler

で2.1倍の性能UP

Shuffle命令併用

で3.7倍の性能UP

x2.1

x3.7

(41)

Read-Only(RO) Cache

TEX

Texture API

CUDA Arrays

一般的なRead-Onlyキャッ

シュとして使用可能

Kepler以降

コンパイラに指示

41

DRAM

L2 cache

SM

SMEM

L1

Read

_only

TEX

(42)

__global__ kernel( int* output, int* input ) {

...

output[idx] = ... + input[idx + delta] + ...; ...

}

...

}

２つの使い方

組み込み関数:

__ldg()

型修飾子:

const restrict

...

output[idx] = ... + __ldg( &input[idx + delta] ) + ...; ...

}

__global__ kernel( int* output, const int* __restrict__ input ) {

...

(43)

RO Cacheの効果

Himeno BMT

19ポイント・ステンシル

テストコード

共有メモリを使用せずに

CUDA化

43

jacobi_kernel( ..., float* p, ... );

(44)

RO Cacheの効果事例(Himeno BMT)

0 10 20 30 40 50 60 70 80 90 100 Without RO cache

(Fermi: C2075) Without RO cache(Kepler: K20) With RO cache(Kepler: K20)

GF LOP S

GFLOPS (Himeno BMT)

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

25%性能UP

25%

(45)

Hyper-Q

Work Distributor 32 active grids Stream Queue Mgmt C B A R Q P Z Y X

Grid Management Unit

Pending & Suspended Grids 1000s of pending grids SMX SMX SMX SMX SM SM SM SM Work Distributor 16 active grids Stream Queue Mgmt C B A Z Y X R Q P

Fermi

Kepler

より多くのカーネルを

同時実行可能に

45 CUDA Generated Work

(46)

Without Hyper-Q (Fermi)

最多16同時実行

制限

: 同時実行できるのはストリーム端のカーネル

Kernel P, Q, R

Kernel A, B, C

Kernel X, Y, Z

Stream 1 Stream 2 Stream 3

Single Hardware Work Queue

(47)

With Hyper-Q (Kepler)

最多32同時実行

(偽の)ストリーム依存性から開放

47

Kernel P, Q, R

Kernel A, B, C

Kernel X, Y, Z

Stream 1 Stream 2 Stream 3

A—B—C

P—Q—R

X—Y—Z

(48)

小カーネル同時実行テストコード

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]>>>( ... ); }

(49)

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

49

部分的に同時実行

シングルハードウェアキューの制約

Tesla C2075

(50)

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

全カーネル(ストリーム)を同時実行

これまでより簡単に同時実行が可能に

(51)

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 );

データ転送とカーネル実行のオーバーラップ

3つの処理をオーバーラップ可能

データ転送

(Host to Device)

カーネル実行

データ転送

(Device to Host)

51

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 );

(52)

パイプライン化

パイプラインコードも

Hyper-Qで効率化

データ転送とカーネル実行のオーバーラップ

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;

(53)

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

53

パイプライン化前

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

Tesla C2075

(54)

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

パイプライン後

データ転送とカーネル実行、相応の時間でオーバーラップ

カーネル実行の間に

隙間

(55)

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

完全にオーバーラップ

カーネル実行の間に空き無し

Hyper-Qの効果

55

Tesla K20

(56)

Dynamic Parallelismとは？

GPUからカーネルを起動する仕組み

Dynamically

実行時のデータ値に基づくカーネル起動

Simultaneously 複数スレッドから同時に起動

Independently スレッド毎に独自グリッドで起動

CPU GPU CPU GPU

(57)

CPU

GPU

CPU

GPU

Dynamic Parallelismの動作イメージ

GPUが自律的にに動作

CPUがきめ細かく制御

(58)

Dynamic Parallelism コードサンプル

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