Пример оптимизации оптимизации CUDA

Question

Я пытаюсь реализовать оптимизацию на CUDA Reduction и должен был успешно завершиться до части 6. Спасибо всем вам за помощь. Чтобы получить полное представление о CUDA, мне также нужно закончить окончательную оптимизацию, как указано в слайде № 31, известном как каскадирование алгоритма.

Идея состоит в том, чтобы иметь 512 элементов на поток и суммировать их все до начала выполнения.

Я пробовал с подходом, где каждый мой поток имел доступ к непрерывным 512 номерам из памяти. К сожалению, его производительность была наихудшей. Я угадываю причину быть банковскими конфликтами, но до сих пор не понял точно. Может ли кто-нибудь из вас предложить причину такого поведения?

Я также отправляю образец кода, предоставленный Nvidia ниже.

unsigned int tid = threadIdx.x; 
unsigned int i = blockIdx.x*(blockSize*2) + threadIdx.x; 
unsigned int gridSize = blockSize*2*gridDim.x; 
sdata[tid] = 0; 
while (i < n) { 
    sdata[tid] += g_idata[i] + g_idata[i+blockSize]; 
    i += gridSize; 
} 
__syncthreads(); 

Несколько параметров остались неопределенными. Я могу сделать вывод, что blockSize равно числу потоков на блок. Но я не могу сделать вывод о важности переменной «gridSize». Каков подходящий способ доступа к памяти, чтобы мы получили лучшую производительность? Является ли это примером плавного доступа?

Благодарим вас за помощь и комментарий ниже, если у вас есть другие вопросы.

Answer 1

Это пример объединенного доступа. Лучший gridDim зависит от вашего оборудования. В зависимости от регистров на поток и максимальных потоков на один блок это значение должно быть некоторым множителем количества многопроцессоров, доступных на вашем оборудовании. Если ваша проблема достаточно велика, 8-кратный многопроцессорный счет является хорошим выбором для Kepler и 16x для Maxwell.

Answer 2

Скажите, что у вас есть blockDim.x = blockSize = 256 нитей на блок и gridDim.x = 32 блоков в сетке, и вы хотите уменьшить большой массив g_idata[8,192,000].

Тогда у вас есть 8192 темы. Давайте использовать

thread[x][y], x=0..31, y=0..255 

для представления этих потоков.

thread[x][y] Каждый загружается

g_idata[iter*512*x+y] and g_idata[iter*512*x+256+y], iter = 0 .. 999 

к совместно используемой памяти sdata.

Для каждой итерации iter все 8192 threads[x][y] загрузит gridSize = 16384 элементов из памяти графического процессора.

Это объединенный доступ к памяти, и это правильный способ доступа к памяти графического процессора.

Однако ваш путь, где каждый thread[x] читает data[i*x*512 .. i*(x+1)*512-1], i=0..., не является хорошим способом. На самом деле это самый неэффективный способ доступа к памяти GPU.