cuda - CUDAの共有メモリへのローカルポインタ

Question

スレッドブロック内の各スレッドに共有メモリへの独自のポインタを持たせるにはどうすればよいですか？私はそのようなポインタの宣言のいくつかの例を見つけました：

int __shared__ *p;
__shared__ int array[256];

p = &array[threadId];

これは正しいですか、それとも別の方法がありますか？

Answer 1

いいえ、それは正しい方法ではありません。このサンプルコードでは、pが共有されているため、ブロック内のすべてのスレッドが同じメモリにアクセスしようとします。threadIdブロックを持つ一意のスレッドインデックスである場合は、次のように実行できます。

int *p;
__shared__ int array[256];

p = &array[threadId];

この場合、コンパイラはレジスタまたはスレッドローカルメモリのいずれかを使用arrayして、ブロック内の各スレッドの静的共有メモリ割り当てに要素の一意のアドレスを格納します。

Answer 2

あなたが正しいです。より良い方法は、共有メモリの動的割り当てです。例は仲間とし​​てです：

void __global__ test(){
extern __shared__ int s[];
int *p = &s[xx];

}

...
test<<<x,y, shared memory length>>>();
...