Объединяется ли неполный глобальный доступ к памяти?

Question

Это объединилось, если n < warpSize?

// In kernel
int x;
if (threadId < n)
  x = globalMem[threadId];

Такая ситуация возникает в последней итерации цикла, если некоторые N неделима warpSize, Должен ли я выполнять эти ситуации и распределять память устройства только по warpSize или это слилось как есть?

Answer 1

Если 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.