目次 >> CUDA >> 複数のGPUを使う

CUDA - 複数のGPUを使う

一枚のグラフィックカードに複数の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
である。

一次元拡散方程式を2つのGPUで解く

まず、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

今回はあくまで、使い方の例なので、実行時間は考慮に入れていないが、非常に重い処理などをさせる場合、有効なテクニックである。