シェアードメモリ

シェアードメモリについて

シェアードメモリとは†

シェアードメモリ(shared memory)はブロック内のスレッドが共用して使えるオンチップメモリで， グローバルメモリに比べて非常に高速です．シェアードメモリはバンクと呼ばれるメモリモジュールに分割されており， スレッド間のバンク・コンフリクトがなければレジスタアクセスと同じ速さで処理できます．

確保方法†

シェアードメモリ変数，配列はCUDAカーネル上で宣言します．
静的な確保(コンパイル時に配列の長さが決まる)の場合，

  1

__shared__ float x[16];

動的な確保(実行時に配列の長さが決まる)の場合は，カーネル呼び出し側で <<< ... >>> の第3引数で指定します．

  1
  2
  3
  4
  5
  6
  7
  8
  9
 10
 11
 12
 13
 14
 15
 16
 17
 18

__global__
void KernelFunc()
{
    extern __shared__ float x[];
 
           ...
}
 
__host__
void HostFunc(int shared_size)
{
       ...
 
    KernelFunc<<< grid, block, sizeof(float)*shared_size >>>();
 
       ...
}

シェアードメモリの大きさ†

デバイスプロパティを取得し，sharedMemPerBlockで確認できます．

  1
  2
  3
  4

    cudaDeviceProp dev;
    cutilSafeCall(cudaGetDeviceProperties(&dev, 0));
 
    printf(" shared memory / block : %d (KB)\n", dev.sharedMemPerBlock/1024);

バンク・コンフリクト(Programing guide 5.1.2.5節参照)†

広いメモリバンド幅を達成するためにシェアードメモリはバンクと呼ばれるサイズが同じメモリモジュールに分割されています． これらのバンクは同時並列にアクセスすることができ，例えば， n回のメモリアクセス場所が全て異なるバンクであり，これらが同時にリード or ライトされた場合， 単一メモリモジュールの場合と比べてn倍広いバンド幅となります． しかし，もしもリクエストしたメモリアドレスが同じバンク上にあった場合， バンク・コンフリクトが発生し，そのアクセスは直列に実行されます．