GPUDirect の現状整理 multi-gpu に取組むために G-DEP チーフエンジニア河井博紀 (kawai@gdep.jp)
名称 : 日本 GPU コンピューティングパートナーシップ (G-DEP) 所在 : 東京都文京区本郷 7 丁目 3 番 1 号東京大学アントレプレナープラザ, 他工場 URL http://www.gdep.jp アライアンスパートナー コアテクノロジーパートナー NVIDIA JAPAN ソリューションパートナー株式会社エルザジャパン みずほ情報総研株式会社 株式会社アーク情報システム 日本ネスト株式会社株式会社システム計画研究所 有限会社イワタシステムサポート サーヴァンツインターナショナル株式会社 株式会社ソフテック 株式会社サイコム プロメテック ソフトウェア株式会社 トーワ電機株式会社
小型 PC から ノード内 multi-gpu 高性能ワークステーション 1 2 3 4 複数ノード multi-gpu InfiniBand GPU クラスタまで 1 4 8 12 16
Outline mult i-gpu プログラミング概論 GPUDirect とは? GPUDirect の現状整理
We assume CUDA 4.0 or later Fermi or later Tesla 64-bit Linux / 64-bit Windows with TCC driver
Why multi-gpu? 性能を Scaling-up させるため メモリを増やして 大きな問題に取組むため
ノード内 Tesla K20 2 枚で Game of Life 各セルは alive( =1) or dead( =0) 誕生 Rules 過疎死 Grid Size:N N 境界条件 :periodic b.c. 過密死
ノード内 Tesla K20 2 枚で Game of Life Tesla 0 Grid Size:N N 境界条件 :periodic b.c. Tesla 1
ノード内 Tesla K20 2 枚で Game of Life // Tesla 0 を操作 cudasetdevice( 0 ); kernel<<< >>>( ); Tesla 0 // Tesla 1 を操作 cudasetdevice( 1 ); kernel<<< >>>( ); Tesla 1
ノード内 Tesla K20 2 枚で Game of Life 境界 Tesla 0 Tesla 0 の境界の更新には Tesla 1 の境界が必要 Tesla 1
ノード内 Tesla K20 2 枚で Game of Life 境界 Tesla 0 ghost 領域 隣の境界を copy Tesla 1
ノード内 Tesla K20 2 枚で Game of Life 境界 ghost を参照する事で境界も update できる ghost 領域 隣の境界を copy Tesla 0 Tesla 1
ノード内 Tesla K20 2 枚で Game of Life 境界 ghost を参照する事で境界も update できる Tesla 0 ghost 領域 Tesla 1
ノード内 Tesla K20 2 枚で Game of Life 境界 ghost を参照する事で境界も update できる Tesla 0 ghost 領域 Tesla 1
ノード内 Tesla K20 2 枚で Game of Life 境界 ghost を参照する事で境界も update できる Tesla 0 ghost 領域 Tesla 1
ノード内 Tesla K20 2 枚で Game of Life 1 update Tesla 0 2 copy して ghost を更新 1 update Tesla 1
もう少し工夫できます
ノード内 Tesla K20 2 枚で Game of Life copy 時間の 隠蔽 が重要に Kernel 処理 と データ送受信 は同時に行える naive update copy update copy 隠蔽 境界 update copy 内部 update 境界 update copy 内部 update 短縮!
ノード内 Tesla K20 2 枚で Game of Life 1 境界 update 2 内部 update 1 境界 update 1 境界 update 2 内部 update 2 同時に copy して ghost 更新 Tesla 0 1 境界 update Tesla 1
ノード内 Tesla K20 2 枚で Game of Life K20 1 K20 2 (naive) Grid Size:2048 2048 更新回数 :10 万回境界条件 :periodic b.c. K20 2 ( 隠蔽 ) 0 10 20 30 40 50 Processing Time [sec]
ノード内 Tesla K20 2 枚で Game of Life for ( n=0; n<num_steps; n++) { for ( i=0; i<2; i++) { cudasetdevice( i ); update_bnd<<<, stream_up[i]>>>( ); cudaeventrecord( event[i], stream_up[i] ); update_inner<<<, stream_up[i]>>>( ); } 境界 update kernel 内部 update kernel } for ( i=0; i<2; i++) { cudastreamwaitevent( stream_cpy[i], event[i] ); cudamemcpyasync(, stream_cpy[i] ); } for ( i=0; i<2; i++) cudamemcpyasync(, stream_cpy[i] ); for ( i=0; i<2; i++) { cudasetdevice( i ); cudadevicesynchronize(); } 境界の update 完了まで cudamemcpyasync を停止 双方向 cudamemcpyasync 2 セット ここで同期
Outline GPUDirect とは? GPUDirect の現状整理
GPUDirect とは 異なる Tesla 間のデータ転送 を高速に行う機能の事です
隠蔽できるのに 恩恵はあるのか?
Tesla の枚数 境界 内部 境界 内部 copy copy 境界 内部 境界 内部 copy copy 境界 内部 copy 境界内部 copy 隠蔽できなくなった
GPUDirect の恩恵はあるのか? Teslaの枚数を増やす可能性を考慮して使っておくべき 原理的に隠蔽出来ない時は ある 隠蔽してる暇がない時も ある
Outline GPUDirect の現状整理
GPUDirect には現在 version 1 と version 2 があります
GPUDirect ver.1 InfiniBand クラスタ専用 = ノードをまたぐ高速転送 GPUDirect ver.2 = ノード内の高速転送
Summary GPUDirect 用途ダイレクト転送? CUDA 使用方法正式リリース ver.1 IB クラスタノード間 済 ver.2 ノード内済 転送を隠蔽できない / してない時に GPUDirect は有用
GPUDirect ver.2 はノード内のダイレクト転送です
No GPUDirect ver.2 With GPUDirect ver.2 CPU メインメモリ CPU メインメモリ Tesla 0 遠回り Direct! Tesla 1 Tesla 0 Tesla 1 M/B M/B
GPUDirect ver.2 を使用するには cudadeviceenablepeeraccess() を宣言するだけ
GPUDirect ver.2 の実装 with GPUDirect ver.2 cudasetdevice(0); cudadeviceenablepeeraccess(1, 0); cudamalloc(&p0, size); cudasetdevice(1); Tesla 1 への GPUDv2を有効化 CPU cudamalloc(&p1, size); Direct! メインメモリ *p0 Tesla 0 cudamemcpy(p1, p0, size, cudamemcpydefault); // UVA (CUDA 4.0) により異なる GPU 間の転送が可能 M/B *p1 Tesla 1
GPUDirect ver.2 の実装 cudasetdevice(0); // cudadeviceenablepeeraccess(1, 0); cudamalloc(&p0, size); cudasetdevice(1); cudamalloc(&p1, size); no GPUDirect ver.2 CPU Fallback メインメモリ *p0 Tesla 0 cudamemcpy(p1, p0, size, cudamemcpydefault); // UVA (CUDA 4.0) により異なる GPU 間の転送が可能 // 但し Nsight のタイムラインによるとパイプライン処理されており cudamemcpy 2 回よりは高速 M/B *p1 Tesla 1
Bandwidth [GB/s] cudamemcpy 6 5 4 with GPUDv2 no GPUDv2 (fallback) Fallback wins 3 2 1 latency with GPUDv2: 11μs no GPUDv2: 20μs int: 200 万個 8MB single: 200 万個 double: 100 万個 0 4B 16B 64B 256B 1KB 4KB 16KB 64KB 256KB 1MB 4MB 16MB 64MB 256MB 転送データサイズ CentOS 6.2, NV Driver 319.32, CUDA5, K20c, X79Chipset Performance may vary based on system HDW and config.
Bandwidth [GB/s] 10 9 8 7 6 5 4 3 2 1 0 latency 双方向 cudamemcpyasync with GPUDv2 no GPUDv2 with GPUDv2: 14μs no GPUDv2: 31μs (fallback) Fallback wins 8B 32B 128B 512B 2KB 8KB 32KB 128KB 512KB 2MB 8MB 32MB 128MB 512MB 転送データサイズ int: 800 万個 32MB single: 800 万個 double: 400 万個 CentOS 6.2, NV Driver 319.32, CUDA5, K20c, X79Chipset Performance may vary based on system HDW and config.
再度 Tesla K20 2 枚で Game of Life K20 1 K20 2 (naive) no GPUDv2 with GPUDv2 11% 高速化 Grid Size:2048 2048 更新回数 :10 万回境界条件 :periodic b.c. K20 2 ( 隠蔽 ) no GPUDv2 with GPUDv2 隠蔽できているので変わらず 0 10 20 30 40 50 Processing Time [sec]
Summary GPUDirect 用途ダイレクト転送? CUDA 使用方法正式リリース ver.1 IB クラスタノード間 済 ver.2 ノード内 Yes 4.0 cudadeviceenablepeeraccess() 同じIOHの必要あり 済 転送を隠蔽できない / してない時に GPUDirect は有用
GPUDirect ver.1 はマーケティング ネームです
GPUDirect ver.1 ダイレクト転送ではありません メインメモリをしっかり通ります
InfiniBand InfiniBand 用 buffer CPU メインメモリ rank0 の M/B *Ad Tesla 1 Tesla 用 buffer InfiniBand クラスタの 1 ノード
MPI_Send naive なコード CPU rank0 の M/B メインメモリ *A *A *Ad Tesla 1 InfiniBand クラスタの 1 ノード // 後述の CUDA-Aware MPI を除いて // MPI 関数にデバイスポインタは渡せない if ( myrank == 0 ) { cudamemcpy( A, Ad, ); MPI_Send( A, ); } else if ( myrank == 1 ) { } MPI_Recv( B, ); cudamemcpy( Bd, B, );
MPI_Send CPU rank0 の M/B メインメモリ *A *Ad Tesla 1 InfiniBand クラスタの 1 ノード GPUDirect ver.1 を使用 cudamallochost( &A, size ); cudamallochost( &B, size ); if ( myrank == 0 ) { cudamemcpy( A, Ad, ); MPI_Send( A, ); } else if ( myrank == 1 ) { } MPI_Recv( B, ); cudamemcpy( Bd, B, );
MPI_Send CPU rank0 の M/B メインメモリ パイプライン処理 *Ad Tesla 1 InfiniBand クラスタの 1 ノード CUDA-Aware MPI を使用 MVAPICH2 ver. 1.8 以上 OpenMPI ver. 1.7 以上 Cray, IBM if ( myrank == 0 ) { MPI_Send( Ad, ); } else if ( myrank == 1 ) { } MPI_Recv( Bd, ); でのみ可能 デバイスポインタを渡せる
MPI_Isend CPU rank0 の M/B メインメモリ パイプライン処理 *Ad Tesla 1 InfiniBand クラスタの 1 ノード CUDA-Aware MPI を使用 MVAPICH2 ver. 1.8 以上 OpenMPI ver. 1.7 以上 Cray, IBM if ( myrank == 0 ) { MPI_Isend( Ad, ); } else if ( myrank == 1 ) { } MPI_Irecv( Bd, ); でのみ可能 デバイスポインタを渡せる
Bandwidth [GB/s] 3.5 3 2.5 2 CUDA-Aware MPI with GPUDv1 no GPUDv1 MPI_Send/Recv 1.5 1 0.5 latency CUDA-Aware : 16μs with GPUDv1: 22μs no GPUDv1: 16μs 0 4B 16B 64B 256B 1KB 4KB 16KB 64KB 256KB 1MB 4MB 16MB 64MB 256MB 転送データサイズ CentOS 6.2, NV Driver 319.32, CUDA5, QDR, K20c, X79Chipset Performance may vary based on system HDW and config.
Bandwidth [GB/s] 6 5 4 CUDA-Aware MPI with GPUDv1 no GPUDv1 双方向 MPI_Isend/Irecv 3 2 1 latency CUDA-Aware : 16μs with GPUDv1: 22μs no GPUDv1: 17μs 0 8B 32B 128B 512B 2KB 8KB 32KB 128KB 512KB 2MB 8MB 32MB 128MB 512MB 転送データサイズ CentOS 6.2, NV Driver 319.32, CUDA5, QDR, K20c, X79Chipset Performance may vary based on system HDW and config.
G-DEP は CUDA-Aware MPI をインストールした GPU クラスタの販売を開始しております Drawer InfiniBand Switch 計算ノード 計算ノード 計算ノード さらに大規模なクラスタも販売しております 是非お気軽に kawai@gdep.jp までお問い合わせ下さい 計算ノード Gigabit Switch
Summary GPUDirect 用途ダイレクト転送? CUDA 使用方法正式リリース ver.1 IB クラスタノード間 No 3.1 中継メモリを pinned に CUDA-Aware MPI が最速 済 ver.2 ノード内 Yes 4.0 cudadeviceenablepeeraccess() 同じ IOH の必要あり 済 転送を隠蔽できない / してない時に GPUDirect は有用
GPUDirect ver.3 はまだ正式リリースされてません
ノードをまたぐダイレクト転送
Summary GPUDirect 用途ダイレクト転送? CUDA 使用方法正式リリース ver.1 IB クラスタノード間 No 3.1 中継メモリを pinned に CUDA-Aware MPI が最速 済 ver.2 ノード内 Yes 4.0 cudadeviceenablepeeraccess() 同じ IOH の必要あり 済 ver.3 IB クラスタノード間 Yes 5.0? 2013 Q4 転送を隠蔽できない / してない時に GPUDirect は有用