SM - GTC Japan 2016

Jetson TX1 開発キットで始める
CUDAプログラミング
菱洋エレクトロ CUDAエンジニア石井琢人
1
イントロダクション
Agenda
GPUアーキテクチャとCUDAプログラミングモデル
CUDAプログラミングの基礎
ライブラリを用いたCUDA画像処理
2
イントロダクション
3
菱洋エレクトロについて
1961年設立
半導体・システム情報機器・組み込み製品の
販売及びサポートを行う
代表取締役会長：小川贒八郎
資本金：136億7,200万円
従業員：600人
本社：東京都中央区築地
4
GAMING
DESIGN
ENTERPRISE
VIRTUALIZATION
PC
DATA CENTER
GeForce | Quadro
Tesla | GRID
HPC & CLOUD
SERVICE PROVIDERS
AUTONOMOUS
MACHINES
MOBILE
Tegra | SHIELD
THE WORLD LEADER IN VISUAL COMPUTING
5
NVIDIA GPU ロードマップ
72
Volta
60
SGEMM / W
48
Pascal
36
24
Maxwell
12
Kepler
Fermi
Tesla
0
2008
2010
2012
2014
2016
2018
6
CUDA / GPUコンピューティング
 CUDA
– Compute Unified Device Architecture
– Linux/Windows/MacOS X (+Android)で動作
– CUDA Toolkit 最新バージョンは 8.0
 GPUコンピューティング
– GPUによる汎用コンピューティング
– いわゆるGPGPUとほぼ同義
– GPU = Graphics Processing Unit
7
CUDA開発者数の推移
6.4億
3700万
CUDA GPUs
CUDA GPUs
350万
15万
CUDA Downloads
CUDA Downloads
… 77,000
Supercomputing
77
Supercomputing
Teraflops
Teraflops
950
60
Universities
Teaching
Universities
Teaching
77,500
4,000
Academic Papers
Academic Papers
2008
2016
8
Tesla GPU Accelerators Family
World’s Fastest and Most Efficient HPC Accelerators
GPUs
Single Precision
Peak (SGEMM)
Double Precision
Peak (DGEMM)
Memory
Size
Memory
Bandwidth
(ECC off)
PCIe Gen
System
Solution
K80
5.6 TF
1.87 TF
24 GB
480 GB/s
Gen 3
Server
K40
4.29 TF
(3.22TF)
1.43 TF
(1.33 TF)
12 GB
288 GB/s
Gen 3
Server +
Workstation
M40
6.8 TF
0.2 TF
12 / 24
GB
288 GB/s
Gen 3
Server
P100
9.3 TF
4.7 T
12 / 16
GB
720 / 540
GB/s
Gen 3
Server
P40
11.76 TF
0.37 TF
24 GB
346 GB/s
Gen 3
Server
9
Quadro Professional Graphics Family
CUDA Cores
M6000 K6000 M5000
M4000
M2000
K620
K420
3072
2880
2048
CC
5.2
3.5
5.2
5.2
5.2
5.0
3.0
Single Precision
6.8 TFLOPs
5.2 TFLOPs
4.3 TFLOPs
2.6 TFLOPs
1.8 TFLOPs
0.8 TFLOPs
0.3 TFLOPs
PCIe Gen
1664
768
384
3.0
192
2.0
Memory Size
24 GB
12 GB
8 GB
8 GB
4 GB
2 GB
1 GB
Memory BW
317 GB/s
288 GB/s
211 GB/s
192 GB/s
106 GB/s
29 GB/s
29 GB/s
4x DP*
4x DP*
Slots + Display
Connectors
4x DP* + 1x DVI
2x DP* + 2x DVI
Max Resolution
Max Displays
4096 x 2160
4
Pro Features
Board Power
4x DP* + 1x DVI
4
4
225 W
150 W
* DisplayPort 1.2 multi-streaming can be used to drive multiple displays from a single DP connector
DP* + DVI
3840 x 2160
4
4
SDI, SYNC, STEREO, MOSAIC, NVIEW
250W
DP* + DVI
4
4
MOSAIC, NVIEW
120 W
75 W
41 W
41 W
10
Quadro GPUs (Compute Capability)
P6000
P5000
CC
6.1
6.1
Single Precision
12 TFLOPs
8.9 TFLOPs
CUDA Cores
3840
PCIe Gen
2560
3.0
Memory Size
24 GB
16 GB
Memory BW
433 GB/s
288 GB/s
Slots + Display
Connectors
4x DP* + 1x DVI
4x DP* + 1x DVI
Max Resolution
Max Displays
Pro Features
Board Power
7680 x 4320
4
4
SLI, SYNC, STEREO, MOSAIC,
NVIEW
250W
225 W
* DisplayPort 1.2 multi-streaming can be used to drive multiple displays from a single DP connector
11
TEGRA X1
256-Core Mobile Super Chip
BRIDGING THE GAP
Maxwell
Kepler
Advancements
Fermi
Tesla
Tegra X1
Tegra K1
GEFORCE ARCHITECTURE
Tegra 4
Tegra 3
MOBILE ARCHITECTURE
13
NVIDIA TEGRA X1
- The New Level of Mobile Performance
MAXWELL GPU, 256 CORES,
500 GFLOPS
CUDA
25.6 GB/S BANDWIDTH
VIDEO IMAGE COMPOSITOR (VIC)
14
プロセッサと製品シリーズ
Tesla
共通のGPUアーキテクチャから各製品
シリーズに合わせてコア数を調整
全製品で同じプログラムが動作する
Quadro
用途に応じて製品を選定
GeForce
Tegra
15
GPUアーキテクチャと CUDAプログラミングモデル
16
GPUのハードウエア構造について
GPU Diagram
内容
Compute Capability
CUDAのプログラミングモデルについて
Grid, Block, Warp, Thread
17
GPUの構造
Tegra X1
Giga Thread Engine
GPU
Giga Thread Engine
SM
SM
L2 Cache
Memory Interface
処理をSMに割り振る
SM (Streaming Multiprocessor)
並列プロセッサ
L2 Cache
R/W可能な二次キャッシュ
Memory Interface
CPUと物理的に共有するメモリへのアクセス
18
GPUの構造
一般的なGPUカード
Giga Thread Engine
GPU
処理をSMに割り振る
PCI Express
Giga Thread Engine
SM
SM
SM
SM
SM (Streaming Multiprocessor)
…
並列プロセッサ
L2 Cache
R/W可能な二次キャッシュ
L2 Cache
DRAM
DRAM
すべてのSMとPCI Expressから
アクセス可能なメモリ
PCI Express
PC(ホスト)との接続インターフェース
19
Tegra X1のGPU構造
20
Streaming Multiprocessor (SM)
GPU内部の並列プロセッサ
SMの単位で処理を実行
CUDA coreは単体で動作しない
Maxwell : 128 CUDA cores / SMM
21
SMアーキテクチャとCompute Capability
SM
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
22
Compute Capability
GPUコアアーキテクチャのバージョン
CUDA GPUs : https://developer.nvidia.com/cuda-gpus
アーキテクチャは進化する
より高効率の命令実行
省消費電力
23
Compute Capability
Architecture
Compute
Capability
3.0
Kepler
Maxwell
Pascal
3.2
特徴・機能
発表
192 cores/SMX, 1536 cores (Max)
2012/3
Tegra K1
2013/1
Single / Double = 24 : 1
192 cores/SMX, 2880 cores (Max)
3.5
Single / Double = 3 : 1
Hyper-Q, Dynamic Parallelism
2012/11
3.7
Tesla K80
2014/11
5.0
128 cores/SMM
2014/2
Single / Double = 32 : 1
5.2
96 KB shared memory
2014/9
5.3
Tegra X1, FP16 support
2015/1
Tesla P100, HBM2, 64 cores/SM
2016/4
128 cores/SMM
2016/5
6.0
6.1
Single / Double = 2 : 1
Single / Double = 32 : 1
24
Tesla GPUs (Compute Capability)
デバイス名
コア数
ピーク演算性能
単精度/倍精度(FLOPS)
メモリバンド幅
GB/sec
消費電力
W
2880
4.29 T / 1.43 T
288
235
2496 x 2
5.6 T / 1.87 T
480
300
3072
6.8 T / 0.2 T
288
250
3584
9.3 T / 4.7 T
720 / 540
250
3840
11.76 T / 0.37 T
346
250
Compute Capability 3.5
Tesla K40
Compute Capability 3.7
Tesla K80
Compute Capability 5.2
Tesla M40
Compute Capability 6.0
Tesla P100
Compute Capability 6.1
Tesla P40
25
Quadro GPUs (Compute Capability)
CUDA Cores
M6000 K6000 M5000
M4000
M2000
K620
K420
3072
2880
2048
CC
5.2
3.5
5.2
5.2
5.2
5.0
3.0
Single Precision
6.8 TFLOPs
5.2 TFLOPs
4.3 TFLOPs
2.6 TFLOPs
1.8 TFLOPs
0.8 TFLOPs
0.3 TFLOPs
PCIe Gen
1664
768
384
3.0
192
2.0
Memory Size
12GB
12 GB
8 GB
8 GB
4 GB
2 GB
1 GB
Memory BW
317 GB/s
288 GB/s
211 GB/s
192 GB/s
106 GB/s
29 GB/s
29 GB/s
Slots + Display
Connectors
4x DP* + 1x DVI
2x DP* + 2x DVI
4x DP* + 1x DVI
Max Resolution
Max Displays
4x DP*
4096 x 2160
4
Pro Features
Board Power
4x DP*
4
4
225 W
150 W
* DisplayPort 1.2 multi-streaming can be used to drive multiple displays from a single DP connector
DP* + DVI
3840 x 2160
4
4
SLI, SYNC, STEREO, MOSAIC, NVIEW
250W
DP* + DVI
4
4
MOSAIC, NVIEW
120 W
75 W
41 W
41 W
26
Quadro GPUs (Compute Capability)
P6000
P5000
CC
6.1
6.1
Single Precision
12 TFLOPs
8.9 TFLOPs
CUDA Cores
3840
PCIe Gen
2560
3.0
Memory Size
24 GB
16 GB
Memory BW
433 GB/s
288 GB/s
Slots + Display
Connectors
4x DP* + 1x DVI
4x DP* + 1x DVI
Max Resolution
Max Displays
Pro Features
Board Power
7680 x 4320
4
4
SLI, SYNC, STEREO, MOSAIC,
NVIEW
250W
225 W
* DisplayPort 1.2 multi-streaming can be used to drive multiple displays from a single DP connector
27
CUDAプログラミングモデル
異なるSMアーキテクチャで同じプログラムが動作
GPUハードウェアとソフトウェアの対応
並列処理のみ実行可能
並列化不可能な逐次処理には向かない
100万スレッド以上での超並列動作 (Massively Parallel)
28
GPUは並列処理部分を担当
並列化可能な
逐次実行部分
ボトルネック
GPU
アプリケーションタスク
CPU
29
Jetson TX1の構成
SOC
GPU
ARM CPU
(4コア)
Giga Thread Engine
制御
SM
L2 Cache
DRAM
CPU側アドレス空間
転送
GPU側アドレス空間
30
一般的な装置構成
CPUにつながった外部演算装置
PC
GPU
CPU
Giga Thread Engine
制御
PCIe
ホスト側DRAM
SM
SM
SM
SM
…
L2 Cache
DRAM
転送
31
典型的な実行例
CPU
プログラム
開始
データ
転送
GPUはCPUからの制御で
動作する
CUDA
カーネル
実行
完了
待ち
GPU
GPUでの演算
データ
転送
入力データはCPU→GPU
へ転送
結果はGPU→CPUへ転送
GPU上に常駐する
プログラムはない
32
カーネル実行の階層
GPU
CPU
データ
転送
Block1
Thread
Thread
Thread
Thread
Block2
Thread
Thread
Thread
Thread
Grid
CPUから呼び出される
Blockにより構成される
Block
Threadにより構成される
Thread数は、Grid内で一定
…
…
CUDA
カーネル
実行依頼
Grid
Block0
Thread
Thread
Thread
Thread
Thread
Block n
Thread
Thread
Thread
Thread
最小の実行単位
33
Block は SM内で実行
プログラミングモデル
Grid
Block0
GPU
Block1
Block2
SM
Block3
SM
Block4
SM
Block5
Block6
…
Block N
SM
※ Blockの実行順序は保証されない。特定のSMへのBlock割り当てはできない
34
Block は SM内で実行
Blockは1つのSMに割り当てられる
複数のSMにまたがらない
（SM内では複数Blockが実行される場合もある）
Block内のThread間でSMのリソースを活用可能
各々のBlockは、独立に、非同期に処理を実行する
実行順序の保証はない
Block間の通信・同期は行わない
（Block内のThread間では同期可能）
35
Streaming Multiprocessor on Maxwell
128 Cores/SMM
Compute Capability 5.3
SMX (簡略化しています)
レジスタファイル
レジスタ 64 K個 (256 KB)
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
SFU
SFU
SFU
SFU
SFU
SFU
SFU
SFU
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
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
SFU
SFU
SFU
SFU
SFU
SFU
SFU
SFU
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
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
SFU
SFU
SFU
SFU
SFU
SFU
SFU
SFU
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
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
SFU
SFU
SFU
SFU
SFU
SFU
SFU
SFU
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
Texture/L1
Cache
48 KB
Shared Mem.
64 KB
http://docs.nvidia.com/cuda/cuda-c-programming-guide/#features-and-technical-specifications
SFU
Special Function Unit
LD/ST
Memory Load/Store
36
Warp
並列実行の最小単位
Warp : 32 GPU スレッド
1命令を Warp (32スレッド)が、並列に処理
SIMT (Single Instruction Multiple Thread)
Thread
Core
Core
Core
Core
CUDA cores
SM
…
Thread
Thread
Thread
Thread
Thread
…
32 GPU Thread
Warp
…
Warp
Block
Warp
SW
1命令を
32並列実行
Core
37
CUDAプログラム実行の概要
SM
CPU
Grid
Block
Warp (32 Thread)
Warp (32 Thread)
Warp (32 Thread)
Block
Block
Grid
CPUからの呼び出し単位
Blockで構成される
Block
SM上の実行単位
Grid内でスレッド数固定
Warp
一命令を32並列で実行
SM
GPUの並列プロセッサ
38
CUDAプログラミングの基礎
39
CUDAプログラミングの基礎
ホストプログラミングとカーネルプログラミング
ホストプログラミング
メモリ転送、カーネルの実行
カーネルプログラミング
GPU上の関数の実装
40
デバイス・メモリ構成
ホスト
GPU (デバイス)
SM(X)
CPU
デバイスメモリ
(グローバルメモリ)
ホストメモリ
CPU側：ホスト
ホストメモリ
ホストコード
GPU側：デバイスグローバルメモリデバイスコード (カーネル)
41
CUDA ホストプログラミング
メモリのアロケーション、解放
cudaMalloc()/cudaFree()
メモリコピー
cudaMemcpy()
カーネルの呼び出し
特殊な構文
同期
cudaDeviceSynchronize()
終了処理
cudaDeviceReset()
42
cudaMalloc() / cudaFree()
グローバルメモリの確保/解放
cudaError_t cudaMalloc(void ∗∗ devPtr, size_t size)
cudaError_t cudaFree(void *);
例：
float *dptr;
/* float型、1024個の要素分のデバイスメモリをアロケート */
cudaMalloc((void**)&dptr, sizeof(float) * 1024);
/* 解放 */
cudaFree(dptr);
43
cudaMemcpy()
srcからdstへcountバイトの要素を転送
cudaError_t
cudaMemcpy (void ∗ dst, const void ∗ src, size_t count,
enum cudaMemcpyKind kind)
例：
float src[1024] = {…..}
float *ddst;
cudaMalloc((void**)&ddst, sizeof(float) * 1024);
cudaMemcpy(ddst, src, sizeof(float) * 1024,
cudaMemcpyHostToDevice);
44
cudaMemcpy()
cudaMemcpyKind – 転送方向の指定
enum cudaMemcpyKind
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice
cudaMemcpyHostToHost
cudaMemcpyDefault
45
カーネル呼び出し
カーネルを呼び出す特殊構文
kernelName<<<GridDim, BlockDim>>>(引数);
GridDim : グリッド中のブロック数
BlockDim : ブロックあたりのスレッド数
引数は、複数個指定可能
例：
sample<<<1, 256>>>(x, y, z);
46
cudaDeviceSynchronize()
カーネル処理が終了するのを待機する同期API
cudaError_t cudaDeviceSynchronize (void)
例：
someKernel<<<xxx, yyy>>>(a, b, c);
cudaDeviceSynchronize();
47
cudaDeviceReset()
アプリケーションの終了時に呼び出す
cudaError_t cudaDeviceReset (void)
役割：
使用したリソースの解放
プロファイリングデータの回収
cuda-memcheckを用いた、デバイスメモリのリークチェック
48
cudaError_t
エラーチェックに使用するCUDAのエラー型
成功時は、cudaSuccessを返す
エラーの場合、内容を確認
const char∗ cudaGetErrorString (cudaError_t error)
エラーを説明する文字列を返す
49
CUDAカーネル
GPUで実行される関数の構文
__global__
void myKernel(int a, float *pb, …) {
/* device code */
}
通常のC/C++の構文が使用可能
50
CUDAカーネル関数の修飾子
__global__ , __device__ , __host__
__global__
ホストから呼び出し可能なデバイス側の関数
戻り値はvoid固定
__device__
デバイスから呼び出し可能なデバイス側の関数
__host__
ホストから呼び出し可能なホスト側の関数
通常のC/C++関数と同義
51
プログラム例
配列の和の計算
c[i] = a[i] + b[i]
メモリの取り扱い
基本的なカーネルの実装
52
配列の和
メモリの扱い
ホスト
float *a, *b, *c をアロケート
GPU
float *da, *db, *dc をアロケート (デバイスメモリ)
*a, *bに値を設定
ホスト->デバイス転送 a-> da, b->db
カーネル実行依頼
カーネル
dc[i] = da[i] + db[i]
ホスト <- デバイス転送 c <- dc
結果表示・検証
float *da, *db, *dc を開放 (デバイスメモリ)
float *a, *b, *c を開放
53
配列の和 : ホストコード
int main() {
static const int size= 256 * 100;
int memSize = sizeof(float) * size;
float *a, *b, *c, *da, *db, *dc; /* ホストもデバイスもメモリは同じポインタ型 */
/* ホスト側メモリの確保と値の初期化(略)*/
/* GPU側メモリをアロケート */
cudaMalloc(&da, memSize); cudaMalloc(&db,memSize); cudaMalloc(&dc, memSize);
cudaMemcpy(da, a, memSize, cudaMemcpyHostToDevice); /* メモリ転送(Host→Device) */
cudaMemcpy(db, b, memSize, cudaMemcpyHostToDevice);
/* カーネル(addArrayKernel)をここで呼ぶ */
}
cudaMemcpy(c, dc, memSize, cudaMemcpyDeviceToHost); /* メモリ転送(Host←Device) */
/* 表示などの処理 (略) */
cudaFree(da); cudaFree(db); cudaFree(dc);
free(a); free(b); free(c);
cudaDeviceReset(); /* 最後に実行 */
54
並列化
カーネル設計
1 スレッドで、一つの要素の和を計算
複数のブロックに分割
1 Blockあたりの最大スレッド数は1024 (図はBlockあたり4スレッドとした)
Thread ID :
0
1
2
3
4
Block[0]
a[i]
b[i]
0
1
2
5
6
7
8
Block[1]
3
4
5
6
10
11
12
Block[2]
7
+ + + +
+ + + +
15 14 13 12
11 10
9
9
8
8
14
15
Block[3]
10 11
12 13 14 15
+ + + +
+ + + +
7
9
13
6
5
4
3
2
0
1
c[i]
55
Global ID
GPUスレッドの通し番号
Global ID、Grid内で一意
blockDim.x * blockIdx.x + threadIdx.x
threadIdx.x
Global ID
0
1
2
3
Thread番号、Block内で一意
blockIdx.x
Block番号、Grid内で一意
blockDim.x
Block内のスレッド数
blockIdx
0
blockIdx
4
5
6
7
1
threadIdx
0
1
2
3
Thread
Thread
Thread
Thread
threadIdx
0
1
2
3
Thread
Thread
Thread
Thread
56
カーネル実装
__global__
void addArrayKernel(float *dc, const float *da, const float *db, int size) {
/* Global IDを算出 */
int globalID = blockDim.x * blockIdx.x + threadIdx.x;
if (globalID < size) { /* 範囲チェック */
/* 自スレッド担当の要素のみ、処理 */
dc[globalID] = da[globalID] + db[globalID];
}
}
57
カーネル呼び出し
並列度の指定
ブロック数でスケールする。ブロックごとのスレッド数は一定
⇒ 切り上げの割り算
/* gridDim * blockDim個のスレッドを起動する */
int blockDim = 256;
int gridDim = (size + blockDim – 1) / blockDim;
addArrayKernel<<<gridDim, blockDim>>>(dc, da, db, size);
58
配列の和
ホストコード
int main() {
static const int size= 256 * 100;
int memSize = sizeof(float) * size;
float *a, *b, *c, *da, *db, *dc; /* ホストもデバイスもメモリは同じポインタ型 */
/* ホスト側メモリの確保と値の初期化(略)*/
/* GPU側メモリをアロケート */
cudaMalloc(&da, memSize); cudaMalloc(&db,memSize); cudaMalloc(&dc, memSize);
cudaMemcpy(da, a, memSize, cudaMemcpyHostToDevice); /* メモリ転送(Host→Device) */
cudaMemcpy(db, b, memSize, cudaMemcpyHostToDevice);
int blockDim = 256; int gridDim = (size + blockDim – 1) / blockDim;
addArrayKernel<<<gridDim, blockDim>>>(dc, da, db, size);
// cudaDeviceSynchronize(); /* 同期。今回は、必須ではない。 */
}
cudaMemcpy(c, dc, memSize, cudaMemcpyDeviceToHost); /* メモリ転送(Host←Device) */
/* 表示などの処理 (略) */
cudaFree(da); cudaFree(db); cudaFree(dc);
free(a); free(b); free(c);
cudaDeviceReset();
59
ライブラリを用いたCUDA画像処理
60
CUDAのライブラリを
用いた画像処理
NVIDIAの提供するライブラリ
VisionWorks
デモ
61
https://developer.nvidia.com/gpu-accelerated-libraries
NVIDIAの提供するライブラリ
GPU-Accelerated Libraries
nvGRAPH
cuBLAS
NPP
cuDNN
TensorRT
Thrust
CUDA Math Library
cuSPARSE
cuFFT
cuRAND
62
NPP
NVIDIA Performance Primitives library
画像・動画・信号処理向けにGPUに最適化された関数群からなるライブラリ
CPUにおけるIntel IPPのような位置づけ
CPUに対して5～10倍程度の処理性能
CUDA Toolkitに付属
63
Image from https://developer.nvidia.com/npp
VisionWorks
VisionWorksの目的
高パフォーマンスでロバストなコンピュータビジョン・プリミティブの提供
Tegra組み込みプラットフォームでのCVアプリケーション開発を容易に
プロダクトサイクルにおけるプロトタイプの作成を素早く
64
VisionWorksの概要
VisionWorks Software Stack
コンピュータビジョン・アプリケーション
VisionWorksサンプルアプリケーション
Feature Tracking, Optical Flow, etc.
VisionWorks
API Extensions
VisionWorks
Framework Extensions
OpenVX API
OpenVX Framework
VisionWorks Toolkit
ユーザー
CUDA Acceleration Framework
Khronos
Tegra K1/X1, Kepler/Maxwell GPU
NVIDIA
65
Portable, Power-efficient Vision Processing
Khronosによって策定されたコンピュータビジョン向けの標準規格
OpenVXの主な目標は，
1. CV向けのプリミティブと画像及びデータフォーマットのサブセットの定義
2. ヘテロジニアス・アーキテクチャにおけるアクセラレーションの実現
3. 異なるアーキテクチャ間におけるポータビリティの提供
66
OpenVX Graphs – The Key to Efficiency
color
convert
処理の流れをグラフとしてまとめることでホストプロセッサ
とのやり取りを減らす
Gaussian
pyramid
メモリ転送も含めて自動的に最適化
アクセラレータ内部の高速なメモリやキャッシュにデータ
を保持できる
pyr -1
pyr 0
pts -1
pts 0
pyrLK optical flow
67
VisionWorks Framework
CV data containers
image, pyramid, array, matrix, scalar, LUT, threshold など
CV Primitives
画像フィルタ，画像の演算や解析，特徴抽出など
User extensibility
ユーザー定義カーネル (CPU/GPU)
68
VisionWorks Primitives
OpenVX
Primitives
NVIDIA
Extensions
69
VisionWorks のインストール
JetPack
JetPackに同梱
ホストマシンにも同時にインストール可能
for WindowsはBeta版が利用可能
https://developer.nvidia.com/embedded/visionworks
70
VisionWorks のインストール
VisionWorks Samples のインストール
サンプルのインストールスクリプトを実行
/usr/share/visionworks/sources/install-samples.sh <install-path>
<install-path>にVisionWorks-1.5-Samplesが展開される
同様にSFMとTrackingのサンプルもそれぞれインストール
/usr/share/visionworks-sfm/sources/install-samples.sh <install-path>
/usr/share/visionworks-tracking/sources/install-samples.sh <install-path>
71
OpenVX API と CUDA API
NVXCU
VisionWorks 1.5からCUDA API (NVXCU) 追加，オブジェクトを透過的に扱えるようになった
NVXCU
OpenVX
アプリケーション
VisionWorks
(context)
Image
Object
vx_graph graph
NVXCU API
user CUDA kernel
Image
Object
参照
参照
Graph
Object
ブロッキングなしにVXプリミティブと
ユーザー定義カーネルを実行可能に
72
プリミティブの実行 (OpenVX API)
イミディエイトモードとグラフモード
イミディエイトモード
OpenCVのように関数呼び出しごとにブロッキングされる使い方
グラフモード
グラフに処理を登録することでブロッキングを減らし最適化の機会を増やす使い方
73
プリミティブの実行 (OpenVX API)
Immediate Mode
OpenCVの関数のような使い方で呼び出しごとにブロッキングされる
プロトタイプとして作る際に便利
プレフィックスとして‘vxu’がつく
// 3x3 box filter
vxuBox3x3( context, src0, tmp );
// 2画像の絶対値差分
vxuAbsDiff( context, tmp, src1, dst );
74
プリミティブの実行 (OpenVX API)
Graph Mode
処理の実行前にワークロードがわかる
動画のストリーム処理に適している
vx_graph graph = vxCreateGraph( context );
// ノードを作ってあらかじめグラフをチェックする
// グラフに問題があればここでエラーを検出
vxBox3x3Node( graph, src0, tmp );
vxAbsDiffNode( graph, tmp, src1, dst );
vxVerifyGraph( graph );
// グラフの処理を実行
vxProcessGraph( graph );
75
データオブジェクトへのアクセス (OpenVX API)
Semi-opaque Objects
vxAccessImagePatch( img, &rect, 0, &addr, &ptr,
VX_READ_AND_WRITE );
// ‘ptr’ を使ってデータへアクセス
VisionWorks
(context)
アプリケーション
vx_uint8 *ptr
vxCommitImagePatch( img, &rect, 0, &addr, ptr );
// ‘ptr’ はすでに無効
vxBox3x3( img, dst );
vx_image img
参照
pixels
データへのポインタは一時的にしか利用できない
VisionWorksがメモリを管理し必要な時にのみ同期
やCPU-GPU間のデータ転送を行う
76
デモ
77
デモ内容
ぼかし
動画ファイル
グレースケール
OpenCVとVisionWorks(NVXCU)で実装
ヒストグラム平均化
repository : https://bitbucket.org/tishii_ry/vx_gtcj_demo
エッジ抽出
78
参考資料・書籍
CUDA C Programming Guide
http://docs.nvidia.com/cuda/index.html#programming-guides
NVIDIAによるCUDAに関する公式リファレンス
Parallel for All
http://devblogs.nvidia.com/parallelforall/
NVIDIAのエンジニアによるGPU関連技術の紹介など
NVIDIA Japan Facebook
https://ja-jp.facebook.com/NVIDIAGPUComputing
NVIDIA Japanのブログ
CUDA C プロフェッショナルプログラミング
Professional CUDA C Programmingの邦訳書
NVIDIA Japanのエンジニアによる監修
各種メモリの特徴やストリームの使い方まで網羅
79
Thank you
80

Download Report