Разделяемая (shared) память является очень эффективным средством оптимизации за счет очень быстрого доступа (в 100 раз быстрее чем глобальная память). Однако, при неправильном использовании ее возможны конфликты банков, которые существенно замедляют быстродействие. В данной статье пойдет речь о том, как эти конфликты возникают, и как их избежать.
Конфликты возникают, когда 2 или более потоков из одного варпа (warp) (для устройств версии 2.0) или половины варпа (для устройстве версии 1.3 и ниже) осуществляют доступ к байтам, которые принадлежат разным 32 битным словам, находящимся в одном банке памяти. В случае конфликта доступ осуществляется последовательно. Количество потоков, обращающихся к банку, называется степенью конфликта. Если степень конфликта N, то доступ осуществляется в N раз медленнее, чем если бы конфликта не было.
На устройствах версии 1.x конфликта можно избежать, если несколько потоков осуществляют доступ к одному и тому же слову, принадлежащему одному и тому же банку, и только если этот запрос одиночный — в данном случае задействуется механизм широковещательного доступа.
На устройствах версии 2.x таких запросов может быть несколько и осуществятся они будут параллельно (разные потоки могут осуществлять доступ к разным байтам слова).
При 64 битном доступе конфликт банков возникает только если 2 или более потоков из любой из половин варпа осуществляют доступ по адресам, принадлежащим одному и тому же банку.
При 128 битном доступе как правило возникают конфликты банков второй степени.
Доступ разрядностью больше, чем 32 разбивается на запросы разрядностью 32, 64 и 128 бит.
Память распределяется по банкам таким образом, что каждое 32 битное слово в последовательности, последовательно назначается одному из 32 банков случае устройства версии 2.0 и 16 банков в случае устройства версии 1.3 и ниже. Соответственно номер банка можно рассчитать по следующей формуле:
Номер банка = (Адрес в байтах/4)%32 — для устройства версии 2.0
Номер банка = (Адрес в байтах/4)%16 — для устройства версии 1.x
Для устройств версии 1.x
1. 8 и 16 битный доступ
В данном примере первые 4 байта находятся в одном банке, поэтому первые 4 потока будут конфликтовать при доступе
Проблема решается добавлением избыточных данных (padding) и изменение схемы доступа:
Для 16-битного доступа:
В данном примере первые 2 шорта находятся в одном банке, поэтому первые 2 потока будут конфликтовать при доступе
Проблема решается аналогично 8-битному доступу:
2. 32-х битный доступ
Для данного типа доступа конфликты банков менее очевидны, но могут возникнуть при, например, такой схеме доступа:
В этом случае 0-й и 8-й поток читают из 0 и 1 банков соответственно, создавая таким образом конфликт 2-й степени.
Решить эту проблему можно к примеру так:
Для устройств версии 2.0
Из-за особенностей широковещательного доступа, 8 и 16 битные схемы доступа на данных устройствах не вызывают конфликтов банков, однако, конфликт может возникнуть в следующем случае:
Конфликт возникает, если s — четная. Если s — нечетная, но конфликтов не возникает.
Конфликты можно отследить, если воспользоваться макросом CUT_BANK_CHECKER( array, index), входящим в состав CUDA Utility Toolkit. Для этого необходимо пользоваться этим макросом для доступа к памяти и выполнять приложение в режиме эмуляции. При завершении приложения, будет напечатан отчет о конфликтах.
Например вот так:
Также, для отслеживания конфликтов можно пользоваться профайлером. Данная информация отображается в разделе warp serialize. Данный счетчик показывает количество варпов, которым необходимо сериализовывать свой доступ при адресации константной или разделяемой памяти, другими словами, этот счетчик показывает конфликты банков.
В заключение отмечу, что наиболее эффективная методика устранения конфликтов банков — это разработка схем доступа, которая минимизирует их возникновение и последующий анализ приложения профайлером (что никогда не лишнее).
Как возникают конфликты разделяемой памяти
Конфликты возникают, когда 2 или более потоков из одного варпа (warp) (для устройств версии 2.0) или половины варпа (для устройстве версии 1.3 и ниже) осуществляют доступ к байтам, которые принадлежат разным 32 битным словам, находящимся в одном банке памяти. В случае конфликта доступ осуществляется последовательно. Количество потоков, обращающихся к банку, называется степенью конфликта. Если степень конфликта N, то доступ осуществляется в N раз медленнее, чем если бы конфликта не было.
Механизм широковещательного доступа
На устройствах версии 1.x конфликта можно избежать, если несколько потоков осуществляют доступ к одному и тому же слову, принадлежащему одному и тому же банку, и только если этот запрос одиночный — в данном случае задействуется механизм широковещательного доступа.
На устройствах версии 2.x таких запросов может быть несколько и осуществятся они будут параллельно (разные потоки могут осуществлять доступ к разным байтам слова).
Особенности доступа на устройствах версии 2.0
При 64 битном доступе конфликт банков возникает только если 2 или более потоков из любой из половин варпа осуществляют доступ по адресам, принадлежащим одному и тому же банку.
При 128 битном доступе как правило возникают конфликты банков второй степени.
Доступ разрядностью больше, чем 32 разбивается на запросы разрядностью 32, 64 и 128 бит.
Как память распределяется по банкам
Память распределяется по банкам таким образом, что каждое 32 битное слово в последовательности, последовательно назначается одному из 32 банков случае устройства версии 2.0 и 16 банков в случае устройства версии 1.3 и ниже. Соответственно номер банка можно рассчитать по следующей формуле:
Номер банка = (Адрес в байтах/4)%32 — для устройства версии 2.0
Номер банка = (Адрес в байтах/4)%16 — для устройства версии 1.x
Примеры доступа к памяти, вызывающие конфликты
Для устройств версии 1.x
1. 8 и 16 битный доступ
__shared__ char shmem8[32];
char data = shmem8[threadIdx.x];
В данном примере первые 4 байта находятся в одном банке, поэтому первые 4 потока будут конфликтовать при доступе
Проблема решается добавлением избыточных данных (padding) и изменение схемы доступа:
__shared__ char shmem8[32*4];
char data = shmem8[threadIdx.x*4];
Для 16-битного доступа:
__shared__ short shmem16[32];
short data = shmem16[threadIdx.x];
В данном примере первые 2 шорта находятся в одном банке, поэтому первые 2 потока будут конфликтовать при доступе
Проблема решается аналогично 8-битному доступу:
__shared__ short shmem16[32*2];
short data = shmem16[threadIdx.x*2];
2. 32-х битный доступ
Для данного типа доступа конфликты банков менее очевидны, но могут возникнуть при, например, такой схеме доступа:
__shared__ int shmem32[64];
int data1 = shmem32[threadIdx.x*2];
int data2 = shmem32[threadIdx.x*2+1];
В этом случае 0-й и 8-й поток читают из 0 и 1 банков соответственно, создавая таким образом конфликт 2-й степени.
Решить эту проблему можно к примеру так:
__shared__ int shmem32_1[32];
__shared__ int shmem32_2[32];
int data1 = shmem32_1[threadIdx.x];
int data2 = shmem32_2[threadIdx.x];
Для устройств версии 2.0
Из-за особенностей широковещательного доступа, 8 и 16 битные схемы доступа на данных устройствах не вызывают конфликтов банков, однако, конфликт может возникнуть в следующем случае:
__shared__ int shared[64];
int data = shared[threadIdx.x*s];
Конфликт возникает, если s — четная. Если s — нечетная, но конфликтов не возникает.
Отслеживание конфликтов банков
NVIDIA Banck Checker
Конфликты можно отследить, если воспользоваться макросом CUT_BANK_CHECKER( array, index), входящим в состав CUDA Utility Toolkit. Для этого необходимо пользоваться этим макросом для доступа к памяти и выполнять приложение в режиме эмуляции. При завершении приложения, будет напечатан отчет о конфликтах.
Например вот так:
__shared__ int shared[64];
int data = CUT_BANK_CHECKER(shared, threadIdx.x*s);
CUDA Profiler
Также, для отслеживания конфликтов можно пользоваться профайлером. Данная информация отображается в разделе warp serialize. Данный счетчик показывает количество варпов, которым необходимо сериализовывать свой доступ при адресации константной или разделяемой памяти, другими словами, этот счетчик показывает конфликты банков.
Заключение
В заключение отмечу, что наиболее эффективная методика устранения конфликтов банков — это разработка схем доступа, которая минимизирует их возникновение и последующий анализ приложения профайлером (что никогда не лишнее).
