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

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), возможно, вы попытаетесь сохранить копию этого значения в каждом регистре потока (используя локальная переменная).