Twitterに投稿 はてなブックマークに追加 Yahoo!ブックマークに追加 Google Bookmarksに追加

目次 >> CUDA >> 非同期処理

CUDA - 非同期処理

先回までは、GPU上のみでデータをどのように処理するかについて検討してきた。
この場合、GPUで計算している間、CPUはアイドル状態、すなわち何もせずに遊んでいる状態である。
このCPUのアイドル時間を有効に活用としようというのが今回の目的である。
まずは、1次元の拡散方程式を解く次のコードを基準として用いる。まず、diff1d.cuは、

#include <stdio.h>
#include <cutil.h>
#include "diff1d_kernel.cu"
int main( int argc, char** argv) 
{
    //デバイスの初期化
    CUT_DEVICE_INIT(argc, argv);

    //タイマーを作成して計測開始
    unsigned int timer = 0;
    CUT_SAFE_CALL( cutCreateTimer( &timer));
    CUT_SAFE_CALL( cutStartTimer( timer));

    //メインメモリ上にfloat型のデータを100個生成する
    float* h_idata = (float*) malloc(sizeof( float) * 100);
    //初期条件をセット
    for( int i = 0; i < 50; i++) 
        h_idata[i] = 0;
    for( int i = 50; i < 100; i++) 
        h_idata[i] = 1;

    //デバイス上（ビデオカードのこと）にも同じくfloat型100個分のメモリを確保する
    float* d_idata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_idata, sizeof( float) * 100 ));

    dim3  grid( 1, 1, 1);
    //100は100個並列であるため
    dim3  threads(100, 1, 1);
    
    //メインメモリからデバイスのメモリにデータを転送する
    CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, sizeof( float) * 100 , cudaMemcpyHostToDevice) );

    for (int t=0;t<100;t++)
    {
        //ここでGPUを使った計算が行われる
        diff1dKernel<<< grid, threads, sizeof( float) * 200 >>>( d_idata);
        //デバイスからメインメモリ上に実行結果をコピー
        CUDA_SAFE_CALL( cudaMemcpy( h_idata, d_idata, sizeof( float) * 100, cudaMemcpyDeviceToHost) );
    }
    //タイマーを停止しかかった時間を表示
    CUT_SAFE_CALL( cutStopTimer( timer));
    printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutDeleteTimer( timer));
    //各種メモリを解放
    free( h_idata);
    CUDA_SAFE_CALL(cudaFree(d_idata));

    //終了処理
    CUT_EXIT(argc, argv);
    return 0;
}

カーネル部分（diff1d_kernel.cu）は、

__global__ void diff1dKernel( float* g_idata) 
{
    extern __shared__ char sharedmem[]; 
    float * sdata = (float *) sharedmem; 
    float * sdata2 = (float *) (sharedmem+sizeof(float)*100);

    // スレッドIDを取得
    const unsigned int tid = threadIdx.x;

    //グローバルメモリから入力データの読み込み
    sdata[tid] = g_idata[tid];
    __syncthreads();

    //ここで計算を行う
    const float Dfu=1;
    const float dt=0.2;
    const float dx=1;
    float Dfudtdx2=Dfu*dt/(dx*dx);
    for (int n=0;n<100000;n++)
    {
        if (tid==0)
            sdata2[tid]=sdata[tid]+(2*sdata[tid+1]-2*sdata[tid])*Dfudtdx2;
        else if (tid==99)
            sdata2[tid]=sdata[tid]+(2*sdata[tid-1]-2*sdata[tid])*Dfudtdx2;
        else
            sdata2[tid]=sdata[tid]+(sdata[tid-1]+sdata[tid+1]-2*sdata[tid])*Dfudtdx2;
        __syncthreads();
        sdata[tid]=sdata2[tid];
    }

    //グローバルメモリに結果を書き込む
    g_idata[tid] = sdata[tid];
}

コンパイルは、

nvcc diff1d.cu -lcutil32 -lkernel32 -o diff1d.exe

で行った。
当方で3回実行した際の平均実行時間は7374.1msであった。

さて、次にまずこのコードに、非同期処理を使わずに、単純にCPUで実行するコードを加えてみる。

#include <stdio.h>
#include <cutil.h>
#include "diff1d_kernel.cu"
int main( int argc, char** argv) 
{
    //デバイスの初期化
    CUT_DEVICE_INIT(argc, argv);

    //タイマーを作成して計測開始
    unsigned int timer = 0;
    CUT_SAFE_CALL( cutCreateTimer( &timer));
    CUT_SAFE_CALL( cutStartTimer( timer));

    //メインメモリ上にfloat型のデータを100個生成する
    float* h_idata = (float*) malloc(sizeof( float) * 100);
    //初期条件をセット
    for( int i = 0; i < 50; i++) 
        h_idata[i] = 0;
    for( int i = 50; i < 100; i++) 
        h_idata[i] = 1;

    //デバイス上（ビデオカードのこと）にも同じくfloat型100個分のメモリを確保する
    float* d_idata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_idata, sizeof( float) * 100 ));

    dim3  grid( 1, 1, 1);
    //100は100個並列であるため
    dim3  threads(100, 1, 1);
    
    //メインメモリからデバイスのメモリにデータを転送する
    CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, sizeof( float) * 100 , cudaMemcpyHostToDevice) );

    for (int t=0;t<100;t++)
    {
        //ここでGPUを使った計算が行われる
        diff1dKernel<<< grid, threads, sizeof( float) * 200 >>>( d_idata);
        //デバイスからメインメモリ上に実行結果をコピー
        CUDA_SAFE_CALL( cudaMemcpy( h_idata, d_idata, sizeof( float) * 100, cudaMemcpyDeviceToHost) );
        
        //CPUを使ってデータの合計値の100000倍の値を求め表示する
        float sum=0;
        for (int n=0;n<100000;n++)
        {
            for (int i=0;i<100;i++)
            {
                sum+=h_idata[i];
            }
        }
        printf("%f\n",sum);
    }
    //タイマーを停止しかかった時間を表示
    CUT_SAFE_CALL( cutStopTimer( timer));
    printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutDeleteTimer( timer));
    //各種メモリを解放
    free( h_idata);
    CUDA_SAFE_CALL(cudaFree(d_idata));

    //終了処理
    CUT_EXIT(argc, argv);
    return 0;
}

赤字の部分が追加したCPUで実行するコードである。コードの内容には特に実用的な意味はない。
3回の実行時間の平均は11546.53msであった。GPUの実行を部分を取り除いたコード、

#include <stdio.h>
#include <cutil.h>
#include "diff1d_kernel.cu"
int main( int argc, char** argv) 
{
    //デバイスの初期化
    CUT_DEVICE_INIT(argc, argv);

    //タイマーを作成して計測開始
    unsigned int timer = 0;
    CUT_SAFE_CALL( cutCreateTimer( &timer));
    CUT_SAFE_CALL( cutStartTimer( timer));

    //メインメモリ上にfloat型のデータを100個生成する
    float* h_idata = (float*) malloc(sizeof( float) * 100);
    //初期条件をセット
    for( int i = 0; i < 50; i++) 
        h_idata[i] = 0;
    for( int i = 50; i < 100; i++) 
        h_idata[i] = 1;

    //デバイス上（ビデオカードのこと）にも同じくfloat型100個分のメモリを確保する
    float* d_idata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_idata, sizeof( float) * 100 ));

    dim3  grid( 1, 1, 1);
    //100は100個並列であるため
    dim3  threads(100, 1, 1);
    
    //メインメモリからデバイスのメモリにデータを転送する
    CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, sizeof( float) * 100 , cudaMemcpyHostToDevice) );

    for (int t=0;t<100;t++)
    {
        //CPUを使ってデータの合計値の100000倍の値を求め表示する
        float sum=0;
        for (int n=0;n<100000;n++)
        {
            for (int i=0;i<100;i++)
            {
                sum+=h_idata[i];
            }
        }
        printf("%f\n",sum);
    }
    //タイマーを停止しかかった時間を表示
    CUT_SAFE_CALL( cutStopTimer( timer));
    printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutDeleteTimer( timer));
    //各種メモリを解放
    free( h_idata);
    CUDA_SAFE_CALL(cudaFree(d_idata));

    //終了処理
    CUT_EXIT(argc, argv);
    return 0;
}

このコードの3階の平均実行時間は、4189.17msであり、4189.17（CPUのみ）+7374.1ms（GPUのみ）がおよそ11546.53ms（GPU+CPU）となっているのでそれぞれ順にシリアルに実行されていることがわかる。
次はいよいよ、GPUとCPUをパラレルに実行するコードである。

#include <stdio.h>
#include <cutil.h>
#include "diff1d_kernel.cu"
int main( int argc, char** argv) 
{
    //デバイスの初期化
    CUT_DEVICE_INIT(argc, argv);

    //タイマーを作成して計測開始
    unsigned int timer = 0;
    CUT_SAFE_CALL( cutCreateTimer( &timer));
    CUT_SAFE_CALL( cutStartTimer( timer));

    //メインメモリ上にfloat型のデータを100個生成する
    float* h_idata = (float*) malloc(sizeof( float) * 100);
    //初期条件をセット
    for( int i = 0; i < 50; i++) 
        h_idata[i] = 0;
    for( int i = 50; i < 100; i++) 
        h_idata[i] = 1;

    //デバイス上（ビデオカードのこと）にも同じくfloat型100個分のメモリを確保する
    float* d_idata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_idata, sizeof( float) * 100 ));

    dim3  grid( 1, 1, 1);
    //100は100個並列であるため
    dim3  threads(100, 1, 1);
    
    //メインメモリからデバイスのメモリにデータを転送する
    CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, sizeof( float) * 100 , cudaMemcpyHostToDevice) );

    for (int t=0;t<100;t++)
    {
        cudaEvent_t fin;
        CUDA_SAFE_CALL( cudaEventCreate(&fin));
        //ここでGPUを使った計算が行われる
        diff1dKernel<<< grid, threads, sizeof( float) * 200 >>>( d_idata);
        cudaEventRecord(fin, 0);

        //CPUを使ってデータの合計値の100000倍の値を求め表示する
        //この部分がGPUの計算と平行して行われる。
        float sum=0;
        for (int n=0;n<100000;n++)
        {
            for (int i=0;i<100;i++)
            {
                sum+=h_idata[i];
            }
        }
        printf("%f\n",sum);
        //diff1dKernelが終わるまで待つ
        while( cudaEventQuery(fin) == cudaErrorNotReady ){}
        //デバイスからメインメモリ上に実行結果をコピー
        CUDA_SAFE_CALL( cudaMemcpy( h_idata, d_idata, sizeof( float) * 100, cudaMemcpyDeviceToHost) );
        CUDA_SAFE_CALL( cudaEventDestroy(fin) );
    }
    //タイマーを停止しかかった時間を表示
    CUT_SAFE_CALL( cutStopTimer( timer));
    printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutDeleteTimer( timer));
    //各種メモリを解放
    free( h_idata);
    CUDA_SAFE_CALL(cudaFree(d_idata));

    //終了処理
    CUT_EXIT(argc, argv);
    return 0;
}

赤字で書いた部分が、今回追加した部分である。もともと、__global__で始まるカーネルの実行は非同期で行われる。
すなわち、カーネルの実行（<<<>>>のある行）は終了する前に、すぐに次の行に実行が移る。
cudaEventRecord(fin, 0);も非同期で行われるが、これは前に非同期処理のカーネルの実行が終わったときに実行される。
以上の2つの行が読み込まれた後、GPUが計算中に、CPUの計算が実行される。
cudaEventRecord(fin, 0)が実行されたかどうかは、cudaEventQuery(fin) == cudaErrorNotReadyで判定できる。
具体的には、while文でGPUの計算が終わるのを待っている。

さて、このコードの実行時間であるが、3回の平均は7395.93msであった。これは基準としたGPUのみのコードの平均時間7374.1msとほぼ同じであり、CPUの計算が平行して行われたことがわかる。

なお、非同期で行われないcudaMemcpyをCPU計算部分より前に持ってくると、

       diff1dKernel<<< grid, threads, sizeof( float) * 200 >>>( d_idata);
       CUDA_SAFE_CALL( cudaMemcpy( h_idata, d_idata, sizeof( float) * 100, cudaMemcpyDeviceToHost) );//GPUコードが終わるまで先に進まない

この部分で、カーネルの実行が終わるのを待つため、非同期で実行できない。

なおより実用的なコードとしては例えば、次のような、GPU実行中に結果をファイルに書き込むなどがある。

#include <stdio.h>
#include <cutil.h>
#include "diff1d_kernel.cu"
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));

    //メインメモリ上にfloat型のデータを100個生成する
    float* h_idata = (float*) malloc(sizeof( float) * 100);
    //初期条件をセット
    for( int i = 0; i < 50; i++) 
        h_idata[i] = 0;
    for( int i = 50; i < 100; i++) 
        h_idata[i] = 1;

    //デバイス上（ビデオカードのこと）にも同じくfloat型100個分のメモリを確保する
    float* d_idata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_idata, sizeof( float) * 100 ));

    dim3  grid( 1, 1, 1);
    //100は100個並列であるため
    dim3  threads(100, 1, 1);
    
    //メインメモリからデバイスのメモリにデータを転送する
    CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, sizeof( float) * 100 , cudaMemcpyHostToDevice) );

    for (int t=0;t<100;t++)
    {
        cudaEvent_t fin;
        CUDA_SAFE_CALL( cudaEventCreate(&fin));
        //ここでGPUを使った計算が行われる
        diff1dKernel<<< grid, threads, sizeof( float) * 200 >>>( d_idata);
        cudaEventRecord(fin, 0);

        //実行結果を非同期でファイルに書き込む
        for (int i=0;i<100;i++)
        {
            fprintf(fp,"%f\t",h_idata[i]);
        }
        fprintf(fp,"\n");
        //diff1dKernelが終わるまで待つ
        while( cudaEventQuery(fin) == cudaErrorNotReady ){}
        //デバイスからメインメモリ上に実行結果をコピー
        CUDA_SAFE_CALL( cudaMemcpy( h_idata, d_idata, sizeof( float) * 100, cudaMemcpyDeviceToHost) );
        CUDA_SAFE_CALL( cudaEventDestroy(fin) );
    }
    //タイマーを停止しかかった時間を表示
    CUT_SAFE_CALL( cutStopTimer( timer));
    printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutDeleteTimer( timer));
    //各種メモリを解放
    free( h_idata);
    CUDA_SAFE_CALL(cudaFree(d_idata));

    //終了処理
    fclose(fp);
    CUT_EXIT(argc, argv);
    return 0;
}


本文中のFC4はFedora ProjectのFedora Core 4を、FC5はFedora Core 5を、FC6はFedora Core 6をopenSUSEはNovellのSUSE Linux OSSを表します。Fedora7以降は、単にFedora7、Fedora8、Fedora9、Fedora10、Fedora11、Fedora12、Fedora13、Fedora14、Fedora15と表示しています。Ubuntuは、必要に応じて10.04、11.04のようにバージョン番号をつけて区別しています。MandrivaはMandriva Linuxを表します。

ここに登場するドメイン名やIPアドレスなどはフィクションです。実在の人物・団体等とは一切関係がありません。
実際に使用する際は、各自の環境に合わせて書き換えてください。
もし何か間違いなどありましたらこちらからご連絡ください
リンクに許可は不要です。
Copyright (C) 2011 Chikuma Engineering Co., Ltd. All Rights Reserved.