Локальный указатель на общую память в 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>>>(); 
...