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 и существуют, лучше использовать концепцию неизменяемых данных и сохранять результаты расчетов в новых объектах, которые и передавать на следующие этапы расчетов.
комментарии (20)