GPU コンピューティング⼊⾨ & ハンズオン

Agenda•Introduction to GPU Computing• GPU Computing• CUDA•

ハンズオン• Lab 1• Lab 2

GPU Computing

GPU コンピューティング
Low latency + High throughput

CPUGPUアプリケーション実⾏

アプリケーション‧コードGPU

Do i=1,N

CPU

計算の重い部分
End do

逐次部分はCPU上で実⾏並列部分はGPUで実⾏GPU の構造
NVIDIA A100

⼤量の CUDA コア
並列性の抽出が鍵64 CUDA core / SM
108 Streaming Multiprocessor (SM) / chipGPU Applications

https://www.nvidia.com/en-us/gpu-accelerated-applications/GPU Applications

https://www.nvidia.com/en-us/gpu-accelerated-applications/1000 以上のアプリケーションが GPU に対応

Leading Applications

アプリを GPU 対応する⽅法Application

Library

OpenACCCUDA

GPU 対応ライブラリにチェンジ

既存コードにディレクティブを挿⼊

主要処理を CUDA で記述

簡単に開始

簡単に加速

⾼い⾃由度

NVIDIA Libraries

cuSPARSE

cuSOLVER

AmgX

cuFFT

cuRAND

nvGRAPHThrustPerformance
Primitives

cuDNN

TensorRT

NCCL

cuBLASPartner Libraries

Computer Vision

Audio and Video

Matrix, Signal and Image

Linear Algebra

Math, Signal and Image

Graph

Sparse direct solvers

Linear AlgebraLinear Algebra

Computational
GeometrySparse Iterative
Methods

Real-time visual
simulation

アプリを GPU 対応する⽅法

Application

Library

OpenACCCUDA

GPU 対応ライブラリにチェンジ

既存コードにディレクティブを挿⼊

主要処理を CUDA で記述

簡単に開始

簡単に加速

⾼い⾃由度

SAXPY (Y = A*X + Y)

OpenMP

OpenACC

void saxpy(int n,float a,float *x,float *restrict y){#pragma omp parallel forfor (int i = 0; i < n; ++i)y[i] += a*x[i];}void saxpy(int n,float a,float *x,float *restrict y){#pragma acc parallel copy(y[:n]) copyin(x[:n])for (int i = 0; i < n; ++i)y[i] += a*x[i];}...
saxpy(N, 3.0, x, y);...

...
saxpy(N, 3.0, x, y);...

アプリを GPU 対応する⽅法Application

Library

OpenACCCUDA

GPU 対応ライブラリにチェンジ

既存コードにディレクティブを挿⼊

主要処理を CUDA で記述

簡単に開始

簡単に加速

⾼い⾃由度

SAXPY (Y = A*X + Y)

CUDA

CPUvoid saxpy(int n, float a,float *x, float *y){for (int i = 0; i < n; ++i)y[i] += a*x[i];}__global__ void saxpy(int n, float a,float *x, float *y){
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)y[i] += a*x[i];
}

...

...
saxpy(N, 3.0, x, y);...

cudaMemcpy(d_x, x, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, y, size, cudaMemcpyHostToDevice);
saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y);cudaDeviceSynchronize();
cudaMemcpy(y, d_y, size, cudaMemcpyDeviceToHost);...

アプリを GPU 対応する⽅法

CuPy
Application

LibraryOpenACC

CUDA

GPU 対応ライブラリにチェンジ

既存コードにディレクティブを挿⼊

主要処理を CUDA で記述

簡単に開始

簡単に加速

⾼い⾃由度

SAXPY (Y = A*X + Y)

CPU

CuPy

import numpy as np

import numpy as np
import cupy as cp

def saxpy(a, x, y):return a * x + y

def saxpy(a, x, y):return a * x + y

...

...

x = np.arange(N, dtype=...)
y = np.arange(N, dtype=...)

x = np.arange(N, dtype=...)
y = np.arange(N, dtype=...)

y = saxpy(a, x, y)...

d_x = cp.asarray(x)d_y = cp.asarray(y)d_y = saxpy(a, d_x, d_y)y = cp.asnumpy(d_y)
...

CUDA

§ プログラミングモデル§ アーキテクチャGPU コンピューティング
CPUGPU
CPU Memory

GPU Memory
PCI• ⾼スループット指向のプロセッサ• 分離されたメモリ空間

GPU プログラム

CUDA

CPU
void saxpy(int n, float a,float *x, float *y){for (int i = 0; i < n; ++i)y[i] += a*x[i];}__global__ void saxpy(int n, float a,float *x, float *y){
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)y[i] += a*x[i];
}

...

...
saxpy(N, 3.0, x, y);...

cudaMemcpy(d_x, x, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, y, size, cudaMemcpyHostToDevice);
saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y);cudaDeviceSynchronize();
cudaMemcpy(y, d_y, size, cudaMemcpyDeviceToHost);...

GPU 実⾏の基本的な流れ

CPUGPU• GPU は、CPU からの制御で動作

⼊⼒データ
転送
GPU カーネル
投⼊• ⼊⼒データ: CPU から GPU に転送 (H2D)• GPU カーネル: CPU から投⼊
GPU 上で演算

同期出⼒データ転送• 出⼒データ: GPU から CPU に転送 (D2H)*GPU カーネル : GPU 上で実⾏される関数GPU プログラム

CUDA

CPU
void saxpy(int n, float a,float *x, float *y){for (int i = 0; i < n; ++i)y[i] += a*x[i];}__global__ void saxpy(int n, float a,float *x, float *y){
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)y[i] += a*x[i];}⼊⼒データ転送
カーネル起動
...
saxpy(N, 3.0, x, y);
同期
...

...
cudaMemcpy(d_x, x, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, y, size, cudaMemcpyHostToDevice);
saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y);cudaDeviceSynchronize();
cudaMemcpy(y, d_y, size, cudaMemcpyDeviceToHost);...

出⼒データ転送GPU カーネルCPU
void saxpy(int n, float a,float *x, float *y){}for (int i = 0; i < n; ++i)y[i] += a*x[i];...
saxpy(N, 3.0, x, y);
...

CUDA
__global__ void saxpy(int n, float a,float *x, float *y){
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)y[i] += a*x[i];Global スレッド ID}...
saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y);...

• GPU カーネル: 1 つの GPU スレッドの処理内容を記述• 基本: 1 つの GPU スレッドが、1 つの配列要素を担当

Execution Configuration
ブロック数とブロックサイズ

スレッド ID__global__ void saxpy(int n, float a,float *x, float *y){
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)y[i] += a*x[i];}ブロックサイズ...
saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y);
...

ブロック数

ブロックサイズ

ブロック数 x ブロックサイズ ≧ 配列要素数

ブロック ID

スレッド階層スレッド、ブロック、グリッド

y[i] = a*x[i] + y[i]x[]0127 128

255 256

383 384y[]

0127 128

255 256

383 3840127 1280255 256
127
0383
127

スレッド
(global)(local)ブロック0

ブロック2ブロック1

グリッド• ブロックサイズ (スレッド数/ブロック) は、カーネル毎に設定可能• 推奨: 128 or 256 スレッド

3840ブExecution Configuration
ブロック数とブロックサイズ__global__ void saxpy(int n, float a,float *x, float *y){
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)y[i] += a*x[i];}...
N/256,64, 128
25664 >>>(N, 3.0, d_x, d_y);
saxpy<<< N/
N/128,...

ブロック数

ブロックサイズ

ブロック数 x ブロックサイズ = 配列要素数

2D 配列の GPU カーネル例__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]){Global スレッド ID (x)int i = threadIdx.x + blodkDim.x * blockIdx.x;
int j = threadIdx.y + blodkDim.y * blockIdy.y;if ( i < N && j < N )Global スレッド ID (y)C[i][j] = A[i][j] + B[i][j];}ブロックサイズ (x,y)...
dim3 sizeBlock( 16,32, 1664,8 );4
dim3 numBlocks( N/sizeBlock.x, N/sizeBlock.y );MatAdd<<< numBlocks, sizeBlock >>>(A, B, C);...

ブロック数 (x,y)• ブロックサイズ (ブロック形状) は、1D~3D で表現可能

ブロック マッピング、スレッド マッピング(0,0)

dim3 sizeBlock(16,16)

(0,0)

(1,0)

(0,0)

dim3 sizeBlock(32,8)

(0,2)

(1,1)

(1,2)

(1,0)

(0,1)

(31,7)
(1,1)

(0,2)

(1,2)

(0,3)

(1,3)

(0,4)

(1,4)

(2,0)

(15,15)
(0,1)

(0,0)

(2,1)

(2,2)

ブロック ID (blockIdx)

スレッド ID (threadIdx)

CUDA§ プログラミングモデル§ アーキテクチャGPU アーキテクチャ概要•

PCI I/F•

ホスト接続インタフェース•

Giga Thread Engine• SM に処理を割り振るスケジューラ•

DRAM I/F (HBM2e)•

全 SM、PCI I/F からアクセス可能なメモリ (デバイスメモリ, フレームバッファ)•

L2 cache (40 MB)• 全 SM からアクセス可能な R/Wキャッシュ•

SM (Streaming Multiprocessor)• 「並列」プロセッサ、A100 : 108

SM (Streaming Multi-processor)

••

•
•

演算ユニット• INT32: 64 個
• FP32: 64 個
• FP64: 32 個
• TensorCore: 4 個Other units• LD/ST, SFU, etc
レジスタ (32 bit): 64K 個共有メモリ/L1 キャッシュ: 192 KB

GPU カーネル実⾏の流れ• CPU が、GPU にグリッドを投⼊• 具体的な投⼊先は、Giga Thread Engine

グリッド
ブロック
スレッド
スレッド

ブロック

ブロックを SM に割り当て• 各ブロックは、互いに独⽴に実⾏• ブロック間では同期しない、実⾏順序の保証なし• 1 つのブロックは複数 SM にまたがらない• 1 つの SM に複数ブロックが割り当てされることはある
グリッド
ブロック
ブロック
ブロック
ブロック

GPU カーネル実⾏の流れ• SM 内のスケジューラが、スレッドを CUDA コアに投⼊

グリッド
ブロック
スレッド
スレッド

ブロック

GPU カーネル実⾏の流れ• SM 内のスケジューラが、ワープを CUDA コアに投⼊• ワープ: 32 スレッドの塊• ブロックをワープに分割、実⾏可能なワープを、空 CUDA コアに割り当てる
グリッド
ブロック
ワープ
ワープ
ブロック

ワープの CUDA コアへの割り当て• ワープ内の 32 スレッドは、同じ命令を同期して実⾏• 各ワープは、互いに独⽴して実⾏(Single Instruction Multiple Threads)• 同じブロック内のワープは、明⽰的に同期可能 (__syncthreads())
グリッド
ブロック
ワープ
ワープ
ブロック
ワープ

SIMT

GPU アーキの変化を問題としないプログラミングモデル

Kepler, CC3.5
192 cores /SM

Pascal, CC6.0
64 cores /SM
Maxwell, CC5.0
128 cores /SM

Ampere, CC8.0
64 cores /SM