GPUDirectの現状整理
名称: 日本GPUコンピューティングパートナーシップ
(G-DEP)
所在: 東京都文京区本郷7丁目3番1号 東京大学アントレプレナープラザ， 他工場
URL
http://www.gdep.jp
アライアンスパートナー
コアテクノロジーパートナー
NVIDIA JAPAN
ソリューションパートナー
株式会社エルザジャパン、みずほ情報総研株式会社、
株式会社アーク情報システム、日本ネスト株式会社
株式会社システム計画研究所、有限会社イワタシステムサポート、
サーヴァンツインターナショナル株式会社、株式会社ソフテック、
小型PCから…
高性能ワークステーション、
InfiniBand GPUクラスタ まで。
ノード内 multi-GPU
複数ノード multi-GPU
1 2 3 4
1
4
8
12
16
• • • • • • • • • • • •Outline
mult
i-GPU
プログラミング概論
GPUDirect
とは？
We assume …
・Fermi or later Tesla
・CUDA 4.0 or later
Why multi-GPU ?
・性能を Scaling-up させるため
・各セルは alive( =1) or dead( =０)
・Rules
誕生
過疎死
過密死
ノード内
Tesla K20 ×
2枚で
“Game of Life”
Grid Size：N
×
N
Tesla 0
×
Tesla 0
Tesla 1
cudaSetDevice( 0 );
kernel<<<…>>>(…);
cudaSetDevice( 1 );
kernel<<<…>>>(…);
// Tesla 0 を操作
// Tesla 1 を操作
Tesla 0
Tesla 0 の境界の更新にはTesla 1の境界が必要
境界
Tesla 0
Tesla 1
境界
隣の境界をcopy
“
ghost領域
”
Tesla 0
隣の境界をcopy
境界
ghost を参照する事で
境界も update できる。
“
ghost領域
”
Tesla 0
Tesla 1
境界
ghost を参照する事で
境界も update できる。
“
ghost領域
”
Tesla 0
境界
“
ghost領域
”
ghost を参照する事で
境界も update できる。
Tesla 0
Tesla 1
境界
“
ghost領域
”
ghost を参照する事で
境界も update できる。
Tesla 0
① update
① update
②
copy
して
ghost
を更新
copy
時間の “隠蔽” が重要に
・「Kernel処理」と「データ送受信」は同時に行える。
update
copy
△
naive
境界
update
◎
隠蔽
update
copy
内部
update
境界
update
内部
update
Tesla 0
Tesla 1
② 内部
update
② 同時に
、
copyしてghost更新
①
境界
update
①
境界
update
② 内部
update
①
境界
update
①
更新回数：10 万回
Grid Size：2048
×
2048
K20 ×1
K20 ×2
(naive)
K20 ×2
(隠蔽)
cudaSetDevice( i );
update_bnd<<<…, stream_up[i]>>>(…);
for ( i=0; i<2; i++) {
}
cudaEventRecord( event[i], stream_up[i] );
update_inner<<<…, stream_up[i]>>>(…);
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 ( n=0; n<NUM_STEPS; n++)
{
cudaSetDevice( i );
for ( i=0; i<2; i++) {
cudaDeviceSynchronize();
境界
の
update完了まで、
cudaMemcpyAsync
を停止
内部 update kernel
境界 update kernel
双方向 cudaMemcpyAsync
× 2セット
ここで同期
Outline
GPUDirect
とは？
GPUDirect
とは、
「異なる
Tesla
間のデータ転送」
隠蔽できるのに、
恩恵はあるのか？
copy
copy
境界 内部
copy
copy
copy
copy
隠蔽できなく
なった…
境界
境界
境界
境界
境界
内部
内部
内部
内部
内部
・
・
・
Teslaの枚数
・隠蔽してる暇がない時も、ある。
・原理的に隠蔽出来ない時は、ある。
・Teslaの枚数を増やす可能性を考慮して
使っておくべき。
Outline
GPUDirect
には現在、
「version 1」と「version 2」
があります。
InfiniBandクラスタ専用
ノードをまたぐ高速転送
GPUDirect ver.1
GPUDirect ver.2
ノード内の高速転送
＝
＝
GPUDirect
用途
ダイレクト転送？
CUDA
使用方法
正式リリース
ver.1
IBクラスタ
ノード間
済
ver.2
ノード内
済
GPUDirect ver.2
CPU
メインメモリ
Tesla 0
Tesla 1
CPU
メインメモリ
Tesla 0
Tesla 1
No
GPUDirect ver.2
With
GPUDirect ver.2
cudaDeviceEnablePeerAccess()
を宣言するだけ。
GPUDirect ver.2
CPU
メインメモリ
Tesla 0
Tesla 1
cudaDeviceEnablePeerAccess(1, 0);
cudaMemcpy(p1, p0, size, cudaMemcpyDefault);
*p0
*p1
cudaSetDevice(0);
cudaMalloc(&p0, size);
cudaSetDevice(1);
cudaMalloc(&p1, size);
Direct !
with GPUDirect ver.2
GPUDirect ver.2 の実装
// ※ UVA (CUDA ≧ 4.0) により異なるGPU間の転送が可能
Tesla “1” への
GPUDv2を有効化
// cudaDeviceEnablePeerAccess(1, 0);
cudaSetDevice(0);
cudaMalloc(&p0, size);
cudaSetDevice(1);
cudaMalloc(&p1, size);
M/B
CPU
メインメモリ
Tesla 0
Tesla 1
*p0
*p1
no GPUDirect ver.2
Fallback …
// ※ UVA (CUDA ≧ 4.0) により異なるGPU間の転送が可能
// ※ 但し、Nsightのタイムラインによるとパイプライン処理
されており、cudaMemcpy ×2回よりは高速
GPUDirect ver.2 の実装
1
2
3
4
5
6
Bandwi
dth [GB/
s]
cudaMemcpy
with GPUDv2
no GPUDv2
Fallback wins
〜
8MB
single: 200万個
double: 100万個
int:
200万個
latency
with GPUDv2: 11μs
no GPUDv2: 20μs
(fallback)
Fallback wins
0
1
2
3
4
5
6
7
8
9
10
8B
32B
128B
512B
2KB
8KB
32KB 128KB 512KB
2MB
8MB
32MB 128MB 512MB
Ban
dwi
dth [GB/
s]
双方向 cudaMemcpyAsync
with GPUDv2
no GPUDv2
〜
32MB
single: 800万個
double: 400万個
int:
800万個
latency
with GPUDv2: 14μs
再度
、
Tesla K20 ×
2枚で
Game of Life
更新回数：10 万回
Grid Size：2048
×
2048
K20 ×1
K20 ×2
(naive)
K20 ×2
(隠蔽)
with GPUDv2
with GPUDv2
11%高速化
隠蔽できている
ので変わらず
no GPUDv2
no GPUDv2
GPUDirect
用途
ダイレクト転送？
CUDA
使用方法
正式リリース
ver.1
IBクラスタ
ノード間
済
ver.2
ノード内
Yes
4.0〜
cudaDeviceEnablePeerAccess()
※ 同じIOHの必要あり
済
Summary
・転送を隠蔽できない/してない時に
GPUDirectは有用。
GPUDirect ver.1
GPUDirect ver.1
ダイレクト転送ではありません。
メインメモリ
Tesla 1
*Ad
Tesla
用
buffer
InfiniBand
用
buffer
CPU
InfiniBand
rank0
の
M/B
メインメモリ
Tesla 1
*Ad
MPI_Send
{
if ( myrank == 0 )
cudaMemcpy( A, Ad, … );
MPI_Send( A, … );
naive なコード
// 後述の “CUDA-Aware MPI” を除いて
// MPI関数にデバイスポインタは渡せない
*A
*A
CPU
else if ( myrank == 1 )
}
{
cudaMemcpy( Bd, B, … );
MPI_Recv( B, … );
}
CPU
メインメモリ
Tesla 1
MPI_Send
{
if ( myrank == 0 )
cudaMemcpy( A, Ad, … );
MPI_Send( A, … );
cudaMallocHost( &A, size );
cudaMallocHost( &B, size );
GPUDirect ver.1を使用
*Ad
*A
else if ( myrank == 1 )
}
{
MPI_Recv( B, … );
rank0
の
M/B
CPU
メインメモリ
Tesla 1
*Ad
OpenMPI
ver. 1.7 以上
Cray, IBM
でのみ
可能
CUDA-Aware MPIを使用
MVAPICH2 ver. 1.8 以上
{
if ( myrank == 0 )
else if ( myrank == 1 )
MPI_Send( Ad, … );
}
{
MPI_Recv( Bd, … );
}
デバイスポインタ
パイプライン処理
MPI_Send
CPU
メインメモリ
Tesla 1
*Ad
OpenMPI
ver. 1.7 以上
Cray, IBM
でのみ
可能
CUDA-Aware MPIを使用
MVAPICH2 ver. 1.8 以上
{
if ( myrank == 0 )
else if ( myrank == 1 )
MPI_Isend( Ad, … );
}
{
パイプライン処理
MPI_Isend
0
0.5
1
1.5
2
2.5
3
3.5
4B
16B
64B
256B
1KB
4KB
16KB
64KB 256KB 1MB
4MB
16MB 64MB 256MB
Bandwi
dth [GB/
s]
MPI_Send/Recv
CUDA-Aware MPI
with GPUDv1
no GPUDv1
latency
with GPUDv1: 22μs
no GPUDv1: 16μs
CUDA-Aware : 16μs
1
2
3
4
5
6
Bandwi
dth [GB/
s]
双方向 MPI_Isend/Irecv
CUDA-Aware MPI
with GPUDv1
no GPUDv1
latency
with GPUDv1: 22μs
no GPUDv1: 17μs
CUDA-Aware : 16μs
Drawer InfiniBand Switch 計算ノード 計算ノード 計算ノード 計算ノード Gigabit Switch
