一枚のグラフィックカードに複数のGPUを載せたGeforce 9800 GX2や、SLIを使って複数のグラフィックカードを搭載したシステムでは、これらを同時に使って処理能力を高めることができる。
今回は、2つのGPUを使って一次元拡散方程式を解いてみる。
複数のGPUを使う場合、pthreadなどを使い、GPUの数だけスレッドを作成する。
GPU間で直接データをやりとりする方法はないので、それぞれ、メインメモリを介してデータのやりとりをする。
1つのGPUを使った並列処理はOpenMPに近いものがあるが、一方複数のGPUを使った処理はMPIのような処理となる。
今回使ったシステムは、
Opteron 2.0 GHz
NVIDIA Geforce 9800 GX2
OpenSUSE 10.2
である。
まず、diff1dmgpu.cuであるが、
#include <stdio.h> #include <cutil.h> #include <multithreading.h> //GPUの数(このプログラムのではGPUの数は2つの場合のみ考慮している) #define MAX_CPU_THREAD 2 #include "diff1dmgpu_kernel.cu" //ケーブル長 #define len 100 //グローバル変数 float exchange0,exchange1; float* v_h; float* v_d[MAX_CPU_THREAD]; //スレッド間の同期を取るためのバリア pthread_barrier_t barrier; //GPUの数だけ下記の関数が実行される static CUT_THREADPROC gpuThread(int * dev) { //デバイス(GPU)の選択 CUDA_SAFE_CALL(cudaSetDevice(*dev)); //GPU内にメモリの確保 CUDA_SAFE_CALL( cudaMalloc( (void**) &v_d[*dev], sizeof( float) * (len/2+100) )); //メインメモリからGPUメモリへデータの転送 if (*dev==0) { CUDA_SAFE_CALL( cudaMemcpy( &v_d[*dev][0], &v_h[0], sizeof( float) * len/2 , cudaMemcpyHostToDevice) ); CUDA_SAFE_CALL( cudaMemcpy( &v_d[*dev][len/2], &v_h[len/2], sizeof( float) * 1 , cudaMemcpyHostToDevice) ); } else { CUDA_SAFE_CALL( cudaMemcpy( &v_d[*dev][1], &v_h[len/2], sizeof( float) * len/2 , cudaMemcpyHostToDevice) ); CUDA_SAFE_CALL( cudaMemcpy( &v_d[*dev][0], &v_h[len/2-1], sizeof( float) * 1 , cudaMemcpyHostToDevice) ); } //拡散方程式のメインループ for (int tn=0;tn<100;tn++) { //拡散方程式を解くカーネル diff1dKernel<<< 1, 51 >>>(v_d[*dev]); //GPU間でのデータの交換 if (*dev==0) { CUDA_SAFE_CALL( cudaMemcpy( &exchange0, &v_d[*dev][len/2-1], sizeof(float) *1, cudaMemcpyDeviceToHost)); } else { CUDA_SAFE_CALL( cudaMemcpy( &exchange1, &v_d[*dev][1], sizeof(float) * 1, cudaMemcpyDeviceToHost)); } pthread_barrier_wait( &barrier ); if (*dev==0) { CUDA_SAFE_CALL( cudaMemcpy(&v_d[*dev][len/2], &exchange1, sizeof(float) * 1, cudaMemcpyHostToDevice)); } else { CUDA_SAFE_CALL( cudaMemcpy(&v_d[*dev][0], &exchange0, sizeof(float) * 1, cudaMemcpyHostToDevice)); } } //結果のメインメモリへの書き込み if (*dev==0) { CUDA_SAFE_CALL( cudaMemcpy( &v_h[0], &v_d[*dev][0], sizeof( float) * len/2 , cudaMemcpyDeviceToHost) ); } else { CUDA_SAFE_CALL( cudaMemcpy( &v_h[len/2], &v_d[*dev][1], sizeof( float) * len/2 , cudaMemcpyDeviceToHost) ); } //GPUメモリの解放 CUDA_SAFE_CALL(cudaFree(v_d[*dev])); CUT_THREADEND; } int main( int argc, char** argv) { //デバイスの初期化 CUT_DEVICE_INIT(argc, argv); //結果書き込み用ファイルのオープン FILE *fp=fopen("result.txt","w"); //タイマーを作成して計測開始 unsigned int timer = 0; CUT_SAFE_CALL( cutCreateTimer( &timer)); CUT_SAFE_CALL( cutStartTimer( timer)); //バリアを初期化 pthread_barrier_init( &barrier, 0, MAX_CPU_THREAD ); //GPUの分だけスレッドを確保 int s_gpuCount; CUDA_SAFE_CALL(cudaGetDeviceCount(&s_gpuCount)); int threadIds[MAX_CPU_THREAD]; printf("%d GPUs found\n", s_gpuCount); CUTThread * threads = (CUTThread *)malloc(sizeof(CUTThread) * s_gpuCount); //メインメモリ上にfloat型のデータを100(len)個生成する v_h = (float*) malloc(sizeof( float) * len); //初期条件をセット for( int i = 0; i < len/2; i++) v_h[i] = 0; for( int i = len/2; i < len; i++) v_h[i] = 1; //実行結果を表示 for (int i=0;i<len;i++) { fprintf(fp,"%f\t",v_h[i]); } fprintf(fp,"\n"); //デバイスの数だけスレッドを開始 for(int i = 0; i < s_gpuCount; i++) { threadIds[i] = i; threads[i] = cutStartThread((CUT_THREADROUTINE)gpuThread, (void *)&threadIds[i]); } //すべてのスレッドが終了するのを待つ cutWaitForThreads(threads, s_gpuCount); //実行結果を表示 for (int i=0;i<len;i++) { fprintf(fp,"%f\t",v_h[i]); } fprintf(fp,"\n"); //タイマーを停止しかかった時間を表示 CUT_SAFE_CALL( cutStopTimer( timer)); printf("Processing time: %f (ms)\n", cutGetTimerValue( timer)); CUT_SAFE_CALL( cutDeleteTimer( timer)); //各種メモリを解放 free( v_h); fclose(fp); return 0; }
スレッド間の同期はバリアでとっている。
一方、diff1dmgpu_kernel.cuは、
__global__ void diff1dKernel( float* v) { __shared__ float vold[256]; // スレッドIDを取得 const unsigned int tid = threadIdx.x; //グローバルメモリから入力データの読み込み vold[tid] = v[tid]; __syncthreads(); //ここで計算を行う const float Dfu=1; const float dt=0.2; const float dx=1; float Dfudtdx2=Dfu*dt/(dx*dx); if (tid==0) v[tid]=vold[tid]+(2*vold[tid+1]-2*vold[tid])*Dfudtdx2; else if (tid==50) v[tid]=vold[tid]+(2*vold[tid-1]-2*vold[tid])*Dfudtdx2; else v[tid]=vold[tid]+(vold[tid-1]+vold[tid+1]-2*vold[tid])*Dfudtdx2; }
上記のプログラムを下記のようにしてコンパイルする。
$ nvcc diff1dmgpu.cu -o diff1dmgpu -I/usr/local/NVIDIA_CUDA_SDK/common/inc /usr/local/NVIDIA_CUDA_SDK/lib/libcutil.a
今回はあくまで、使い方の例なので、実行時間は考慮に入れていないが、非常に重い処理などをさせる場合、有効なテクニックである。
最終更新日