Особенности глобальной памяти.

Как уже было отмечено, глобальная память – это основная память на GPU. Она является основным местом для размещения и хранения большого объема данных для обработки ядрами. Глобальная является самой медленной из всех типов памяти (она плохо работает с обращениями по случайному адресу, не имеет кэша, т.е. каждый раз данные надо считывать заново).

Переменные, объявленные спецификатором _global_, размещаются в глобальной памяти. Память также можно выделить динамически, вызвав функцию cudaMalloc на хосте. Из устройства эту функцию вызывать нельзя, так как распределением памяти должна заниматься программа, работающая на CPU. Данные с хоста можно отправлять в устройство вызовом функции cudaMemcpy.

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

Выравнивание типа данных позволяет скомпилировать запрос в глобальную память в одну команду 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);

};

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

__device__ int4 data[512];

__global__ void initData()

{

int idx = threadIdx.x

data[idx] = make_int4(idx, idx, idx, 0);

};

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

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

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

struct vector3

{

float x;

float y;

float z;

};

int main()

{

printf("%i\n", sizeof(vector3));

return 0;

};

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

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

struct __align__(16) vector3

{

float x;

float y;

float z;

};

int main()

{

printf("%i\n", sizeof(vector3));

return 0;

};

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

Обращение к глобальной памяти осуществляется одновременно из половины варпа (16 потоков). Также существуют особые принципы доступа, позволяющие значительно ускорить работу с глобальной памятью. Графический процессор имеет возможность объединять запросы в блоки по 32, 64 и 128 байт (coalescing global memory accesses). Максимальное объединение – это когда все запросы одной полвины варпа удается объединить в один большой запрос (транзакцию) на чтение непрерывного блока памяти. Возможности объединения запросов различаются в зависимости от вычислительных возможностей конкретного GPU. Условия возникновения объединенного запроса к глобальной памяти в зависимости от версии аппаратной архитектуры графического адаптера приведены в табл. 3.5.

Таблица 3. 5.

Условия возникновения объединенного запроса

Сompute Сapability 1.0, 1.1	Сompute Сapability 1.2 и выше
Нити обращаются: к 32-битовым словам, создавая 64-байтовый блок к 64-битовым словам, создавая 128-байтовый блок. Все 16 слов лежат в пределах блока, k-ая нить обращается к k-му слову блока	Нити обращаются: к 8-битовым словам, создающим один 16-байтовый сегмент к 16-битовым словам, создающим один 32-байтовый сегмент к 32-битовым словам, создающим один 64-байтовый сегмент Получающийся сегмент выровнен по своему размеру