Microsoft PowerPoint - 先端GPGPUシミュレーション工学特論(web).pptx

Transcription

1 複数 GPU の利用 長岡技術科学大学電気電子情報工学専攻出川智啓

2 今回の内容 並列アーキテクチャと並列処理の分類 OpenMP 複数 GPU の利用 GPU Direct によるデータ通信 939

3 複数の GPU を利用する目的 Grouse の 1 ノードには 4 台の GPU を搭載 Tesla M2050 1T FLOPS/ 台 3 GB/ 台 4 台全てを使う事で期待できる性能 GPU を 1 台だけ使った時の 4 倍高速化 GPU1 台では収まらないサイズの問題を計算 940

4 複数の GPU を利用する目的 4 台全てを使う事で期待できる性能 GPU を 1 台だけ使った時の 4 倍高速化 高速化は達成可能 複数の GPU を使うための手続きやデータの移動が必要 GPU 間のデータ交換が必要になると速度が低下 GPU1 台では収まらないサイズの問題を計算 利用できるメモリ容量は GPU の台数に比例 1 台では実行できない大規模な問題も実行可能 941

5 複数の GPU が利用できる環境 1 台の計算機に複数の GPU が搭載 複数の GPU を搭載するために大きな筐体が必要 1 台の PC で完結するので, プログラム中でネットワーク通信を行う必要がない 帯域の狭い PCI Ex バスを複数の GPU で取り合うため, データ移動の効率は悪い PCI Ex バス 942

6 複数の GPU が利用できる環境 1 台の GPU が搭載された計算機をネットワークで接続 小さい筐体のPC( キューブ型 PCなど ) を並べて作成 各 PCにCPUやメモリが必要 複数 GPUの利用にはネットワーク通信が必須 PCI Ex PCI Ex PCI Ex PCI Ex 高速ネットワーク,LAN 943

7 複数の GPU が利用できる環境 複数の GPU が搭載された複数の計算機をネットワークで接続 スーパーコンピュータ等の大規模計算環境 性能の大半は GPU が占める 2~4CPUs 2~4GPUs 2~4CPUs 2~4GPUs 2~4CPUs 2~4GPUs 944 高速ネットワーク

8 複数の GPU が利用できる環境 CUDA における複数 GPU 利用 新しいバージョン GPU 間通信を行う命令が用意されたり, 複数の GPU を利用する数値計算ライブラリが登場 ( 古いバージョンと比べれば ) 利用しやすくなっている 古いバージョン 複数の GPU を利用するための機能は未搭載 従来から並列計算に利用されていたライブラリと併用 OpenMP や MPI など 945

9 複数 GPU の利用 (CUDA 3.2 以前 ) CUDA の制約 1 スレッドは複数の GPU を利用できない GPU は必ず CPU の 1 スレッド ( もしくは 1 プロセス ) に割当て 4 台の GPU を利用するためには 4 スレッド (4 プロセス ) 必要 OpenMP や MPI を利用 並列計算を行うプログラムを作成 スレッドやプロセス番号を基に利用する GPU を決定, 処理を実行 従来の並列計算の素直な拡張 946

10 複数 GPU の利用 (CUDA 4.0 以降 ) CUDA の制約の緩和 1 スレッドから複数の GPU を利用 ( 切替 ) 可能 2 台の GPU 間でデータ通信を行う命令の追加 GPU Direct v2.0 peer to peer 複数 GPU 間で統一されたメモリ空間を利用可能 Unified Virtual Addressing 947

11 複数 GPU の利用 (CUDA 4.0 以降 ) 複数 GPU を利用するプログラムの作成が容易 cudasetdevice() でカレントデバイスを指定 CUDA の命令はそのカレントデバイスに対して発行 別のデバイスを利用する際は cudasetdevice() でカレントデバイスを切替 複数の GPU へのデータコピー, 複数の GPU でカーネルの起動が容易 948

12 複数 GPU 化の例題 ベクトル和 1.OpenMP を利用する方法 従来の並列計算の拡張 2.CUDA4.0 の機能を利用する方法 949

13 並列アーキテクチャの分類 システムの特徴付け, プロセッサの分類 Flynn の分類 並列アーキテクチャのグループ分け データの処理と命令の並列性に着目 SISD 単一命令単一データ SIMD 単一命令複数データ MISD 複数命令単一データ MIMD 複数命令複数データ 950

14 Single Instruction Multiple Data streams 単一命令複数データ 複数のまとまったデータに対して同じ演算を同時に実行 命令は一つ, その命令が同時に多くのデータに対して適用されるアーキテクチャ 数値シミュレーションに最適 数学のベクトルや配列計算の概念に一致 ベクトルプロセッサとも呼ばれる GPU もここに分類 データ 命令 A 0 B 0 A 1 B 1 + A 2 B 2 A 3 B 3 951

15 プロセス OS から資源 ( コア, メモリ, 外部記憶など ) を割り当てられ, 実行状態 ( または待機状態 ) にあるプログラム システムプロセス OS の実行に関係するプログラム ユーザプロセス ユーザ権限で実行されているプログラム 952

16 マルチプロセス 複数のプロセスが存在し, 並列に実行 プロセスが一つのみ プロセスが二つ以上 シングルプロセスマルチプロセス マルチプロセスに対応した OS が必要 現在の OS はマルチプロセスに対応 シングルコア CPU 一つでもマルチプロセスが可能 OS が複数のプロセスを切替 複数のプロセスが並列に実行されているように見せる 953

17 マルチプロセス 各プロセスに専用のメモリ領域を割り当て CPUやメモリは複数のプログラムに割り当てられる プログラムはCPUやメモリを独占しているように振る舞う プロセス A スレッド CPU メモリ OS メモリ プロセス B スレッド メモリ 954

18 スレッド プログラムの処理の最小実行単位 プロセス内で複数のスレッドが存在 1プロセスに一つのスレッドのみ 1プロセスに二つ以上のスレッド シングルプロセス シングルスレッドマルチスレッド シングルスレッド マルチスレッド マルチプロセス シングルスレッド マルチスレッド 955

19 マルチスレッド 一つのプロセスに二つ以上のスレッドが存在 一つのプロセスに専用のメモリ領域を割当 プロセス内の複数のスレッドはメモリ領域を共有 プロセス A スレッド スレッド CPU メモリ OS メモリプロセスB スレッドスレッド メモリ 956

20 並列計算機システム 並列処理の基本 処理を何らかの方法で分割 分割した処理をプロセッサ ( やコア ) に割り当て同時に処理 並列計算機システム 複数のプロセッサをもつ 主にメモリに違いがある 共有メモリシステム 分散メモリシステム ハイブリッドシステム 957

21 共有メモリシステム 複数のプロセッサがメモリ空間を共有 分割した処理は各プロセッサ上で並列的に処理 共有されたメモリ空間上の変数は全ての CPU( やコア ) からアクセス ( 読み書き ) 可能 他からアクセスされない変数を持つことも可能 CPU CPU CPU CPU メモリ 958

22 並列処理の分類 タスク並列 独立なタスクを異なる CPU コアで同時に実行 データ並列 独立なタスクが処理するデータを分割し, 異なる CPU コアがデータを参照し, 処理を実行 Embarrassingly parallel (perfectly parallel) 各 CPU コアが同じタスクを異なるパラメータで実行 GPU が各ピクセルの色を決定し, ディスプレイに描画する処理 あるタスクに対してパラメータの影響を調査するような問題 959

23 データ並列 独立な処理 A,B,C が取り扱うデータを分割して実行 逐次処理 コア 処理 A 処理 B 処理 C 並列処理 コア 1 処理 A 処理 B 処理 C コア 2 処理 A 処理 B 処理 C 高速化 960

24 OpenMP 共有メモリシステムでの並列処理に利用 標準化されたオープンな規格 OpenMP をサポートしているコンパイラであれば同じ書き方が可能 並列化したい箇所をコンパイラに指示 ディレクティブ コンパイラが対応していなければコメントとして扱われる 修正が最小限で済み, 共通のソースコードで管理 961

25 OpenMP 並列に処理を実行させる箇所に指示句 ( ディレクティブ ) を挿入 for 文の並列化 ディレクティブを一行追加 (#pragma omp ~) #pragma omp parallel for for(int i=0; i<n; i++) C[i] = A[i] + B[i] 962

26 逐次 ( 並列化前 ) プログラム #include<stdio.h> #include<stdlib.h> #define N (1024*1024) #define Nbytes (N*sizeof(float)) int main(){ float *a,*b,*c; int i; 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; 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]); return 0; 963

27 並列化プログラム #include<stdio.h> #include<stdlib.h> #define N (1024*1024) #define Nbytes (N*sizeof(float)) int main(){ float *a,*b,*c; int i; 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 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]); return 0; 964

28 並列化プログラム #include<stdio.h> #include<stdlib.h> #define N (1024*1024) #define Nbytes (N*sizeof(float)) int main(){ float *a,*b,*c; int 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; #pragma omp for 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]); return 0; 965

29 コンパイル コンパイル時にコンパイルオプションを付与 fopenmp fopenmp を付けるとディレクティブを処理 fopenmp を付けないとディレクティブは無視される 966

30 処理の並列化 データ並列 for ループをスレッドの数だけ分割 タスク並列もできる スレッド 1 スレッド 2 コア 1 for(i=0; i<n/2 1; i++) c[i] = a[i] + b[i]; a[i] b[i] コア 2 for(i=n/2; i<n; i++) c[i] = a[i] + b[i]; c[i] 967

31 OpenMP の指示文 並列処理制御 OpenMP で並列処理を行う領域の定義 並列実行領域 (Parallel Region) 構文 ワークシェアリング (Work sharing) 構文 同期制御 OpenMP 並列領域内でのデータアクセス, 命令実行の同期 データ属性制御 並列領域内で利用されるデータの属性を定義 その他 968

32 並列実行領域 (Parallel Region) 構文 parallel 構文 parallel 構文で指示された領域では指定されたスレッドが並列に処理を実行 全てのスレッドが同じ処理を実行 #pragma omp parallel //{ < ここに括弧を書くとエラー { 複数のスレッドが起動され, ここに書いてある処理を実行全てのスレッドが同じ処理を実行 969

33 ワークシェアリング (Work sharing) 構文 for 構文 parallel 構文で指定された並列実行領域内で利用 直後の for ループを各スレッドに分割して並列処理を実行 for( 初期化 ; 継続条件 ; 再初期化 ) で構成される for ループが対象 全てのスレッドが処理を終了するまで他のスレッドは待機 #pragma omp parallel { #pragma omp for for(i=0; i<n; i++){ for ループを自動的に分割して各スレッドが実行 全てのスレッドが処理を終了するまで待機 970

34 データ属性制御 shared 指示節 parallel 構文で指定された並列実行領域内での変数の取り扱いを指示 指定した変数を全てのスレッドで共有 int data; #pragma omp parallel shared(data) { 全てのスレッドが data を共有あるスレッドが data を変更すると, 他のスレッドが参照する data の値も変更 971

35 データ属性制御 private 指示節 parallel 構文で指定された並列実行領域内での変数の取り扱いを指示 指定した変数のコピーを全てのスレッドが個別に保持 コピーされた変数の値は引き継がれない 0もしくは未定義 for 構文でワークシェアされたfor 文のループカウンタは private int data; #pragma omp parallel private(data) { 全てのスレッドが data のコピーを個別に保持 972

36 データ属性制御 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 をまとめて総和を計算 973

37 OpenMP ランタイム API omp_set_num_threads 並列実行領域のスレッド数を指定 並列実行領域の直前で呼出 omp_get_num_threads omp_set_num_threads(4); #pragma omp parallel { 4 スレッドが起動し, 並列実行領域内の処理を実行 並列実行領域内で実行されているスレッド数を返す omp_set_num_threads(12); #pragma omp parallel { printf("%d n",omp_get_num_threads());//12 と表示 974

38 OpenMP ランタイム API omp_get_thread_num 並列実行領域内で実行されているスレッドに割り振られた番号を返す omp_get_wtime omp_set_num_threads(4); #pragma omp parallel { printf("%d n",omp_get_thread_num()); // 呼び出したスレッドに応じて 0~3 のいずれかを表示 倍精度浮動小数点で時間 ( 秒単位 ) を返す 全ての OS で実際の実行時間を取得可能 double time_start = omp_get_wtime();... double time_end = omp_get_wtime(); printf(" 実行時間 %f sec n",time_end time_start); 975

39 OpenMP を利用して複数 GPU を利用 #include<stdio.h> #include<stdlib.h> #include<omp.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_omp.cu 976

40 OpenMP を利用して複数 GPU を利用 int main(){ float *a,*b,*c; float *host_c = (float *)malloc(nbytes); int th; omp_set_num_threads(gpus); #pragma omp parallel private(th,a,b,c) shared(host_c) { th = omp_get_thread_num(); cudasetdevice(th); cudamalloc((void **)&a,nbytes/gpus); cudamalloc((void **)&b,nbytes/gpus); cudamalloc((void **)&c,nbytes/gpus); init<<< NB/GPUs, NT >>>(a,b,c); add<<< NB/GPUs, NT >>>(a,b,c); cudamemcpy(&host_c[n/gpus * th], c, Nbytes/GPUs, cudamemcpydevicetohost); cudafree(a); cudafree(b); cudafree(c); double sum=0; for(int i=0;i<n;i++) sum+=host_c[i]; printf("%f n",sum/n); free(host_c); return 0; vectoradd_omp.cu 977

41 コンパイル コンパイルのみ実行し, オブジェクトファイルを作成 OpenMP ディレクティブが記述されたソース (.c,.cpp) は cc/gcc/g++ でコンパイル CUDA のカーネルが記述されたソース (.cu) は nvcc でコンパイル #pragma... cc fopenmp c cpu.c cpu.c cpu.o gpu.cu global void... nvcc c gpu.cu gpu.o 978

42 リンク オブジェクトファイルをリンクして実行ファイルを作成 OpenMP のランタイムライブラリ (libgomp) を追加 cpu.o nvcc cpu.o gpu.o lgomp a.out gpu.o 979

43 コンパイル コンパイルのみ実行し, オブジェクトファイルを作成 OpenMP ディレクティブと CUDA のカーネルが混在する場合 ソースファイルが一つだけなら c オプションは不要 nvcc でコンパイルからリンクまで実行できる gpu.cu global void... #pragma... nvcc c Xcompiler fopenmp gpu.cu gpu.o 980

44 リンク オブジェクトファイルをリンクして実行ファイルを作成 OpenMP のランタイムライブラリ (libgomp) を追加 ソースファイルが一つだけなら c オプションは不要 nvcc でコンパイルからリンクまで実行できる nvcc gpu.o lgomp gpu.o a.out 981

45 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 init (float *a, float *b, float *c){ int i = blockidx.x*blockdim.x + threadidx.x; global void add (float *a, float *b, float *c){ int i = blockidx.x*blockdim.x + threadidx.x; c[i] = a[i] + b[i]; a[i] = 1.0; b[i] = 2.0; c[i] = 0.0; vectoradd_multigpu.cu 982

46 1 スレッドで複数の GPU を利用 int main(){ float *a[gpus],*b[gpus],*c[gpus]; int dev; 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); 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); 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); cudafree(a[dev]); cudafree(b[dev]); cudafree(c[dev]); return 0; vectoradd_multigpu.cu 983

47 GPU Direct v2.0 2 個の GPU 間でデータを移動 ホストメモリを経由せず,GPU 間で peer to peer 通信 ホストメモリを経由した従来のデータコピー メモリ メモリ メモリ CPU GPU1 GPU2 PCI Ex 984

48 GPU Direct v2.0 2 個の GPU 間でデータを移動 ホストメモリを経由せず,GPU 間で peer to peer 通信 ホストメモリを経由した従来のデータコピー メモリ メモリ メモリ CPU GPU1 GPU2 PCI Ex 985

49 GPU Direct v2.0 2 個の GPU 間でデータを移動 ホストメモリを経由せず,GPU 間でpeer to peer 通信 同じIO Hubに接続されていることが条件 GPU Direct による peer to peer コピー メモリ メモリ メモリ CPU GPU1 GPU2 PCI Ex 986

50 GPU Direct v2.0 利用できるデバイス Fermi 以降の GPU 64 bit Linux もしくは Windows で動作 Windows の場合は Tesla GPU のみで利用可能 Linux の場合は GeForce でも利用可能 987

51 GPU Direct v2.0 を利用したデータ交換 手順 GPU を 2 個選択 GPU が peer to peer 通信可能かを確認 peer to peer 通信を個別に有効化 1 台目から 2 台目,2 台目から 1 台目 データをコピー 988

52 GPU Direct v2.0 を利用したデータ交換 GPU の選択 cudasetdevice( 使用 GPU 番号 ) cudasetdevice(1); GPU1 が使われる GPU が peer to peer 通信可能かを確認 cudadevicecanaccesspeer( 結果, アクセス元 GPU, アクセス先 GPU) cudadevicecanaccesspeer(&result,device0,device1); device0からdevice1にpeer 通信できるかをresultに書き込む 989

53 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 個分コピー 990

54 GPU Direct による P2P 通信 #include<stdio.h> #include<stdlib.h> #include<math.h> #define N #define Bytes (N*sizeof(float)) void p2p (float *, float *, float *, float *); 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; //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); 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 991

55 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 992

56 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 993

57 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); //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); 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 994

58 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 995

59 Unified Virtual Addressing 複数の GPU のメモリアドレスと,cudaHostAlloc() で確保した CPU のメモリアドレスを統一的に管理 40 bit メモリ空間を構成 異なる GPU のメモリや CPU のメモリの区別が不要 異なる GPU に置かれたメモリを参照可能 GPU direct を利用してアクセス cudadeviceenablepeeraccess を利用 記述は楽になるが性能は出ない 996

60 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 init (float *a, float *b, float *c){ int i = blockidx.x*blockdim.x + threadidx.x; // カーネルは変更なし global void add (float *a, float *b, float *c){ int i = blockidx.x*blockdim.x + threadidx.x; c[i] = a[i] + b[i]; a[i] = 1.0; b[i] = 2.0; c[i] = 0.0; vectoradd_uva.cu 997

61 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 998

62 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 999

63 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 1000

64 実行結果 配列の要素数 N=2 20 カーネル 単一 GPU 実行時間 [ms] UVA 初期化 init ベクトル和 add Unified Virtual Addressing を利用して P2P アクセスすると 10~20 倍以上の時間を要する 大量のデータにアクセスせず, 限定的な利用が重要 データが少なく, わざわざコピーするまでもない場合等 1001