Конфликт банка с общей памятью GPU

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.Также есть большой раздел в GeForce Gems 3 на эту тему - «39.2.3 Предотвращение конфликтов банков».

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

ОК, так что ваш первый пример :

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

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

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

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

N = 64

Тема 0: запись в 0, 31, 63 Тема 1: записьна 1, 32

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

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

Answer 2

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

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

---

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

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

Это может быть верно для старых графических процессоров, но для недавних графических процессоров с cc> = 2.x они не вызывают банковские конфликты, эффективно из-за механизма широковещания ( link ). Следующая цитата взята из РУКОВОДСТВО ПО ПРОГРАММИРОВАНИЮ CUDA C (v8.0.61) G3.3. Общая память .

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

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

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