VOLTA ARCHITECTURE DEEP DIVE 成瀬彰, シニアデベロッパーテクノロジーエンジニア, 2017/12/12
|
|
|
- いとは よしくに
- 4 years ago
- Views:
Transcription
1 VOLTA ARCHITECTURE DEEP DIVE 成瀬彰, シニアデベロッパーテクノロジーエンジニア, 2017/12/12
2 TESLA V100 の概要 Volta Architecture Improved NVLink & HBM2 Volta MPS Improved SIMT Model Tensor Core Most Productive GPU Efficient Bandwidth Inference Utilization New Algorithms 125 Programmable TFLOPS Deep Learning Deep Learning と HPC 両方に最適な GPU 2
3 Images per Second Images per Second VOLTA DL 性能を大幅に向上 トレーニング 2.4x faster インファレンス TensorRT - 7ms Latency 3.7x faster P100 V100 P100 V100 FP32 Tensor コア FP16 Tensor コア (*) DL モデルは ResNet50 3
4 P100 に対する相対性能 VOLTA HPC 性能を大きく向上 HPC アプリケーション性能 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
5 VOLTA 米国トップスパコンのエンジン SUMMIT AI Exascale Today Performance Leadership 200 PF Accelerated Science ACME DIRAC FLASH GTC HACC LSDALTON NAMD 20 PF NUCCOR NWCHEM QMCPACK RAPTOR SPECFEM XGC 3+EFLOPS Tensor Ops 10X Perf Over Titan 5-10X Application Perf Over Titan 5
6 TESLA V100 トランジスタ数 :21B 815 mm 2 80 SM 5120 CUDA コア 640 Tensor コア HBM2 16 GB, 900 GB/s NVLink 300 GB/s *full GV100 chip contains 84 SMs 6
7 GPU ピーク性能比較 : P100 vs v100 P100 V100 性能 UP トレーニング性能 10 TOPS 125 TOPS 12x インファレンス性能 21 TFLOPS 125 TOPS 6x FP64/FP32 5/10 TFLOPS 7.8/15.6 TFLOPS 1.5x HBM2 バンド幅 720 GB/s 900 GB/s 1.2x NVLink バンド幅 160 GB/s 300 GB/s 1.9x L2 キャッシュ 4 MB 6 MB 1.5x L1 キャッシュ 1.3 MB 10 MB 7.7x 7
8 STREAM: Triad- Delivered GB/s HBM2 メモリ 使用効率 UP 実効バンド幅 1.5 倍 V100 measured on pre-production hardware. HBM2 stack P100 V100 76% 95% 使用効率 8
9 VOLTA NVLINK P100 V100 リンク数 4 6 バンド幅 / リンク 40 GB/s 50 GB/s トータルバンド幅 160 GB/s 300 GB/s (*) バンド幅は双方向 DGX1V 9
10 NEW SM MICROARCHITECTURE 10
11 VOLTA GV100 SM GV100 FP32ユニット 64 FP64ユニット 32 INT32ユニット 64 Tensorコア 8 レジスタファイル 256 KB 統合 L1 共有メモリ 128 KB Activeスレッド 2048 (*) SM あたり 11
12 VOLTA GV100 SM 生産性の向上 命令セットを一新スケジューラを2 倍命令発行機構をシンプルに L1キャッシュの大容量 高速化 SIMTモデルの改善テンソル計算の加速 最もプログラミングの簡単な SM 12
13 VOLTA TENSOR コア 13
14 TENSOR コア混合精度行列計算ユニット 4x4 の行列の積和演算を 1 サイクルで計算する性能 (128 演算 / サクル ) 行列の FMA (Fused Multiply-Add) 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 14
15 VOLTA TENSOR 演算 入力 :FP16 フル精度 FP16 乗算 FP32 加算 出力 :FP32 FP16 に変換 more products FP16 FP16 16bit 16bit 32bit + FP32 FP16 FP32 FP16 加算もサポート ( インファレンス用 ) 15
16 TENSOR コアの使われ方 16x16 の行列の積和演算を Warp レベル (32 スレッド ) で協調実行 Warp (32 スレッド ) 32 スレッドで同期 Tensor コアを使用 16x16 行列の行列積を実行 32 スレッドで同期 16
17 TENSOR コアの使い方 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 cublas, cudnn, TensorRT Volta 向けに最適化されたフレームワーク ライブラリ } wmma::store_matrix_sync(d, Cmat, 16, wmma::row_major); CUDA C++ Warp レベル行列演算テンプレート 17
18 Tensor コア用の行列積アルゴリズムの選択 計算型を指定 (Tensor コアの場合は 加算の計算精度 ) CUBLAS: TENSOR コアの使い方 cublasgemmex で行列積 cublascreate( &handle ); cublassetmathmode( handle, CUBLAS_TENSOR_OP_MATH ); algo = CUBLAS_GEMM_DFALT_TENSOR_OP; cublasgemmex( handle, transa, transb, m, n, k, alpha, A, CUDA_R_16F, lda, B, CUDA_R_16F, ldb, beta, C, CUDA_R_16F, ldc, CUDA_R_32F, algo ); Tensor コア使用モードを選択 入力行列 A,B のデータ型を指定 出力行列 C のデータ型を指定 18
19 CUBLAS: TENSOR コアの使い方 cublasgemmex で行列積 入力行列 A,B のデータ型 出力行列 C のデータ型 加算型 FP16 FP16 FP32 ( 標準的な用途?) FP16 FP16 FP16 FP16 で加算 ( インファレンス ) FP16 FP32 FP32 FP32 で出力 FP32 FP32 FP32 FP32 データのまま Tensor コア使用 他 API でも使用可 : cublassgemmex, cublashgemm, cublashgemmbatched, cublashgemmstridebatched 19
20 CUBLAS: TENSOR コアの実効性能 P100 FP32 vs. V100 Tensor コア 最大 9 倍の性能向上 20
21 TFLOPS CUBLAS: TENSOR コアの実効性能 V100 同士で比較 : FP32 vs. Tensor コア Tensor コア (FP16 加算 ) Tensor コア (FP32 加算 ) FP32 と比べて 最大で 6 倍以上の性能 UP (FP32 加算の場合 ) matrix size (M=N=K) FP32 FP16 TensorCore (FP32 add) TensorCore (FP16 add) CUDA cublasgemmex() 使用 21
22 平均 1.2 TENSOR コアの計算精度 FP32 の計算結果に近い アプリケーション依存 誤差範囲 Tensor コアの演算結果は FP16 と比べて FP32 との誤差が小さい FP32 TensorCore FP16 内積長 行列 A: 指数分布 (activation) 行列 B: 正規分布 (weight) ( 平均 0.0, 分散 1.0) 内積長 : 万サンプル 誤差区間 : 99% 22
23 CUDNN: TENSOR コアの使い方 Convolution cudnncreate( &handle ); cudnncreatetensordescriptor( &cudnnidesc ); cudnncreatetensordescriptor( &cudnnodesc ); cudnncreatefilterdescriptor( &cudnnfdesc ); cudnncreateconvolutiondescriptor( &cudnnconvdesc ); cudnnsetconvolutionnddescriptor( cudnnconvdesc, ); cudnnsetconvolutionmathtype( cudnnconvdesc, CUDNN_TENSOR_OP_MATH ); algo = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM; cudnnconvolutionforward( handle, alpha, cudnnidesc, dev_i, cudnnfdesc, dev_f, cudnnconvdesc, algo, workspace, workspacesize, beta, cudnnodesc, dev_o ); Input Weight Output Tensor コア使用モードを選択 Tensor コア対応の Convolution アルゴリズム選択 23
24 CUDNN: TENSOR コアの使い方 Convolution: Tensor コア対応アルゴリズム Forward w CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED x y BackwardData w CUDNN_CONVOLUTION_BWD_DATA_ALGO_1 CUDNN_CONVOLUTION_BWD_DATA_ALGO_WINOGRAD_NONFUSED dx dy BackwardFilter CUDNN_CONVOLUTION_BWD_FILTER_ALGO_1 CUDNN_CONVOLUTION_BWD_FILTER_ALGO_WINOGRAD_NONFUSED x dw dy 24
25 CUDNN: TENSOR コアの実効性能 Pascal FP32 vs. V100 Tensor コア Convolution 層の性能比較 25
26 INDEPENDENT THREAD SCHEDULING 26
27 VOLTA GV100 SM 生産性の向上 命令セットを一新スケジューラを2 倍命令発行機構をシンプルに L1キャッシュの大容量 高速化 SIMTモデルの改善テンソル計算の加速 最もプログラミングの簡単な SM 27
28 WARP の実装 Pascal まで Program Counter (PC) and Stack (S) Warp(32 スレッド ) 毎に PC は 1 つ 28
29 diverge reconverge PASCAL: WARP 実行モデル if (threadidx.x < 4) { A; パス1 B; } else { X; パス2 Y; } Warp 内で複数パスに分岐した場合 一方のパスが完了するまで もう一方のパスは実行されない パス 1 X; Y; パス 2 A; B; Time 29
30 diverge reconverge PASCAL: WARP 実行モデル 分岐したパス間の同期は NG if (threadidx.x < 4) { A; syncwarp(); パス1 B; } else { X; syncwarp(); Y; } パス 2 パス 1 A; B; X; Y; パス 2 Time 30
31 スレッド間で通信するプログラム Pascal Lock-Free アルゴリズムであれば OK 他スレッドを待つのは NG 31
32 PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S WARP の実装 Pascal まで Program Counter (PC) and Stack (S) Warp(32 スレッド ) 毎に PC は 1 つ Volta Convergence Optimizer スレッド毎に PC を管理 個別にスケジューリングが可能 32
33 diverge synchronize VOLTA: 拡張 WARP 実行モデル Thread Independent Scheduling 分岐したパス間で 同期が可能! if (threadidx.x < 4) { パス1 A; X; syncwarp(); Y; B; } else { X; パス2 syncwarp(); パス1 パス1 Y; } A; B; syncwarp(); パス 2 パス 2 Time ( 注意 ) 同じワープの別スレッドが 同じサイクルに別インストラクションの実行は出来ない 33
34 スレッド間で通信するプログラム Pascal Volta Lock-Free アルゴリズムであれば OK 他スレッドを待つのは NG Starvation Free アルゴリズムであれば OK 他スレッドを待っても OK 34
35 STARVATION FREE アルゴリズムの例 双方向リンクリスト device void insert_after(node *a, Node *b) { Node *c; lock(a); lock(a->next); c = a->next; a->next = b; b->prev = a; b->next = c; c->prev = b; Doubly-Linked List with Fine Grained Lock a->next A C B c->prev } unlock(c); unlock(a); 35
36 STARVATION FREE アルゴリズムの例双方向リンクリスト device void insert_after(node *a, Node *b) { Node *c; lock(a); lock(a->next); c = a->next; a->next = b; b->prev = a; b->next = c; c->prev = b; Doubly-Linked List with Fine Grained Lock a->next A C B c->prev } unlock(c); unlock(a); *Not shown: lock() implementation 36
37 STARVATION FREE アルゴリズムの例双方向リンクリスト device void insert_after(node *a, Node *b) { Node *c; lock(a); lock(a->next); c = a->next; Doubly-Linked List with Fine Grained Lock A C a->next = b; b->prev = a; b->next = c; c->prev = b; a->next b->prev B b->next c->prev } unlock(c); unlock(a); 37
38 STARVATION FREE アルゴリズムの例双方向リンクリスト device void insert_after(node *a, Node *b) { Node *c; lock(a); lock(a->next); c = a->next; Doubly-Linked List with Fine Grained Lock A C a->next = b; b->prev = a; b->next = c; c->prev = b; a->next b->prev B b->next c->prev } unlock(c); unlock(a); Pascal で このプログラムを実行するのは危険アルゴリズムを Lock-free に変える必要がある 38
39 STARVATION FREE アルゴリズムの例双方向リンクリスト device void insert_after(node *a, Node *b) { Node *c; lock(a); lock(a->next); c = a->next; Doubly-Linked List with Fine Grained Lock A C a->next = b; b->prev = a; b->next = c; c->prev = b; a->next b->prev B b->next c->prev } unlock(c); unlock(a); Volta は最大 16 万スレッドを同時起動できるので あるスレッドが lock 獲得で停滞しても 他のスレッドが処理を進められる 39
40 VOLTA: 拡張 SIMT モデル CPU GPU (Pascal) GPU (Volta) データ並列 SIMD SIMT SIMT スレッド並列 ( タスク並列 ) MIMD SIMT (lock-free) SIMT Pascal まで スレッド並列のプログラムは アルゴリズムを lock-free に変更する必要 Volta アルゴリズム変更なく (or 少なく ) GPU で実行可能に 40
41 L1 CACHE AND SHARED MEMORY 41
42 VOLTA GV100 SM 生産性の向上 命令セットを一新スケジューラを2 倍命令発行機構をシンプルに L1キャッシュの大容量 高速化 SIMTモデルの改善テンソル計算の加速 最もプログラミングの簡単な SM 42
43 PASCAL の L1 キャッシュと共有メモリ Pascal SM Load/Store Units 共有メモリ 64 KB 短遅延 L1 キャッシュ 24 KB ストリーミング : キャッシュミスでもノンブロッキング L2 キャッシュ 4 MB 43
44 VOLTA: L1 キャッシュと共有メモリの統合 Pascal SM Load/Store Units Volta SM Load/Store Units 共有メモリ 64 KB 短遅延 L1 キャッシュ and 共有メモリ 128 KB L1 キャッシュ 24 KB ストリーミング L2 キャッシュ 4 MB L2 キャッシュ 6 MB 44
45 VOLTA: L1 キャッシュと共有メモリの統合 Volta: ストリーミング L1 キャッシュ Volta: 共有メモリ ノンブロッキング短い遅延 4 倍以上のバンド幅 5 倍以上の容量 L1 キャッシュとストレージを共用最大 96KB まで設定可能 ( カーネル毎 ) L1 Volta SM Load/Store Units L1キャッシュ共有メモリ L1 キャッシュ and 共有メモリ 32 KB 128 KB 128 KB 96 KB cudafuncsetattribute( func, cudafuncattributepreferredsharedmemorycarveout, cudasharedmemcarveoutmaxl1 cudasharedmemcarveoutmaxshared ); ); L2 キャッシュ 6 MB 45
46 L1 キャッシュで 共有メモリ使用時相応の性能を Volta L1 キャッシュ キャッシュ L1 キャッシュ使用時の性能 ( 平均 ) 共有メモリ使用で最適化した場合が基準 簡単に使える ( ソースコード変更不要 ) 90% 以上のケースで同等の性能 共有メモリ スレッド間の協調が必要なとき Atomicsが高速 安定した性能 70% 93% Pascal Volta 46
47 Volta: ストリーミング L1 キャッシュ VOLTA: L2 キャッシュの改善 Volta SM Load/Store Units ノンブロッキング短い遅延 4 倍以上のバンド幅 5 倍以上の容量 L1 キャッシュ and 共有メモリ 128 KB Volta: 共有メモリ L1 キャッシュとストレージを共用最大 96KB まで設定可能 ( カーネル毎 ) L2 キャッシュ, 6 MB 47
48 Atomcs throughput (Gops) L2 ATOMICS 性能の改善 Volta 最大 2 倍のスループット向上 Pascal Array size (MiB) no conflict (P100) no conflict (V100) AtomicAdd(FP32) 256M threads アクセスパターン : 規則的, ランダム random (P100) random (V100) 48
49 SCHEDULER 49
50 VOLTA GV100 SM 生産性の向上 命令セットを一新スケジューラを2 倍命令発行機構をシンプルに L1キャッシュの大容量 高速化 SIMTモデルの改善テンソル計算の加速 最もプログラミングの簡単な SM 50
51 SM: PASCALとVOLTAの相違 Volta ワープスケジューラ: 2 4 FP32ユニット/スケジューラ: Pascal
52 PASCAL のスケジューラ Pascal GP100: 1 個のスケジューラに 2 個のディスパッチャー 各ディスパッチャーが 16CUDA コアを担当 スケジューラは 1 サイクルに 1 回 Warp を選択 どちらかのディスパッチャーに渡す 各ディスパッチャーは 2 サイクルに 1 回 16CUDA コアに命令を投入 投入された命令は 2 サイクル使って 32 スレッドの処理を実行 32 スレッド = 16 CUDA コア x 2 サイクル 52
53 VOLTA のスケジューラ Volta GV100: 1 個のスケジューラに 1 個のディスパッチャー 各ディスパッチャーが 16 個の FP32 ユニットと INT ユニットを担当 スケジューラは 1 サイクルに 1 回 Warp を選択 ディスパッチャーに渡す ディスパッチャーは 1 サイクルに 1 回 16 個の FP32 or INT ユニットに 命令を投入 投入された命令は 2 サイクル使って 32 スレッドの処理を実行 53
54 INT Warp #2 FP32 Warp #1 VOLTA のスケジューラ Volta GV100: 1 個のスケジューラに 1 個のディスパッチャー 各ディスパッチャーが 16 個の FP32 ユニットと INT ユニットを担当 スケジューラは 1 サイクルに 1 回 Warp を選択 ディスパッチャーに渡す ディスパッチャーは FP32とINT1の同時実行が可能サイクルに1 回 16 個のFP32 or INTユニットに 命令を投入 投入された命令は 2 サイクル使って 32 スレッドの処理を実行 54
55 UNIFIED MEMORY 55
56 PASCAL のユニファイド メモリ GPU CPU Unified Memory GPU GPU CPU CPU Page Migration Engine GPU GPU CPU CPU Memory Memory GPU に最適な状態 CPU に最適な状態 56
57 VOLTA のユニファイド メモリ (CPU と PCI で接続 ) GPU CPU Unified Memory GPU GPU CPU CPU Page Migration Engine GPU GPU CPU CPU Memory GPU に最適な状態 アクセスカウンタの導入 より適切なタイミングで Page Migration Memory CPU に最適な状態 57
58 VOLTA のユニファイド メモリ (CPU と NVLINK で接続 ) GPU CPU Unified Memory GPU GPU CPU CPU Page Migration Engine GPU GPU CPU CPU Memory GPU に最適な状態 アクセスカウンタの導入 NVLINK の新機能 (Coherent, ATS, Atomics) Memory CPU に最適な状態 58
59 ユニファイド メモリの状況 少ない労力で 高い性能を Performance vs no Unified Memory OpenACC on P100 PGI の OpenACC コンパイラは ユニファイド メモリをサポート ( コンパイラ オプション ) SPEC ACCEL ベンチマーク 15 個の平均性能 ( データ移動を手動で最適化した場合との比較 ) PCIe: 86% 86% 明示的にデータを移動する効果 NVLINK: 91% Unified Memory Automatic data movement for allocatables PGI 17.1 Compilers OpenACC SPEC ACCEL 1.1 performance measured March, SPEC and the benchmark name SPEC ACCEL are registered trademarks of the Standard Performance Evaluation Corporation. 59
60 ロードマップ : UNIFIED SYSTEM ALLOCATOR 標準の malloc() で ユニファイド メモリが使えるようになる CUDA コード with System Allocator void sortfile(file *fp, int N) { char *data; // Allocate memory using any standard allocator data = (char *) malloc(n * sizeof(char)); fread(data, 1, N, fp); sort<<<...>>>(data,n,1,compare); OS サポートが必要 HMM Linux Kernel Module Linux kernel 4.14 にマージ CPU と GPU 間のデータ移動は 透過的に行われる ( ユニファイド メモリと同様 ) use_data(data); } // Free the allocated memory free(data); 60
61 CUDA MULTI-PROCESS SCHEDULING 61
62 GPU 上のマルチ プロセスのスケジューリング 背景 A B C CPU Processes 時分割スケジューリング GPU 使用時間帯を プロセスに配分 あるタイミングで GPU を使用しているプロセスは 一つ 各プロセスの排他実行を重視 マルチ プロセスサービス 同じ時間帯に 複数プロセスの同時 GPU 使用を許す 全プロセスで考えたときのスループットを重視 62
63 時分割スケジューリング A B C CPU Processes A GPU Execution Pascal GP100 Timeslice 1 63
64 時分割スケジューリング A B C A B C CPU Processes A B GPU Execution Pascal GP100 Pascal GP100 Timeslice 2 64
65 時分割スケジューリング CPU Processes A B C A B C A B C A Pascal GP100 B Pascal GP100 Pascal GP100 C Timeslice 3 GPU Execution 65
66 時分割スケジューリング 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 各プロセスの GPU 利用率が低ければ 当然 GPU 利用率は低いまま 66
67 PASCAL: マルチ プロセスサービス (MPS) A B C CUDA MPS 各プロセスの GPU 使用率は低くても 同時に GPU リソースを使用することで トータルで GPU 使用率を高めることができる Software Work Submission CUDA MULTI-PROCESS SERVICE (demon) CPU Processes GPU Execution Default ではオフ メモリ保護に制限 ( 他プロセスのメモリを壊す可能性 ) Limited Isolation A B Pascal GP100 C 67
68 VOLTA: マルチ プロセスサービス (MPS) Volta で MPS 改善 : ハードウェアでメモリ保護 ( 安全 ) Hardware Accelerated Work Submission A B C CUDA MULTI-PROCESS SERVICE CONTROL CPU Processes GPU Execution カーネル起動遅延の短縮 カーネル起動スループットの改善 VOLTA MULTI-PROCESS SERVICE スケジューラー分割による QoS 向上 ( 性能安定 ) Hardware Isolation A B C 対応プロセス数の増加 (Pascal:16 Volta:48) Volta GV100 68
69 Resnet50 Images/sec, 7ms latency VOLTA MPS: インファレンス事例 大 Batch サイズを使えないケースでも 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 69
70 TESLA V100 の概要 Volta Architecture Improved NVLink & HBM2 Volta MPS Improved SIMT Model Tensor Core Most Productive GPU Efficient Bandwidth Inference Utilization New Algorithms 125 Programmable TFLOPS Deep Learning Deep Learning と HPC 両方に最適な GPU 70
71
Images per Second Images per Second VOLTA: ディープラーニングにおける大きな飛躍 ResNet-50 トレーニング 2.4x faster ResNet-50 推論 TensorRT - 7ms レイテンシ 3.7x faster P100 V100 P10
NVIDIA TESLA V100 CUDA 9 のご紹介 森野慎也, シニアソリューションアーキテクト (GPU-Computing) NVIDIA Images per Second Images per Second VOLTA: ディープラーニングにおける大きな飛躍 ResNet-50 トレーニング 2.4x faster ResNet-50 推論 TensorRT - 7ms レイテンシ
CUDA 9 AND MORE 成瀬彰, シニアデベロッパーテクノロジーエンジニア, 2017/12/12
CUDA 9 AND MORE 成瀬彰, シニアデベロッパーテクノロジーエンジニア, 2017/12/12 CUDA 9 の概要 VOLTA に対応 ライブラリの高速化 Tesla V100 Volta アーキテクチャ Tensor コア NVLink Independent スレッドスケジューリング cublas ( 主に DL 向け ) NPP ( 画像処理 ) cufft ( 信号処理 ) cusolver
VOLTA AND TURING: ARCHITECTURE Akira Naruse, Developer Technology, 2018/9/14
VOLTA AND TURING: ARCHITECTURE AND PERFORMANCE OPTIMIZATION Akira Naruse, Developer Technology, 2018/9/14 VOLTA AND TURING: ARCHITECTURE Akira Naruse, Developer Technology, 2018/9/14 For HPC and Deep Learning
Slides: TimeGraph: GPU Scheduling for Real-Time Multi-Tasking Environments
計算機アーキテクチャ第 11 回 マルチプロセッサ 本資料は授業用です 無断で転載することを禁じます 名古屋大学 大学院情報科学研究科 准教授加藤真平 デスクトップ ジョブレベル並列性 スーパーコンピュータ 並列処理プログラム プログラムの並列化 for (i = 0; i < N; i++) { x[i] = a[i] + b[i]; } プログラムの並列化 x[0] = a[0] + b[0];
07-二村幸孝・出口大輔.indd
GPU Graphics Processing Units HPC High Performance Computing GPU GPGPU General-Purpose computation on GPU CPU GPU GPU *1 Intel Quad-Core Xeon E5472 3.0 GHz 2 6 MB L2 cache 1600 MHz FSB 80 GFlops 1 nvidia
ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014
ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014 コンカレントな処理の実行 システム内部の複数の処理を 平行に実行する CPU GPU メモリ転送 カーネル実行 複数のカーネル間 ストリーム GPU 上の処理キュー カーネル実行 メモリ転送の並列性 実行順序 DEFAULT STREAM Stream : GPU
hpc141_shirahata.pdf
GPU アクセラレータと不揮発性メモリ を考慮した I/O 性能の予備評価 白幡晃一 1,2 佐藤仁 1,2 松岡聡 1 1: 東京工業大学 2: JST CREST 1 GPU と不揮発性メモリを用いた 大規模データ処理 大規模データ処理 センサーネットワーク 遺伝子情報 SNS など ペタ ヨッタバイト級 高速処理が必要 スーパーコンピュータ上での大規模データ処理 GPU 高性能 高バンド幅 例
GPU チュートリアル :OpenACC 篇 Himeno benchmark を例題として 高エネルギー加速器研究機構 (KEK) 松古栄夫 (Hideo Matsufuru) 1 December 2018 HPC-Phys 理化学研究所 共通コードプロジェクト
GPU チュートリアル :OpenACC 篇 Himeno benchmark を例題として 高エネルギー加速器研究機構 (KEK) 松古栄夫 (Hideo Matsufuru) 1 December 2018 HPC-Phys 勉強会 @ 理化学研究所 共通コードプロジェクト Contents Hands On 環境について Introduction to GPU computing Introduction
ムーアの法則後の世界 年間のマイクロプロセッサのトレンド トランジスタ数 ( 千単位 ) 年率 1.1 倍 シングルスレッド性能 年率 1.5 倍 Original data up t
エヌビディアが加速する AI 革命 エヌビディア合同会社 エンタープライズマーケティング本部長林憲一 1 ムーアの法則後の世界 10 7 40 年間のマイクロプロセッサのトレンド 10 6 10 5 10 4 トランジスタ数 ( 千単位 ) 年率 1.1 倍 10 3 10 2 シングルスレッド性能 年率 1.5 倍 1980 1990 2000 2010 2020 Original data up
VOLTA TENSOR コアで 高速かつ高精度に DL モデルをトレーニングする方法 成瀬彰, シニアデベロッパーテクノロジーエンジニア, 2017/12/12
VOLTA TENSOR コアで 高速かつ高精度に DL モデルをトレーニングする方法 成瀬彰, シニアデベロッパーテクノロジーエンジニア, 2017/12/12 アジェンダ Tensorコアとトレーニングの概要 混合精度 (Tensorコア) で FP32と同等の精度を得る方法 ウェイトをFP16とFP32を併用して更新する ロス スケーリング DLフレームワーク対応状況 ウェイトをFP16で更新する
GPU GPU CPU CPU CPU GPU GPU N N CPU ( ) 1 GPU CPU GPU 2D 3D CPU GPU GPU GPGPU GPGPU 2 nvidia GPU CUDA 3 GPU 3.1 GPU Core 1
GPU 4 2010 8 28 1 GPU CPU CPU CPU GPU GPU N N CPU ( ) 1 GPU CPU GPU 2D 3D CPU GPU GPU GPGPU GPGPU 2 nvidia GPU CUDA 3 GPU 3.1 GPU Core 1 Register & Shared Memory ( ) CPU CPU(Intel Core i7 965) GPU(Tesla
1 OpenCL OpenCL 1 OpenCL GPU ( ) 1 OpenCL Compute Units Elements OpenCL OpenCL SPMD (Single-Program, Multiple-Data) SPMD OpenCL work-item work-group N
GPU 1 1 2 1, 3 2, 3 (Graphics Unit: GPU) GPU GPU GPU Evaluation of GPU Computing Based on An Automatic Program Generation Technology Makoto Sugawara, 1 Katsuto Sato, 1 Kazuhiko Komatsu, 2 Hiroyuki Takizawa
1 GPU GPGPU GPU CPU 2 GPU 2007 NVIDIA GPGPU CUDA[3] GPGPU CUDA GPGPU CUDA GPGPU GPU GPU GPU Graphics Processing Unit LSI LSI CPU ( ) DRAM GPU LSI GPU
![1 GPU GPGPU GPU CPU 2 GPU 2007 NVIDIA GPGPU CUDA[3] GPGPU CUDA GPGPU CUDA GPGPU GPU GPU GPU Graphics Processing Unit LSI LSI CPU ( ) DRAM GPU LSI GPU](/thumbs/89/99132402.jpg "1 GPU GPGPU GPU CPU 2 GPU 2007 NVIDIA GPGPU CUDA[3] GPGPU CUDA GPGPU CUDA GPGPU GPU GPU GPU Graphics Processing Unit LSI LSI CPU ( ) DRAM GPU LSI GPU")
GPGPU (I) GPU GPGPU 1 GPU(Graphics Processing Unit) GPU GPGPU(General-Purpose computing on GPUs) GPU GPGPU GPU ( PC ) PC PC GPU PC PC GPU GPU 2008 TSUBAME NVIDIA GPU(Tesla S1070) TOP500 29 [1] 2009 AMD
VXPRO R1400® ご提案資料
Intel Core i7 プロセッサ 920 Preliminary Performance Report ノード性能評価 ノード性能の評価 NAS Parallel Benchmark Class B OpenMP 版での性能評価 実行スレッド数を 4 で固定 ( デュアルソケットでは各プロセッサに 2 スレッド ) 全て 2.66GHz のコアとなるため コアあたりのピーク性能は同じ 評価システム
MATLAB® における並列・分散コンピューティング ~ Parallel Computing Toolbox™ & MATLAB Distributed Computing Server™ ~
MATLAB における並列 分散コンピューティング ~ Parallel Computing Toolbox & MATLAB Distributed Computing Server ~ MathWorks Japan Application Engineering Group Takashi Yoshida 2016 The MathWorks, Inc. 1 System Configuration
システムソリューションのご紹介
HP 2 C 製品 :VXPRO/VXSMP サーバ 製品アップデート 製品アップデート VXPRO と VXSMP での製品オプションの追加 8 ポート InfiniBand スイッチ Netlist HyperCloud メモリ VXPRO R2284 GPU サーバ 製品アップデート 8 ポート InfiniBand スイッチ IS5022 8 ポート 40G InfiniBand スイッチ
GPGPUクラスタの性能評価
2008 年度理研 HPC シンポジウム第 3 世代 PC クラスタ GPGPU クラスタの性能評価 2009 年 3 月 12 日 富士通研究所成瀬彰 発表の概要 背景 GPGPU による高速化 CUDA の概要 GPU のメモリアクセス特性調査 姫野 BMT の高速化 GPGPU クラスタによる高速化 GPU Host 間のデータ転送 GPU-to-GPU の通信性能 GPGPU クラスタ上での姫野
23 Fig. 2: hwmodulev2 3. Reconfigurable HPC 3.1 hw/sw hw/sw hw/sw FPGA PC FPGA PC FPGA HPC FPGA FPGA hw/sw hw/sw hw- Module FPGA hwmodule hw/sw FPGA h
23 FPGA CUDA Performance Comparison of FPGA Array with CUDA on Poisson Equation (lijiang@sekine-lab.ei.tuat.ac.jp), (kazuki@sekine-lab.ei.tuat.ac.jp), (takahashi@sekine-lab.ei.tuat.ac.jp), (tamukoh@cc.tuat.ac.jp),
CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン
CUDA 画像処理入門 エヌビディアジャパン CUDA エンジニア森野慎也 GTC Japan 2014 CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン RGB Y( 輝度 ) 変換 カラー画像から グレイスケールへの変換 Y = 0.299 R + 0.587
GPGPU
GPGPU 2013 1008 2015 1 23 Abstract In recent years, with the advance of microscope technology, the alive cells have been able to observe. On the other hand, from the standpoint of image processing, the
修士論文
AVX を用いた倍々精度疎行列ベクトル積の高速化 菱沼利彰 1 藤井昭宏 1 田中輝雄 1 長谷川秀彦 2 1 工学院大学 2 筑波大学 1 目次 1. 研究背景 目的 2. 実装, 実験環境 3. 実験 - 倍々精度ベクトル演算 - 4. 実験 - 倍々精度疎行列ベクトル積 - 5. まとめ 多倍長精度計算フォーラム 2 目次 1. 研究背景 目的 2. 実装, 実験環境 3. 実験 - 倍々精度ベクトル演算
1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境 Lin
Windows で始める CUDA 入門 GTC 2013 チュートリアル エヌビディアジャパン CUDA エンジニア森野慎也 1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境
WP _v1.1 目次 NVIDIA Tesla V100 GPU アーキテクチャ概論... 1 Tesla V100: AI コンピューティングと HPC の主戦力... 3 主な機能... 3 AI および HPC 向けの究極のパフォーマンス... 7 NVIDIA GPU -
NVIDIA TESLA V100 GPU アーキテクチャ 世界最先端のデータセンター GPU WP-08608-001_v1.1 August 2017 WP-08608-001_v1.1 目次 NVIDIA Tesla V100 GPU アーキテクチャ概論... 1 Tesla V100: AI コンピューティングと HPC の主戦力... 3 主な機能... 3 AI および HPC 向けの究極のパフォーマンス...
GPU n Graphics Processing Unit CG CAD
GPU 2016/06/27 第 20 回 GPU コンピューティング講習会 ( 東京工業大学 ) 1 GPU n Graphics Processing Unit CG CAD www.nvidia.co.jp www.autodesk.co.jp www.pixar.com GPU n GPU ü n NVIDIA CUDA ü NVIDIA GPU ü OS Linux, Windows, Mac
4 倍精度基本線形代数ルーチン群 QPBLAS の紹介 [index] 1. Introduction 2. Double-double algorithm 3. QPBLAS 4. QPBLAS-GPU 5. Summary 佐々成正 1, 山田進 1, 町田昌彦 1, 今村俊幸 2, 奥田洋司
![4 倍精度基本線形代数ルーチン群 QPBLAS の紹介 [index] 1. Introduction 2. Double-double algorithm 3. QPBLAS 4. QPBLAS-GPU 5. Summary 佐々成正 1, 山田進 1, 町田昌彦 1, 今村俊幸 2, 奥田洋司](/thumbs/97/131936032.jpg "4 倍精度基本線形代数ルーチン群 QPBLAS の紹介 [index] 1. Introduction 2. Double-double algorithm 3. QPBLAS 4. QPBLAS-GPU 5. Summary 佐々成正 1, 山田進 1, 町田昌彦 1, 今村俊幸 2, 奥田洋司")
4 倍精度基本線形代数ルーチン群 QPBLAS の紹介 [index] 1. Introduction 2. Double-double algorithm 3. QPBLAS 4. QPBLAS-GPU 5. Summary 佐々成正 1, 山田進 1, 町田昌彦 1, 今村俊幸 2, 奥田洋司 3 1 1 日本原子力研究開発機構システム計算科学センター 2 理科学研究所計算科学研究機構 3 東京大学新領域創成科学研究科
Microsoft PowerPoint - GPU_computing_2013_01.pptx
GPU コンピューティン No.1 導入 東京工業大学 学術国際情報センター 青木尊之 1 GPU とは 2 GPGPU (General-purpose computing on graphics processing units) GPU を画像処理以外の一般的計算に使う GPU の魅力 高性能 : ハイエンド GPU はピーク 4 TFLOPS 超 手軽さ : 普通の PC にも装着できる 低価格
本文ＡＬＬ.indd
Intel Xeon プロセッサにおける Cache Coherency 時間の性能測定方法河辺峻田口成美古谷英祐 Intel Xeon プロセッサにおける Cache Coherency 時間の性能測定方法 Performance Measurement Method of Cache Coherency Effects on an Intel Xeon Processor System 河辺峻田口成美古谷英祐
GPU CUDA CUDA 2010/06/28 1
GPU CUDA CUDA 2010/06/28 1 GPU NVIDIA Mark Harris, Optimizing Parallel Reduction in CUDA http://developer.download.nvidia.com/ compute/cuda/1_1/website/data- Parallel_Algorithms.html#reduction CUDA SDK
Microsoft Word - HOKUSAI_system_overview_ja.docx
HOKUSAI システムの概要 1.1 システム構成 HOKUSAI システムは 超並列演算システム (GWMPC BWMPC) アプリケーション演算サーバ群 ( 大容量メモリ演算サーバ GPU 演算サーバ ) と システムの利用入口となるフロントエンドサーバ 用途の異なる 2 つのストレージ ( オンライン ストレージ 階層型ストレージ ) から構成されるシステムです 図 0-1 システム構成図
Microsoft Word - Dolphin Expressによる10Gbpソケット通信.docx
Dolphin Express による 10Gbps ソケット通信 Dolphin Express は 標準的な低価格のサーバを用いて 強力なクラスタリングシステムが構築できる ハードウェアとソフトウェアによる通信用アーキテクチャです 本資料では Dolphin Express 製品の概要と 実際にどの程度の性能が出るのか市販 PC での実験結果をご紹介します Dolphin Express 製品体系
スライド 1
知能制御システム学 画像処理の高速化 OpenCV による基礎的な例 東北大学大学院情報科学研究科鏡慎吾 swk(at)ic.is.tohoku.ac.jp 2007.07.03 リアルタイム処理と高速化 リアルタイム = 高速 ではない 目標となる時間制約が定められているのがリアルタイム処理である.34 ms かかった処理が 33 ms に縮んだだけでも, それによって与えられた時間制約が満たされるのであれば,
rank ”«‘‚“™z‡Ì GPU ‡É‡æ‡éŁÀŠñ›»
rank GPU ERATO 2011 11 1 1 / 26 GPU rank/select wavelet tree balanced parenthesis GPU rank 2 / 26 GPU rank/select wavelet tree balanced parenthesis GPU rank 2 / 26 GPU rank/select wavelet tree balanced
GTC Japan, 2018/09/14 得居誠也, Preferred Networks Chainer における 深層学習の高速化 Optimizing Deep Learning with Chainer
GTC Japan, 2018/09/14 得居誠也, Preferred Networks Chainer における 深層学習の高速化 Optimizing Deep Learning with Chainer Chainer のミッション Deep Learning とその応用の研究開発を加速させる 環境セットアップが速い すぐ習熟 素早いコーディング 実験の高速化 結果をさっと公開 論文化
GPUコンピューティング講習会パート１
GPU コンピューティング (CUDA) 講習会 GPU と GPU を用いた計算の概要 丸山直也 スケジュール 13:20-13:50 GPU を用いた計算の概要 担当丸山 13:50-14:30 GPU コンピューティングによる HPC アプリケーションの高速化の事例紹介 担当青木 14:30-14:40 休憩 14:40-17:00 CUDA プログラミングの基礎 担当丸山 TSUBAME の
N08
CPU のキモチ C.John 自己紹介 英語きらい 絵かけない 人の話を素直に信じない CPUにキモチなんてない お詫び 予告ではCとC# とありましたがやる気と時間の都合上 C++のみを対象とします 今日のネタ元 MSDN マガジン 2010 年 10 月号 http://msdn.microsoft.com/ja-jp/magazine/cc850829.aspx Windows と C++
熊本大学学術リポジトリ Kumamoto University Repositor Title GPGPU による高速演算について Author(s) 榎本, 昌一 Citation Issue date Type URL Presentation
熊本大学学術リポジトリ Kumamoto University Repositor Title GPGPU による高速演算について Author(s) 榎本, 昌一 Citation Issue date 2011-03-17 Type URL Presentation http://hdl.handle.net/2298/23539 Right GPGPU による高速演算について 榎本昌一 東京大学大学院工学系研究科システム創成学専攻
倍々精度RgemmのnVidia C2050上への実装と応用
.. maho@riken.jp http://accc.riken.jp/maho/,,, 2011/2/16 1 - : GPU : SDPA-DD 10 1 - Rgemm : 4 (32 ) nvidia C2050, GPU CPU 150, 24GFlops 25 20 GFLOPS 15 10 QuadAdd Cray, QuadMul Sloppy Kernel QuadAdd Cray,
PowerPoint プレゼンテーション
Dell PowerEdge C6320 スケーラブルサーバアプライアンス 仮想化アプライアンスサーバ 最新のプロセッサを搭載したサーバプラットフォーム vsmp Foundation によるサーバ仮想化と統合化の適用 システムはセットアップを完了した状態でご提供 基本構成ではバックプレーン用のスイッチなどが不要 各ノード間を直接接続 冗長性の高いバックプレーン構成 利用するサーバプラットフォームは
HPEハイパフォーマンスコンピューティング ソリューション
HPE HPC / AI Page 2 No.1 * 24.8% No.1 * HPE HPC / AI HPC AI SGIHPE HPC / AI GPU TOP500 50th edition Nov. 2017 HPE No.1 124 www.top500.org HPE HPC / AI TSUBAME 3.0 2017 7 AI TSUBAME 3.0 HPE SGI 8600 System
tabaicho3mukunoki.pptx
1 2 はじめに n 目的 4倍精度演算より高速な3倍精度演算を実現する l 倍精度では足りないが4倍精度は必要ないケースに欲しい l 4倍精度に比べてデータサイズが小さい Ø 少なくともメモリ律速な計算では4倍精度よりデータ 転送時間を減らすことが可能 Ø PCIeやノード間通信がボトルネックとなりやすい GPUクラスタ環境に有効か n 研究概要 l DD型4倍精度演算 DD演算 に基づく3倍精度演算
CUDA 連携とライブラリの活用 2
1 09:30-10:00 受付 10:00-12:00 Reedbush-H ログイン GPU 入門 13:30-15:00 OpenACC 入門 15:15-16:45 OpenACC 最適化入門と演習 17:00-18:00 OpenACC の活用 (CUDA 連携とライブラリの活用 ) CUDA 連携とライブラリの活用 2 3 OpenACC 簡単にGPUプログラムが作成できる それなりの性能が得られる
名称 : 日本 GPU コンピューティングパートナーシップ (G-DEP) 所在 : 東京都文京区本郷 7 丁目 3 番 1 号東京大学アントレプレナープラザ, 他工場 URL アライアンスパートナー コアテクノロジーパートナー NVIDIA JAPAN ソリュ
GPUDirect の現状整理 multi-gpu に取組むために G-DEP チーフエンジニア河井博紀 (kawai@gdep.jp) 名称 : 日本 GPU コンピューティングパートナーシップ (G-DEP) 所在 : 東京都文京区本郷 7 丁目 3 番 1 号東京大学アントレプレナープラザ, 他工場 URL http://www.gdep.jp アライアンスパートナー コアテクノロジーパートナー
スライド 1
GPU クラスタによる格子 QCD 計算 広大理尾崎裕介 石川健一 1.1 Introduction Graphic Processing Units 1 チップに数百個の演算器 多数の演算器による並列計算 ~TFLOPS ( 単精度 ) CPU 数十 GFLOPS バンド幅 ~100GB/s コストパフォーマンス ~$400 GPU の開発環境 NVIDIA CUDA http://www.nvidia.co.jp/object/cuda_home_new_jp.html
Microsoft PowerPoint - SWoPP2010_Shirahata
GPU を考慮した MapReduce の タスクスケジューリング 白幡晃一 1 佐藤仁 1 松岡聡 1 2 3 1 東京工業大学 2 科学技術振興機構 3 国立情報学研究所 大規模データ処理 情報爆発時代における 大規模データ処理 気象 生物学 天文学 物理学など様々な科学技術計算での利用 MapReduce 大規模データ処理のためのプログラミングモデルデ スケーラブルな並列データ処理 GPGPU
<4D F736F F F696E74202D2091E63489F15F436F6D C982E682E992B48D8291AC92B489B F090CD2888F38DFC E B8CDD8
Web キャンパス資料 超音波シミュレーションの基礎 ~ 第 4 回 ComWAVEによる超高速超音波解析 ~ 科学システム開発部 Copyright (c)2006 ITOCHU Techno-Solutions Corporation 本日の説明内容 ComWAVEの概要および特徴 GPGPUとは GPGPUによる解析事例 CAE POWER 超音波研究会開催 (10 月 3 日 ) のご紹介
PowerPoint プレゼンテーション
vsmp Foundation スケーラブル SMP システム スケーラブル SMP システム 製品コンセプト 2U サイズの 8 ソケット SMP サーバ コンパクトな筐体に多くのコアとメモリを実装し SMP システムとして利用可能 スイッチなし構成でのシステム構築によりラックスペースを無駄にしない構成 将来的な拡張性を保証 8 ソケット以上への拡張も可能 2 システム構成例 ベースシステム 2U
TSUBAME2.0 における GPU の 活用方法 東京工業大学学術国際情報センター丸山直也第 10 回 GPU コンピューティング講習会 2011 年 9 月 28 日
TSUBAME2.0 における GPU の 活用方法 東京工業大学学術国際情報センター丸山直也第 10 回 GPU コンピューティング講習会 2011 年 9 月 28 日 目次 1. TSUBAMEのGPU 環境 2. プログラム作成 3. プログラム実行 4. 性能解析 デバッグ サンプルコードは /work0/gsic/seminars/gpu- 2011-09- 28 からコピー可能です 1.
はじめに AI は 感染症の早期発見と治療法の探索 交通事故死の削減 事故発生前の重大なインフラ欠陥の発見など 人類が直面する複雑な > 問題を解決するのに役立てられています AI とディープラーニング利用における 2 つの大きな課題は パフォーマンスの最大化と 絶え間なく変化する基盤技術の管理です
技術概要 NVIDIA GPU CLOUD ディープラーニングソフトウェア 最適化されたディープラーニングコンテナーのガイド はじめに AI は 感染症の早期発見と治療法の探索 交通事故死の削減 事故発生前の重大なインフラ欠陥の発見など 人類が直面する複雑な > 問題を解決するのに役立てられています AI とディープラーニング利用における 2 つの大きな課題は パフォーマンスの最大化と 絶え間なく変化する基盤技術の管理です
Microsoft Word ●IntelクアッドコアCPUでのベンチマーク_吉岡_ _更新__ doc
2.3. アプリ性能 2.3.1. Intel クアッドコア CPU でのベンチマーク 東京海洋大学吉岡諭 1. はじめにこの数年でマルチコア CPU の普及が進んできた x86 系の CPU でも Intel と AD がデュアルコア クアッドコアの CPU を次々と市場に送り出していて それらが PC クラスタの CPU として採用され HPC に活用されている ここでは Intel クアッドコア
NUMAの構成
GPU のプログラム 天野 アクセラレータとは? 特定の性質のプログラムを高速化するプロセッサ 典型的なアクセラレータ GPU(Graphic Processing Unit) Xeon Phi FPGA(Field Programmable Gate Array) 最近出て来た Deep Learning 用ニューロチップなど Domain Specific Architecture 1GPGPU:General
並列・高速化を実現するための 高速化サービスの概要と事例紹介
第 4 回 AVS 可視化フォーラム 2019 並列 高速化を実現するための 高速化サービスの概要と事例紹介 株式会社アーク情報システム営業部仮野亮ソリューション技術部佐々木竜一 2019.08.30 はじめに アーク情報システムの紹介 高速化サービスとは? 事例紹介 コンサルティングサービスについて アーク情報システムの紹介 設立 資本金 :1987 年 10 月 :3 億 600 万円 従業員数
スライド 1
Nehalem 新マイクロアーキテクチャ スケーラブルシステムズ株式会社 はじめに 現在も続く x86 マイクロプロセッサマーケットでの競合において Intel と AMD という 2 つの会社は 常に新しい技術 製品を提供し マーケットでのシェアの獲得を目指しています この技術開発と製品開発では この 2 社はある時は 他社に対して優位な技術を開発し 製品面での優位性を示すことに成功してきましたが
<4D F736F F F696E74202D F A282BD94BD959C89F A4C E682528D652E707074>
発表の流れ SSE を用いた反復解法ライブラリ Lis 4 倍精度版の高速化 小武守恒 (JST 東京大学 ) 藤井昭宏 ( 工学院大学 ) 長谷川秀彦 ( 筑波大学 ) 西田晃 ( 中央大学 JST) はじめに 4 倍精度演算について Lisへの実装 SSEによる高速化 性能評価 スピード 収束 まとめ はじめに クリロフ部分空間法たとえば CG 法は, 理論的には高々 n 回 (n は係数行列の次元数
b4-deeplearning-embedded-c-mw
ディープラーニングアプリケーション の組み込み GPU/CPU 実装 アプリケーションエンジニアリング部町田和也 2015 The MathWorks, Inc. 1 アジェンダ MATLAB Coder/GPU Coder の概要 ディープニューラルネットワークの組み込み実装ワークフロー パフォーマンスに関して まとめ 2 ディープラーニングワークフローのおさらい Application logic
NVIDIA Tesla K20/K20X GPU アクセラレータ アプリケーション パフォーマンス テクニカル ブリーフ
NVIDIA Tesla K20/K20X GPU アクセラレータ アプリケーション パフォーマンス テクニカル ブリーフ K20 GPU2 個に対するスピードアップ NVIDIA は Fermi アーキテクチャ GPU の発表により パフォーマンス エネルギー効率の両面で飛躍的な性能向上を実現し ハイパフォーマンスコンピューティング (HPC) の世界に変革をもたらしました また 実際に GPU
Microsoft PowerPoint - suda.pptx
GPU の HWアーキテクチャと高性能化手法 須田礼仁 ( 東京大学 ) 2011/03/22 GPU 高性能プログラミング GPU のハードウェアを理解する CUDA のソフトウェアを理解する CUDA でプログラムを書くのは難しくないが, CUDA で高速なプログラムを書くのは難しい どうすれば遅くなるかを理解する! 効果が大きいものから順に説明します 1 高性能プログラミングの手順 1. 現在のコードの,
HP Workstation 総合カタログ
HP Workstation Z HP 6 Z HP HP Z840 Workstation P.9 HP Z640 Workstation & CPU P.10 HP Z440 Workstation P.11 17.3in WIDE HP ZBook 17 G2 Mobile Workstation P.15 15.6in WIDE HP ZBook 15 G2 Mobile Workstation
Microsoft Word - nvsi_050110jp_netvault_vtl_on_dothill_sannetII.doc
Article ID: NVSI-050110JP Created: 2005/10/19 Revised: - NetVault 仮想テープ ライブラリのパフォーマンス検証 : dothill SANnetⅡSATA 編 1. 検証の目的 ドットヒルシステムズ株式会社の SANnetll SATA は 安価な SATA ドライブを使用した大容量ストレージで ディスクへのバックアップを行う際の対象デバイスとして最適と言えます
(速報) Xeon E 系モデル 新プロセッサ性能について
( 速報 ) Xeon E5-2600 系モデル新プロセッサ性能について 2012 年 3 月 16 日 富士通株式会社 2012 年 3 月 7 日 インテル社より最新 CPU インテル Xeon E5 ファミリー の発表がありました この最新 CPU について PC クラスタシステムの観点から性能検証を行いましたので 概要を速報いたします プロセッサインテル Xeon プロセッサ E5-2690
160311_icm2015-muramatsu-v2.pptx
Linux におけるパケット処理機構の 性能評価に基づいた NFV 導 の 検討 村松真, 川島 太, 中 裕貴, 林經正, 松尾啓志 名古屋 業 学 学院 株式会社ボスコ テクノロジーズ ICM 研究会 2016/03/11 研究 的 VM 仮想 NIC バックエンド機構 仮想化環境 仮想スイッチ パケット処理機構 物理環境 性能要因を考察 汎 IA サーバ NFV 環境に適したサーバ構成を検討
Microsoft PowerPoint - CCS学際共同boku-08b.ppt
マルチコア / マルチソケットノードに おけるメモリ性能のインパクト 研究代表者朴泰祐筑波大学システム情報工学研究科 taisuke@cs.tsukuba.ac.jp アウトライン 近年の高性能 PC クラスタの傾向と問題 multi-core/multi-socket ノードとメモリ性能 メモリバンド幅に着目した性能測定 multi-link network 性能評価 まとめ 近年の高性能 PC
第 1 回ディープラーニング分散学習ハッカソン <ChainerMN 紹介 + スパコンでの実 法 > チューター福 圭祐 (PFN) 鈴 脩司 (PFN)
第 1 回ディープラーニング分散学習ハッカソン チューター福 圭祐 (PFN) 鈴 脩司 (PFN) https://chainer.org/ 2 Chainer: A Flexible Deep Learning Framework Define-and-Run Define-by-Run Define Define by Run Model
スライド 1
東北大学工学部機械知能 航空工学科 2019 年度クラス C D 情報科学基礎 I 14. さらに勉強するために 大学院情報科学研究科 鏡慎吾 http://www.ic.is.tohoku.ac.jp/~swk/lecture/ 0 と 1 の世界 これまで何を学んだか 2 進数, 算術演算, 論理演算 計算機はどのように動くのか プロセッサとメモリ 演算命令, ロード ストア命令, 分岐命令 計算機はどのように構成されているのか
Microsoft PowerPoint - GPGPU実践基礎工学（web）.pptx
GPU のプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 GPU プログラミング環境 (CUDA) GPU プログラムの実行の流れ CUDA によるプログラムの記述 カーネル (GPU で処理する関数 ) の構造 記述方法とその理由 GPU 固有のパラメータの確認 405 GPU(Graphics Processing Unit) とは 画像処理専用のハードウェア 具体的には画像処理用のチップ
Microsoft PowerPoint - GPUシンポジウム _d公開版.ppt [互換モード]
![Microsoft PowerPoint - GPUシンポジウム _d公開版.ppt [互換モード]](/thumbs/102/156675429.jpg "Microsoft PowerPoint - GPUシンポジウム _d公開版.ppt [互換モード]")
200/0/9 数値流体解析の並列効率とその GPU による高速化の試み 清水建設 ( 株 ) 技術研究所 PHAM VAN PHUC ( ファムバンフック ) 流体計算時間短縮と GPU の活用の試み 現 CPUとの比較によりGPU 活用の可能性 現 CPU の最大利用 ノード内の最大計算資源の利用 すべてCPUコアの利用 適切なアルゴリズムの利用 CPU コア性能の何倍? GPU の利用の試み
PowerPoint プレゼンテーション
各種計算機アプリケーション性能比較 目次. はじめに. 行列積計算.QDR 積計算 4.N 体問題計算 5. 多次元積分計算 5. 次元積分計算 5. 次元積分計算 5. 4 次元積分計算 5.4 5 次元積分計算 5.5 6 次元積分計算 平成 6 年度第 四半期 . はじめに 今までと少し性質の異なるグラフィックボードが使用できる様になったので従来のアプリケーションで性能比較を実施しました 主に使用した計算機は以下のものです
PassMark PerformanceTest ™
KRONOS S ライン 性能ベンチマーク オーバークロックモニター OCCT OverClock Checking Tool i7z (A better i7 (and now i3, i5) reporting tool for Linux) KRONOS S800 CATIA Benchmark Aerospace - 8/17 passengers Jet - Mid Fuse DELL Precision
PowerPoint プレゼンテーション
Oracle GRID Center Flash SSD + 最新ストレージと Oracle Database で実現するデータベース統合の新しい形 2011 年 2 月 23 日日本オラクル Grid Center エンジニア岩本知博 進化し続けるストレージ関連技術 高速ストレージネットワークの多様化 低価格化 10GbE FCoE 8Gb FC ディスクドライブの多様化および大容量 / 低価格化
Slide 1
CUDA プログラミングの基本 パート I - ソフトウェアスタックとメモリ管理 CUDA の基本の概要 パート I CUDAのソフトウェアスタックとコンパイル GPUのメモリ管理 パートII カーネルの起動 GPUコードの具体項目 注 : 取り上げているのは基本事項のみです そのほか多数の API 関数についてはプログラミングガイドを ご覧ください CUDA インストレーション CUDA インストレーションの構成
ERDAS IMAGINE における処理速度の向上 株式会社ベストシステムズ PASCO CORPORATION 2015
ERDAS IMAGINE における処理速度の向上 株式会社ベストシステムズ 本セッションの目的 本セッションでは ERDAS IMAGINEにおける処理速度向上を目的として機器 (SSD 等 ) 及び並列処理の比較 検討を行った 1.SSD 及び RAMDISK を利用した処理速度の検証 2.Condorによる複数 PCを用いた並列処理 2.1 分散並列処理による高速化試験 (ERDAS IMAGINEのCondorを使用した試験
PNopenseminar_2011_開発stack
PROFINET Open Seminar 開発セミナー Software Stack FPGA IP core PROFINET 対応製品の開発 2 ユーザ要求要求は多種多様 複雑な規格の仕様を一から勉強するのはちょっと.. できるだけ短期間で 柔軟なスケジュールで進めたい既存のハードウェアを変更することなく PN を対応させたい将来的な仕様拡張に対してシームレスに統合したい同じハードウェアで複数の
GPGPU によるアクセラレーション環境について
GPGPU によるアクセラレーション環境について 長屋貴量 自然科学研究機構分子科学研究所技術課計算科学技術班 概要 GPGPU とは 単純で画一的なデータを一度に大量に処理することに特化したグラフィックカードの演算資源を 画像処理以外の汎用的な目的に応用する技術の一つである 近年 その演算能力は CPU で通常言われるムーアの法則に則った場合とは異なり 飛躍的に向上しており その演算性能に魅力を感じた各分野での応用が広がってきている
Slide 1
電子情報通信学会研究会組込みシステム研究会 (IPSJ-EMB) 2010 年 1 月 28 日 超並列マルチコア GPU を用いた高速演算処理の実用化 NVIDIA Solution Architect 馬路徹 目次 なぜ今 GPU コンピューティングか? CPUの性能向上速度が減速 性能向上 = 並列処理 にGPUコンピューティングが応える CUDAシステムアーキテクチャによる超並列処理の実現
TopSE並行システム はじめに
はじめに 平成 23 年 9 月 1 日 トップエスイープロジェクト 磯部祥尚 ( 産業技術総合研究所 ) 2 本講座の背景と目標 背景 : マルチコア CPU やクラウドコンピューティング等 並列 / 分散処理環境が身近なものになっている 複数のプロセス ( プログラム ) を同時に実行可能 通信等により複数のプロセスが協調可能 並行システムの構築 並行システム 通信 Proc2 プロセス ( プログラム
05-scheduling.ppt
オペレーティングシステム ~ スケジューリング ~ 山田浩史 hiroshiy @ cc.tuat.ac.jp 2014/06/01 復習 : プロセス 実行状態にあるプログラムのこと プログラムの実行に必要なものをひっくるめて指す テキスト領域 データ領域 スタック領域 CPU のレジスタ値 プログラムカウンタ など OS はプロセス単位で管理する メモリ Hard Disk CPU プロセス execute
Introduction Purpose This training course demonstrates the use of the High-performance Embedded Workshop (HEW), a key tool for developing software for
Introduction Purpose This training course demonstrates the use of the High-performance Embedded Workshop (HEW), a key tool for developing software for embedded systems that use microcontrollers (MCUs)
PowerPoint プレゼンテーション
応用数理概論 準備 端末上で cd ~/ mkdir cppwork cd cppwork wget http://271.jp/gairon/main.cpp wget http://271.jp/gairon/matrix.hpp とコマンドを記入. ls とコマンドをうち,main.cppとmatrix.hppがダウンロードされていることを確認. 1 準備 コンパイル c++ -I. -std=c++0x
Microsoft PowerPoint - sales2.ppt
最適化とは何? CPU アーキテクチャに沿った形で最適な性能を抽出できるようにする技法 ( 性能向上技法 ) コンパイラによるプログラム最適化 コンパイラメーカの技量 経験量に依存 最適化ツールによるプログラム最適化 KAP (Kuck & Associates, Inc. ) 人によるプログラム最適化 アーキテクチャのボトルネックを知ること 3 使用コンパイラによる性能の違い MFLOPS 90
untitled
PC murakami@cc.kyushu-u.ac.jp muscle server blade server PC PC + EHPC/Eric (Embedded HPC with Eric) 1216 Compact PCI Compact PCIPC Compact PCISH-4 Compact PCISH-4 Eric Eric EHPC/Eric EHPC/Eric Gigabit
211 年ハイパフォーマンスコンピューティングと計算科学シンポジウム Computing Symposium 211 HPCS /1/18 a a 1 a 2 a 3 a a GPU Graphics Processing Unit GPU CPU GPU GPGPU G
211 年ハイパフォーマンスコンピューティングと計算科学シンポジウム Computing Symposium 211 HPCS211 211/1/18 GPU 4 8 BLAS 4 8 BLAS Basic Linear Algebra Subprograms GPU Graphics Processing Unit 4 8 double 2 4 double-double DD 4 4 8 quad-double
表面RTX入稿
Quadro 2019.04 NVIDIA Quadro NVIDIA Quadro NVIDIA NVIDIA QUADRO BREAKTHROUGH IN EVERY FORM. RTX NVIDIA QUADRO RTX QUADRO RTX FAMILY QUADRO RTX 6000 24 GB 10 Giga Rays/sec QUADRO RTX 4000 8 GB 6 Giga Rays/sec
HP_PPT_Standard_16x9_JP
Autodesk Simulation に最適 HP Z Workstation 最新情報 日本ヒューレット パッカード株式会社ワークステーション市場開発大橋秀樹 HP Workstation 軌跡 新 Z シリーズ初のモバイル製品 2008 年から日本でマーケットシェア No.1 Unix WS や独自グラフィックスなど開発実績
strtok-count.eps
IoT FPGA 2016/12/1 IoT FPGA 200MHz 32 ASCII PCI Express FPGA OpenCL (Volvox) Volvox CPU 10 1 IoT (Internet of Things) 2020 208 [1] IoT IoT HTTP JSON ( Python Ruby) IoT IoT IoT (Hadoop [2] ) AI (Artificial
IntelR Compilers Professional Editions
June 2007 インテル コンパイラー プロフェッショナル エディション Phil De La Zerda 公開が禁止された情報が含まれています 本資料に含まれるインテル コンパイラー 10.0 についての情報は 6 月 5 日まで公開が禁止されています グローバル ビジネス デベロップメント ディレクター Intel Corporation マルチコア プロセッサーがもたらす変革 これまでは
! 行行 CPUDSP PPESPECell/B.E. CPUGPU 行行 SIMD [SSE, AltiVec] 用 HPC CPUDSP PPESPE (Cell/B.E.) SPE CPUGPU GPU CPU DSP DSP PPE SPE SPE CPU DSP SPE 2
![! 行行 CPUDSP PPESPECell/B.E. CPUGPU 行行 SIMD [SSE, AltiVec] 用 HPC CPUDSP PPESPE (Cell/B.E.) SPE CPUGPU GPU CPU DSP DSP PPE SPE SPE CPU DSP SPE 2](/thumbs/88/115445009.jpg "! 行行 CPUDSP PPESPECell/B.E. CPUGPU 行行 SIMD [SSE, AltiVec] 用 HPC CPUDSP PPESPE (Cell/B.E.) SPE CPUGPU GPU CPU DSP DSP PPE SPE SPE CPU DSP SPE 2")
! OpenCL [Open Computing Language] 言 [OpenCL C 言 ] CPU, GPU, Cell/B.E.,DSP 言 行行 [OpenCL Runtime] OpenCL C 言 API Khronos OpenCL Working Group AMD Broadcom Blizzard Apple ARM Codeplay Electronic Arts Freescale
XACCの概要
2 global void kernel(int a[max], int llimit, int ulimit) {... } : int main(int argc, char *argv[]){ MPI_Int(&argc, &argc); MPI_Comm_rank(MPI_COMM_WORLD, &rank); MPI_Comm_size(MPI_COMM_WORLD, &size); dx
Insert your Title here
マルチコア マルチスレッド環境での静的解析ツールの応用 米 GrammaTech 社 CodeSonar によるスレッド間のデータ競合の検出 2013 GrammaTech, Inc. All rights reserved Agenda 並列実行に起因する不具合の摘出 なぜ 並列実行されるプログラミングは難しいのか データの競合 デッドロック どのようにして静的解析ツールで並列実行の問題を見つけるのか?
PowerPoint プレゼンテーション
高性能計算基盤 第 7 回 CA1003: 主記憶共有型システム http://arch.naist.jp/htdocs-arch3/ppt/ca1003/ca1003j.pdf Copyright 2019 奈良先端大中島康彦 1 2 3 4 マルチスレッディングとマルチコア 5 6 7 主記憶空間の数が 複数 か 1 つ か 8 ただしプログラムは容易 9 1 つの主記憶空間を共有する場合 10
AMD/ATI Radeon HD 5870 GPU DEGIMA LINPACK HD 5870 GPU DEGIMA LINPACK GFlops/Watt GFlops/Watt Abstract GPU Computing has lately attracted
DEGIMA LINPACK Energy Performance for LINPACK Benchmark on DEGIMA 1 AMD/ATI Radeon HD 5870 GPU DEGIMA LINPACK HD 5870 GPU DEGIMA LINPACK 1.4698 GFlops/Watt 1.9658 GFlops/Watt Abstract GPU Computing has
Microsoft PowerPoint - ARCEMB08HayashiSlides.ppt [互換モード]
![Microsoft PowerPoint - ARCEMB08HayashiSlides.ppt [互換モード]](/thumbs/93/113153048.jpg "Microsoft PowerPoint - ARCEMB08HayashiSlides.ppt [互換モード]")
演算 / メモリ性能バランスを考慮した CMP 向けオンチップ メモリ貸与法の提案 九州大学 林徹生今里賢一井上弘士村上和彰 1 発表手順 背景 目的 演算 / メモリ性能バランシング 概要 アクセスレイテンシの削減とオーバーヘッド 提案手法の実現方法 着目する命令 (Cell プロセッサへの ) 実装 性能評価 姫野ベンチマーク Susan@MiBench おわりに 2 チップマルチプロセッサ (CMP)
スパコンに通じる並列プログラミングの基礎
2018.09.10 furihata@cmc.osaka-u.ac.jp ( ) 2018.09.10 1 / 59 furihata@cmc.osaka-u.ac.jp ( ) 2018.09.10 2 / 59 Windows, Mac Unix 0444-J furihata@cmc.osaka-u.ac.jp ( ) 2018.09.10 3 / 59 Part I Unix GUI CUI:
TSUBAME2.0におけるGPUの 活用方法
GPU プログラミング 基礎編 東京工業大学学術国際情報センター 1. GPU コンピューティングと TSUBAME2.0 スーパーコンピュータ GPU コンピューティングとは グラフィックプロセッサ (GPU) は グラフィック ゲームの画像計算のために 進化を続けてきた 現在 CPU のコア数は 2~12 個に対し GPU 中には数百コア その GPU を一般アプリケーションの高速化に利用! GPGPU
マルチコア時代の並列プログラミング
マルチコア時代の 並列プログラミング ~ ロックとメモリオーダリング ~ 中村実 nminoru@nminoru.jp http://www.nminoru.jp/~nminoru/ まずは自己紹介を 電機メーカー勤務のエンジニア Java VM 特に並列 GC JIT コンパイラの研究 開発 Java 系雑誌にときどき寄稿 最近は IA-64 と戯れる日々 趣味で Web に細々とプログラミングのメモを綴る日々
Microsoft PowerPoint - GPGPU実践基礎工学（web）.pptx
GPU のメモリ階層 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 GPU のメモリ階層 グローバルメモリ 共有メモリ モザイク処理への適用 コンスタントメモリ 空間フィルタへの適用 577 GPU の主要部品 基盤 GPU( チップ )+ 冷却部品 画面出力端子 電源入力端子 メモリ 特性の把握が重要 電源入力端子 画面出力端子 メモリ チップ PCI Ex 端子 http://www.geforce.com/whats
( CUDA CUDA CUDA CUDA ( NVIDIA CUDA I
GPGPU (II) GPGPU CUDA 1 GPGPU CUDA(CUDA Unified Device Architecture) CUDA NVIDIA GPU *1 C/C++ (nvcc) CUDA NVIDIA GPU GPU CUDA CUDA 1 CUDA CUDA 2 CUDA NVIDIA GPU PC Windows Linux MaxOSX CUDA GPU CUDA NVIDIA
26102 (1/2) LSISoC: (1) (*) (*) GPU SIMD MIMD FPGA DES, AES (2/2) (2) FPGA(8bit) (ISS: Instruction Set Simulator) (3) (4) LSI ECU110100ECU1 ECU ECU ECU ECU FPGA ECU main() { int i, j, k for { } 1 GP-GPU
catalog_quadro_series_2018
Quadro 2018.0 NVIDIA Quadro NVIDIA Quadro NVIDIA NVIDIA QUADRO BREAKTHROUGH IN EVERY FORM. NVIDIA Quadro GV100VoltaGPU32GBHBM2 CUDA5120 32GB Tensor Core60Deep Learning NVLink 2 NVIDIA Quadro GV100GPU PCIe
CPU Levels in the memory hierarchy Level 1 Level 2... Increasing distance from the CPU in access time Level n Size of the memory at each level 1: 2.2
FFT 1 Fourier fast Fourier transform FFT FFT FFT 1 FFT FFT 2 Fourier 2.1 Fourier FFT Fourier discrete Fourier transform DFT DFT n 1 y k = j=0 x j ω jk n, 0 k n 1 (1) x j y k ω n = e 2πi/n i = 1 (1) n DFT
IPSJ SIG Technical Report Vol.2013-HPC-138 No /2/21 GPU CRS 1,a) 2,b) SpMV GPU CRS SpMV GPU NVIDIA Kepler CUDA5.0 Fermi GPU Kepler Kepler Tesla
GPU CRS 1,a),b) SpMV GPU CRS SpMV GPU NVIDIA Kepler CUDA5.0 Fermi GPU Kepler Kepler Tesla K0 CUDA5.0 cusparse CRS SpMV 00 1.86 177 1. SpMV SpMV CRS Compressed Row Storage *1 SpMV GPU GPU NVIDIA Kepler