общего банк память конфликта с массив символов

Question

Я понимаю, банковский конфликт при работе с 4-байтными типами данных, но я задаюсь , если мы получим какой-либо банковский конфликт (4-полосной/8-позиционные?) Со следующим кодом

__shared__ char shared[]; 
foo = shared[threadIdx.x]; 

Приведенный выше код приводит к 4 последовательным нитям в деформации, обращающейся к одному и тому же адресу слова в одном банке.

Будет ли подобный шаблон доступа к памяти привести к конфликту в банке для любого семейства устройств cuda? По-видимому, он работает только для старых карт, но я хочу подтвердить.

Мой вопрос может быть обобщен далее, что, если несколько потоков обращаются к одному и тому же адресу банковского адресата [8-байтовый или 4-байтовый], но для каждого из них требуется небольшая часть. будет ли аппаратное обеспечение обрабатывать такие запросы без каких-либо банковских конфликтов? Спасибо

Answer 1

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

От the documentation:

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

Для устройств, поддерживающих 8-байтовый банковский режим и находящихся в режиме 8-байтового банка, вышеупомянутый механизм вещания обобщается на 64-разрядное выровненное местоположение.

Обратите внимание, что я тщательно изложил свою формулировку. Предположим, у меня есть две такие группы вещания в одном запросе warp. Теперь также предположим, что эти группы обращаются к двум различным местоположениям, но в двух местах в одном банке. Например. группа A - адрес 0, а группа B - адрес 1024. В этом случае все потоки, участвующие в группе A, будут обслуживаться в одной транзакции, и все потоки, участвующие в группе B, будут обслуживаться в одной транзакции, но те две группы будут сериализованы относительно друг друга.

С другой стороны, если группой целенаправленного адресом 0 и группой В целенаправленном адресе 8, те, находятся в отдельных банках, так что все нити в группе А и всех нитей в группе B может обслуживаться в одна транзакция, поскольку разрешено несколько широковещательных рассылок на запрос.