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

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

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

メモリアクセスについては、メニーコアと格闘ブログ:CUDAプログラム:デバイスの構成と、メモリアクセスに概念図を示しました。

これまでは、Hostから直接アクセスできるグローバルメモリを使って計算を行ってきました。グローバルメモリは全てのスレッド(プロセッサ)から読み書きでき、並列化のための特別な処理も少なく容量が多いため、“とりあえずGPGPUを体験したい!”という場合、比較的お手軽に利用できる方法です。

一方、GPUのマルチプロセッサに搭載されている“共有メモリ”領域は、小容量ながらグローバルメモリと比較して“数百倍”高速にスレッドプロセッサからアクセスできるとの説明があります。これを使わない手はありません。


共有メモリには以下のような特徴があります。

  • 極めて低いレイテンシ
  • 実行時、またはコンパイル時に指定
  • 同じスレッドブロック内のすべてのスレッドがアクセス可能
  • 生存期間はスレッドブロック(≒カーネル関数内)


ここからは、これまでのプログラムを基に、共有メモリを使用する方法を試行します。


共有メモリは同じブロック内のスレッドであれば、情報共有することができます。

言いかえれば、ブロックが異なると、情報を共有することが難しくなります。

ファイル:CUDAを使ってプログラムとC言語プログラムの比較その3-04.jpg



[カーネル呼び出し]

ここでは、

  • 計算領域を16x16=256ブロック(2次元))に分割し、
  • それぞれのブロックは16x16=256スレッド(2次元))

で、構成されていると考えて、プログラムしていきます。

ファイル:メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その4)共有メモリ.jpg


以前、メニーコアと格闘ブログ:CUDAで実装熱伝導解析(その2)デバイスが行う処理を呼ぶで、カーネル関数呼び出しについて書きました。

実は、ここで指定したブロック数、スレッド数は2次元で指定することができます。つまり上で図示した形でブロックとスレッドを配置することができます。具体的には以下のように指定します。

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


[カーネル関数部分]

共有メモリを定義するためには、以下のように変数宣言します。

__shared__ float T_sm[18*18];

カーネル関数内で“__shared__”と修飾して、変数宣言するだけです。

この共有メモリに、グローバルメモリ内のデータをコピーすることで、計算に必要なメモリアクセスの時間を短縮することができます。

次回は引き続き、共有メモリを使った計算を実装していきます。


ナン 2009年9月1日 (火) 18:32 (JST)


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

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

<comments hideform="false" />


Comments

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

個人用ツール