メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その2)デバイスが行う処理を呼ぶ

出典: トータル・ディスクロージャ・サイト(事実をありのままに)

前回、デバイスにデータをコピーすることができたので、今回はデバイス側のスレッドを使って計算を行います。

メディア:gfdm2d-0827-04.txt



デバイスの処理は、ホスト側の処理とは別に記述します。

// デバイス側処理定義
__global__ void funcname( 仮引数1、仮引数2... )
{
    ....
}


// ホスト側処理
int main()
{
        ....

    // デバイス側処理呼び出し
    funcname<<<blocksize, threadsize>>>( 引数1、引数2... );

        ....
}


[デバイス側処理について]

デバイス側で行う処理には、関数修飾子を付けます。

  • __global__   :ホストから呼び出すことができる、デバイスで行う処理。“カーネル関数”と呼ばれます。
  • __device__   :デバイス側処理から呼ぶ、別のデバイス側処理。ホストからは呼び出すことができない。


“[ホスト側からの呼び出し]”

ホストからデバイスの処理を呼び出すには、以下の文を使います。

funcname<<<int blocksize, int threadsize>>>( 引数1、引数2... );

  • blocksize    :グリッドを構成するブロックサイズ
  • threadsize   :ブロックを構成するスレッドサイズ

デバイスでは、ここに指定した( blocksize × threadsize )だけの数のスレッドが起動します。


では、実際に使用したコードを抜粋します。 ここでは簡単のため、CPUで計算する場合と同じように、1スレッドで計算するため、blocksize、threadsizeともに1を指定します。

__global__ void solve(double* T,
                     double* T_new,
                     double gamma,
                     int NH_d,
                     int NW_d,
                     int h_mid,
                     double temp_high)
{
    ....
}

int main(){

    ....

       //solve(gamma);
       solve<<<1,1>>>(T_d, T_new_d, gamma, NH, NW, h_mid, temp_high);

   ....

}


このプログラムの指定方法では、起動するスレッド数が 1×1 =1スレッドとなります。つまり普通のCPUでの計算と同じ結果が得られます。 “カーネル関数”はホストから呼び出されると、その内容は、デバイス内全てのスレッドで一斉に実行されます。この、起動するスレッドの数は、カーネル関数呼び出し時の数字(blocksizeとthreadsize)の設定で、デバイスで起動するスレッド数を指定します。


引数については、T_d,T_new_d は、デバイス側のメモリ領域を指していますが、他のgamma,NH,NW,h_mid,temp_highについては、ホスト側のメモリ領域にあるデータを指定しています。配列など、“頻繁にアクセスするデータはあらかじめデバイス側にコピーしておく”といいみたいです。


[実行]

普通のCPUと同じように1スレッドで正しく計算できることを確認します。


計算時間 69m26.366s (=4192.366s)

*格子数                 256×256
*ステップ数(ループ数) 50000ステップ
*データ書き出し         1000ステップに1度

計算時間がかかるものの、デバイス側にデータや処理を渡し、正しく計算結果を得ることができました。今後はこのプログラムを基に複数スレッドを使って高速化していきます。


[今後の展開]

  • 複数のスレッドを起動して並列に計算

  今回、1スレッドですが、初めてGPUに計算をさせました。今後はこれを発展させる形で、GPUに搭載されているスレッドプロセッサを複数使って並列計算を行うよう、プログラムを改変していきます。。

  •  単精度計算と倍精度計算の比較

  GPUは240個の単精度計算ユニットと、3個の倍精度計算ユニットを持っています。GPUで行う計算を単精度の変えた場合どうなるのかについてみていきます。

  • メモリアクセスの方法を変更し、高速化を狙う

  使用するメモリ領域をグローバルメモリからシェアメモリを使用するようにプログラムを変更し、メモリアクセス時間を短縮していきます。


ナン 2009年8月28日 (金) 13:35 (JST)


この記事へのコメントをお寄せください

  • サイトへの書き込みに差し支えございましたら トータルディスクロージャーサイトサポート係へメールをお送りください
  • トータル・ディスクロージャ・サイトに投稿された文章と画像は、すべてその著作権がHPCシステムズ株式会社に帰属し、HPCシステムズ株式会社が著作権を所有することに同意してください。
  • あなたの文章が他人によって自由に編集、配布されることを望まない場合は、投稿を控えてください。
  • コメントを書き込む場合は名前にひらがなを織り交ぜてください。
  • あなたの投稿する文章と画像はあなた自身によって書かれたものであるか、パブリック・ドメインかそれに類する自由なリソースからの複製であることを約束してください。あなたが著作権を保持していない作品を許諾なしに投稿してはいけません!

<comments hideform="false" />


Comments

ノート:メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その2)デバイスが行う処理を呼ぶ

個人用ツール