Есть ли способ установить значение по умолчанию для массива совместно используемой памяти?

Question

Рассмотрим следующий код:

__global__ void kernel(int *something) {
extern __shared__ int shared_array[];     

// Some operations on shared_array here.

}

Можно ли установить целый shared_array в какое-то значение - например, 0 - без явного обращения к каждой ячейке в каком-то потоке?

Answer 1

Нет. Общая память не инициализирована. Вы должны как-то инициализировать это сами, так или иначе...

Из Руководства по программированию CUDA C 3.2, раздел B.2.4.2, параграф 2:

__shared__ переменные не могут иметь инициализацию как часть их объявления.

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

Answer 2

Вы можете эффективно инициализировать совместно используемые массивы параллельно, как это

// if SHARED_SIZE == blockDim.x, eliminate this loop
for (int i = threadIdx.x; i < SHARED_SIZE; i += blockDim.x) 
    shared_array[i] = INITIAL_VALUE;
__syncthreads();

Answer 3

Да, ты можешь. Вы можете указать, что первый поток в блоке устанавливает его, а другой нет, например:

extern __shared__ unsigned int local_bin[]; // Size specified in kernel call

if (threadIdx.x == 0) // Wipe on first thread - include " && threadIdx.y == 0" and " && threadIdx.z == 0"  if threadblock has 2 or 3 dimensions instead of 1.
{
    // For-loop to set all local_bin array indexes to specified value here - note you cannot use cudaMemset as it translates to a kernel call itself
}

// Do stuff unrelated to local_bin here    

__syncthreads(); // To make sure the memset above has completed before other threads start writing values to local_bin.

// Do stuff to local_bin here

В идеале вы должны выполнить как можно больше работы перед вызовом syncthreads, так как это позволяет всем другим потокам выполнять свою работу до завершения memset - очевидно, это имеет значение только в том случае, если работа может иметь совершенно разные времена завершения потока, например, если есть условное ветвление. Обратите внимание, что для цикла "установка" для потока 0 необходимо передать размер массива local_bin в качестве параметра ядру, чтобы вы знали размер итерируемого массива.

Оригинальная концепция источника