Глобальная память NVIDIA CUDA

Глобальная память

В документации CUDA в качестве одного из основных достижений технологии приводится возможность произвольной адресации  глобальной памяти. То есть можно читать из любой ячейки памяти, и писать можно тоже в произвольную ячейку (на GPU это обычно не так).

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

Глобальная память необходима в основном для сохранения результатов работы программы перед отправкой их на хост (в обычную память DRAM). Причина этого в том, что глобальная память — единственный вид памяти, куда можно что-то записывать.

Переменные, объявленные с квалификатором __global__, размещаются в глобальной памяти. Глобальную память также можно выделить динамически, вызвав функцию cudaMalloc(void* mem, int size) на хосте. Из устройства эту функцию вызывать нельзя. Отсюда следует, что распределением памяти должна заниматься программа-хост, работающая на CPU. Данные с хоста можно отправлять в устройство вызовом функции cudaMemcpy:

cudaMemcpy(void* gpu_mem, void* cpu_mem, int size, cudaMemcpyHostToDevice);

Точно таким же образом можно  проделать и обратную процедуру:

cudaMemcpy(void* cpu_mem, void* gpu_mem, int size, cudaMemcpyDeviceToHost);

Этот вызов тоже осуществляется с хоста.

При работе с глобальной памятью  важно соблюдать правило коалесинга (coalescing). Основная идея в том, что  треды должны обращаться к последоваетльным ячейкам памяти, причем 4,8 или 16 байтовым. При этом, самый первый тред должен обращаться по адресу, выровненному на границу соответственно 4,8 или 16 байт. Адреса, возвращаемые cudaMalloc выровнены как минимум по границе 256 байт.

 

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

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

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

 

Пример оптимизации при использовании  встроенных CUDA-типов. Размер типа int3 – 12 байт, доступ к памяти будет не оптимальным: Лучше использовать тип int4 (16 байтов), даже если четвертый компонент вам не нужен: 
В случае работы со структурами необходимо использовать ключевое слово __align__, которое позволяет выравнивать тип по заданному размеру. 
 
Пример выравнивания размера структуры. 
 
До выравнивания размер структуры vector3 составит 12 байт: На консоль выведется число 12. 
 
После выравнивания размер vector3 составит 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’а.  

Константная память

Константная память кэшируется, как  это видно на рис. 4. Кэш существует в единственном экземпляре для одного мультипроцессора, а значит, общий для всех задач внутри блока. На хосте в константную память можно что-то записать, вызвав функцию cudaMemcpyToSymbol. Из устройства константная память доступна только для чтения.

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

#define N 100

__constant__ int gpu_buffer[N];

void host_function()

{   

 int cpu_buffer[N];   

 ...   

 cudaMemcpyToSymbol(gpu_buffer, cpu_buffer, sizeof(int)*N);

}

...

// __global__ означает, что device_kernel - ядро, которое может быть запущено  на GPU

__global__ void device_kernel()

{   

 int a = gpu_buffer[0];   

 int b = gpu_buffer[1] + gpu_buffer[2];   

 // gpu_buffer[3] = a; ОШИБКА! константная память доступна только для чтения

}

 

 

! Константная память (constant memory) является достаточно быстрой из доступных GPU. Отличительной особенностью константной памяти является возможность записи данных с хоста, но при этом в пределах GPU возможно лишь чтение из этой памяти, что и обуславливает её название. Для размещения данных в константной памяти предусмотрен спецификатор __constant__. Если необходимо использовать массив в константной памяти, то его размер необходимо указать заранее, так как динамическое выделение в отличие от глобальной памяти в константной не поддерживается. Для записи с хоста в константную память используется функция cudaMemcpyToSymbol, и для копирования с device’а на хост cudaMemcpyFromSymbol, как видно этот подход несколько отличается от подхода при работе с глобальной памятью.

Так как для константной памяти используется кэш, доступ к ней в  общем случае довольно быстрый. Единственный, но очень большой недостаток константной  памяти заключается в том, что  ее размер составляет всего 64 Kбайт (на все устройство). Из этого следует, что в контекстной памяти имеет смысл хранить лишь небольшое количество часто используемых данных.

3.2.5. Текстурная память

Текстурная память кэшируется (рис. 4). Для каждого мультипроцессора имеется только один кэш, а значит, этот кэш общий для всех задач внутри блока.

Название текстурной памяти (и, к  сожалению, функциональность) унаследовано от понятий «текстура» и «текстурирование». Текстурирование — это процесс наложения текстуры (просто картинки) на полигон в процессе растеризации. Текстурная память оптимизирована под выборку 2D данных и имеет следующие возможности:

• быстрая выборка значений фиксированного размера (байт, слово, двойное или учетверенное слово) из одномерного или двухмерного массива;
• нормализованная адресация числами типа float в интервале [0,1).
• аппаратная линейная или билинейная (в случае 2D) интерполяция соседних значений в случае нормализованной адресации;
• аппаратная обработка выхода за границу массива с использованием двух режимов: clamp и wrap (не путать с пулами потоков warp!).

Текстурная память физически не отделена от глобальной памяти. Выделив  функцией cudaMalloc некий участок глобальной памяти, можно затем указать с  помощью функции cudaBindTexture, что данный участок теперь будет рассматриваться как текстурная память.

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

Но есть и плохая новость. Из текстурной памяти можно читать данные только встроенных в nvcc типов, имеющих размер 1, 2, 4, 8 или 16 байт, и только с помощью специальных функций — tex1D, tex2D или tex1Dfetch, tex2Dfetch. Иначе говоря, нельзя сделать указатель на текстурную память и разыменовать его произвольным образом (например, прочитав какую-либо структуру размером в 26 байт).

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

Теперь подробнее про возможности. Нормализованная адресация — это адресация числом с плавающей точкой от 0 до 1.

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

Интерполяцию можно использовать только в том случае, если текстура выбирается значениями float. При этом допускаются векторизованные типы и типы, отличные от типов с плавающей  точкой.

Например, можно создать текстуру с элементами uchar4, находящимися в интервале [0,255]. Затем можно их выбирать, используя нормализованную адресацию. Результирующим значением будетет слово типа float4, отображенное в интервал [0,1). Здесь используется стандартное математическое обозначение для интервала, где 0 включается, а 1 — нет.

Обработку выхода за границы массива  можно сделать в двух режимах wrap и clamp. При использовании режима clamp происходит «схлопывание» значения до ближайшей границы. В режиме wrap координаты заворачиваются таким образом, что если координата x выходит за границу массива на N, то будет взят N-ый элемент массива. Т.е. фактически всегда берется остаток от деления на размер массива. Следует отметить, что это свойство текстурной памяти весьма полезно для реализации хэш-таблиц.

#define N 1024

 

// намеренно выбран тип uint4, встроенный  в nvcc, так как он имеет максимально 

// возможный размер для текстурной  выборки. Это позволяет сократить  количество  
// выборок

uint4* gpu_memory;

 

texture<uint4, 1, cudaReadModeElementType> texture;

 

void host_function()

{   

 uint cpu_buffer[N];   

 ...   

 cudaMalloc((void**) &gpu_memory, N*sizeof(uint4)); // выделим память в GPU   

 // настройка параемтров текстуры texture    

 texture.addressMode[0] = cudaAddressModeWrap; // режим Wrap   

 texture.addressMode[1] = cudaAddressModeWrap;   

 texture.filterMode = cudaFilterModePoint; // ближайшее значение   

texture.normalized = false; // не использовать нормализованную адресацию    

 cudaBindTexture(0, texture, gpu_memory,N) // отныне эта память будет считаться текстурной   

 cudaMemcpy(gpu_memory,cpu_buffer,N*sizeof(uint4),cudaMemcpyHostToDevice); // копируем данные на GPU

}

...

// __global__ означает, что device_kernel - ядро, которое нужно распараллелить

__global__ void device_kernel()

{   

 uint4 a = tex1Dfetch(texture,0); // можно выбирать данные только таким способом !    

 uint4 b = tex1Dfetch(texture,1);