Pull to refresh

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

Reading time 5 min
Views 22K
Основная тема этой части – оптимизация работы с глобальной памятью при программировании 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 и существуют, лучше использовать концепцию неизменяемых данных и сохранять результаты расчетов в новых объектах, которые и передавать на следующие этапы расчетов.
Tags:
Hubs:
+14
Comments 22
Comments Comments 22

Articles