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

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

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

Ответ 1

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

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

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

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

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

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

Число потоков = 32

N = 64

Тема 0: Будет писать до 0, 31, 63 Тема 1: Будет писать 1, 32

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

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

Ответ 2

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

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

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

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

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

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

Это означает, в частности, что банковские конфликты отсутствуют, если доступ к массиву из char осуществляется следующим образом, например:
   extern __shared__ char shared[];
   char data = shared[BaseIndex + tid];