Конфликты банка общей памяти в CUDA: как память выровнена по банкам

Question

Насколько я понимаю, разделяемая память делится на банки и доступ нескольких потоков к одному элементу данных в одном банке вызовет конфликт (или рассылку).

На данный момент я выделяю довольно большой массив, который концептуально представляет несколько пар из двух матриц:

__shared__ float A[34*N]

Где N - количество пар, первые 16 чисел пары представляют собой одну матрицу, а следующие 18 чисел - вторые.

Дело в том, что доступ к первой матрице не конфликтует, но доступ ко второй имеет конфликты. Эти конфликты неизбежны, однако я думаю, что, поскольку вторая матрица равна 18, все будущие матрицы будут смещены по отношению к банкам, и, следовательно, возникнет больше конфликтов, чем необходимо.

Это правда, если так, как я могу избежать этого?

Каждый раз, когда я распределяю разделяемую память, она начинается в новом банке? Так что потенциально я мог бы сделать

__shared__ Apair1[34]
__shared__ Apair2[34]
...

Есть идеи?

Спасибо

Answer 1

Если ваши пары матриц хранятся смежно, и если вы обращаетесь к элементам линейно по индексу потока, то у вас не будет конфликтов банков общей памяти.

Другими словами, если у вас есть:

A[0]  <- mat1 element1
A[1]  <- mat1 element2
A[2]  <- mat1 element3
A[15] <- mat1 element16
A[16] <- mat2 element1
A[17] <- mat2 element2
A[33] <- mat2 element18

И вы получаете доступ к этому, используя:

float element;
element = A[pairindex * 34 + matindex * 16 + threadIdx.x];

Тогда смежные потоки обращаются к смежным элементам в матрице, и у вас нет конфликтов.

В ответ на ваши комментарии (ниже) кажется, что вы ошибаетесь в своем понимании. Это правда, что существует 16 банков (в нынешних поколениях, 32 в следующем поколении, Fermi), но последовательные 32-битные слова находятся в последовательных банках, то есть адресное пространство чередуется между банками. Это означает, что при условии, что у вас всегда есть индекс массива, который можно разложить до x + threadIdx.x (где x не зависит от threadIdx.x или, по крайней мере, является постоянным для групп из 16 потоков), у вас не возникнет конфликтов банков.

Когда вы получаете доступ к матрицам дальше по массиву, вы все равно получаете к ним доступ непрерывно, и, следовательно, у вас не будет банковских конфликтов. Конфликты между банками возникают только при начале доступа к несмежным элементам.

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

Answer 2

Банки настроены так, что каждый последующий 32 бита находится в следующем банке. Таким образом, если вы объявляете массив с 4-байтовыми числами с плавающей точкой, каждый последующий объект с плавающей точкой в ​​массиве будет находиться в следующем банке (по модулю 16 или 32, в зависимости от вашей архитектуры). Я предполагаю, что у вас есть вычислительная способность 1.x, поэтому у вас есть банк шириной 16.

Если у вас есть массивы 18 и 16, все может быть смешно. Вы можете избежать банковских конфликтов в массиве 16x16, объявив его как

__shared__ float sixteen[16][16+1]

, который позволяет избежать конфликтов банков при доступе к элементам транспонирования с помощью threadIdx.x (как я полагаю, вы делаете, если у вас возникают конфликты). При доступе к элементам, скажем, в первой строке матрицы 16x16 все они будут находиться в 1-м банке. То, что вы хотите сделать, это иметь каждый из них в последовательном банке. Паддинг делает это для вас. Вы обрабатываете массив точно так же, как и раньше, как шестнадцать [строка] [столбец], или аналогично для уплощенной матрицы, как шестнадцать [строка * (16 + 1) + столбец], если хотите.

Для случая 18x18, когда вы получаете доступ в транспонировании, вы двигаетесь с равномерной скоростью. Ответ снова заключается в том, чтобы дополнить 1.

__shared__ float eighteens[18][18+1]

Так что теперь, когда вы получите доступ к транспонированию (скажем, к элементам в первом столбце), он получит доступ как (18 + 1)% 16 = 3, и вы получите доступ к банкам 3, 6, 9, 12, 15, 2, 5, 8 и т. Д., Чтобы не возникало конфликтов.

Конкретный сдвиг выравнивания из-за наличия матрицы размера 18 не проблема, потому что начальная точка массива не имеет значения, это только порядок, в котором вы обращаетесь к нему. Если вы хотите сгладить массивы, которые я предложил выше, и объединить их в 1, это нормально, если вы обращаетесь к ним аналогичным образом.