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)>>>(...);
となる.