CUDA: Работа с памятью. Часть II.

    Основная тема этой части – оптимизация работы с глобальной памятью при программировании GPU.

    У GPU есть ряд особенностей, игнорирование которых может стоить многократной потери производительности при использовании глобальной памяти. Но если учесть все тонкости, то можно получить действительно эффективные CUDA-программы.

    Приступаем.


    Что не так с глобальной памятью?


    Объем глобальной памяти самый большой из всех типов памяти, но в тоже время эта память – самая медлительная по техническим характеристикам: скорости считывания и записи.

    В предыдущей части я рассматривал пример транспонирования матрицы. Для повышения производительности использовался буфер разделяемой памяти, что позволило увеличить производительность почти в четыре раза. Но было достаточно страно видеть это увеличение при лишнем посреднике. Секрет же кроется в правильном обращении к глобальной памяти.

    Можно выделить два способа оптимизации в работе с глобальной памятью: выравнивание размеров используемых типов и использование объединенных запросов.

    Выравнивание размеров используемых типов


    Выравнивание типа данных позволяет скомпилировать запрос в глобальную память в одну команду GPU, в противном случае компилятор сгенерирует дополнительный код, что может значительно понизить производительность. Для оптимальной производительности тип данных должен иметь размерность 4, 8 или 16 байт.

    Если размер типа не соответствует 4, 8 или 16 байтам, то лучше использовать тип большей размерности или произвести выравнивание с помощью ключевого слова __align__(размер выравнивания).

    Пример оптимизации при использовании встроенных CUDA-типов.

    Размер типа int3 – 12 байт, доступ к памяти будет не оптимальным:

    __device__ int3 data[512];

    __global__ void initData()
    {
      int idx = threadIdx.x
      data[idx] = make_int3(idx, idx, idx);
    };

    * This source code was highlighted with Source Code Highlighter.


    Лучше использовать тип int4 (16 байтов), даже если четвертый компонент вам не нужен:

    __device__ int4 data[512];

    __global__ void initData()
    {
      int idx = threadIdx.x
      data[idx] = make_int4(idx, idx, idx, 0);
    };

    * This source code was highlighted with Source Code Highlighter.


    В случае работы со структурами необходимо использовать ключевое слово __align__, которое позволяет выравнивать тип по заданному размеру.

    Пример выравнивания размера структуры.

    До выравнивания размер структуры vector3 составит 12 байт:

    struct vector3
    {
      float x;
      float y;
      float z;
    };

    int main()
    {
      printf("%i\n", sizeof(vector3));
      return 0;
    };


    * This source code was highlighted with Source Code Highlighter.


    На консоль выведется число 12.

    После выравнивания размер vector3 составит 16 байт:

    struct __align__(16) vector3
    {
      float x;
      float y;
      float z;
    };

    int main()
    {
      printf("%i\n", sizeof(vector3));
      return 0;
    };

    * This source code was highlighted with Source Code Highlighter.


    На консоль выведется число 16.

    Использование объединеных запросов


    Куда больший прирост производительности можно получить при объединении большого количества запрос в глобальную память в один (иногда запросы назвают транзакциями). В документации nVidia это назвается coalescing global memory accesses. Но, перед тем, как перейти к непосредственному обсуждению того, что необходимо для объединения запросов в память, необходимо знать пару дополнительных вещей о работе GPU.

    Для контроля исполнения работы нитей GPU использует так называемый warp. С программной точки зрения warp представляет пул нитей. Именно в пределах этого warp’а происходит параллельная работа нитей, которые были запрошены при вызове ядра, именно в warp’е нити могут взаимодействовать между собой. Размер warp’а для всех GPU составляет 32, то есть параллельно в warp’е исполняются только 32 нити. Одновременно на GPU можно запустить несколько warp’ов, это количество определяется размерами доступной регистровой и разделяемой памяти. Другая интересная особенность, что для доступа к памяти используется half-warp, то есть в начале к памяти обращаются первые 16 нитей, а затем вторая половина из 16 нитей. Почему доступ происходи т именно так, я точно сказать не могу, могу лишь предположить, что это связано с первичными задачами GPU – обработкой графики.

    Теперь рассмотрим требования, необходимые для объединения запросов в глобальную память. Не забываем, что обращение к памяти происходит через half-warp.

    Условия необходимые для объединения при обращении в память зависят от версии Compute Capability, я привожу их для версии 1.0 и 1.1, больше подробностей можно узнать в документации от nVidia.
    • Нити должны обращаться либо к 32-битовым словам, давая при этом в результате один 64-байтовый блок (транзакцию), либо к 64-битовым словам, давая при этом один 128-байтовый блок (транзакцию)
    • Если используется обращение к 128-битовым словам, то в результате будет выполнено две транзакции, каждая из которых вернет по 128 байт информации
    • Нити должны обращаться к элементам памяти последовательно, каждой следующей нити должно соответствовать следующее слово в памяти (некоторые нити могут вообще не обращаться к соответствующим словам)
    • Все 16 слов должны быть в пределах блока памяти, к которому выполняется доступ

    Пара примечаний к условиям:
    • Под словами подразумевается любой тип данных, главное – соблюдение нужных размерностей
    • Размерность слов дана в битах, а размерность получаемых блоков данных в байта.



    Рис. 1. Запросы, дающие объединение при обращении к памяти

    На рис. 1 приведены примеры запросов к глобальной памяти, которые дают объединение в одну транзакцию. Слева выполнены все условия: каждый поток из half-warp’а обращается к соответствующему по порядку 32-битному слову, адрес начала памяти выровнен по размеру блока транзакции (16 нитей * 4 байт = 64 байта). Справа приведен пример, когда некоторые потоки из блока вообще не обращаются к соответствующим им словам в памяти.


    Рис. 2. Запросы, не дающие объединение при обращении к памяти

    На рис. 2 приведены примеры, которые не дают объединения при обращении к глобалной памяти. Слева не выполнены условие обращения нитей соответствующим словам в памяти. Справа не выполнено условие по выравниванию адреса памяти по размеру блока. В результате: вместо одной объединеной транзакции получаем по 16 отдельных, по одной на каждый поток half-warp’а.

    Структуры массивов или массивы структур?



    Пару слов стоит уделить вопросу по работе со структурами и как при этом добиться повышения производительности. Если есть необходимость использования массива структур, то лучше создать отдельные массивы компонентов структуры, что позволит уменьшить количество запросов в глобальную память за счет объединений.

    Рассмотрим пример.

    Неэффективная работа с глобальной памятью:

    struct __align__(16) vec3
    {
      float x;
      float y;
      float z;
    };

    __device__ vec3 data[SIZE];

    __global__ void initData()
    {
      int idx = blockDim.x * blockIdx.x + threadIdx.x;
      data[idx].x = idx;
      data[idx].y = idx * 2;
      data[idx].z = idx * 3;
    };

    * This source code was highlighted with Source Code Highlighter.


    Эффективнее использовать отдельные массивы:

    __device__ float x[SIZE];
    __device__ float y[SIZE];
    __device__ float z[SIZE];

    __global__ void initArr()
    {
      int idx = blockDim.x * blockIdx.x + threadIdx.x;
      x[idx] = idx;
      y[idx] = idx * 2;
      z[idx] = idx * 3;
    };

    * This source code was highlighted with Source Code Highlighter.


    В первом случае использования массива векторов для обращения к каждому полю структуры необходим отдельный запрос в память, во втором случае за счет объединения достаточно 3 запросов для каждого half-warp’а. В среднем, этот подход позволяет увеличить производительность в 2 раза.

    Заключение


    В заключение всего выше сказанного хочу дать самый важный совет при работе с памятью в CUDA:

    НИКОГДА НЕ ПЫТАЙТЕСЬ ИЗМЕНЯТЬ ЗНАЧЕНИЕ ОДНОЙ ЯЧЕЙКИ ПАМЯТИ НЕСКОЛЬКИМИ НИТЯМИ ОДНОВРЕМЕННО.

    Это самая частая ошибка в многопоточном программировании. На самом деле CUDA не гарантирует атомарного доступа для каждой нити к определенной области памяти, поэтому результаты могут получиться не совсем такими, как ожидается. Хотя атомарные операции в CUDA и существуют, лучше использовать концепцию неизменяемых данных и сохранять результаты расчетов в новых объектах, которые и передавать на следующие этапы расчетов.
    Поделиться публикацией

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

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

      +2
      Очень полезная информация.

      Оффтоп: мне вот тут подумалось:
      — Вычисления на GPU правильно применять для задач по параллельной обработке более-менее объёмных данных.
      — Точнее всего под это подходит СУБД.
      — А вот круто было бы иметь СУБД, которая работает на GPU, в следствие чего её производительность сильно выше других СУБД на обычных CPU. Такая СУБД имела бы оромный успех у highload проектов, ИМХО.
      — А если бы она ещё бы легко масштабировалась добавлением видеокарт и машин, было бы вообще супер!
        +2
        Идея хорошая, но есть ряд узких моментов, таких как копирование данных с хоста на девайс и обратно, это может стать критичным в СУБД. Использование GPU подразумевает большой объем вычислений, проводимых над данными, поэтому СУБД для GPU должна быть весьма специализированной.
          +1
          > копирование данных с хоста на девайс и обратно, это может стать критичным в СУБД
          Кстати, а во сколько раз копирование с хоста на девайс медленнее копирования из памяти в память внутри хоста?

          > поэтому СУБД для GPU должна быть весьма специализированной.
          Я тут подумал, что логичней сделать какой-нить движок таблик к мусклу/постресу, которые подразумевает хранение таблицы в памяти видеокарты. Тогда можно было бы засунуть туда пару таблиц даже не особо маленьких и выполнять выборку из них со сложными условиями и сортировками. И лишь изредка синхронизировать таблицы с хостом, да и то только если они меняются.

          Я вот сейчас прикинул, найдётся немного задач, которые лучше решались бы таким вот движком таблиц. Однако если бы я на этапе архитектуры имел в виду такой движок, то этих задач было бы на много больше.

          Но в целом действительно выглядит очень уж узкоспециализированно.
            0
            >Кстати, а во сколько раз копирование с хоста на девайс медленнее копирования из памяти в память внутри хоста?
            Страно, что сам этим вопросом не задавался. Протестирую — отпишусь.

            По-второму пункту: количество операций с таблицами должно быть весьма внушительно, чтобы компенсировать издержки синхронизации девайса и хоста. Ждем архитектуру без хоста, которую в Intel так боятся :).
              0
              > количество операций с таблицами должно быть весьма внушительно,
              > чтобы компенсировать издержки синхронизации девайса и хоста.
              Ну это в любом случае подходит для задач, когда нам не надо иметь полностью актуальную таблицу на хосте. Либо, если изменения таблицы достаточно дёшевы, их можно выполнять параллельно на хосте и девайсе. Надо только позаботиться, чтобы они выполнялись абсолютно одинаково.
                0
                $ ./bandwidthTest
                Running on…
                device 0:GeForce 9600 GT
                Quick Mode
                Host to Device Bandwidth for Pageable memory
                .
                Transfer Size (Bytes) Bandwidth(MB/s)
                33554432 2274.1

                Quick Mode
                Device to Host Bandwidth for Pageable memory
                .
                Transfer Size (Bytes) Bandwidth(MB/s)
                33554432 2103.7

                Quick Mode
                Device to Device Bandwidth
                .
                Transfer Size (Bytes) Bandwidth(MB/s)
                33554432 29141.2

                Как видим, разница почти в 15 раз
            0
            Всегда думал, что СУБД почти всегда упираются в IO, а не в CPU.
              0
              Не редко возникают задачи, где СУБД упирается в CPU как раз.
                0
                Как думаете, почему на серверах БД оперативки 16-32гб?)
                В CPU…
                  0
                  Чтобы данные таблиц с диска кешировать и буфер разобранных запросов держать. CPU почти не у дел…
                0
                Я как раз разрабатываю такую СУБД :) Не совсем СУБД конечно, моя разработка основана на sqlite.
                В рамках работы над своей диссертацией. Самому интересно узнать чем это все кончится, планирую чуть позже открыть исходники. Приятно слышать, что у кого то схожие мысли.
                  0
                  У самого есть пара проектов, которые sqlite используют. Очень интересно что у вас получится.
                    0
                    Я постараюсь написать об этом в своем блоге, как только получу более менее работающую бету. Сейчас все разрознено и стоят немного другие цели.
                    0
                    А это будет кроссплатформенно? И под какой лицензией собираетесь выкладывать, если собираетесь?
                      0
                      Думаю да. Лицензия — да такая же что и у sqlite наверное, т.е. никаких ограничений.
                        0
                        Шикарно. Хочу быть как минимум в курсе о Вашей разработке. Вы где-нибудь о ней пишете?
                          0
                          Нет, пока не пишу. Много времени отнимает основная работа. Но как только будет то, что можно потестировать, я напишу администратору этого раздела, так что, думаю, вы не пропустите.
                      0
                      Новости есть?
                    +1
                    Спасибо за статьи :)

                    Интересно, кто-нибудь из разработчиков физических движков (в частности ODE) уже начал примеряться к CUDA? Было бы очень полезно и интересно, насколько выросла бы производительность.

                    С одной стороны, задачи твердотельного моделирования как нельзя лучше подходят для SIMD (по сути сплошные матрицы). С другой — там много мест где принимаются решения на базе условий, так что хз, можно ли разложить это в поток.
                      0
                      У nVidia есть физический движок PhysX, которые базируется на CUDA.
                      По поводу использования CUDA в задачах моделирования — примеры можно посмотреть здесь
                        0
                        Я знаю про physX, но меня интересует именно ODE. Возможно конечно они сделают порт для физикса.
                          0
                          Если физический движок изначально писался под CPU, то сделать версию под GPU, все равно, что с нуля начать проект, разве что интерфейсную часть оставить неизменной для пользователей.

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

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