TensorRTをもっと高性能に使う

TOKYO | OCT. 5, 2016
TensorRTをもっと高性能に使う
森野慎也, シニアCUDAエンジニア, ディープラーニング部, エヌビディアジャパン
TensorRT
ディープラーニングにおける最高の推論パフォーマンス
高性能なフレームワークで、
GPU上の推論エンジンを、容易に開発する。
ディープラーニングによる推論を用いた、
製品展開のためのソリューション
トレーニング済みネットワーク、ターゲットGPUに対して、
推論エンジンの性能を最適化
32-bit、16-bit 演算を用いた推論エンジン
Hyperscale, ADAS, Embedded 向け
https://developer.nvidia.com/tensorrt
HTTP (~10ms)
データセンターにおける GPU Inference Engine
画像分類
オブジェクト
検知
音声認識
---
TensorRT
ニューラルネットワークの最適化
• ネットワークレイヤの融合
• Concatenationレイヤの除去
• カーネル特殊化
• ターゲットプラットフォームに対する
オートチューニング
トレーニング済みの
ニューラルネットワーク
https://developer.nvidia.com/tensorrt
• 最適なテンソルレイアウトの選択
• バッチサイズのチューニング
最適化された
推論
エンジン
TensorRT
ワークフロー
DIGITS
NEURAL
NETWORK
developer.nvidia.com/gpu-inference-engine
OPTIMIZATION
ENGINE
PLAN
EXECUTION
ENGINE
TensorRT @ Parallel ForAll
TensorRTそのものについては、
こちらをご参照ください。
簡潔ですが、情報量も多いです。
https://devblogs.nvidia.com/parallelforall/production-deep-learning-nvidia-gpu-inference-engine/
GPU REST ENGINE
https://github.com/NVIDIA/gpu-rest-engine
CaffeをTensorRTに
入れ替えてみよう！
推論サーバの実装。
Go言語を使用した、HTTPサーバ
Caffeを用いた推論
NVIDIA Dockerを使用したコンテナ化
TensorRTとCaffeの推論性能比較
CaffeとTensorRTの比較
実行例 : 画像から対象を認識
[{"confidence":0.9255,"label":"n02328150 Angora,
Angora rabbit"},
{"confidence":0.0466,"label":"n02326432 hare"},
{"confidence":0.0272,"label":"n02325366 wood rabbit,
cottontail, cottontail rabbit"},
{"confidence":0.0001,"label":"n02097298 Scotch
terrier, Scottish terrier, Scottie"},
{"confidence":0.0001,"label":"n02364673 guinea pig,
Cavia cobaya"}]
CaffeとTensorRTの比較
評価条件
使用したモデル、ファイル
-
Caffeに含まれる、GoogleNetモデル。
https://github.com/BVLC/caffe/tree/master/models/bvlc_googlenet
-
ILSVRC12の学習に用いられた、平均値イメージ。
https://github.com/BVLC/caffe/tree/master/data/ilsvrc12
フレームワーク・ライブラリ
-
Caffe : NVIDIA github上のフォークを使用 (v0.15.9)
-
TensorRT RC1 (CUDA 8.0)
基本となる実装
https://github.com/NVIDIA/caffe/tree/caffe-0.15/examples/cpp_classification
推論(INFERENCE)の処理
処理内容の説明
Caffe版
TensorRT版
画像のデコード
cv::decode
cv::decode
前処理:
Float化・リサイズ・平均値を引く
メモリレイアウト変換
cv::cuda モジュール
CUDA
カスタム実装
推論
Caffe
TensorRT
推論結果から
画像のラベルを取得
CPU処理
CPU処理
少しだけ、実装を、見てみる
本日使用したGPU
Quadro M5000
Compute Capability : 5.2 (GM204)
CUDA Cores : 2048
Memory : 8 GB GDDR5 (256 bit)
TensorRTのご利用時には、
Teslaでの運用をお願いいたします。
評価結果
Caffe版
TensorRT版
速度向上率
前処理
3.0 ms
0.8 ms
x 3.7
推論
7.4 ms
2.4 ms
x 3.1
全体
10.4 ms
3.2 ms
x 3.2
前処理も、最低限の処理に整理し、CUDA化することで、3.7倍 高速化。
TensorRTの使用により、推論速度は、3.1倍、向上
TensorRTから、もっと性能を引き出す！
高効率化の指針
コンカレントな実行
- 複数の推論処理を、並行に実行することで、GPUの使用率をあげる。
バッチサイズの効果
バッチサイズを大きくとる
- 一般に、バッチサイズが大きい方が、演算効率が高い。
コンカレントに実行
cudaStreamを利用する
Thread 1
Thread 2
Thread 3
CUDA Stream 1
CUDA Stream 2
output1
decode1
decode4
output2
decode2
decode5
output3
decode3
infer1
infer3
infer2
スレッドの数と、CUDA Streamの数が、一致しなくとも構わない。
dec…
Infer…
Infer..
ストリームの設定
nvinfer1::IExecutionContext *ctx = …;
cudaStream_t stream;
cudaStreamCreate(&stream);
/* 略 */
ctx->enqueue (batchSize, d_inputs, stream, NULL);
cudaStreamSynchronize(stream);
コンカレントに実行
性能計測結果
5.00
700
同じスレッド数、ストリーム数で実行。
600
8スレッド:8ストリームで飽和。
推論時間 [ms]
500
3.00
400
300
2.00
200
1.00
100
0.00
0
1:1
2:2
4:4
8:8
16:16
スレッド数 : ストリーム数
32:32
スループット [1/s]
4.00
約2倍の性能改善
プロファイラ
シリアル実行時
コンカレント実行時には、複数のカーネルが、
オーバーラップして動作している。
コンカレント実行時
拡大図
カーネルが、オーバーラップして
いる。
推論用のカーネルの並列度は、
低いものがある。
カーネルのオーバーラップ実行が
可能。
複数の推論をコンカレントに実
行することで、効率向上。
WEBサーバと結合して評価
シーケンシャル・コンカレントな推論 スループット比較
推論サーバの
「推論処理スループット」
「処理レイテンシ」
を計測。
推論サーバ
GOの net/http パッケージを使用して実装
推論Server
http:
…
「同時リクエスト数」を変えて、計測。
GOのboomを使用
Localhost経由での計測
(イーサネット通信の影響なし)
Client : boom
同時リクエスト数
1 ~ 32
推論のリクエストに対して、並行実行
推論リクエスト
推論リクエスト
推論リクエスト
推論リクエスト
推論リクエスト
実行待ち
実行待ち
GPU
推論実行
Stream1
推論実行
Stream2
推論実行
Stream3
推論サーバの性能評価
スループット・レイテンシ
シリアルな推論に比べ…
スループット [リクエスト数 / s]
700
シリアル
コンカレント
600
コンカレントに推論を実施することで
約 “650 リクエスト / 秒” を実現。
約 2倍のスループットを達成
500
400
同時リクエスト数 8 で、性能は飽和。
300
レイテンシ(同時リクエスト数8)
200
シリアル
コンカレント
100
0
1
2
4
8
同時(並行)リクエスト数
16
32
22.1 ms
14.2 ms
バッチサイズの効果
一般に、バッチサイズを大きくすることで、
処理性能は向上。
バッチ数が小さい :
メモリアクセス量に対して、
演算量が少ない。効率低い。
TensorRT
(バッチサイズ=1)
TensorRT
(バッチサイズ=4)
バッチ数を増やす :
メモリアクセス量に比して
多くの演算を効率的に実行
rabbit
Rabbit
Dog
Cat
Mouse
高効率
バッチ数の増加による性能改善
Chart Title
3.50
1000
3.00
TensorRTの
推論処理時間のみを計測。
800
2.00
600
1.50
400
1.00
200
0.50
0.00
0
1
2
4
8
16
バッチサイズ
32
64
128
スループット / s
実行時間 [ms]
2.50
バッチサイズごとに、
ネットワークの最適化行う。
バッチサイズの増加により
スループットは、「2.5倍」改善
推論のリクエストに対して、並行実行
複数バッチ数で実行
推論リクエスト
推論リクエスト
実行待ち
推論リクエスト
推論リクエスト
推論リクエスト
推論リクエスト
推論リクエスト
推論リクエスト
推論リクエスト
推論実行
Stream2
推論リクエスト
推論リクエスト
推論リクエスト
推論リクエスト
推論実行
Stream1
推論リクエスト
推論実行
Stream3
推論(バッチ数4)を適用する
推論サーバのスループット
900
シリアル
800
コンカレント
700
バッチ
同時リクエスト数 32にて、
バッチ数4の推論サーバが
高速になる。
スループット [1/s]
600
一方、レイテンシは同程度。
500
400
コンカレント
バッチ(N=4)
300
53.5 ms
49 ms
200
100
0
1
2
4
8
同時リクエスト数
16
32
64
シリアル実行に比して、
2.4倍の性能向上
TensorRTから、もっと性能を引き出す！
まとめ
前処理も最適化する
コンカレントな実行
-
バッチサイズの効果
複数の推論処理を、並行に実行することで、GPUの使用率をあげる。
バッチサイズを大きくとる
- 一般に、GPUでは、バッチサイズが大きい方が、演算効率が高い。
TOKYO | OCT. 5, 2016
THANK YOU SO MUCH !!
JOIN THE CONVERSATION
#GTCJapan