【CUDA】Shared memoryの動的な確保

Shared memoryは,オンチップでキャッシュされないグローバルメモリにアクセスするより100倍速い.
Shared memoryはブロックごとに配置されているおり,同じブロックのスレッドは,Shared memoryにアクセスできる.
そのため,競合を防ぐために__syncthreads()で,同期をとる処置がとられやすい.


このShared memoryは通常静的に宣言される.
たとえばこんな感じ

__shared__ int s[64];

でも,コンパイル時にサイズが分からない場合もある.
その時には,カーネル呼び出し時に指定したサイズ分だけ動的に確保できる.

//host側のカーネルの呼び出し
dynamicReverse<<<1, n, n*sizeof(int)>>>(d_d, n);
//device側
extern __shared__ int s[];

extern 修飾子を使う.

複数のShared memoryを確保したいとき

カーネル呼び出しのときに指定できるのは1つの数だけなので,
複数のshared memoryを使いたいときは,
shared memoryのサイズの合計を指定して,
カーネル内部で切り分ける必要がある.

extern __shared__ int s[];
int *integerData = s;                        // nI ints
float *floatData = (float*)&integerData[nI]; // nF floats
char *charData = (char*)&floatData[nF];      // nC chars

カーネル呼び出しは,

myKernel<<<gridSize, blockSize, nI*sizeof(int)+nF*sizeof(float)+nC*sizeof(char)>>>(...);

となる.

[参考]
Using Shared Memory in CUDA C/C++ | Parallel Forall