cuda - カーネルで共有メモリ変数を宣言する

Question

共有変数の仕組みについて質問があります。

このようにカーネルで共有変数を宣言すると __shared__ int array1[N] 、各アクティブ ブロックのすべての一意の共有メモリに、サイズ N の array1 のインスタンスが含まれるようになりました。つまり、各アクティブ ブロックのすべての共有メモリに N*sizeof(int) バイトが割り当てられたことを意味します。N*sizeof(int) は、計算能力 1.3 の GPU の場合、最大で 16KB でなければなりません。

したがって、上記が正しいと仮定し、次のようにホストで割り当てられた 2D スレッドと 2D ブロックを使用します。

dim3 block_size(22,22);
dim3 grid_size(25,25);

それぞれサイズ N*sizeof(int) の array1 の 25x25 インスタンスがあり、ブロックの各共有メモリにアクセスできるほとんどのスレッドは 22x22 です。これは私の最初の質問であり、回答がありました。

Q: array1 に値を代入すると

array1[0]=1;

次に、すべてのアクティブなブロックがその値を独自の共有メモリに即座に割り当てますか?

Answer 1

各ブロックは常に独自の共有メモリ配列を割り当てます。したがって、25x25 ブロックを起動すると、最終的に共有メモリに 25x25 配列が作成されます。

ただし、すべてのブロックが同時に存在するとは限らないため、これらすべての配列が同時に存在するわけではありません。アクティブなブロックの数は、実行されている GPU の実際のモデルによって異なります。GPU ドライバーは可能な限り多くの起動を試み、余分なブロックは前のブロックが作業を終了した後に実行されます。

N*sizeof(int) の最大値は、カードの計算能力と L1 キャッシュ構成によって異なります。8KB、16KB、32KB、および 48KB の間で変化する可能性があります。

最後の質問に答えるには、各共有配列は、対応するブロックに属するすべてのスレッドから見えます。あなたの場合、各共有配列は、対応する 22x22 スレッドによって表示されます。