GPUコンピューティング入門 - 九州大学情報基盤研究開発センター

GPUコンピューティング入門
2015.08.26
エヌビディア合同会社 CUDAエンジニア村上真奈
Agenda
エヌビディアのGPUについて(20分)
GPUコンピューティングとは?(10分)
OpenACC入門(90分)
CUDA入門(90分)
2
Agenda
エヌビディアのGPUについて(20分)
GPUコンピューティングとは?(10分)
OpenACC入門(90分)
CUDA入門(90分)
3
18
1993年創立
共同創業者、社長兼CEO : ジェンスン・フアン
1999年、銘柄コード「NVDA」で NASDAQ に株式上場
1999年に GPU を発明し、現在までに10億点以上を出荷
2015年会計年度の収益: 46.8億ドル
世界中に 9,300 名の従業員
7,300 件の特許取得済み資産
本社: カリフォルニア州サンタクララ
4
3
ゲーミング
オートモーティブ
エンタープライズ
HPC ＆クラウド
当社は、ビジュアル・コンピューティングが必要不可欠で重要な価値を持つ大規模な市場に特化し、プロセッサのプラット
フォーム、ソフトウェア、システム、サービスを提供しています。当社はPCテクノロジ、データセンター・テクノロジ、モバイル・テ
クノロジの革新に取り組んでいます。そして、当社の発明はさまざまな業界のOEM製品の原動力となっています。
5
6
東京工業大学 TSUBAME 2.5
4,224枚の Tesla K20X
単精度理論性能値で日本 No.1 スパコン
17PFLOPS SP
7
NVIDIA GPU OFFERS TOP LEVEL COMPUTATIONAL
PERFORMANCE WITH HIGH ENERGY EFFICIENVY
From SC TOP500 Nov., 2014
Rank Country Site
System
1
China
National Super Computer
Center in Guangzhou
Tianhe-2 (MilkyWay-2) - TH-IVB-FEP Cluster,
Intel Xeon E5-2692 12C 2.200GHz, TH
Express-2, Intel Xeon Phi 31S1P
2
US
DOE/SC/Oak Ridge
National laboratory
Titan-Cray XK7, Opt. 6274 16C 2.2GHz,
NVIDIA K20x
3
US
DOE/NNSA/LLNL
Sequoia - BlueGene/Q, Power BQC 16C 1.60
GHz, Custom
4
Japan
RIKEN Advanced Institute
for Computational Science
(AICS)
K computer, SPARC64 VIIIfx 2.0GHz, Tofu
interconnect
5
US
DOE/SC/Argonne National Mira - BlueGene/Q, Power BQC 16C 1.60GHz,
Laboratory
Custom
Cores
Rmax
(TFlop/s)
Rpeak
(TFlop/s)
Power
(kW)
3,120,000
33,862.70
54,902.40
17,808
560,640
17,590.00
27,112.50
8,209
1,572,864
17,173.20
20,132.70
7,890
705,024
10,510.00
11,280.40
12,660
786,432
8,586.60
10,066.30
3,945
In GREEN500 the most energy efficient super computers, NVIDIA GPU drives 8 systems out of TOP 10.
8
REAL WORLD EXAMPLE
Rendering 30-second Animation at Renault
9
Deep Learning における GPU の活用
Deep Learning に GPU を活用
Input
Result
110
28%
26%
GPU 対応した Deep Learning 用ツール
16%
60
person
dog
12%
0
0
4
2010
2011
2012
7%
2013
2014
chair
Caffe
Torch
Theano
Cuda-convnet
cuDNN
cuBLAS
10
GPUロードマップ
72
Volta
60
SGEMM / W
48
Pascal
36
24
Maxwell
12
Fermi
Tesla
0
2008
Kepler
2010
2012
2014
2016
2018
11
TESLA KEPLER FAMILY
WORLD’S FASTEST AND MOST EFFICIENT HPC
ACCELERATORS
CFD, BioChemistry, Neural
Networks, High Energy Physiscs,
Graph analytics, Material
Science, BioInformatics, M&E
Weather & Climate, Physics,
BioChemistry, CAE,
Material Science
Image, Signal,
Video, Seismic
Memory
Bandwidth
(ECC off)
GPUs
Single Precision
Peak (SGEMM)
Double Precision
Peak (DGEMM)
Memory
Size
PCIe Gen
System
Solution
K80
8.74 TF
(5.6TF)
2.91TF
(1.87TF)
24 GB
480GB/s
(240GB/s x2)
Gen 3
Server +
Workstation
K40
4.29 TF
(3.22TF)
1.43 TF
(1.33 TF)
12 GB
288 GB/s
Gen 3
Server +
Workstation
K20X
3.95 TF
(2.90 TF)
1.32 TF
(1.22 TF)
6 GB
250 GB/s
Gen 2
Server only
K20
3.52 TF
(2.61 TF)
1.17 TF
(1.10 TF)
5 GB
208 GB/s
Gen 2
Server +
Workstation
K10
4.58 TF
0.19 TF
8 GB
320 GB/s
Gen 3
Server only
12
THE NEW QUADRO FAMILY
M6000 K6000
K5200
K4200
K2200
K620
K420
2880
2304
1344
640
384
192
Single Precision
5.2 TFLOPs
3.1 TFLOPs
2.1 TFLOPs
1.3 TFLOPs
0.8 TFLOPs
0.3 TFLOPs
PCIe Gen
3.0
# CUDA Cores
3072
2.0
Memory Size
12GB
12 GB
8 GB
4 GB
4 GB
2 GB
1 GB
Memory BW
317 GB/s
288 GB/s
192 GB/s
173 GB/s
80 GB/s
29 GB/s
29 GB/s
2x DP* + 2x DVI
2x DP* + DVI
Slots + Display
Connectors
2x DP* + 2x DVI 2x DP* + 2x DVI
4096 x 2160
Max Resolution
Max Displays
4
250W
225 W
*
DP + DVI
DP* + DVI
3840 x 2160
4
4
4
SDI, SYNC, STEREO, MOSAIC, NVIEW
Pro Features
Board Power
4
2x DP* + DVI
150 W
* DisplayPort 1.2 multi-streaming can be used to drive multiple displays from a single DP connector
108 W
4
4
MOSAIC, NVIEW
68 W
45 W
41 W
13
前世代比 3倍の性能
Double Precision FLOPS (DGEMM)
Tesla
M2090
Tesla
K40
CUDA コア数
512
2880
倍精度演算性能
DGEMM
665 G
400 GF
1.43 TF
1.33 TF
単精度演算性能
SGEMM
1.33 TF
0.89 TF
4.29 TF
3.22 TF
メモリバンド幅
178 GB/s
288 GB/s
メモリサイズ
6 GB
12 GB
消費電力
225W
235W
1.33 TFLOPS
1.4
TFLOPS
1.2
1
0.8
0.6
0.4
0.40 TFLOPS
0.2
0
Tesla M2090
Tesla K40
Single Precision FLOPS (SGEMM)
3.5
3.22 TFLOPS
3
TFLOPS
2.5
2
1.5
1
0.89 TFLOPS
0.5
0
Tesla M2090
Tesla K40
14
NVIDIA GPU SCALABLE ARCHITECTURE
FROM SUPER COMPUTER TO MOBILE
Tesla
Tegra
In Super Computers
Quadro
In Work Stations
GeForce
In PCs
Mobile
GPU
In Tegra
17
2015
TEGRA X1 MOBILE SUPERCHIP
256-core Maxwell GPU | 8-core 64-bit CPU | 4Kp60 10-bit H.265/VP9
19
TEGRA X1 OVERVIEW
CPU: Quad ARM Cortex
A57/A53 64/32b CPU that
delivers Performance and
Power Efficiency
GPU: Next Generation 256Core Maxwell GPU that deliver
Class-Leading Performance
and Power Efficiency
End-to-End 4k 60fps Pipeline
that delivers Premium 4K
Experience
Built on 20nm Process
Technology
20
BRIDGING THE GAP
Maxwell
Kepler
Advancements
Fermi
Tesla
Tegra X1
Tegra K1
GEFORCE ARCHITECTURE
Tegra 4
Tegra 3
MOBILE ARCHITECTURE
21
WORLD’S 1ST TERAFLOPS MOBILE PROCESSOR
Tegra X1
1200
Tegra X1 (FP16)
Core i7
1000
GPU
GPU
CPU
800
CPU
GFLOPS
FP16/INT16
600
400
Tegra K1
200
Tegra 4
Tegra 2
Tegra 3
0
TIME
Note: 4790K Core i7, CPU @ 4GHz, GPU 22
@ 350 MHz
Agenda
エヌビディアのGPUについて(20分)
GPUコンピューティングとは?(10分)
OpenACC入門(90分)
CUDA入門(90分)
24
NVIDIA GPU の歴史
2010
Fermi
3 Billion
Transistors
GPU
2012
Kepler
7 Billion
Transistors
統合シェーダ + CUDA
25
GPUの構造
GPU
PCI Express
Giga Thread Engine
SM
SM
SM
SM
…
L2 Cache
DRAM
26
GPUアーキテクチャ概要
PCI I/F
SM
SM
SM
SM
SM
SM
SM
SM
ホスト接続インタフェース
Giga Thread Engine
SMに処理を割り振るスケジューラ
DRAM (384-bit, GDDR5)
SM
SM
Kepler GK110
SM
SM
SM
SM
SM
全SM、PCI I/Fからアクセス可能なメモリ
(デバイスメモリ, フレームバッファ)
L2 cache (1.5MB)
全SMからアクセス可能なR/Wキャッシュ
SM (Streaming Multiprocessor)
「並列」プロセッサ
27
SM (STREAMING MULTIPROCESSOR)
CUDA core
GPUスレッドはこの上で動作
Kepler: 192個
Other units
DP, LD/ST, SFU
Register File (65,536 x 32bit)
Shared Memory/L1 Cache
(64KB)
Read-Only Cache(48KB)
Kepler GK110
28
COMPUTE CAPABILITY
• GPUコアアーキテクチャのバージョン
•
CUDA GPUs : https://developer.nvidia.com/cuda-gpus
• アーキテクチャは進化する
• 高効率の命令実行
• 省消費電力
29
SM ARCHITECTURE VS COMPUTE
CAPABILITY
Instruction Cache
Scheduler
Scheduler
Dispatch
Dispatch
Register File
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Core
Load/Store Units x 16
Special Func Units x 4
Interconnect Network
64K Configurable
Cache/Shared Mem
Uniform Cache
Fermi
CC 2.0 : 32 cores / SM
Kepler
CC 3.5 : 192 cores / SMX
Maxwell
CC 5.0 : 128 cores / SMM
30
GPUコンピューティングとは?
• GPUは何の略?
• Graphics Processing Unit
• 3DCG等、画像データ処理の為のデバイス
• GPUによる汎用コンピューティングのこと
• 計算科学など様々な用途でGPUを利用する
31
ヘテロジニアス・コンピューティング
CPU
逐次処理に最適化
GPU Accelerator
並列処理に最適化
32
GPUアプリケーションの例
画像処理コンピュータビジョン
医療画像
防衛
計算化学
気象
金融工学
バイオ
数値解析
33
GPUアクセラレーションの実現方法
アプリケーション
GPU
ライブラリ
ライブラリを呼び出すだけ
簡単に高速化を実現
OpenACC
ディレクティブ
CUDA
既存コードにディレクティブ
を挿入して高速化
重要なコードをCUDAで記述
最も自由度が高い
34
GPUアクセラレーションの実現方法
簡単
ライブラリ
ライブラリを呼び出すだけで、高速化が可能
ライブラリとして提供されている機能のみ高速化が可能
OpenACC
既存のC言語やFortranのコードにディレクティブを挿入するだけで簡単に高速化
最適化はコンパイラが行う為、細かいチューニングを行う事は出来ない
CUDA
高速化
自由度が最も高く、細かいチューニングが可能
CUDAでのプログラミングを学ぶ必要がある
35
Agenda
エヌビディアのGPUについて(20分)
GPUコンピューティングとは?(10分)
OpenACC入門(90分)
CUDA入門(90分)
36
GPUアクセラレーションの実現方法
アプリケーション
GPU
ライブラリ
ライブラリを呼び出すだけ
簡単に高速化を実現
OpenACC
ディレクティブ
CUDA
既存コードにディレクティブ
を挿入して高速化
重要なコードをCUDAで記述
最も自由度が高い
37
OPENACC
標準的なGPUディレクティブ
シンプル:
ディレクティブを挿入するのみ。コードを変更する事なく高速化
オープン:
OpenACCはマルチコアプロセッサで並列化を行う為のオープン標準
柔軟:
GPU ディレクティブは、高い並列性を保ちつつ同一コードで複数のアーキテクチャに対応可能
38
OpenACCメンバーとパートナー
39
コンパイラとツール
2013年12月～
コンパイラ
2014年1月～
2015年(予定)
OpenACC 2.0対応
デバッグツール
40
簡単に高速
自動車
金融
生命科学
Real-Time Object
Detection
Valuation of Stock Portfolios
using Monte Carlo
Interaction of Solvents and
Biomolecules
Global Manufacturer of Navigation
Systems
Global Technology Consulting Company
University of Texas at San Antonio
40時間で5倍
4時間で2倍
8時間で5倍
41
簡単に始められる
大学関係者の方は無償で使用可能に
下記のサイトからOpenACC toolkitをダウンロード
https://developer.nvidia.com/openacc
PGIコンパイラ/MPI/CUDAなど
一式が簡単にインストール可能
42
実行モデル
アプリケーション・コード
GPU
$acc parallel
計算の
重い部分
CPU
$acc end parallel
並列部分は
GPUコードを生成
逐次部分は
CPUコードを生成
43
OpenACC ディレクティブ
CPU
GPU
コンパイラへシンプルなヒント
Program myscience
... serial code ...
!$acc kernels
do k = 1,n1
do i = 1,n2
... parallel code ...
enddo
enddo
!$acc end kernels
...
End Program myscience
コンパイラがコードを並列化
コンパイラへの
OpenACC
ヒント
並列部はGPUで
逐次処理はCPUで動作
Fortran または C言語
のオリジナルコード
44
OpenMPとOpenACCの比較
OpenMP
CPU
main() {
double pi = 0.0; long i;
#pragma omp parallel for reduction(+:pi)
for (i=0; i<N; i++)
{
double t = (double)((i+0.05)/N);
pi += 4.0/(1.0+t*t);
}
CPUコアに
printf(“pi = %f\n”, pi/N);
計算処理を分散
}
OpenACC
CPU
GPU
main() {
double pi = 0.0; long i;
#pragma acc kernels
for (i=0; i<N; i++)
{
double t = (double)((i+0.05)/N);
pi += 4.0/(1.0+t*t);
}
printf(“pi = %f\n”, pi/N);
}
GPUコアに
計算処理を分散
45
OpenACC ディレクティブ構文
C/C++
#pragma acc 指示行 [節[,]節] …]
{ structured block }
Fortran
!$acc 指示行 [節[,]節] …]
{ structured block }
!$acc end directive
46
OpenACC構文: parallel 指示行
• parallel : 並列に実行される領域を指示行で指定
#pragma acc parallel
for(int i=0;i<n;i++){
a[i] = 0.0;
b[i] = 1.0;
c[i] = 2.0;
}
kernel 1
Kernel(カーネル):
GPU上で実行される
関数
47
OpenACC構文: kernels 指示行
• kernels : 複数のカーネルを作成
#pragma acc kernels
for(int i=0;i<n;i++){
a[i] = 0.0;
b[i] = 1.0;
c[i] = 2.0;
}
#pragma acc kernels
for(int i=0;i<n;i++){
a[i] = b[i] + c[i];
}
kernel 1
Kernel(カーネル):
GPU上で実行される
関数
kernel 2
48
[C tips]: restrict修飾子
コンパイラに対して明示的にrestrict修飾子を指定。ポインタのエイリアシング
を制限
例)
float *restrict ptr
OpenACCコンパイラにrestrict修飾子をつけ変数の独立性を伝える
独立性の保障がないとコンパイラは並列化を行う事が出来ない
49
http://en.wikipedia.org/wiki/Restrict
例:SAXPY (Y=A*X+Y)
Trivial first example
Apply a loop directive
Learn compiler commands
int main(int argc, char **argv)
{
int N = 1<<20; // 1 million floats
if (argc > 1)
N = atoi(argv[1]);
float *x = (float*)malloc(N * sizeof(float));
float *y = (float*)malloc(N * sizeof(float));
*restrict:
#include <stdlib.h>
“yはxのエイリアスでない”と明示的
に指定
void saxpy(int n,
float a,
float *x,
float *restrict y)
{
#pragma acc kernels
for (int i = 0; i < n; ++i)
y[i] = a * x[i] + y[i];
}
for (int i = 0; i < N; ++i) {
x[i] = 2.0f;
y[i] = 1.0f;
}
saxpy(N, 3.0f, x, y);
return 0;
}
50
C言語:SAXPY (Y=A*X+Y)
OpenMP
OpenACC
void saxpy(int n,
float a,
float *x,
float *restrict y)
{
#pragma omp parallel for
for (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);
...
 omp  acc
 データの移動
51
Fortran: SAXPY (Y=A*X+Y)
OpenMP
OpenACC
subroutine saxpy(n, a, X, Y)
real :: a, X(:), Y(:)
integer :: n, i
subroutine saxpy(n, a, X, Y)
real :: a, Y(:), Y(:)
integer :: n, i
!$omp parallel do
do i=1,n
Y(i) = a*X(i)+Y(i)
enddo
!$omp end parallel do
end subroutine saxpy
!$acc parallel copy(Y(:)) copyin(X(:))
do i=1,n
Y(i) = a*X(i)+Y(i)
enddo
!$acc end parallel
end subroutine saxpy
...
call saxpy(N, 3.0, x, y)
...
...
call saxpy(N, 3.0, x, y)
...
52
コンパイルオプション
C:
pgcc –acc -ta=nvidia -Minfo=accel –o saxpy_acc saxpy.c
Fortran:
pgf90 –acc -ta=nvidia -Minfo=accel –o saxpy_acc saxpy.f90
ターゲットに
nvidiaを指定
コンパイラがGPU用のコードを生
成する際の情報を表示する
53
簡単にコンパイル
OpenMP / OpenACC
void saxpy(int n, float a,
float *x,
float *restrict y)
$ pgcc –acc {–ta=nvidia –Minfo=accel saxpy.c
#pragma acc parallel copy(y[:n]) copyin(x[:n])
saxpy:
#pragma
omp parallel for
16, Generating
present_or_copy(y[:n])
for present_or_copyin(x[:n])
(int i = 0; i < n; ++i)
Generating
+= code
a*x[i];
Generatingy[i]
Tesla
} parallelizable
19, Loop is
Accelerator kernel generated
...
19, #pragma
acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
saxpy(N, 3.0, x, y);
...
54
簡単に実行
OpenMP / OpenACC
void saxpy(int n, float a,
float *x,
float *restrict y)
{
$ pgcc -Minfo -acc
saxpy.c
$ nvprof ./a.out
#pragma acc kernels copy(y[:n]) copyin(x[:n])
saxpy:
==10302==
NVPROF
ispresent_or_copy(y[:n])
profiling
process
#pragma
omp parallel
for 10302, command: ./a.out
16, Generating
==10302==
Profiling
application:
for present_or_copyin(x[:n])
(int
i = 0; i ./a.out
< n; ++i)
Generating
==10302==
Profiling
result:
y[i]
+= code
a*x[i];
Generating
Tesla
Time(%)
Calls
Avg
Min
Max Name
} parallelizable
19, Loop Time
is
62.95% Accelerator
3.0358ms
1.5179ms 1.5172ms 1.5186ms [CUDA memcpy HtoD]
kernel2generated
31.48% 19,
1.5181ms
1 1.5181ms
1.5181ms/*1.5181ms
[CUDA
memcpy DtoH]
...
#pragma
acc loop
gang, vector(128)
blockIdx.x
threadIdx.x
*/
5.56% 268.31us
268.31us 268.31us saxpy_19_gpu
saxpy(N, 3.0,1 x,268.31us
y);
...
55
例: ヤコビ反復法
正しい値になるように反復計算を行う。隣接点の平均値で値を更新
連立一次方程式を解く為のオーソドックスな手法
例: 2次元ラプラス方程式: 𝛁𝟐 𝒇(𝒙, 𝒚) = 𝟎
A(i,j+1)
A(i-1,j)
A(i,j)
A(i+1,j)
𝐴𝑘+1 𝑖, 𝑗 =
𝐴𝑘 (𝑖 − 1, 𝑗) + 𝐴𝑘 𝑖 + 1, 𝑗 + 𝐴𝑘 𝑖, 𝑗 − 1 + 𝐴𝑘 𝑖, 𝑗 + 1
4
A(i,j-1)
56
ヤコビ反復法(アルゴリズム)
A(i,j+1)
while ( error > tol ) {
error = 0.0;
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]));
}
}
A(i-1,j)
A(i,j) A(i+1,j)
A(i,j-1)
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
57
並列領域 (OpenMP)
while ( error > tol ) {
error = 0.0;
#pragma omp parallel for shared(m, n, Anew, A)
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
#pragma omp parallel for shared(m, n, Anew, A)
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
58
並列領域 (OpenACC)
while ( error > tol ) {
error = 0.0;
#pragma acc kernels
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
#pragma acc kernels
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
 Parallels と Kernels
— 並列領域を指示
 Parallels
— 並列実行スタート
 Kernels
— 複数のカーネル
59
[PGI tips] コンパイラメッセージ
$ pgcc –acc –ta=nvidia –Minfo=accel jacobi.c
jacobi:
44, Generating copyout(Anew[1:4094][1:4094])
Generating copyin(A[:][:])
Generating Tesla code
45, Loop is parallelizable
46, Loop is parallelizable
Accelerator kernel generated
45, #pragma acc loop gang /* blockIdx.y */
46, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
49, Max reduction generated for error
60
並列領域 (KERNELS CONSTRUCT)
while ( error > tol ) {
error = 0.0;
 Parallels と Kernels
— 並列領域を指示
#pragma acc kernels
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
— 並列走行の開始
$ pgcc –acc
–ta=nvidia
error = max(error, -Minfo=accel
abs(Anew[j][i] -jacobi.c
A[j][i]);
jacobi:
}
} 59, Generating present_or_copyout(Anew[1:4094][1:4094])
 Parallels
 Kernels
— 複数のGPUカーネル
Generating present_or_copyin(A[:][:])
#pragma acc kernels
Generating
code{
for (int
j = 1; j <Tesla
N-1; j++)
for (int
= 1;
i < M-1; i++) {
61,
Loop iis
parallelizable
A[j][i]
= Anew[j][i];
63,
Loop is
parallelizable
}
Accelerator kernel generated
}
61, #pragma acc loop gang /* blockIdx.y */
}
63, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
Max reduction generated for error
61
データ転送 (DATA CLAUSE)
while ( error > tol ) {
error = 0.0;
#pragma acc kernels
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
$ pgcc –acc –ta=nvidia
-Minfo=acc
jacobi.c
A[j-1][i]
+ A[j+1][i])
* 0.25;
jacobi: error = max(error, abs(Anew[j][i] - A[j][i]);
}
59, Generating present_or_copyout(Anew[1:4094][1:4094])
}
Generating present_or_copyin(A[:][:])
Generating
#pragma
acc kernelsTesla code
for
(int
j =is
1; parallelizable
j < N-1; j++) {
61,
Loop
for (int i = 1; i < M-1; i++) {
63, Loop is parallelizable
A[j][i] = Anew[j][i];
Accelerator kernel generated
}
}
61, #pragma acc loop gang /* blockIdx.y */
}
63, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
Max reduction generated for error
62
コードの解析
実行状況を確認。ボトルネックはどの部分か？
実行時間の内訳を調べる
63
-ta=nvidia,time
コンパイルオプションに-ta=nvidia,timeを追加して、コンパイル・実行
/home/OpenACC/C/jacobi.c
Kernel実行:196ms
jacobi NVIDIA devicenum=0
time(us): 4,595,922
44: compute region reached 200 times
46: kernel launched 200 times
grid: [32x4094] block: [128]
データコピー(H->D):1087ms
device time(us): total=196,036 max=1,053 min=931 avg=980
elapsed time(us): total=201,618 max=1,084 min=958 avg=1,008
46: reduction kernel launched 200 times
grid: [1] block: [256]
device time(us): total=39,356 max=206 min=187 avg=196
elapsed time(us): total=42,155 max=227 min=200 avg=210
44: data region reached 200 times
44: data copyin transfers: 800
device time(us): total=1,087,027 max=1,374 min=1,354
avg=1,358
データコピーがボトルネック
53: compute region reached 200 times
55: kernel launched 200 times
64
grid: [32x4094] block: [128]
NVIDIA Visual Profiler (NVVP)を使用
65
NVVPによる解析: データ転送がボトルネック
利用率:低い
1 cycle
GPU
kernel
GPU
kernel
66
計算処理とデータ転送
CPU Memory
データ転送
GPU Memory
PCI
計算オフロード
計算オフロード、データ転送、両方を考慮する必要がある
67
OpenACC構文: データ指示行
• copy ( X )
• copyin(list) + copyout(list)
• copyin ( X )
• アクセラレータ領域に入る際にGPU上に X 用のメモリを確保し、ホストからGPU(デバ
イス)へ X を転送する
• copyout ( X )
• アクセラレータ領域に入る際にGPU上に X 用のメモリを確保し、アクセラレータ領域か
ら出る時にGPU(デバイス)からホストへ X を転送する
• create ( X )
• アクセラレータ領域に入る時にGPU上に X 用のメモリが確保される (転送はされない)
• present ( X )
• アクセラレータ領域に入る時に X が既にデバイス上に存在することを示す
68
OpenACC構文: データ指示行
• pcopy ( X )
• present (X) + copy(X)
• pcopyin ( X )
• present (X) + copyin(X)
• pcopyout ( X )
• present (X) + copyout(X)
• pcreate ( X )
• present (X) + create(X)
69
データ転送 (DATA CLAUSE)
while ( error > tol ) {
error = 0.0;
#pragma acc kernels \
pcopyout(Anew[1:N-2][1:M-2]) pcopyin(A[0:N][0:M])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
#pragma acc kernels \
pcopyout(A[1:N-2][1:M-2]) pcopyin(Anew[1:N-2][1:M-2])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
 copyin (HostGPU)
 copyout (HostGPU)
 copy
 create
 present
 pcopyin
 pcopyout
 pcopy
 pcreate
70
データ転送 (DATA CLAUSE)
while ( error > tol ) {
error = 0.0;
#pragma acc kernels \
pcopy(Anew[:][:]) pcopyin(A[:][:])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
#pragma acc kernels \
pcopy(A[:][:]) pcopyin(Anew[:][:])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
 copyin (HostGPU)
 copyout (HostGPU)
 copy
 create
 present
 pcopyin
 pcopyout
 pcopy
 pcreate
71
過剰なデータ転送
while ( error > tol ) {
error = 0.0;
#pragma acc kernels \
pcopy(Anew[:][:]) pcopyin(A[:][:])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
#pragma acc kernels \
pcopy(A[:][:]) pcopyin(Anew[:][:])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
72
過剰なデータ転送
GPU
Host
while ( error > tol ) {
error = 0.0;
#pragma acc kernels \
pcopy(Anew[:][:]) \
pcopyin(A[:][:])
{
}
#pragma acc kernels \
pcopy(A[:][:]) \
pcopyin(Anew[:][:])
{
}
}
copyin
copyout
copyin
copyout
#pragma acc loop reduction(max:error)
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
73
データ領域 (data construct)
#pragma acc data pcopy(A, Anew)
while ( error > tol ) {
error = 0.0;
#pragma acc kernels pcopy(Anew[:][:]) pcopyin(A[:][:])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
#pragma acc kernels pcopy(A[:][:]) pcopyin(Anew[:][:])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
 copyin (CPUGPU)
 copyout (CPUGPU)
 copy
 create
 present
 pcopyin
 pcopyout
 pcopy
 pcreate
}
74
データ領域 (data CONSTRUCT)
#pragma acc data pcopy(A) create(Anew)
while ( error > tol ) {
error = 0.0;
#pragma acc kernels pcopy(Anew[:][:]) pcopyin(A[:][:])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
#pragma acc kernels pcopy(A[:][:]) pcopyin(Anew[:][:])
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
 copyin (CPUGPU)
 copyout (CPUGPU)
 copy
 create
 present
 pcopyin
 pcopyout
 pcopy
 pcreate
75
適正なデータ転送
GPU
Host
#pragma acc data \
pcopy(A) create(Anew)
while ( error > tol ) {
error = 0.0;
copyin
#pragma acc kernels \
pcopy(Anew[:][:]) \
pcopyin(A[:][:])
{
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
Anew[j][i] = (A[j][i+1] + A[j][i-1] +
A[j-1][i] + A[j+1][i]) * 0.25;
error = max(error, abs(Anew[j][i] - A[j][i]);
}
}
}
#pragma acc kernels \
pcopy(A[:][:]) \
pcopyin(Anew[:][:])
{
for (int j = 1; j < N-1; j++) {
for (int i = 1; i < M-1; i++) {
A[j][i] = Anew[j][i];
}
}
}
}
copyout
76
データ転送の削減 (NVVP)
稼働率:高い
1 cycle
77
GPUアクセラレーションの実現方法
アプリケーション
GPU
ライブラリ
ライブラリを呼び出すだけ
簡単に高速化を実現
OpenACC
ディレクティブ
CUDA
既存コードにディレクティブ
を挿入して高速化
重要なコードをCUDAで記述
最も自由度が高い
78
CUDAとは?
• Compute Unified Device Architectureの略
• NVIDIA GPU上の汎用並列計算プラットフォーム
• Linux・Windows・MacOS X(+Android)で動作
• 現在、7.0が最新
• ※7.5RCも公開中
79
CUDA 開発・実行環境
ライブラリ・ミドルウェア
cuDNN
cuSOLVER
cuRAND
cuFFT
cuBLAS
cuSPARSE
Thrust
NPP
NVRTC
プログラミング言語
C
C++
C++11
Fortran
Java
Python
etc..
NVIDIA-GPUs(ハードウェア)
HyperQ
Dynamic
Parallelism
開発環境
MATLAB
Mathematica
etc..
NVCC
(CUDA compiler)
CUDA-GDB
Profiler
Nsight IDE
GPU
Direct
82
CUDA 開発・実行環境
ライブラリ・ミドルウェア
cuDNN
cuSOLVER
cuRAND
cuFFT
cuBLAS
cuSPARSE
Thrust
NPP
NVRTC
プログラミング言語
C
C++
C++11
Fortran
Java
Python
etc..
NVIDIA-GPUs(ハードウェア)
HyperQ
Dynamic
Parallelism
開発環境
MATLAB
Mathematica
etc..
NVCC
(CUDA compiler)
CUDA-GDB
Profiler
Nsight IDE
GPU
Direct
83
進化するハードウェア NVIDIA-GPUS
84
進化するハードウェア NVIDIA-GPUS
Hyper-Q
Dynamic Parallelism
GPU Direct
85
CUDA 開発・実行環境
ライブラリ・ミドルウェア
cuDNN
cuSOLVER
cuRAND
cuFFT
cuBLAS
cuSPARSE
Thrust
NPP
NVRTC
プログラミング言語
C
C++
C++11
Fortran
Java
Python
etc..
NVIDIA-GPUs(ハードウェア)
HyperQ
Dynamic
Parallelism
開発環境
MATLAB
Mathematica
etc..
NVCC
(CUDA compiler)
CUDA-GDB
Profiler
Nsight IDE
GPU
Direct
87
プログラミング言語
C
C++
CUDA C
CUDA C++(C++11),Thrust
Python
PyCUDA
Fortran
CUDA Fortran
その他
F#, MATLAB, Mathematica, …
88
CUDA 開発・実行環境
ライブラリ・ミドルウェア
cuDNN
cuSOLVER
cuRAND
cuFFT
cuBLAS
cuSPARSE
Thrust
NPP
NVRTC
プログラミング言語
C
C++
C++11
Fortran
Java
Python
etc..
NVIDIA-GPUs(ハードウェア)
HyperQ
Dynamic
Parallelism
開発環境
MATLAB
Mathematica
etc..
NVCC
(CUDA compiler)
CUDA-GDB
Profiler
Nsight IDE
GPU
Direct
90
CUDAライブラリ
cuDNN
ディープニューラルネットワーク
計算用ライブラリ
cuFFT
高速フーリエ変換ライブラリ
cuSOLVER
cuRAND
cuSPARSE
乱数生成ライブラリ
疎行列計算用ライブラリ
cuBLAS
NPP
Thrust
線形代数演算LAPACK用
ライブラリ
線形代数計算用ライブラリ
動画像処理・信号処理用
ライブラリ
Ｃ++テンプレートライブラリ
91
CUDAを使用したソフトウェア
MATLAB
Mathematica
OpenCV
ArrayFire
etc…
Caffe
torch
theano
92
CUDA 開発・実行環境
ライブラリ・ミドルウェア
cuDNN
cuSOLVER
cuRAND
cuFFT
cuBLAS
cuSPARSE
Thrust
NPP
NVRTC
プログラミング言語
C
C++
C++11
Fortran
Java
Python
etc..
NVIDIA-GPUs(ハードウェア)
HyperQ
Dynamic
Parallelism
開発環境
MATLAB
Mathematica
etc..
NVCC
(CUDA compiler)
CUDA-GDB
Profiler
Nsight IDE
GPU
Direct
94
開発環境 DEBUG & ANALYSIS
NVCC
CUDA用コンパイラ
CUDA-GDB
CUDA用デバッガ(Linux,Mac)
CUDA-MEMCHECK
GPUメモリエラーチェックツール
Nsight IDE
Profiler
CUDA統合開発環境(Linux,Windows)
CUDA解析ツール
95
NSIGHT VISUAL STUDIO EDITION
96
ここまでの復習
• CUDAでは、様々なプログラミング言語やライブラリを使う事が可能
• ケースによって最適なものを選択すれば良い
• 既存のライブラリやミドルウェアを有効活用する
• CUDAはロードマップが存在し、進化し続けている
• よりプログラミングしやすく
• パフォーマンスが出やすいように
97
CUDA C/C++アプリケーション入門
今回は、CUDA C/C++で説明します
98
典型的な装置構成
CPUにつながった外部演算装置
GPU
PC
CPU
(数コア)
Giga Thread Engine
制御
PCIe
ホスト側DRAM
SM
SM
SM
SM
…
L2 Cache
DRAM
転送
99
典型的な実行例
プログラム
開始
GPUはCPUからの制御
で動作する。
CPU
データ
転送
CUDA
カーネル
実行
完了
待ち
データ
転送
入力データは
CPU→GPUへと転送。
結果は、
GPU→CPUと転送
GPU
GPUでの演算
GPU上に常駐する
プログラムはない。
100
CUDA C/C++用語
• GPUで実行される関数をカーネル(kernel)と呼ぶ
• CPUで実行されるコードをホストコード、GPUで実行されるコードを
デバイスコードと呼ぶ
• データ並列を表現する為に以下の概念を用いる
• グリッド (grid)
• ブロック (block)
• スレッド (thread)
101
グリッド・ブロック・スレッド
• グリッド (grid)
•
ブロックをまとめた物
• ブロック (block)
スレッドをまとめた物
•
1ブロックあたり最大1024スレッド
• スレッド (thread)
•
カーネルを動作させる最小単位
Thread
Thread
Block1
Thread
Thread
Thread
Thread
Block2
Grid
Thread
Thread
…
•
Thread
Thread
Block0
Thread
Thread
Thread
Thread
Block
n
Thread
Thread
103
グリッド・ブロック・スレッド
GPU
CUDA
Block0
Grid
GPU
Thread
Thread
Thread
Thread
Block1
Thread
Thread
Thread
Thread
Block2
Thread
Thread
Thread
Thread
…
…
Block n
Thread
Thread
Thread
Thread
SM
SM
SM
SM
core
105
カーネル実行の流れ
•
Giga Thread EngineがブロックをSMに割り当てる
Grid
Block0
Giga Thread Engine
Block1
Block2
Block3
Block4
…
BlockN
Block4
107
カーネル実行の流れ
•
SMの中のスケジューラがコアにスレッドを投入する
Grid
Block 0
ワープを投入
Thread
Thread
Thread
Thread
32スレッド単位で
投入
Thread
Thread
Thread
Thread
Thread
Thread
Thread
Thread
Thread
Thread
Thread
Thread
…
BlockN
108
BLOCK は SM上で実行
• Block ⇒ 1 SM
• 複数のSMにまたがらない。
（SM中では、複数Blockが実行される場合もある。）
• Block内部では、SMXのリソースを活用可能
• 各々のBlockは、独立に、非同期に処理を実行する。
• 実行順序の保証はない。
• Block間の通信・同期は行わない。
109
例: 一次元配列の加算
• 配列Aと配列Bの加算結果を配列Cに書き込む
[0] [1]
A
10
1
[2] [3] [4] [5] [6] [7]
8
7
14
13
2
5
[8] [9] [10] [11] [12] [13] [14] [15]
6
15
3
9
12
11
0
4
+ + + + + + + + + + + + + + + +
B
0
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
…
…
=
=
=
=
=
=
=
=
=
=
=
=
=
=
=
=
C
10
2
10
10
18
18
8
12
14
24
13
20
24
24
14
19
…
110
例: 一次元配列の加算(CPU)
• 配列の0番から逐次加算していく
CPU
C[0] = A[0] + B[0];
C[1] = A[1] + B[1];
C[2] = A[2] + B[2];
for(int i=0
C[3]; =i<nMatrixSize
A[3] + B[3]; ; i++)
{
C[4]= =A[i]
A[4]
+ B[4];
C[i]
+ B[i];
C[5] = A[5] + B[5];
}
…
111
例: 一次元配列の加算(GPU)
[0] [1]
A
B
10
1
[2] [3] [4] [5] [6] [7]
8
7
14
13
2
5
[8] [9] [10] [11] [12] [13] [14] [15]
6
15
3
9
12
11
0
4
+ + + + + + + + + + + + + + + + …
0
1
2
5
3
4
6
7
8
9 10 11 12 13 14 15
=
=
=
=
=
=
=
=
=
=
=
=
=
=
=
=
C
…
13
20
24
24
10
2
10
10
18
18
8
12
14
24
T0
T1
T2 T3
T4
T5
T6
T7
T8
T9 T10 T11 T12 T13
Block0
Block1
Block2
14
19
………
…
…
BlockN
112
ブロックIDとスレッドID
ブロックIDとスレッドIDから、インデックス(グローバルID)を生成する
インデックスを用いて各スレッドから、グローバルメモリへアクセスする
index = blockDim.x * blockIdx.x + threadIdx.x;
×
+
2
1
6
8
Thread
Block
0
1
2
3
4
5
6
7
8
9 10 11
0
1
2
3
4
5
0
1
2
3
0
11
4
5
12 13 14 15 16 17
0
1
2
3
4
5
2
113
例: 一次元配列の加算(GPU)
ホスト側(CPU)
int main(int argc,char** argv){
…
MatrixAdd<<<N, M>>>(C,A,B);
…
}
デバイス側(GPU)
__global__ void
MatrixAdd(float* C,const float* A,const float* B){
int i = blockDim.x * blockIdx.x + threadIdx.x;
C[i] = A[i] + B[i];
}
114
GPU側メモリ
の確保
復習:典型的な実行例
GPUはCPUからの制御
で動作する。
CPU
データ
転送
CUDA
カーネル
実行
完了
待ち
データ
転送
入力データは
CPU→GPUへと転送。
結果は、
GPU→CPUと転送
GPU
GPUでの演算
GPU上に常駐する
プログラムはない。
115
ホスト側から呼び出すAPI
cudaMalloc
GPU上のDRAM(グローバルメモリ)にメモリの確保を行う
cudaFree
cudaMallocで取得したメモリの解放を行う
cudaMemcpy
CPU->GPU、GPU->GPU、GPU->CPUのメモリ転送を行う
cudaDeviceSynchronize
CUDAカーネルが終了するまで待つ
116
cudaMemcpy()
メモリは、「ホスト」「デバイス」の二種類
enum cudaMemcpyKind
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice
cudaMemcpyHostToHost
(cudaMemcpyDefault : UVA)
117
サンプルコード(ホスト)
int main() {
…… 略……
int matrixSize= 256 * 100;
float *A, *B, *C;
cudaMalloc(&A,sizeof(float)*matrixSize);
cudaMalloc(&B,sizeof(float)*matrixSize);
cudaMalloc(&C,sizeof(float)*matrixSize);
cudaMemcpy(A,hA, sizeof(float)*matrixSize, cudaMemcpyHostToDevice);
cudaMemcpy(B,hB, sizeof(float)*matrixSize, cudaMemcpyHostToDevice);
MatrixAdd<<<matrixSize/256, 256>>>(C, A, B, matrixSize);
cudaDeviceSynchronize();
cudaMemcpy(hC, C, sizeof(float)*matrixSize, cudaMemcpyDeviceToHost);
cudaFree(A); cudaFree(B); cudaFree(C);
…… 略……
}
118
サンプルコード(デバイス)
__global__ void
MatrixAdd(float* C,const float* A,const float* B,const int size){
int i = blockDim.x * blockIdx.x + threadIdx.x;
if( i < size){
C[i] = A[i] + B[i];
}
}
119
例:RGB->YUV変換を考える
𝑌
0.299
𝑈 = −0.169
𝑉
0.500
0.587
−0.331
−0.419
0.114
0.500
−0.081
𝑅
𝐺
𝐵
121
最適化の為に理解する事
1. GPUのメモリ構造
2. スレッド(thread)構成と占有率(Occupancy)
122
最適化の為に理解する事
1. GPUのメモリ構造
2. スレッド(thread)構成と占有率(Occupancy)
123
GPUのメモリ階層
アクセスが速い
SM
SMEM
Threads
L1
Read
TEX
only
L2 cache
DRAM
アクセスが遅い
124
GPU上のメモリ・キャッシュ・レジスタ
GPU内部の記憶域
Cache
Global Memory
GPU上のDRAM。すべてのSMからアクセス
可能。
Local Memory
Threadスコープのメモリ。GPU上のDRAM。 L1(Keplerのみ)、L2
スレッド内部の配列、レジスタスピル時の
記憶域。
L2
Shared Memory
SM内部のメモリ。Blockスコープでアクセス。なし。手動管理のキャッ
シュとして用いる場合あ
低レイテンシのRead/Write。
スレッド間のデータ共有
り。
Texture Memory
テクスチャユニット経由でアクセスするメ
モリ。
L1(Texture)、L2
Read-only Data Cache Read OnlyでアクセスできるGlobal Memory。L1(Texture)、L2
Constant Memory
Registers
定数を収めるメモリ。ブロードキャストア
クセスに特化。
SM内部のキャッシュ
SM内部のレジスタ。演算可能。
なし
125
READ-ONLY(RO) CACHE
SM
TEX
Threads
Texture API
SMEM
L1
L2 cache
Read
TEX
only
CUDA Arrays
一般的なRead-Onlyキャッシュ
として使用可能
Kepler以降
コンパイラに指示
DRAM
126
12
RO DATA CACHE 使い方
• 型修飾子: const __restrict__ を付ける
__global__ kernel( int* output, const int* __restrict__
input )
input )
{
...
output[idx] = ... + input[idx + delta] + ...;
...
}
127
GLOBAL MEMORY
SM
SMEM
• GPU上のメモリの中で最もポピュラーなメモリ
• メモリサイズは大きく
• アクセスコストは高い
Threads
L1
Read
TEX
only
L2 cache
Global
DRAM
Memory
128
コアレスアクセス
連続するスレッドは連続するメモリアクセスになるように
デバイスメモリは32byte,64byte,128byteの単位でロード・ストア
※CUDA7.0現在
thread0
thread1
thread2
thread3
thread4
160
128
Device Memory
thread5
・・・・・・
192
129
コアレスアクセス
連続するスレッドは連続するメモリアクセスになるように
デバイスメモリは32byte,64byte,128byteの単位でロード・ストア
※CUDA7.0現在
thread0
thread1
thread2
thread3
thread4
160
128
Device Memory
thread5
・・・・・・
192
130
コアレスアクセス
連続するスレッドは連続するメモリアクセスになるように
デバイスメモリは32byte,64byte,128byteの単位でロード・ストア
※CUDA7.0現在
thread0
thread1
thread2
thread3
thread4
160
128
Device Memory
thread5
・・・・・・
192
131
パディングを考慮したメモリの確保
x方向の先頭アドレスが32byteの倍数になるようにパディング
例: RGB 24byte
height
padding = 32 – (3*width%32)
width
padding
132
2次元メモリ確保・転送API
cudaMallocPitch
widthバイトのメモリを、height行分、取得する
行は、パディングを考慮したpitchバイトで整列する
cudaMemcpy2D
cudaMallocPitchで取得したパディングを考慮したメモリ(Dst)に、
Srcのメモリ(パディングなし)をコピーする
133
サンプルコード
uchar4 *src, *dImage;
size_t spitch, dPitch, dPitchInPixel;
// ピッチつきで、メモリをアロケート
cudaMallocPitch(&dImage, &dPitch, width * sizeof(uchar4), height);
dPitchInPixel = dPitch / sizeof(uchar4);
// ピッチを変換しつつ、ホスト→デバイスへと、メモリ転送
cudaMemcpy2D(dImage, dPitch, src, sPitch, width * sizeof(uchar4),
height,cudaMemcpyHostToDevice);
134
最適化の為に理解する事
1. GPUのメモリ構造
2. スレッド(thread)構成と占有率(Occupancy)
135
復習:一次元配列の加算
__global__ void MatrixAdd(float *A, const float *B,const float *C)
{
グローバルID
int i = threadIdx.x + blodkDim.x * blockIdx.x;
if ( i >= N || j >=N ) return;
C[i][i] = A[i][j] + B[i][j];
}
総スレッド数
1ブロックあたりのスレッド数
...
MatrixAdd<<< N/128, 128>>>(A, B, C);
...
136
復習:ブロックIDとスレッドID
ブロックIDとスレッドIDから、インデックス(グローバルID)を生成する
インデックスを用いて各スレッドから、グローバルメモリへアクセスする
index = blockDim.x * blockIdx.x + threadIdx.x;
×
+
2
1
6
8
Thread
Block
0
1
2
3
4
5
6
7
8
9 10 11
0
1
2
3
4
5
0
1
2
3
0
11
4
5
12 13 14 15 16 17
0
1
2
3
4
5
2
137
ブロックIDとスレッドID(二次元)
BLOCK * blockIdx.x + threadIdx.x;
BLOCK =BLOCK
index_x
blockDim.x
・・・・・・ (M,N-2) (M,N-1) (M,N)
Index_y = blockDim.y * blockIdx.y + threadIdx.y;
BLOCK
(1,1)
BLOCK
(0,0)
BLOCK BLOCK
(0,1)
(0,2)
thread thread
(1,0) (1,1)
thread
thread BLOCK
(0,1)
・・・・・・(0,0) (0,N)
thread
(15,15)
・・・・・・
・・・
BLOCK
(1,0)
・・・・・・
thread BLOCK
・・・・・・
(15,0) (M-1,N)
・・・・・・
BLOCK
(M,0)
thread
・・・・ (0,15)
138
二次元配列の加算
__global__ void MatrixAdd(float A[N][N], float *B[N][N], float *C[N][N])
{
int i = threadIdx.x + blodkDim.x * blockIdx.x;
int j = threadIdx.y + blodkDim.y * blockIdy.y;
if ( i >= N || j >=N ) return;
C[i][i] = A[i][j] + B[i][j];
}
1ブロックあたり16*16=256スレッド
...
dim3 sizeBlock( 16, 16 );
dim3 numBlocks( N/sizeBlock.x, N/sizeBlock.y );
MatrixAdd<<< numBlocks, sizeBlock >>>(A, B, C);
...
139
例:RGB->YUV変換を考える
1スレッドで1pixelぶんの処理を行う
ピクセルの数だけスレッドを作成
例) 1920*1080 = 2,073,600 スレッド
3840*2160 = 8,294,400 スレッド
140
例:RGB->YUV変換を考える
・・・・・・
thread7
thread6
thread5
thread4
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
Height
if ((x < w) && (y < h)) {
//Global Memory(Src)から4byte ロード
uchar4 uRGB = gSrc[index];
//Global Memory(Dst)へ変換後の値を4byteストア
gDst[idx] = RGB2YUV(uRGB.x, uRGB.y, uRGB.z);
thread3
thread2
}
thread1
thread0
Width
141
ブロックサイズの決定
x = BlockDim.x * BlockIdx.x + threadIdx.x
(0<= x < width)
y = BlockDim.x * BlockIdx.x + threadIdx.x
(0<= y < height)
グリッド・ブロックサイズの例 )
• 960 threads / block
• 128 threads / block
• 32 threads / block
?
height
width
142
ブロックサイズの決定
占有率を100%にする
ブロックサイズ(ブロック辺りのスレッドの数)は少ない方が良い
ブロックは横長の方が良い
143
占有率(OCCUPANCY)とは?
• マルチプロセッサで同時に実行されるワープの数を同時に実行でき
るワープの最大数で除算したもの
144
BLOCKDIMの決定 (占有率から)
項目
値
最大のBlock数 / SMX
16
最大のThread数 / SMX
2048
最大のThread数 / Block
1024
SMXあたり、2048 Thread走らせたい。
Occupancy (占有率) = 100 %
Occupancy = 100 % を満たす、Blockあたりのスレッド数は、
2048 Thread / 16 Block = 128 Thread / Block
2048 Thread / 8 Block = 256 Thread / Block
2048 Thread / 4 Block = 512 Thread / Block
2048 Thread / 2 Block = 1024 Thread / Block
145
BLOCKDIMの決め方(BLOCKの粒度から)
Grid = 4096 Thread の実行例を考えてみる
Block : 256 Thread、1024 Threadで比較
3 SMX / GPU、1 SMXあたり 1 Blockが実行可能とする
SMX 0
SMX 1
SMX 2
Block
Block
Block
Block
Block
Block
Block
Block
Block
Block
Block
Block
Block
Block
Block
Block
t
256 Thread / Block
SMX 0
SMX 1
SMX 2
Block
Block
Block
Block
t
1024 Thread / Block
Blockサイズは小さいほうが得 → 128 Threads / Block
146
復習:カーネル実行の流れ
•
Giga Thread EngineがブロックをSMに割り当てる
Grid
Block0
Giga Thread Engine
Block1
Block2
Block3
Block4
…
BlockN
Block4
148
復習:カーネル実行の流れ
•
SMの中のスケジューラがコアにスレッドを投入する
Grid
Block 0
Thread
Thread
Thread
Thread
32スレッド単位で
ワープを投入
投入
Thread
Thread
Thread
Thread
Thread
Thread
Thread
Thread
Thread
Thread
Thread
Thread
…
BlockN
149
ワープ(WARP) : 並列実行の最少単位
-ワープ(Warp) : 32 GPU スレッド
1命令を Warp (32スレッド)が、並列に処理
SIMT (Single Instruction Multiple Thread)
Thread
Core
Core
Core
Core
CUDA cores
Thread
Thread
Thread
Thread
Thread
SMX
…
32 GPU Thread
Warp
…
Warp
Block
Warp
SW
1命令を
32並列実行
Core
150
BLOCKDIMの決め方 (SMXの構造から)
Warp Scheduler x 4 :
1 clockあたり、4 Warpに対する命令発行
Blockのサイズは、 128 Thread の倍数が望ましい。
(128 Thread = 32 Thread/Warp x 4 Warp)
152
タイルは横長がよい
タイルの横幅は、32(Warpの幅)の倍数がよい。
32より小さい場合、16、もしくは、8 を使う。
Thread :
0 1 2 3 4 5 6 7
8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31
Memory :
threadIdx.x
153
RGB→Y変換時のバンド幅 : TESLA K20C
blockDim.y
1
1
2
4
8
16
32
64
128
256
512
1024
1.4
2.6
4.8
8.4
13.4
17.7
20.7
20.7
20.7
20.5
19.1
2
2.8
5.2
9.6
16.7
26.3
40.4
41.7
41.6
41.0
37.6
-
4
8
blockDim.x
16
32
64
5.6 11.2 22.1
Occupancy
10.4 20.6 40.7
19.2 37.8 74.0
33.3 69.6 115.0
60.6 106.7 115.0
81.1 103.9 110.9
79.8 99.0 83.5
75.6 75.3
60.3
-
blockDim.x < 8
43.9
<77.9
100
119.4
117.9
114.3
86.9
-
128
256
512 1024
78.5 119.8 119.3 115.4 87.7
%119.8 119.4 115.3 87.4
118.2 114.2 87.3
111.9 87.1
87.2
-値: バンド幅
(GB/sec)
(ECC
off)
- Tesla K20c
154
RGB->YUV変換(ホスト)
/* value、radixで割って、切り上げる */
int divRoundUp(int value, int radix) {
return (value + radix – 1) / radix;
}
/* gridDim, blockDimを、２次元(x, y方向)に初期化 */
dim3 blockDim(128, 1);
/* divRoundUp()は、切り上げの割り算 */
dim3 gridDim(divRoundUp(width, blockDim.x),divRoundUp(height, blockDim.y));
RGB2YUV<<<gridDim, blockDim>>>(dDst, dSrc, …);
155
RGB->YUV変換(デバイス)
__device__ inline uchar4
rgb_2_yuv(unsigned char R, unsigned char G, unsigned char B){
float fY,fU,fY; unsigned char uY,uU,uV;
fY = 0.299f * value.x + 0.587f * value.y + 0.114f * value.z;
uY = (unsigned char)min(255, (int)Y);
… U と Y の処理は省略 …
make_uchar4(uY, uU, uV, 0);
}
__global__ void RGB2YUV (uchar4 *gDst, const uchar4 *gSrc, int w, int h){
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
if ((x < w) && (y < h)) {
int index = y * width + x;
//Global Memory(Src)から4byte ロード
uchar4 uRGB = gSrc[index];
//Global Memory(Dst)へ変換後の値を4byteストア
gDst[idx] = rgb_2_yuv(uRGB.x, uRGB.y, uRGB.z);
}
156
まとめ
• グローバルメモリはコアレスアクセスする
•
二次元の場合はcudaMallocPitchを使う事でメモリアライメントを考慮
したメモリ確保が可能
• メモリのLoadのみの場合はRead Only Data Cacheを活用
• 占有率(Occupancy)とBlock内のスレッド構成を意識
•
Blockサイズは、128が適当 (単純なカーネルの場合)
•
Blockの横幅は、32の倍数。無理な場合、16, 8を選択。
(4 byte / pixelの場合)
157
158
159
160
Appendix.
CUDAダウンロードサイト
https://developer.nvidia.com/cuda-toolkit
OpenACC toolkit
https://developer.nvidia.com/openacc
OpenACCオンライン講座
http://info.nvidianews.com/GettingStartedwithPGIOpenACCCompiler_RegLanding-Page.html
GPUコンピューティング Facebookページ
https://www.facebook.com/NVIDIAGPUComputing
161
Thankyou
162
Thank you
173

Download Report