Как разделить общее значение между потоками в данном блоке?

Question

У меня есть ядро, которое для каждого потока в данном блоке вычисляет цикл for с разным количеством итераций. Я использую буфер размером N_BLOCKS для хранения количества итераций, необходимых для каждого блока. Следовательно, каждый поток в данном блоке должен знать количество итераций, характерных для его блока.

Однако я не уверен, какой способ является лучшим (говоря о производительности), чтобы прочитать значение и распределить его среди всех других потоков. Я вижу только один хороший способ (скажите, пожалуйста, есть ли что-то лучше): сохраните значение в разделяемой памяти и прочитайте каждый поток. Например:

__global__ void foo( int* nIterBuf )
{
   __shared__ int nIter;

   if( threadIdx.x == 0 )
      nIter = nIterBuf[blockIdx.x];

   __syncthreads();

   for( int i=0; i < nIter; i++ )
      ...
} 

Есть ли другие лучшие решения? Мое приложение будет использовать много данных, поэтому я хочу добиться максимальной производительности.

Спасибо!

Answer 1

Значения, доступные только для чтения, которые одинаковы для всех потоков в блоке, вероятно, лучше всего хранить в массивах __constant__. В некоторых архитектурах CUDA, таких как Fermi (SM 2.x), если вы объявляете аргумент массива или указателя с помощью ключевого слова C ++ const и обращаетесь к нему равномерно в пределах блока (то есть индекс зависит только от blockIdx, а не threadIdx), то компилятор может автоматически выдавать ссылку на постоянную память.

Преимущество постоянной памяти в том, что она проходит через выделенный кеш, поэтому она не загрязняет L1, и если объем данных, к которым вы обращаетесь, на блок относительно невелик, после первого доступа в каждом блоке вы всегда должен попадать в кеш после начального обязательного промаха в каждом блоке потока.

Вам также не нужно будет использовать какую-либо общую память или переносить из глобальной в общую память.

Answer 2

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

Если чтение этих данных из разделяемой памяти на каждой итерации замедляет работу, и у вас все еще есть доступные регистры (см. Вычислительные возможности и спецификации вашего GPU), вы, возможно, можете попытаться сохранить копию этого значения в регистре каждого потока (используя локальная переменная).