Конфликт банка с общей памятью графического процессора

Question

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

shared_a[threadIdx.x]=global_a[threadIdx.x]

это простое действие приводит к банковскому конфликту?

Предположим теперь, что размер массива больше, чем количество потоков, поэтому я сейчас использую это для копирования глобальной памяти в общую память:

tid = threadIdx.x;
for(int i=0;tid+i<N;i+=blockDim.x)
     shared_a[tid+i]=global_a[tid+i];

приведенный выше код приводит к банковскому конфликту?

Answer 1

Лучший способ проверить это - профилировать код с помощью "Compute Visual Profiler"; это идет с инструментарием CUDA. Также есть большой раздел в GPU Gems 3 на эту тему - "39.2.3 Как избежать банковских конфликтов".

" Когда несколько потоков в одной и той же деформации обращаются к одному и тому же банку, возникает конфликт банков, если все потоки деформации не обращаются к одному и тому же адресу в одном и том же 32-битном слове " - во-первых, существует 16 банков памяти по 4 байта каждый. Таким образом, по существу, если у вас есть какой-либо поток в памяти для чтения в пол варпа из тех же 4 байтов в банке общей памяти, у вас будут конфликты банков, сериализация и т. Д.

Итак, ваш первый пример:

Сначала предположим, что ваши массивы, например, типа int (32-битное слово). Ваш код сохраняет эти целые числа в разделяемой памяти, через любой полусоюз, который K-нить сохраняет в банке K-й памяти. Так, например, нить 0 первой половины деформации сохранит в shared_a[0] который находится в первом банке памяти, поток 1 сохранит в shared_a[1] каждая половина деформации имеет 16 потоков, которые отображаются в 16 банках по 4 байта. В следующей половине деформации первый поток теперь сохранит свое значение в shared_a[16], который снова находится в первом банке памяти. Так что, если вы используете 4-байтовое слово типа int, float и т. Д., То ваш первый пример не приведет к конфликту банков. Если вы используете 1-байтовое слово, такое как char, в первой половине потоков деформации 0, 1, 2 и 3 сохранят свои значения в первом банке разделяемой памяти, что приведет к конфликту банков.

Второй пример:

Опять же, все это будет зависеть от размера используемого вами слова, но для примера я буду использовать 4-байтовое слово. Итак, глядя на первую половину основы:

Количество потоков = 32

N = 64

Поток 0: запись в 0, 31, 63 Поток 1: запись в 1, 32

Все потоки в половине деформации выполняются одновременно, поэтому запись в общую память не должна вызывать конфликты банков. Я должен дважды проверить это, хотя.

Надеюсь, это поможет, извините за огромный ответ!

Answer 2

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

Профилирование этого кода с помощью NVIDIA Visual Profiler показывает, что при размере элемента меньше 32 и кратном 4 (4, 8, 12, ..., 28) последовательный доступ к общей памяти не приводит к конфликту банков. Однако размер элемента 32 приводит к конфликту банков.

---

Ответ от Ljdawson содержит устаревшую информацию:

... Если вы используете 1-байтовое слово, такое как char, в первой половине основы деформации все 0, 1, 2 и 3 сохранят свои значения в первом банке разделяемой памяти, что приведет к конфликту банков.

Это может быть верно для старых графических процессоров, но для недавних графических процессоров с cc >= 2.x они не вызывают банковские конфликты, эффективно из-за механизма вещания ( ссылка). Следующая цитата взята из Руководства по программированию CUDA C (v8.0.61) G3.3. Общая память.

Запрос на совместную память для деформации не приводит к конфликту банков между двумя потоками, которые обращаются к любому адресу в пределах одного и того же 32-разрядного слова (даже если два адреса попадают в один и тот же банк): в этом случае для доступа для чтения слово транслируется запрашивающим потокам (в одной транзакции может транслироваться несколько слов), а для доступа к записи каждый адрес записывается только одним из потоков (какой поток выполняет запись не определено).

Это, в частности, означает, что нет банковских конфликтов, если к массиву символов обращаются следующим образом, например:

   extern __shared__ char shared[];
   char data = shared[BaseIndex + tid];