メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その4)共有メモリ(あえて失敗例)

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

今回は、グローバルメモリにコピーしたデータを、ブロックの共有メモリに分配する方法について書きます。

結果的に計算は失敗していますが、その顛末をご覧頂いた上で、アドバイス等いただければと思います。

メディア:gfdm2d.txt


[カーネル呼び出し]

ここでは、境界条件の計算のために、1格子点分“太く”領域を確保しました。このご利益は後半記載します。

カーネル呼び出しを以下のようにします。カーネル呼び出しについては前回も記載しましたが、違う部分が分かるでしょうか。

solve<<<dim3(16,16),dim3(18,18)>>>(T_d, T_new_d, gamma_d, NH_d, NW_d, h_mid_d, temp_high_d);

ブロックの数は16×16=256、1ブロックを形成するスレッドの数は18×18になりました。イメージとしては以下のようになります。


ファイル:CUDAで実装熱伝導解析(その4)共有メモリ(あえて失敗例).jpg


赤枠部分がブロックを表します。枠内に記載した、0,1,16,17…は、計算領域の要素のインデックスを表します。


[スレッドにおいて共有メモリにアクセス]

今、赤枠で示した ブロック(インデックス(0,0)) の中の、18×18=324スレッドが起動したところを考えます。

これまで同様、カーネルの内容はすべてのスレッド(ここでは324スレッド)で実行されます。


まず、グローバルメモリから共有メモリに、データをコピーします。

共有メモリの定義の仕方は、“__shred__”修飾子を付けて変数を定義します。

   __shared__ float T_sm[324];

定義した共有メモリに、グローバルメモリからデータをコピーします。

ここで、スレッドのインデックスは、threadIdx.x,threadIdx.yにより知ることができますが、ブロックが変わるとこのインデックスはリセットされます。つまり、グローバルメモリと共有メモリで、インデックスを対応づける必要があります。

ここでは、“あるブロック”に属する“あるスレッド”のインデックスが、全体から見るとどこを指すのかを、thidxという変数名で以下の式から割り出しました。

   thidx = blockIdx.x*4128 + blockIdx.y*16 + threadIdx.x*258 + threadIdx.y;

共有メモリにデータをコピーするのは、変数同士の代入文により行います。

   T_sm[threadIdx.x.18+threadIdx.y] = T[thidx];

これでグローバルメモリTから共有メモリT_smに値をコピーできました。


[条件分岐を減らす努力]

それでは実際の計算について見てみます。

以前の計算と同様、1格子点1スレッドで計算を行うことには変わりありません。直行等間隔格子で差分法を行うので、ある点と、その周囲4点の情報を使用します。

ここで、“ある点”が計算領域の境界(上下左右の端だったり、角だったり)だった場合、“断熱条件”を表現するために、一つ内側の温度を、境界の外側の温度として使用します。

このとき、“もし今計算しているところが、領域の境界だったら”という条件分岐が登場しますが、“GPGPUを使用する計算は、条件分岐が多発すると処理が遅くなる”ということですので、if文の数を減らすため、はじめから一つ内側の温度を、外側(1格子点分太くしたところ)にコピーしておき、余計なif文なしに全部の格子点を一律に計算できる方法をとりました。これが、計算領域の範囲を1格子余分にとった理由です。

ここまでくれば、ある点と、その周囲4点の情報を共有メモリから取得して、計算を実行するだけです。

上記の理由から、計算する範囲は赤枠いっぱいではなく、青枠(一つ内側の格子点)について、一斉に計算を行います。計算した結果は計算した端から、再びグローバルメモリに書き込みます。


[同期がとれない罠]

私、ここに至るまで、“ブロック同士で同期がとれる”ものだと、思い込んでいました。

その結果がこのとおり、見事にブロックごとに同期がとれていません。


ファイル:CUDAで実装熱伝導解析(その4)共有メモリ(あえて失敗例)1.jpg


次回はブロック同士で同期を取り、実行速度のスピードアップを図ります。


ナン 2009年9月4日 (金) 17:52 (JST)


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

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

<comments hideform="false" />


Comments

ノート:メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その4)共有メモリ(あえて失敗例)

個人用ツール