複数GPUの利用 - 長岡技術科学大学

第14回複数GPUの利用
長岡技術科学大学
電気電子情報工学専攻
出川智啓
今回の内容

並列アーキテクチャと並列処理の分類

OpenMP

複数GPUの利用

GPU Directによるデータ通信
1040
先端GPGPUシミュレーション工学特論
2015/07/16
複数のGPUを利用する目的

grouseの1ノードには4台のGPUを搭載




Tesla M2050
1T FLOPS/台
3 GB/台
4台全てを使う事で期待できる性能

GPUを1台だけ使った時の4倍高速化

GPU1台では収まらないサイズの問題を計算
1041
先端GPGPUシミュレーション工学特論
2015/07/16
複数のGPUを利用する目的

4台全てを使う事で期待できる性能

GPUを1台だけ使った時の4倍高速化
 高速化は達成可能
 複数のGPUを使うための手続きやデータの移動が必要
 GPU間のデータ交換が必要になると速度が低下

GPU1台では収まらないサイズの問題を計算
 利用できるメモリ容量はGPUの台数に比例
 1台では実行できない大規模な問題も実行可能
1042
先端GPGPUシミュレーション工学特論
2015/07/16
複数のGPUが利用できる環境

1台の計算機に複数のGPUが搭載



複数のGPUを搭載するために大きな筐体が必要
1台のPCで完結するため，プログラム中でネットワーク通信
を行う必要がない
帯域の狭いPCI‐Exバスを複数のGPUで取り合うため，デー
タ移動の効率は悪い
PCI‐Exバス
1043
先端GPGPUシミュレーション工学特論
2015/07/16
複数のGPUが利用できる環境

1台のGPUが搭載された計算機をネットワークで接続



小さい筐体のPC（キューブ型PCなど）を並べて作成
PCごとにCPUやメモリが必要
複数GPUの利用にはネットワーク通信が必須
PCI‐Ex
PCI‐Ex
PCI‐Ex
PCI‐Ex
高速ネットワーク,LAN
1044
先端GPGPUシミュレーション工学特論
2015/07/16
複数のGPUが利用できる環境

複数のGPUが搭載された複数の計算機をネットワー
クで接続


スーパーコンピュータ等の大規模計算環境
性能の大半はGPUによって実現
2~4CPUs
2~4GPUs
2~4CPUs
2~4GPUs
2~4CPUs
2~4GPUs
・・・
高速ネットワーク
1045
先端GPGPUシミュレーション工学特論
2015/07/16
複数のGPUが利用できる環境


CUDAにおける複数GPU利用

新しいバージョン
 GPU間通信を行う命令が用意されたり，複数のGPUを利用
する数値計算ライブラリが登場
 （古いバージョンと比べれば）利用しやすくなっている

古いバージョン
 複数のGPUを利用するための機能は未搭載
従来から並列計算に利用されていたライブラリと併用

1046
OpenMPやMPIなど
先端GPGPUシミュレーション工学特論
2015/07/16
複数GPUの利用（CUDA 3.2以前）

CUDAの制約




OpenMPやMPIを利用



1スレッドは複数のGPUを利用できない
GPUは必ずCPUの1スレッド（もしくは1プロセス）に割当て
4台のGPUを利用するためには4スレッド（4プロセス）必要
並列計算を行うプログラムを作成
スレッドやプロセス番号を基に利用するGPUを決定，処理を
実行
従来の並列計算の素直な拡張
1047
先端GPGPUシミュレーション工学特論
2015/07/16
複数GPUの利用（CUDA 4.0以降）

CUDAの制約の緩和

1スレッドから複数のGPUを利用（切替）可能

2台のGPU間でデータ通信を行う命令の追加
 GPU Direct v2.0 peer‐to‐peer

複数GPU間で統一されたメモリ空間を利用可能
 Unified Virtual Addressing
1048
先端GPGPUシミュレーション工学特論
2015/07/16
複数GPUの利用（CUDA 4.0以降）

複数GPUを利用するプログラムの作成が容易

cudaSetDevice()でカレントデバイスを指定
 CUDAの命令はそのカレントデバイスに対して発行

別のデバイスを利用する際はcudaSetDevice()でカレン
トデバイスを切替

複数のGPUへのデータコピー，複数のGPUでカーネルの起
動が容易
1049
先端GPGPUシミュレーション工学特論
2015/07/16
複数GPU化の例題

ベクトル和
1.OpenMPを利用する方法

従来の並列計算の拡張
2.CUDA4.0の機能を利用する方法
1050
先端GPGPUシミュレーション工学特論
2015/07/16
並列アーキテクチャの分類

システムの特徴付け，プロセッサの分類

Flynnの分類






1051
並列アーキテクチャのグループ分け
データの処理と命令の並列性に着目
SISD
SIMD
MISD
MIMD
単一命令単一データ
単一命令複数データ
複数命令単一データ
複数命令複数データ
先端GPGPUシミュレーション工学特論
2015/07/16
Single Instruction Multiple Data streams

単一命令複数データ




数値シミュレーションに最適
命令
数学のベクトルや配列計算の概念
に一致


複数のまとまったデータに対して同じ演算を同時に実行
命令は一つ，その命令が同時に多くのデータに対して適用
されるアーキテクチャ
ベクトルプロセッサとも呼ばれる
データ
GPUもここに分類
A0
B0
A1
B1
A2
A3
1052
先端GPGPUシミュレーション工学特論
+
B2
B3
2015/07/16
プロセス

OSから資源（コア，メモリ，外部記憶など）を割り当て
られ，実行状態（または待機状態）にあるプログラム

システムプロセス


OSの実行に関係するプログラム
ユーザプロセス

1053
ユーザ権限で実行されているプログラム
先端GPGPUシミュレーション工学特論
2015/07/16
マルチプロセス

複数のプロセスが存在し，並列に実行



シングルプロセス
マルチプロセス
マルチプロセスに対応したOSが必要


プロセスが一つのみ
プロセスが二つ以上
現在のOSはマルチプロセスに対応
シングルコアCPU一つでもマルチプロセスが可能


1054
OSが複数のプロセスを切替
複数のプロセスが並列に実行されているように見せる
先端GPGPUシミュレーション工学特論
2015/07/16
マルチプロセス

各プロセスに専用のメモリ領域を割り当て


CPUやメモリは複数のプログラムに割り当てられる
プログラムはCPUやメモリを独占しているように振る舞う
プロセスA
スレッド
メモリ
CPU
メモリ
OS
プロセスB
スレッド
メモリ
1055
先端GPGPUシミュレーション工学特論
2015/07/16
スレッド

プログラムの処理の最小実行単位

プロセス内で複数のスレッドが存在


シングルスレッド
マルチスレッド
1プロセスに一つのスレッドのみ
1プロセスに二つ以上のスレッド
シングルプロセス
シングルスレッド
マルチスレッド
マルチプロセス
シングルスレッド
マルチスレッド
1056
先端GPGPUシミュレーション工学特論
2015/07/16
マルチスレッド


一つのプロセスに二つ以上のスレッドが存在
一つのプロセスに専用のメモリ領域を割当

プロセス内の複数のスレッドはメモリ領域を共有
プロセスA
スレッド
スレッド
メモリ
CPU
メモリ
OS
プロセスB
スレッド
スレッド
メモリ
1057
先端GPGPUシミュレーション工学特論
2015/07/16
並列計算機システム

並列処理の基本



処理を何らかの方法で分割
分割した処理をプロセッサ（やコア）に割り当て同時に処理
並列計算機システム





1058
複数のプロセッサをもつ
主にメモリに違いがある
共有メモリシステム
分散メモリシステム
ハイブリッドシステム
先端GPGPUシミュレーション工学特論
2015/07/16
共有メモリシステム



複数のプロセッサがメモリ空間を共有
分割した処理は各プロセッサ上で並列的に処理
共有されたメモリ空間上の変数は全てのCPU（やコ
ア）からアクセス（読み書き）可能

他からアクセスされない変数を持つことも可能
CPU
CPU
CPU
CPU
メモリ
1059
先端GPGPUシミュレーション工学特論
2015/07/16
並列処理の分類

タスク並列


データ並列


独立なタスクを異なるCPU・コアで同時に実行
独立なタスクが処理するデータを分割し，異なるCPU・コアがデータ
を参照し，処理を実行
Embarrassingly parallel (perfectly parallel)

1060
各CPU・コアが同じタスクを異なるパラメータで実行
 GPUが各ピクセルの色を決定し，ディスプレイに描画する処理
 あるタスクに対してパラメータの影響を調査するような問題
先端GPGPUシミュレーション工学特論
2015/07/16
データ並列
独立な処理A,B,Cが取り扱うデータを分割して実行


逐次処理
処理A
コア

処理B
処理C
並列処理
コア1
処理A
処理B
処理C
コア2
処理A
処理B
処理C
高速化
1061
先端GPGPUシミュレーション工学特論
2015/07/16
OpenMP

共有メモリシステムでの並列処理に利用

標準化されたオープンな規格

OpenMPをサポートしているコンパイラであれば同じ
書き方が可能

並列化したい箇所をコンパイラに指示

ディレクティブ
コンパイラが対応していなければコメントとして扱われる
修正が最小限で済み，共通のソースコードで管理
1062
先端GPGPUシミュレーション工学特論


2015/07/16
OpenMP

並列に処理を実行させる箇所に指示句（ディレクティ
ブ）を挿入

for文の並列化

ディレクティブを一行追加(#pragma omp ～)
#pragma omp parallel for
for(int i=0; i<N; i++)
C[i] = A[i] + B[i]
1063
先端GPGPUシミュレーション工学特論
2015/07/16
逐次（並列化前）プログラム
#include<stdio.h>
#include<stdlib.h>
#define N (1024*1024)
#define Nbytes (N*sizeof(float))
for(i=0; i<N; i++)
c[i] = a[i] + b[i];
for(i=0; i<N; i++)
printf("%f+%f=%f¥n",
a[i],b[i],c[i]);
int main(){
float *a,*b,*c;
int i;
return 0;
a = (float *)malloc(Nbytes);
b = (float *)malloc(Nbytes);
c = (float *)malloc(Nbytes);
}
for(i=0; i<N; i++){
a[i] = 1.0;
b[i] = 2.0;
c[i] = 0.0;
}
1064
先端GPGPUシミュレーション工学特論
2015/07/16
並列化プログラム
#include<stdio.h>
#include<stdlib.h>
#define N (1024*1024)
#define Nbytes (N*sizeof(float))
for(i=0; i<N; i++)
c[i] = a[i] + b[i];
for(i=0; i<N; i++)
printf("%f+%f=%f¥n",
a[i],b[i],c[i]);
int main(){
float *a,*b,*c;
int i;
return 0;
a = (float *)malloc(Nbytes);
b = (float *)malloc(Nbytes);
c = (float *)malloc(Nbytes);
#pragma omp parallel for
for(i=0; i<N; i++){
a[i] = 1.0;
b[i] = 2.0;
c[i] = 0.0;
}
#pragma omp parallel for
1065
}
先端GPGPUシミュレーション工学特論
2015/07/16
並列化プログラム
#include<stdio.h>
#include<stdlib.h>
#define N (1024*1024)
#define Nbytes (N*sizeof(float))
}
#pragma omp for
for(i=0; i<N; i++)
c[i] = a[i] + b[i];
}
int main(){
float *a,*b,*c;
int i;
for(i=0; i<N; i++)
printf("%f+%f=%f¥n",
a[i],b[i],c[i]);
a = (float *)malloc(Nbytes);
b = (float *)malloc(Nbytes);
c = (float *)malloc(Nbytes);
#pragma omp parallel
{
#pragma omp for
for(i=0; i<N; i++){
a[i] = 1.0;
b[i] = 2.0;
c[i] = 0.0;
1066
return 0;
}
先端GPGPUシミュレーション工学特論
2015/07/16
コンパイル

コンパイル時にコンパイルオプションを付与
‐fopenmp

‐fopenmpを付けるとディレクティブを処理

‐fopenmpを付けないとディレクティブは無視される
1067
先端GPGPUシミュレーション工学特論
2015/07/16
処理の並列化

データ並列


コア0
forループをスレッドの数だけ分割
タスク並列もできる
for(i=0; i<N/2‐1; i++)
c[i] = a[i] + b[i];
スレッド0
スレッド1
＋
＋
a[i]
＋
＋
b[i]
コア1
for(i=N/2; i<N; i++)
c[i] = a[i] + b[i];
c[i]
1068
先端GPGPUシミュレーション工学特論
2015/07/16
OpenMPの指示文

並列処理制御




同期制御


OpenMP並列領域内でのデータアクセス，命令実行の同期
データ属性制御


OpenMPで並列処理を行う領域の定義
並列実行領域（Parallel Region）構文
ワークシェアリング（Work sharing）構文
並列領域内で利用されるデータの属性を定義
その他
1069
先端GPGPUシミュレーション工学特論
2015/07/16
並列実行領域（Parallel Region）構文

parallel構文


parallel構文で指示された領域では指定されたスレッド
が並列に処理を実行
全てのスレッドが同じ処理を実行
#pragma omp parallel //{ <‐ここに括弧を書くとエラー
{
複数のスレッドが起動され，ここに書いてある処理を実行
全てのスレッドが同じ処理を実行
}
1070
先端GPGPUシミュレーション工学特論
2015/07/16
ワークシェアリング（Work sharing）構文

for構文




parallel構文で指定された並列実行領域内で利用
直後のforループを各スレッドに分割して並列処理を実行
for(初期化;継続条件;再初期化)で構成されるforルー
プが対象
全てのスレッドが処理を終了するまで他のスレッドは待機
#pragma omp parallel
{
#pragma omp for
for(i=0; i<N; i++){
forループを自動的に分割して各スレッドが実行
}
全てのスレッドが処理を終了するまで待機
}
1071
先端GPGPUシミュレーション工学特論
2015/07/16
データ属性制御

shared指示節


parallel構文で指定された並列実行領域内での変数の
取り扱いを指示
指定した変数を全てのスレッドで共有
int data;
#pragma omp parallel shared(data)
{
全てのスレッドがdataを共有
あるスレッドがdataを変更すると，他のスレッドが参照する
dataの値も変更
}
1072
先端GPGPUシミュレーション工学特論
2015/07/16
データ属性制御

private指示節




parallel構文で指定された並列実行領域内での変数の
取り扱いを指示
指定した変数のコピーを全てのスレッドが個別に保持
コピーされた変数の値は引き継がれない
 0もしくは未定義
for構文でワークシェアされたfor文のループカウンタは
private
int data;
#pragma omp parallel private(data)
{
全てのスレッドがdataのコピーを個別に保持
}
1073
先端GPGPUシミュレーション工学特論
2015/07/16
データ属性制御

reduction指示節


値の総和や最大値などを求めるfor文の並列化に利用
結果を保持する変数のコピーが各スレッドに作成され，並
列処理の最後でまとめられる


reduction(op:変数） opは+,‐,*,&,|,^,&&,||のいずれか
変数 = 変数 op 変数もしくは値
int sum=0, c[N];
#pragma omp parallel for reduction(+:sum)
for(i=0; i<N; i++){
sum += c[i];
forループを分割し，各スレッドがsumを個別に保持して
総和計算を実行
}
各スレッドのsumをまとめて総和を計算
1074
先端GPGPUシミュレーション工学特論
2015/07/16
OpenMPランタイムAPI

omp_set_num_threads


並列実行領域のスレッド数を指定
並列実行領域の直前で呼出
omp_set_num_threads(4);
#pragma omp parallel
{
4スレッドが起動し，並列実行領域内の処理を実行
}

omp_get_num_threads

並列実行領域内で実行されているスレッド数を返す
omp_set_num_threads(12);
#pragma omp parallel
{
printf("%d¥n",omp_get_num_threads());//12と表示
}
1075
先端GPGPUシミュレーション工学特論
2015/07/16
OpenMPランタイムAPI

omp_get_thread_num

並列実行領域内で実行されているスレッドに割り振られた番号を返す
omp_set_num_threads(4);
#pragma omp parallel
{
printf("%d¥n",omp_get_thread_num());
//呼び出したスレッドに応じて0~3のいずれかを表示
}

omp_get_wtime


倍精度浮動小数点で時間（秒単位）を返す
全てのOSで実際の実行時間を取得可能
double time_start = omp_get_wtime();
...
double time_end
= omp_get_wtime();
printf("実行時間 %f sec¥n",time_end‐time_start);
1076
先端GPGPUシミュレーション工学特論
2015/07/16
OpenMPを利用して複数GPUを利用
#include<stdio.h>
#include<stdlib.h>
#include<omp.h>
__global__ void add
(float *a, float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
#define N (1024*1024*2)
#define Nbytes (N*sizeof(float))
#define NT 256
#define NB (N/NT)
#define GPUs 4
c[i] = a[i] + b[i];
}
__global__ void init
(float *a, float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
a[i] = 1.0;
b[i] = 2.0;
c[i] = 0.0;
}
vectoradd_omp.cu
1077
先端GPGPUシミュレーション工学特論
2015/07/16
OpenMPを利用して複数GPUを利用
int main(){
float *a,*b,*c;
float *host_c
= (float *)malloc(Nbytes);
int th;
cudaMemcpy(&host_c[N/GPUs * th], c,
Nbytes/GPUs, cudaMemcpyDeviceToHost);
cudaFree(a);
cudaFree(b);
cudaFree(c);
omp_set_num_threads(GPUs);
#pragma omp parallel private(th,a,b,c) }
double sum=0;
shared(host_c)
for(int i=0;i<N;i++)
{
sum+=host_c[i];
th = omp_get_thread_num();
cudaSetDevice(th);
printf("%f¥n",sum/N);
cudaMalloc((void **)&a,Nbytes/GPUs);
cudaMalloc((void **)&b,Nbytes/GPUs);
free(host_c);
cudaMalloc((void **)&c,Nbytes/GPUs);
return 0;
}
init<<< NB/GPUs, NT >>>(a,b,c);
add<<< NB/GPUs, NT >>>(a,b,c);
vectoradd_omp.cu
1078
先端GPGPUシミュレーション工学特論
2015/07/16
コンパイル

コンパイルのみ実行し，オブジェクトファイルを作成


OpenMPディレクティブが記述されたソース（.c, .cpp）は
cc/gcc/g++でコンパイル
CUDAのカーネルが記述されたソース（.cu）はnvccでコン
パイル
#pragma...
$ cc ‐fopenmp ‐c cpu.c
cpu.c
cpu.o
__global__ void ...
$ nvcc ‐c gpu.cu
gpu.cu
1079
gpu.o
先端GPGPUシミュレーション工学特論
2015/07/16
リンク

オブジェクトファイルをリンクして実行ファイルを作成

OpenMPのランタイムライブラリ(libgomp)を追加
cpu.o
$ nvcc cpu.o gpu.o ‐lgomp
a.out
gpu.o
1080
先端GPGPUシミュレーション工学特論
2015/07/16
コンパイル

コンパイルのみ実行し，オブジェクトファイルを作成


OpenMPディレクティブとCUDAのカーネルが混在する場合
ソースファイルが一つだけなら‐cオプションは不要
 nvccでコンパイルからリンクまで実行できる
$ nvcc ‐c ‐Xcompiler ‐fopenmp gpu.cu
gpu.cu
1081
__global__ void ...
gpu.o
#pragma...
先端GPGPUシミュレーション工学特論
2015/07/16
リンク

オブジェクトファイルをリンクして実行ファイルを作成


OpenMPのランタイムライブラリ(libgomp)を追加
ソースファイルが一つだけなら‐cオプションは不要
 nvccでコンパイルからリンクまで実行できる
nvcc gpu.o ‐lgomp
gpu.o
1082
a.out
先端GPGPUシミュレーション工学特論
2015/07/16
1スレッドで複数のGPUを利用
#include<stdio.h>
#include<stdlib.h>
#define N (1024*1024*2)
#define Nbytes (N*sizeof(float))
#define NT 256
#define NB (N/NT)
#define GPUs 4
__global__ void add
(float *a, float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
c[i] = a[i] + b[i];
}
__global__ void init
(float *a, float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
a[i] = 1.0;
b[i] = 2.0;
c[i] = 0.0;
}
vectoradd_multigpu.cu
1083
先端GPGPUシミュレーション工学特論
2015/07/16
1スレッドで複数のGPUを利用
int main(){
float *a[GPUs],*b[GPUs],*c[GPUs];
int dev;
for(dev=0;dev<GPUs;dev++){
cudaSetDevice(dev);
cudaMemcpy(&host_c[dev*N/GPUs],
c[dev], Nbytes/GPUs,
cudaMemcpyDeviceToHost);
}
double sum=0;
for(int i=0;i<N;i++)sum+=host_c[i];
printf("%f¥n",sum/N);
free(host_c);
for(dev=0;dev<GPUs;dev++){
cudaSetDevice(dev);
cudaMalloc((void **)&a[dev],Nbytes/GPUs);
cudaMalloc((void **)&b[dev],Nbytes/GPUs);
cudaMalloc((void **)&c[dev],Nbytes/GPUs);
}
for(dev=0;dev<GPUs;dev++){
cudaSetDevice(dev);
for(dev=0;dev<GPUs;dev++){
cudaSetDevice(dev);
cudaFree(a[dev]);
cudaFree(b[dev]);
cudaFree(c[dev]);
}
return 0;
init<<<NB/GPUs,NT>>(a[dev],b[dev],c[dev]);
add<<<NB/GPUs,NT>>>(a[dev],b[dev],c[dev]);
}
float *host_c = (float *)malloc(Nbytes);
1084
}
vectoradd_multigpu.cu
先端GPGPUシミュレーション工学特論
2015/07/16
GPU Direct v2.0

2個のGPU間でデータを移動

ホストメモリを経由せず，GPU間でpeer to peer通信
ホストメモリを経由した従来のデータコピー
メモリ
メモリ
メモリ
CPU
GPU1
GPU2
PCI‐Ex
1085
先端GPGPUシミュレーション工学特論
2015/07/16
GPU Direct v2.0

2個のGPU間でデータを移動

ホストメモリを経由せず，GPU間でpeer to peer通信
ホストメモリを経由した従来のデータコピー
メモリ
メモリ
メモリ
CPU
GPU1
GPU2
PCI‐Ex
1086
先端GPGPUシミュレーション工学特論
2015/07/16
GPU Direct v2.0

2個のGPU間でデータを移動


ホストメモリを経由せず，GPU間でpeer to peer通信
同じIO‐Hubに接続されていることが条件
GPU Directによるpeer to peerコピー
メモリ
メモリ
メモリ
CPU
GPU1
GPU2
PCI‐Ex
1087
先端GPGPUシミュレーション工学特論
2015/07/16
GPU Direct v2.0

利用できるデバイス

Fermi以降のGPU

64 bit LinuxもしくはWindowsで動作

Windowsの場合はTesla GPUのみで利用可能

Linuxの場合はGeForceでも利用可能
1088
先端GPGPUシミュレーション工学特論
2015/07/16
GPU Direct v2.0を利用したデータ交換

手順

GPUを2個選択

GPUがpeer to peer通信可能かを確認

peer to peer通信を個別に有効化
 1台目から2台目，2台目から1台目の双方向を有効化

データをコピー
1089
先端GPGPUシミュレーション工学特論
2015/07/16
GPU Direct v2.0を利用したデータ交換

GPUの選択

cudaSetDevice(使用GPU番号)
cudaSetDevice(2); //GPU2を選択

GPUがpeer to peer通信可能かを確認

cudaDeviceCanAccessPeer(結果, アクセス元GPU，アクセス先GPU)
cudaDeviceCanAccessPeer(&result,2,3);
//GPU2からGPU3へPeer通信できるかをresultに書き込む
//resultが1ならPeer通信可能，0なら不可能
1090
先端GPGPUシミュレーション工学特論
2015/07/16
GPU Direct v2.0を利用したデータ交換

Peer to peer通信を個別に有効化

cudaDeviceEnablePeerAccess(アクセス先GPU,0);
cudaSetDevice(2);
cudaDeviceEnablePeerAccess(3,0);
//GPU2から3へのpeer通信を有効化

データをコピー

cudaMemcpyPeer(コピー先GPU, コピー先変数, コピー元GPU, コ
ピー元変数, 変数サイズ);
cudaMemcpyPeer(3, b, 2, a, sizeof(float)*100);
//GPU2の変数aからGPU3の変数bへfloat型100個分コピー
1091
先端GPGPUシミュレーション工学特論
2015/07/16
GPU DirectによるP2P通信
#include<stdio.h>
#include<stdlib.h>
#include<math.h>
#define N 65536
#define Bytes (N*sizeof(float))
void p2p (float *, float *, float *, float *);
//a,bをそれぞれa_cp,b_cpにコピー
p2p(a,b,a_cp,b_cp);
//正しくコピーできているかチェック
int num_error=0;
for(int i=0;i<N;i++)
if(abs(a[i]‐a_cp[i])>0.1f)
num_error++;
printf("error = %d¥n",num_error);
int main(void){
float *a = (float *)malloc(Bytes);
float *b = (float *)malloc(Bytes);
float *a_cp=(float *)malloc(Bytes);
float *b_cp=(float *)malloc(Bytes);
for(int i=0;i<N;i++){
a[i]=i+1;
b[i]=2*(i+1);
a_cp[i]=0;
b_cp[i]=0;
}
1092
num_error=0;
for(int i=0;i<N;i++)
if(abs(b[i]‐b_cp[i])>0.1f)
num_error++;
printf("error = %d¥n",num_error);
return 0;
}
peer2peer.cu
先端GPGPUシミュレーション工学特論
2015/07/16
GPU DirectによるP2P通信
void p2p(float *a, float *b, float *a_cp, float *b_cp){
float *dev0_a, *dev0_b;
float *dev1_a, *dev1_b;
cudaDeviceProp deviceProp;
int dev0, dev1, canAccess0To1, canAccess1To0;
//使用するGPUを選択．dev0が1台目，dev1が2台目
dev0 = 2;
dev1 = 3;
//peer to peer通信が可能なGPUかチェックする
//1台目から2台目へP2P通信が可能かチェック
cudaDeviceCanAccessPeer(&canAccess0To1, dev0, dev1);
//2台目から1台目へP2P通信が可能かチェック
cudaDeviceCanAccessPeer(&canAccess1To0, dev1, dev0);
printf("dev 0 To 1 P2P access = %d¥n
dev 1 To 0 P2P access = %d¥n",canAccess0To1,canAccess1To0);
peer2peer.cu
1093
先端GPGPUシミュレーション工学特論
2015/07/16
GPU DirectによるP2P通信
//2台のGPUがお互いにP2P通信できるならデータコピーを実行
if(canAccess0To1 == 1 && canAccess1To0 == 1){
//1台目のGPUがUVAをサポートしているかをチェック
cudaSetDevice(dev0);
cudaGetDeviceProperties(&deviceProp, dev0);
printf("device %d supports Unified Virtual Addressing : %d¥n"
,dev0, deviceProp.unifiedAddressing);
//1台目から2台目へのP2P通信を有効化．2個目の引数は必ず0
cudaDeviceEnablePeerAccess(dev1,0);
//2台目のGPUがUVAをサポートしているかをチェック
cudaSetDevice(dev1);
cudaGetDeviceProperties(&deviceProp, dev1);
printf("device %d supports Unified Virtual Addressing : %d¥n"
,dev1, deviceProp.unifiedAddressing);
//2台目から1台目へのP2P通信を有効化．2個目の引数は必ず0
cudaDeviceEnablePeerAccess(dev0,0);
peer2peer.cu
1094
先端GPGPUシミュレーション工学特論
2015/07/16
GPU DirectによるP2P通信
//1台目のGPUで使うメモリを確保し，変数aをCPUからGPUへコピー（dev0_bは未初期化）
cudaSetDevice(dev0);
cudaMalloc((void **)&dev0_a, Bytes);
cudaMalloc((void **)&dev0_b, Bytes);
cudaMemcpy(dev0_a, a, Bytes, cudaMemcpyHostToDevice);//dev0_aを初期化
//2台目のGPUで使うメモリを確保し，変数bをCPUからGPUへコピー（dev1_aは未初期化）
cudaSetDevice(dev1);
cudaMalloc((void **)&dev1_a, Bytes);
cudaMalloc((void **)&dev1_b, Bytes);
cudaMemcpy(dev1_b, b, Bytes, cudaMemcpyHostToDevice);//dev1_bを初期化
cudaSetDevice(dev0);
//2台目のGPU(dev1)にあるdev1_aへ，1台目のGPU(dev0)のdev0_aをBytes分コピー
cudaMemcpyPeer(dev1_a, dev1, dev0_a, dev0, Bytes);
//1台目のGPU(dev0)にあるdev0_bへ，2台目のGPU(dev1)のdev1_bをBytes分コピー
cudaMemcpyPeer(dev0_b, dev0, dev1_b, dev1, Bytes);
peer2peer.cu
1095
先端GPGPUシミュレーション工学特論
2015/07/16
GPU DirectによるP2P通信
//1台目のGPU(dev0)にあるdev0_bを，CPUへコピー（結果の確認用）
cudaSetDevice(dev0);
cudaMemcpy(b_cp,dev0_b, Bytes, cudaMemcpyDeviceToHost);
//2台目のGPU(dev1)にあるdev1_aを，CPUへコピー（結果の確認用）
cudaSetDevice(dev1);
cudaMemcpy(a_cp,dev1_a, Bytes, cudaMemcpyDeviceToHost);
//1台目のGPU(dev0)で確保したメモリをクリア
cudaSetDevice(dev0);
cudaFree(dev0_a);
cudaFree(dev0_b);
cudaDeviceDisablePeerAccess(dev1);//dev1へのPeer通信を無効化
//2台目のGPU(dev1)で確保したメモリをクリア
cudaSetDevice(dev1);
cudaFree(dev1_a);
cudaFree(dev1_b);
cudaDeviceDisablePeerAccess(dev0);//dev0へのPeer通信を無効化
}
}
peer2peer.cu
1096
先端GPGPUシミュレーション工学特論
2015/07/16
Unified Virtual Addressing

複数のGPUのメモリアドレスとcudaHostAlloc()で
確保したCPUのメモリアドレスを統一的に管理



40 bitメモリ空間を構成
異なるGPUのメモリやCPUのメモリの区別が不要
異なるGPUに置かれたメモリを参照可能

GPU directを利用してアクセス


1097
cudaDeviceEnablePeerAccessを利用
記述は楽になるが性能は出ない
先端GPGPUシミュレーション工学特論
2015/07/16
UVAを利用したベクトル和
#include <stdio.h>
#include <stdlib.h>
#define N (1024*1024*1)
#define Nbytes (N*sizeof(float))
#define NT 256
#define NB (N/NT)
//カーネルは変更なし
__global__ void add
(float *a, float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
c[i] = a[i] + b[i];
}
//カーネルは変更なし
__global__ void init
(float *a, float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
a[i] = 1.0;
b[i] = 2.0;
c[i] = 0.0;
}
vectoradd_uva.cu
1098
先端GPGPUシミュレーション工学特論
2015/07/16
UVAを利用したベクトル和
int main(void){
int dev0,dev1;
float *a,*b,*c;
//使用するGPUを選択．dev0が1台目，dev1が2台目
dev0 = 2;
dev1 = 3;
int canAccess1To0=0;
//2台目から1台目へP2P通信が可能かチェック
cudaDeviceCanAccessPeer(&canAccess1To0, dev1, dev0);
printf("dev 1 To 0 P2P access = %d¥n",canAccess1To0);
//2台目から1台目へP2P通信が可能ならif文の中を実行
if(canAccess1To0 == 1){
//1台目のGPUでメモリを確保（変数a,b,cは1台目のメモリに存在）
cudaSetDevice(dev0);
cudaMalloc( (void **)&a, Nbytes);
cudaMalloc( (void **)&b, Nbytes);
cudaMalloc( (void **)&c, Nbytes);
vectoradd_uva.cu
1099
先端GPGPUシミュレーション工学特論
2015/07/16
UVAを利用したベクトル和
//2台目のGPUを利用するようにデバイスを切替
cudaSetDevice(dev1);
//2台目から1台目へのP2P通信を有効化．2個目の引数は必ず0
cudaDeviceEnablePeerAccess(dev0, 0);
//1台目のGPUのメモリを参照して2台目のGPUで初期化とベクトル和を実行
init<<<NB, NT>>>(a, b, c);
add<<<NB, NT>>>(a, b, c);
//2台目のGPUが1台目のGPUのメモリをCPUへコピー
float *host_c = (float *)malloc(Nbytes);
cudaMemcpy(host_c, c, Nbytes, cudaMemcpyDeviceToHost);
int i;double sum=0;
for(i=0;i<N;i++) sum+=host_c[i];
printf("%f¥n",sum/N);
free(host_c);
//利用するGPUを1台目に切り替えて確保したメモリをクリア
cudaSetDevice(dev0);
cudaFree(a);
cudaFree(b);
cudaFree(c);
vectoradd_uva.cu
}
1100
先端GPGPUシミュレーション工学特論
2015/07/16
UVAを利用したベクトル和
else{//P2P通信が不可能なら1台のGPUで実行
cudaMalloc( (void **)&a, Nbytes);
cudaMalloc( (void **)&b, Nbytes);
cudaMalloc( (void **)&c, Nbytes);
init<<<NB, NT>>>(a, b, c);
add<<<NB, NT>>>(a, b, c);
float *host_c = (float *)malloc(Nbytes);
cudaMemcpy(host_c, c, Nbytes, cudaMemcpyDeviceToHost);
int i;double sum=0;
for(i=0;i<N;i++)sum+=host_c[i];
printf("%f¥n",sum/N);
free(host_c);
cudaFree(a);
cudaFree(b);
cudaFree(c);
}
return 0;
}
vectoradd_uva.cu
1101
先端GPGPUシミュレーション工学特論
2015/07/16
実行結果

配列の要素数
N=220
実行時間[ms]
カーネル


1102
単一GPU
UVA
初期化init
0.108
2.28
ベクトル和add
0.113
1.30
Unified Virtual Addressingを利用してP2Pアクセス
すると10~20倍以上の時間を要する
大量のデータにアクセスせず，限定的な利用が重要
 データが少なく，わざわざコピーするまでもない場合等
先端GPGPUシミュレーション工学特論
2015/07/16

Download Report