本地指針指向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>>>(); 
...