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