Объединяется ли неполный глобальный доступ к памяти?
Это объединилось, если n < warpSize?
// In kernel
int x;
if (threadId < n)
x = globalMem[threadId];
Такая ситуация возникает в последней итерации цикла, если некоторые N неделима warpSize, Должен ли я работать с этими ситуациями и распределять память устройства только делится на warpSize или это слилось как есть?
Решение
Если threadId вычисляется правильно, как описано в руководство по программированию cuda — нить иерархии, чем этот доступ будет объединен — это будет иметь место для threadId = threadIdx.x,
Для разных вычислительных архитектур объединение памяти немного отличается. Более подробную информацию можно найти на приложение G к руководству по программированию cuda.
В общем, вы можете сказать, что глобальные обращения к памяти объединяются, если ваши потоки захватывают последовательные элементы в памяти, начиная с адреса элемента, к которому обращается ваш первый поток.
Предположим, у вас есть массив с плавающей точкой.
float array[]
и ваш доступ к памяти выглядит таким образом
array[threadIdx.x == 0, threadId.x == 1, threadIdx.x == 2, ..., threadIdx.x == 31]
чем ваш доступ будет ограничен.
Но если вы обращаетесь к памяти таким образом (чередование, например)
array[threadIdx.x == 0, NONE, threadId.x == 1, NONE, threadIdx.x == 2, ..., NONE, threadIdx.x == 31]
чем ваш доступ не слился (NONE означает, что этот элемент массива не доступен ни одному потоку)
В первом случае вы получаете 128 последовательных байтов памяти. Во втором случае вы получаете 256 байтов. Во втором случае требуется две деформации для загрузки памяти из глобальной памяти вместо одной деформации для первого случая. Но в обоих случаях для следующих вычислений требуется только 32 элемента с плавающей точкой (128 байтов). Таким образом, ваша общая нагрузка в этом простом случае снизится с 1,0 до 0,5.
Другие решения
Других решений пока нет …