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

目次 >> CUDA >> メモリアクセス

CUDA - メモリアクセス

CUDAは、コンピュータ上のRAM、グラフィックカード上のRAMなどいくつかのメモリをもつ。
それぞれのメモリは、アクセス速度やアクセスできる範囲などが違う。
具体的には、コンピュータ上のメインメモリには直接アクセスはできない。そのため、cudaMemcpyを使って、メインメモリからグラフィックカード上のメモリへ転送してやる必要がある。
一方、グラフィックカード上のメモリにもいくつか種類があり、一つはグローバルメモリであるが、これは同一グリッド内であれば、どのブロック、どのスレッドからもアクセスできる。
一方、シェアードメモリは同一ブロック内のスレッドからしかアクセスできないものの、グローバルメモリに比べて、きわめて高速にアクセスできる。
今回は、これらの違いがどの程度実行速度に影響するか調べてみた。なお今回使用しているのは、比較的ローエンドのGeforce8400GSである。
まずは、上記のプログラムから、ファイル書き込み部分を取り除き、計算量を100倍にして計算してみた。具体的なコードは、

#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<10000;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は元のプログラムと同等である。結果は1186.88msである。

次に、再びシェアードメモリを使わずに、直接グローバルメモリを使って計算させた場合である。コードは、

#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<10000;t++)
    {
        //ここでGPUを使った計算が行われる
        diff1dKernel<<< grid, threads, sizeof( float) * 100 >>>( 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;
}

CUDA_SAFE_CALL( cudaMemcpy( h_idata, d_idata, sizeof( float) * 100, cudaMemcpyDeviceToHost) );の行であるが、一時値を格納するためにfloat*100個分のみ確保した。

diff1d_kernel.cuはg_idataを使って直接計算している。

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

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

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

この場合の、5回の実行時間の平均は、4218.34msであった。3.5倍以上遅くなっている。

つぎに、呼び出し回数は変えずに(元のプログラムと同じ100回)、GPUコードの部分のループを100倍にして計算量を100倍にしてみた。これにより、デバイスからメインメモリ上に実行結果をコピーする回数や、グローバルメモリからシェアードメモリにコピーする回数が少なくなり、純粋な計算速度に近づくことができる。

#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;
}

これの実行結果であるが、5階の実行の平均は767.79msであった。これは、オリジナルのものに比べて1.55倍高速である。逆に言えばかなりの部分がメモリの読み書きに使われていたことがわかる。

なおシェアードメモリを使わずに、直接グローバルメモリを使って計算させた場合で、GPUコード内で計算量を100倍にしたものであるが、プログラムの5回の実行時間の平均値は、3620.55msであった。やはり、3.05倍遅くなっている。

まどめ
nvidiaのプログラミングガイドにもあるように、シェアードメモリにコピーしてから計算させるのがもっとも高速になることがわかる。


最終更新日


本文中の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は、必要に応じて17.04、18.04のようにバージョン番号をつけて区別しています。

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