Понимание конфликтов банков разделяемой (shared) памяти в NVIDIA CUDA

    Разделяемая (shared) память является очень эффективным средством оптимизации за счет очень быстрого доступа (в 100 раз быстрее чем глобальная память). Однако, при неправильном использовании ее возможны конфликты банков, которые существенно замедляют быстродействие. В данной статье пойдет речь о том, как эти конфликты возникают, и как их избежать.

    Как возникают конфликты разделяемой памяти



    Конфликты возникают, когда 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. Данный счетчик показывает количество варпов, которым необходимо сериализовывать свой доступ при адресации константной или разделяемой памяти, другими словами, этот счетчик показывает конфликты банков.

    Заключение



    В заключение отмечу, что наиболее эффективная методика устранения конфликтов банков — это разработка схем доступа, которая минимизирует их возникновение и последующий анализ приложения профайлером (что никогда не лишнее).

    Похожие публикации

    AdBlock похитил этот баннер, но баннеры не зубы — отрастут

    Подробнее
    Реклама

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

      0
      Спасибо за статью.
      Сами мы с этими банками уже разобрались, но если бы ваша статья раньше появилась, то у нас бы меньше времени ушло на разруливание конфликтов.
        +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];
        добавление одного неиспользуемого столбца избавляет нас от конфликтов вообще в данном случае.
        (возможно я где-то напутал со строками/столбцами, но идея надеюсь понятна).
          –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)
            0
            Спасибо за замечания, однако они все отображены в тексте так и ли иначе.

            То что начинать надо с паттерна я написал в заключении. статья не про общую оптимизацию, а про конфликты банков.
            Про конфликты на Ферми с маленькими типами также описано. Цитата: «Из-за особенностей широковещательного доступа, 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
            Я немного недопонял пример:

            __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?
              0
              Всего 16 банков. Притом, 0-ой элемент массива попадает в 0-ой банк, 1->1,… 15->15, 16->0.
              Отсюда следует, что 0-ой поток читает 0-ой элемент из 0-го банка и 8-й поток читает 16-й элемент из 0-го банка.

            Только полноправные пользователи могут оставлять комментарии. Войдите, пожалуйста.

            Самое читаемое