Как загрузить значения во внешний общий массив

Question

при загрузке массивов из глобальной памяти в общую память переменные в общей памяти не делают то, что я хочу.

template<class T> 
__global__ void kernel(T *t1,T *t2) 
{ 
    int tid=threadIdx.x; 
    extern __shared__ T array1[]; 
    extern __shared__ T array2[]; 
    array1[tid]=t1[tid];//copy (1) 
    array2[tid]=t2[tid];//copy (2) 
    __syncthreads(); 
} 

все получается, что array1 [tid] = array2 [tid] = t2 [tid]. при замене места копирования (1) и копирования (2), результатом является array1 [tid] = array2 [tid] = t1 [tid]. только когда я удаляю «extern», результат получается тем, что я хочу (array1 [tid] = t1 [tid], array2 [tid] = t2 [tid]). Может ли кто-нибудь объяснить, почему?

Answer 1

Это:

extern __shared__ T array1[]; 
extern __shared__ T array2[]; 

не будет работать так, как вы думаете.

Эти указатели (array1 и array2) указывают на то же местоположение.

Если вы хотите иметь несколько массивов с использованием динамически распределенной общей памяти, вы должны следовать инструкциям, приведенным в programming guide. Что-то вроде этого:

extern __shared__ T array1[]; 
T * array2 = array1 + size_of_array_1; 

должно работать.

И обязательно пройти распределение по размерам в параметрах запуска ядра, что является достаточным байт для обоих размеров array1 и размера array2