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. Также есть отличный раздел в Графические камни 3 На этом - "39.2.3 Избегание банковских конфликтов".

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

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

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

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

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

Количество потоков = 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, в первой половине Warp Threads 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];