Как стать автором
Обновить

Комментарии 6

Спасибо за статью.
Сами мы с этими банками уже разобрались, но если бы ваша статья раньше появилась, то у нас бы меньше времени ушло на разруливание конфликтов.
еще способ:
пусть у нас есть двумерный буффер, доступ из последовательных потоков (полу-)варпа к которому осуществляется по столбцам, а не по строкам (транспонированный доступ)
__shared__ int buf[N][N]; // N = 16 для 1.x, N = 32 для 2.x
int v = buf[threadIdx.x][threadIdx.y]
очевидно что потоки одного (полу-)варпа будут обращаться к одному столбцу, что вызовет конфликт N-ого порядка.
для предотвращения этого можно использовать след. прием:
__shared__ int buf[N][N+1];
добавление одного неиспользуемого столбца избавляет нас от конфликтов вообще в данном случае.
(возможно я где-то напутал со строками/столбцами, но идея надеюсь понятна).
Все же оптимизацию CUDA-программы надо начинать не с конфликтов банков, а с определением паттерна доступа к глобальной памяти. Затем приоритетным этапом является определение мест нерационального бранчинга, и только после этого можно сосредоточиться на тюнинге на уровне регистров, банков и их конфликтов. Это потому, что если вы вычистите все конфликты, но не будут соблюдены правила объединенных запросов к глобальной памяти, то быстродействие программы будет отличаться от максимально-возможного в несколько раз.

Теперь немного замечаний по тексту:

1. Конфликт банков на Fermi гораздо сложнее вызвать, особенно при работе с маленькими типами char и short. Можно любым числом потоков адресовать один банк (разные его байты, но в рамках одного слова).

2. При необходимости обрабатывать один байт на поток на архитектурах до Fermi можно использовать т.н. bit-twiddling hack, который заключается в подмене threadIdx.x на такую пермутацию, которая позволяет обходить конфликт банков. Идея заключается в произведении циклического сдвига в младших 4 (Например для пермутации линейного блока из 64 потоков в группы по 16):

__device__ DEVICEINLINE int permuteThreads8u(int x)
{
return (x >> 4) + ((x & 0xF) << 2);
}

3. Счетчик warp serialize показывает именно количество сериализаций варпов, случившихся в железе по факту исполнения. Но складывается он не только из конфликта банков. Например, любое ветвление (и в частности те, про которые пишется в branching и divergent branching) вызывает одну сериализацию. Также есть менее значительные (подвластные программисту) явления, вызывающие нарастание этого счетчика. Вообще, счетчики профилировщика рекомендуется оценивать в динамике, а не конкретные их значения. Т.е. лучще уменьшать плохие счетчики (uncoalesced, divergent branch) и увеличивать хорошие (occupancy, coalesced, cache hit rate)
Спасибо за замечания, однако они все отображены в тексте так и ли иначе.

То что начинать надо с паттерна я написал в заключении. статья не про общую оптимизацию, а про конфликты банков.
Про конфликты на Ферми с маленькими типами также описано. Цитата: «Из-за особенностей широковещательного доступа, 8 и 16 битные схемы доступа на данных устройствах не вызывают конфликтов банков, однако, конфликт может возникнуть в следующем случае...».

Warp serialize не связан с бранчами, как бы логично это не казалось, читайте документацию:

warp serialize: If two addresses of a memory request fall in the same memory bank, there is a bank conflict and the access has to be serialized. This counter gives the number of thread warps that serialize on address conflicts to either shared or constant memory.

Спасибо!
Я немного недопонял пример:

__shared__ int shmem32[64];
int data1 = shmem32[threadIdx.x*2];
int data2 = shmem32[threadIdx.x*2+1];
В этом случае 0-й и 8-й поток читают из 0 и 1 банков соответственно, создавая таким образом конфликт 2-й степени.

Каким образом потоки 0 и 8 читая из разных банков создадут bank conflict?
Всего 16 банков. Притом, 0-ой элемент массива попадает в 0-ой банк, 1->1,… 15->15, 16->0.
Отсюда следует, что 0-ой поток читает 0-ой элемент из 0-го банка и 8-й поток читает 16-й элемент из 0-го банка.
Зарегистрируйтесь на Хабре , чтобы оставить комментарий

Публикации

Истории