Images per Second Images per Second VOLTA: ディープラーニングにおける大きな飛躍 ResNet-50 トレーニング 2.4x faster ResNet-50 推論 TensorRT - 7ms レイテンシ 3.7x faster P100 V100 P10

Size: px

Start display at page:

Download "Images per Second Images per Second VOLTA: ディープラーニングにおける大きな飛躍 ResNet-50 トレーニング 2.4x faster ResNet-50 推論 TensorRT - 7ms レイテンシ 3.7x faster P100 V100 P10"

さゆりみのしま
8 years ago
Views:

1 NVIDIA TESLA V100 CUDA 9 のご紹介森野慎也, シニアソリューションアーキテクト (GPU-Computing) NVIDIA

2 Images per Second Images per Second VOLTA: ディープラーニングにおける大きな飛躍 ResNet-50 トレーニング 2.4x faster ResNet-50 推論 TensorRT - 7ms レイテンシ 3.7x faster P100 V100 P100 V100 FP32 Tensor コア FP16 Tensor コア V100 measured on pre-production hardware.

E5-2690 v4, 2.6GHz, w/ 1X Tesla P100 or V100.

3 Relative to Tesla P100 EXASCALE への歩み US で最も強力なスーパーコンピュータを Volta の演算性能で実現 Volta HPC Application Performance System Config Info: 2X Xeon E v4, 2.6GHz, w/ 1X Tesla P100 or V100. V100 measured on pre-production hardware. Summit Supercomputer 200+ PetaFlops ~3,400 Nodes 10 Megawatts

4 TESLA V100 のご紹介 Volta アーキテクチャ改善された NVLink と HBM2 Volta MPS 改善された SIMT モデル Tensor コア最も生産性の高い GPU 広帯域バンド幅推論での活用新しいアルゴリズムプログラマブルなディープラーニング演算エンジンディープラーニングと HPC における最も高速で最も生産性の高い GPU

5 *full GV100 chip contains 84 SMs TESLA V100 21B transistors 815 mm 2 80 SM 5120 CUDA Cores 640 Tensor Cores 16 GB HBM2 900 GB/s HBM2 300 GB/s NVLink

6 PASCAL / VOLTA GPU の性能比較 P100 V100 Ratio トレーニングの高速化 10 TOPS 120 TOPS 12x 推論の高速化 21 TFLOPS 120 TOPS 6x FP64/FP32 5/10 TFLOPS 7.5/15 TFLOPS 1.5x HBM2 バンド幅 720 GB/s 900 GB/s 1.2x NVLink バンド幅 160 GB/s 300 GB/s 1.9x L2 Cache 4 MB 6 MB 1.5x L1 Caches 1.3 MB 10 MB 7.7x

7 STREAM: Triad- Delivered GB/s 新しい HBM2 メモリアーキテクチャ 1.5x のバンド幅を実現 V100 measured on pre-production hardware. HBM2 stack P100 76% のメモリバンド幅効率 V100 95% のメモリバンド幅効率

8 VOLTA NVLINK 300GB/sec リンク数 : 4 本 6 本伝送速度 ( 片方向 ) : GB/sec

9 刷新された SM マイクロアーキテクチャ

10 VOLTA GV100 SM GV100 FP32 units 64 FP64 units 32 INT32 units 64 Tensor Cores 8 Register File Unified L1/Shared memory 256 KB 128 KB Active Threads 2048

11 VOLTA GV100 SM 生産性のために刷新された設計大容量高速な L1 キャッシュテンソル演算の加速化完全に新しいISA スケジューラを倍増簡素化された命令発行ロジック改善された SIMT モデル = もっとも容易にプログラミングできるSM

12 容易なプログラミングでディープラーニングを推進する CUDA で実装されたディープラーニングの手法ディープラーニングの手法 AlexNet Batch Normalization NCCL WinoGrad FlowNet Sparsely Gated Mixture of Experts Billion-Scale Similarity Search (FAISS) bit SGD FFT Convolutions cudnn Persistent RNNs Phased LSTM? 新しいソルバー新しいレイヤー新しいスケーリング手法古い手法を用いた新しいアプリケーションなどなど

13 再録 : PASCAL の L1 とシェアードメモリ Pascal SM Load/Store ユニットシェアードメモリ 64 KB 低レイテンシ L1$ 24 KB ストリーミング : その場で必要になるデータをキャッシュキャッシュミスを許す L2$ 4 MB

14 UNIFYING KEY TECHNOLOGIES Pascal SM Load/Store ユニット Volta SM Load/Store ユニットシェアードメモリ 64 KB 低いレイテンシ L1$ とシェアードメモリ 128 KB L1$ 24 KB ストリーミング L2$ 4 MB L2$ 6 MB

15 VOLTA L1 AND シェアードメモリ Volta ストリーミング L1$ : SM Load/Store ユニット処理中のキャッシュミスを許すキャッシュヒット時のレイテンシが低い 4 倍のバンド幅 5 倍の容量 L1$ and シェアードメモリ 128 KB Volta シェアードメモリ : L1 キャッシュと統合された記憶域最大 96 KB まで構成可能 L2$ 6 MB

16 小さくなったシェアードメモリとの性能差 with the GV100 L1 cache Cache: vs shared Directed testing: shared in global 簡単に使用できる 90% 以上の場合で十分な性能 Shared: vs cache より高速なアトミクスより多くのバンク性能を予測しやすいシェアードメモリによる恩恵 ( 平均 ) 70% Pascal 93% Volta

17 VOLTA TENSOR コア

18 TENSOR コア混合精度行列演算 4x4 行列 A 0,0 A 0,1 A 0,2 A 0,3 B 0,0 B 0,1 B 0,2 B 0,3 C 0,0 C 0,1 C 0,2 C 0,3 D = A 1,0 A 1,1 A 1,2 A 1,3 B 1,0 B 1,1 B 1,2 B 1,3 C 1,0 C 1,1 C 1,2 C 1,3 A 2,0 A 2,1 A 2,2 A 2,3 B 2,0 B 2,1 B 2,2 B 2,3 C 2,0 C 2,1 C 2,2 C 2,3 A 3,0 A 3,1 A 3,2 A 3,3 B 3,0 B 3,1 B 3,2 B 3,3 C 3,0 C 3,1 C 3,2 C 3,3 FP16 or FP32 FP16 FP16 FP16 or FP32 D = AB + C

19 VOLTA TENSOR OPERATION FP16 入力積 (Full precision) FP32 加算器で加算結果を FP32 に変換他の積算結果 F16 F16 + F32 F32 推論のために FP16 加算モードもサポート

$USING TENSOR CORES device void tensor_op_16_16_16( float *d, half *a, half *b, float *c) { wmma::fragment<matrix_a, >$ wmma::load_matrix_sync(bmat, b, 16); wmma::fill_fragment(cmat, 0.

wmma::load_matrix_sync(bmat, b, 16); wmma::fill_fragment(cmat, 0.

20 USING TENSOR CORES device void tensor_op_16_16_16( float *d, half *a, half *b, float *c) { wmma::fragment<matrix_a, > Amat; wmma::fragment<matrix_b, > Bmat; wmma::fragment<matrix_c, > Cmat; wmma::load_matrix_sync(amat, a, 16); wmma::load_matrix_sync(bmat, b, 16); wmma::fill_fragment(cmat, 0.0f); wmma::mma_sync(cmat, Amat, Bmat, Cmat); NVIDIA cudnn, cublas, TensorRT Volta に最適化されたフレームワークとライブラリ } wmma::store_matrix_sync(d, Cmat, 16, wmma::row_major); CUDA C++ Warp レベル行列演算

21 Relative Performance A GIANT LEAP FOR DEEP LEARNING 9.3x faster cublas 混合精度演算 ( 入力 FP16, 演算結果 FP32) 行列積 (M=N=K=2048) P100 (CUDA 8) (CUDA 9) V100 Tensor コア V100 measured on pre-production hardware.

22 マルチプロセス実行時のスケジューリング

23 GPU 上のマルチプロセススケジューリング背景 A B C CPU プロセススケジューリングはタイムスライス単一プロセス実行の場合スループットが最適化されるマルチプロセスサービス (MPS) マルチプロセスのスループットが最適化される

24 マルチプロセスの実行はタイムスライス A B C CPU Processes A GPU Execution Pascal GP100 Timeslice 1

25 マルチプロセスの実行はタイムスライス A B C A B C CPU Processes A B GPU Execution Pascal GP100 Pascal GP100 Timeslice 2

26 マルチプロセスの実行はタイムスライス CPU Processes A B C A B C A B C A Pascal GP100 B Pascal GP100 Pascal GP100 C Timeslice 3 GPU Execution

27 マルチプロセスの実行はタイムスライス CPU Processes A B C A B C A B C A B C Pascal GP100 Pascal GP100 Pascal GP100 Timeslice 1 Timeslice 2 Timeslice 3 GPU Execution プロセスがアイソレートされているそれぞれのプロセスで最高性能を発揮することができる

28 PASCAL マルチプロセスサービス A B C CUDA マルチプロセスサービス : 小さなジョブ間で演算リソースを共有し GPU の利用率を改善処理の実行依頼 CUDA MULTI-PROCESS SERVICE CPU Processes GPU Execution プロセス間では隔離されていない A B C Pascal GP100 Opt-in: プロセス間の隔離は限定的プロセスを束ねることでピークスループットを実現

29 VOLTA マルチプロセスサービス Volta における改善 : ハードウエアで高速化された処理の実行依頼 A B C CUDA MULTI-PROCESS SERVICE CONTROL CPU Processes 処理実行のレイテンシを削減 GPU Execution 処理実行のスループットを改善 VOLTA MULTI-PROCESS SERVICE プロセス単位でのスケジューリングより安定したパフォーマンス Pascal に比べ 3 倍のクライアントハードウエアレベルのプロセス隔離 A B C Volta GV100

30 Resnet50 Images/sec, 7ms latency VOLTA MPS の推論時性能バッチ処理を行わなくとも高速な推論 60% of perf with batching 7x faster V100 measured on pre-production hardware. Single Volta Client, No Batching, No MPS Multiple Volta Clients, No Batching, Using MPS Volta with Batching System

31 TESLA V100 のご紹介 Volta アーキテクチャ改善された NVLink と HBM2 Volta MPS 改善された SIMT モデル Tensor コア最も生産性の高い GPU 広帯域バンド幅推論での活用新しいアルゴリズムプログラマブルなディープラーニング演算エンジン他の V100 の新機能 : 2x L2 アトミクス, int8, 新しいメモリモデル, コピーエンジンページマイグレーション, などなどディープラーニングと HPC における最も高速で最も生産性の高い GPU

VOLTA ARCHITECTURE DEEP DIVE 成瀬彰, シニアデベロッパーテクノロジーエンジニア, 2017/12/12

VOLTA ARCHITECTURE DEEP DIVE 成瀬彰, シニアデベロッパーテクノロジーエンジニア, 2017/12/12 TESLA V100 の概要 Volta Architecture Improved NVLink & HBM2 Volta MPS Improved SIMT Model Tensor Core Most Productive GPU Efficient Bandwidth