NVIDIA TESLA V100 CUDA 9 のご紹介 森野慎也, シニアソリューションアーキテクト (GPU-Computing) NVIDIA
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.
Relative to Tesla P100 EXASCALE への歩み US で最も強力なスーパーコンピュータを Volta の演算性能で実現 Volta HPC Application Performance System Config Info: 2X Xeon E5-2690 v4, 2.6GHz, w/ 1X Tesla P100 or V100. V100 measured on pre-production hardware. Summit Supercomputer 200+ PetaFlops ~3,400 Nodes 10 Megawatts
TESLA V100 のご紹介 Volta アーキテクチャ 改善された NVLink と HBM2 Volta MPS 改善された SIMT モデル Tensor コア 最も生産性の高い GPU 広帯域バンド幅 推論での活用 新しいアルゴリズム プログラマブルなディープラーニング演算エンジン ディープラーニングと HPC における 最も高速で 最も生産性の高い GPU
*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
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
STREAM: Triad- Delivered GB/s 新しい HBM2 メモリアーキテクチャ 1.5x のバンド幅を実現 V100 measured on pre-production hardware. HBM2 stack P100 76% のメモリバンド幅効率 V100 95% のメモリバンド幅効率
VOLTA NVLINK 300GB/sec リンク数 : 4 本 6 本伝送速度 ( 片方向 ) : 20 25 GB/sec
刷新された SM マイクロアーキテクチャ
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
VOLTA GV100 SM 生産性のために刷新された設計 大容量 高速な L1 キャッシュテンソル演算の加速化完全に新しいISA スケジューラを倍増簡素化された命令発行ロジック改善された SIMT モデル = もっとも容易にプログラミングできるSM
容易なプログラミングでディープラーニングを推進する CUDA で実装されたディープラーニングの手法ディープラーニングの手法 AlexNet Batch Normalization NCCL WinoGrad FlowNet Sparsely Gated Mixture of Experts Billion-Scale Similarity Search (FAISS) 2012 2013 2014 2015 2016 2017 2018 1-bit SGD FFT Convolutions cudnn Persistent RNNs Phased LSTM? 新しいソルバー 新しいレイヤー 新しいスケーリング手法 古い手法を用いた新しいアプリケーション などなど
再録 : PASCAL の L1 とシェアードメモリ Pascal SM Load/Store ユニット シェアードメモリ 64 KB 低レイテンシ L1$ 24 KB ストリーミング : その場で必要になるデータをキャッシュキャッシュミスを許す L2$ 4 MB
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
VOLTA L1 AND シェアードメモリ Volta ストリーミング L1$ : SM Load/Store ユニット 処理中のキャッシュミスを許すキャッシュヒット時のレイテンシが低い 4 倍のバンド幅 5 倍の容量 L1$ and シェアードメモリ 128 KB Volta シェアードメモリ : L1 キャッシュと統合された記憶域最大 96 KB まで構成可能 L2$ 6 MB
小さくなったシェアードメモリとの性能差 with the GV100 L1 cache Cache: vs shared Directed testing: shared in global 簡単に使用できる 90% 以上の場合で十分な性能 Shared: vs cache より高速なアトミクス より多くのバンク 性能を予測しやすい シェアードメモリによる恩恵 ( 平均 ) 70% Pascal 93% Volta
VOLTA TENSOR コア
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
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, > 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 レベル行列演算
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.
マルチプロセス実行時のスケジューリング
GPU 上のマルチプロセススケジューリング 背景 A B C CPU プロセス スケジューリングはタイムスライス 単一プロセス実行の場合スループットが最適化される マルチプロセスサービス (MPS) マルチプロセスのスループットが最適化される
マルチプロセスの実行はタイムスライス A B C CPU Processes A GPU Execution Pascal GP100 Timeslice 1
マルチプロセスの実行はタイムスライス A B C A B C CPU Processes A B GPU Execution Pascal GP100 Pascal GP100 Timeslice 2
マルチプロセスの実行はタイムスライス CPU Processes A B C A B C A B C A Pascal GP100 B Pascal GP100 Pascal GP100 C Timeslice 3 GPU Execution
マルチプロセスの実行はタイムスライス 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 プロセスがアイソレートされている それぞれのプロセスで最高性能を発揮することができる
PASCAL マルチプロセスサービス A B C CUDA マルチプロセスサービス : 小さなジョブ間で 演算リソースを共有し GPU の利用率を改善 処理の実行依頼 CUDA MULTI-PROCESS SERVICE CPU Processes GPU Execution プロセス間では隔離されていない A B C Pascal GP100 Opt-in: プロセス間の隔離は限定的 プロセスを束ねることでピークスループットを実現
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
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
TESLA V100 のご紹介 Volta アーキテクチャ 改善された NVLink と HBM2 Volta MPS 改善された SIMT モデル Tensor コア 最も生産性の高い GPU 広帯域バンド幅 推論での活用 新しいアルゴリズム プログラマブルなディープラーニング演算エンジン 他の V100 の新機能 : 2x L2 アトミクス, int8, 新しいメモリモデル, コピーエンジンページマイグレーション, などなど ディープラーニングと HPC における 最も高速で 最も生産性の高い GPU
www.nvidia.com/dli