Хабр Курсы для всех
РЕКЛАМА
Практикум, Хекслет, SkyPro, авторские курсы — собрали всех и попросили скидки. Осталось выбрать!
Хабр доступен 24/7 благодаря поддержке друзей
Комментарии 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];
добавление одного неиспользуемого столбца избавляет нас от конфликтов вообще в данном случае.
(возможно я где-то напутал со строками/столбцами, но идея надеюсь понятна).
пусть у нас есть двумерный буффер, доступ из последовательных потоков (полу-)варпа к которому осуществляется по столбцам, а не по строкам (транспонированный доступ)
__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)
Теперь немного замечаний по тексту:
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.
Спасибо!
То что начинать надо с паттерна я написал в заключении. статья не про общую оптимизацию, а про конфликты банков.
Про конфликты на Ферми с маленькими типами также описано. Цитата: «Из-за особенностей широковещательного доступа, 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.
Спасибо!
Я немного недопонял пример:
Каким образом потоки 0 и 8 читая из разных банков создадут bank conflict?
__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?
Зарегистрируйтесь на Хабре, чтобы оставить комментарий
Понимание конфликтов банков разделяемой (shared) памяти в NVIDIA CUDA