CUDA 中指向共享内存的本地指针答案

【问题标题】：Local pointer to shared memory in CUDACUDA 中指向共享内存的本地指针
【发布时间】：2021-05-30 10:18:35
【问题描述】：

如何使线程块中的每个线程都有自己的指向共享内存的指针？我找到了一些声明此类指针的示例：

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

p = &array[threadId];

这是对的还是有其他方法？

【问题讨论】：

标签： cuda

【解决方案1】：

不，这不是正确的方法。在该示例代码中，p 是共享的，因此这意味着块中的每个线程都将尝试访问同一块内存。如果 threadId 是该块的唯一线程索引，您可以这样做：

int *p;
__shared__ int array[256];

p = &array[threadId];

在这种情况下，编译器将使用寄存器或线程本地内存为块中的每个线程在静态共享内存分配array 中存储元素的唯一地址。

【讨论】：

谢谢，我以为我在 [forums.nvidia.com/index.php?showtopic=35294] 中看到了它，但是有像 int *__shared__ p; 这样的声明它应该是存储在线程本地内存中并指向共享内存的指针。对吗？
不，这是不正确的。您可以放心地忽略该线程中的所有内容 - 从 2008 年开始，CUDA 是新的，编译器是原始的，没有多少人了解语言扩展的细节以及它们如何映射到硬件。
如何强制 p 存储在寄存器中？
@spurra，通过将 p 声明为局部变量。

【解决方案2】：

你是对的。更好的方法是动态分配共享内存。一个例子是：

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

}

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

【讨论】：