メニーコアと格闘ブログ:CUDAプログラム:グローバルメモリとシェアメモリを比較

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

ここではマルチプロセッサに搭載されている“シェアメモリ”にデータをコピーし、シェアメモリ上のデータを使って計算する方法を紹介します。

参考:GPGPU勉強会 1次元線形移流拡散方程式 CUDA プログラム 2[1]


[プログラムad2.cu]

__global__
void
integration_kernel0(float *u, float *du, int nx, float nu, float dx)
{
  __shared__ float sm[BLOCKDIM+2]; //①
  int i = threadIdx.x;
  int n = (BLOCKDIM-2) * blockIdx.x + i;
  if (n<=nx+1)
    sm[i] = u[n]; //②
  //__syncthreads();
  if (n<=nx && i>0 && i<BLOCKDIM-1)
    du[n-1] = -(sm[i+1]-sm[i-1])/(2*dx) + nu * (sm[i+1]-2*sm[i]+sm[i-1])/(dx*dx); //③
}


[CUDAプログラム例: ad1.cu] から変更が加わったのは、integration_kernel0です。

シェアメモリ上の配列として sm を定義しています(①)

そして sm にグローバルメモリ上の配列 u のデータをコピーしています(②)

後はこの sm の値を使って計算するだけ(③)。簡単です。



[結果]

計算量を100倍(NX =1000000)にして計算時間を比較してみました。タイムコマンドによる出力です。

・CPUのみの計算: 172 sec

・GPU,GlobalMEM: 7.49 sec

・GPU,ShareMEM: 6.28 sec


CPU、GPUの比較をすると段違いに高速化していますが、グローバルメモリとシェアメモリの比較となると、1.2倍程度の速度上昇にとどまっています。もっともっとメモリアクセスが頻発する計算だったら効果が発揮できると思うんですが…。


ナン 2009年8月21日 (金) 16:08 (JST)


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

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

<comments hideform="false" />


Comments

ノート:メニーコアと格闘ブログ:CUDAプログラム:グローバルメモリとシェアメモリを比較

個人用ツール