メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その3)デバイスによる並列計算

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

前回は1スレッドのみですが、デバイス側のスレッドを使った計算を行いました。

これを基にして、複数のスレッドを使った並列計算ができるよう、プログラムに改造を加えていきます。


必要な改造は

  • スレッド番号取得
  • 各スレッドの振舞定義

CUDAデバイス並列計算ソースコード メディア:Gfdm2d-0827-04.txt


[スレッド番号取得]

カーネル関数を呼び出すと、GPU側では、“同じ内容のカーネル関数”がスレッド数分動作します。


GPUで複数起動するスレッドには、スレッド同士が区別できるようにIDが振られます。スレッドにとって、今自分がIDいくつで起動しているのかを知るためには、以下の命令を実行します。

    int thidx = blockDim.x * blockIdx.x + threadIdx.x;

これで得られたIDを使えば、スレッド毎の振舞の内容を変えることができます。例えば、“スレッドIDが0~255番であれば計算領域の右側境界部分の特別な計算を行う”、“スレッドIDが256~511版は左側境界部分の特別な計算”、“スレッドIDが512~であれば、普通の計算”といった具合です。


次に、カーネル関数を呼び出す方法は、<<<dim nb, dim nt>>>の中身を変更します。この引数により、起動するスレッドの数を変えることができます。

  • nb:1グリッドを形成するブロック数
  • nt:1ブロックを形成するスレッドの数

というこうとです。この意味については後日記載します。ここではnb=256、nt=256として話を進めます。

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


[各スレッドの振舞定義]

計算領域の左側境界の計算を例に、これまでの計算方法との違いを表しまします。

“1格子点を1スレッドが計算する”という観点で考えると、ソースコードは以下のように書き換えられます。

   // 左端境界の計算
   //i = 0;
   //for(j=1 ; j<(NH_d)-1 ; j++){

       //Tw = T[(i+1)*(NH_d)+j]; // insulation .
       //Te = T[(i+1)*(NH_d)+j];
       //Ts = T[i*(NH_d)+(j-1)];
       //Tn = T[i*(NH_d)+(j+1)];
       //T_new[i*(NH_d)+j] = gamma*(Tw + Te + Ts + Tn - 4.0*T[i*(NH_d)+j]) + T[i*(NH_d)+j];
   //}

   if( thidx<NH_d-1 ){
       Tw = T[thidx+NH_d];
       Te = T[thidx+NH_d];
       Ts = T[thidx-1];
       Tn = T[thidx+1];
       T_new[thidx] = gamma*(Tw + Te + Ts + Tn - 4.0*T[thidx]) + T[thidx];
   }
他の計算についても同様

コメントアウトしてある部分が逐次用のソースコード、そして、その下に記載してあるのが、スレッド並列用の書き方で同内容の処理を表したものです。

スレッド並列用では、スレッド番号によって振舞を変えるため、先頭にIF文を入れました。この例では、“スレッド番号が NH_d-1以下の場合”には、境界条件として、TwとTeが同じ値を使います(断熱条件だから)。

また、“1格子点を1スレッドが計算する”なので、スレッド番号=座標と考えてi,jによるループをなくしました。

各スレッドは自分の計算が終わり次第、T_newに新しい温度を書きます。

ここで、同期関数syncthreads を呼んでいます。これにより、“全てのスレッドの計算終了を確認”してから、全格子点のデータを一斉更新します。


[実行時間について]

計算時間 9.605s

  • 格子数 256×256
  • ステップ数(ループ数) 100000ステップ
  • データ書き出し 1000ステップに1度
  • GPU NVIDIA Tesla C1060
  • CPU Intel(R) Xeon(R) E5540 @ 2.53GHz
  • MEM 24738040 kB(≒24GB)


メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その0)で1CPUでの計算時間を評価しましたが、それと比較するとおよそ10倍の速度向上率となりました。

102.295s / 9.605s = 10.65


また、メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その2)デバイスが行う処理を呼ぶにて、1スレッドによる計算を評価した結果と比較すると、

(4192s×2) / 9.605s = 872.955


[今後の展開]

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

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

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

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

ナン 2009年8月31日 (月) 11:42 (JST)


目次

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

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

<comments hideform="false" />


Comments

guest said ...

見逃していたらすみませんが、計算時間の測定方法と測定した範囲を載せてもらえませんか?

--guest 2009年9月2日 (水) 02:28 (UTC)

ナン said ...

コメントありがとうございます。

今のところtimeコマンドにより、プログラムの最初から最後まで時間を計測しています。

今後、gettimeofday()関数など使って、時間を計測してみたいと思います。

--ナン 2009年9月2日 (水) 04:30 (UTC)

個人用ツール