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

Я пытаюсь понять, как возникают конфликты банков.
если у меня есть массив размером 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];

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


c++ cuda gpgpu bank-conflict
person scatman    schedule 09.12.2010    source источник

Ответы (2)


arrow_upward
15
arrow_downward

Лучший способ проверить это — профилировать ваш код с помощью «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

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

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

person Ljdawson    schedule 09.12.2010
comment
на самом деле для второй части поток 0 запишет в 0,32, а поток 1 запишет в 1,33 и так далее..... пока последний поток 31 не запишет в 31,63. но спасибо за первую часть поста, очень познавательно - person scatman; 09.12.2010
comment
Отредактировано, чтобы отразить ваш комментарий, отвечает ли это на ваш вопрос? - person Ljdawson; 09.12.2010
comment
Обратите внимание, что на устройствах sm_20 и более поздних версиях имеется 32 банка, и доступ должен учитываться для каждого варпа, а не для полуварпа. - person Tom; 09.12.2010
comment
Что касается второй части вопроса, вы правы, конфликтов банков нет, поскольку поток 0 записывает в 0, 32, 64, а поток 1 в 1, 33, 65 и т. д. (небольшое исправление из вашего ответа). Это обычно пишется как for (int i = tid ; i < N ; i += blockDim.x) shared_a[i] = global_a[i]; - person Tom; 09.12.2010
comment
Так что же происходит, когда варп из 32 потоков (речь идет о картах sm_20 и выше) пытается записать массив из 32 целых чисел (2 байта на int = половина слова) в разделяемую память? Приведет ли это к конфликтам банков (таким образом, к сериализации)? - person ; 09.07.2014
comment
Извините за комментарий к старому вопросу. Я случайно наткнулся на этот вопрос по ссылке, и это смутило меня в случае с 1-байтовым элементом. Я думаю, что этот ответ нуждается в некотором обновлении, поскольку для последних графических процессоров последовательный доступ в 1 байт к общей памяти не приводит к конфликту банков. Я разместил ответ. Спасибо за прочтение! - person nglee; 01.06.2017
comment
@Madhatter Это зависит от того, распространяется ли доступ каждого потока на каждый банк или нет. В случае последовательного доступа к разделяемой памяти это не вызовет конфликта банков. - person nglee; 02.06.2017

arrow_upward
3
arrow_downward

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

Профилирование этого кода с помощью 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-битного слова (даже если два адреса попадают в один и тот же банк): в этом случае для доступа для чтения слово передается запрашивающим потокам (в одной транзакции может передаваться несколько слов), а для доступа на запись каждый адрес записывается только одним из потоков (какой поток выполняет запись, не определено).

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

   extern __shared__ char shared[];
   char data = shared[BaseIndex + tid];
person nglee    schedule 18.04.2017