Очень важный вопрос как теоретический, так и с точки зрения практической части — это работа с памятью в CUDA. Ведь для написания программ, которые будут задействовать память GPU нужно понимать основные команды работы с памятью и знать ее иерархическую структуру.
Физически память GPU разделяется на DRAM (dynamic random access memory — динамическая память с произвольным доступом) и на мультипроцессорную, размещенную непосредственно на самом графическом процессоре.
В CUDA имеется всего шесть типов памяти:
- Регистры;
- Локальная память;
- Разделяемая память;
- Глобальная память;
- Текстурная память;
- Константная память.
Наиболее простым видом памяти является регистровая память. Каждый потоковый мультипроцессор содержит 8192 или 16384 32-битовых регистра.
На этапе компиляции они распределяются между тредами блока, тем самым влияя на количество исполняемых блоков одним мультипроцессором.
Каждый тред получает в свое монопольное пользование некоторое количество регистров, которые доступны как на чтение, так и на запись (read/write). Тред не имеет доступа к регистрам других тредов, но свои регистры доступны ей на протяжении выполнения данного ядра.
Поскольку регистры расположены непосредственно в потоковом мультипроцессоре, то они обладают максимальной скоростью доступа.
Если имеющихся регистров не хватает, то для размещения локальных данных (переменных) треда используется так называемая локальная память, размещенная в DRAM. Поэтому доступ к локальной памяти характеризуется очень высокой латентностью – от 400 до 600 тактов.
Локальная память — это область в глобальной памяти, выделенная компилятором для хранения локальных значений потоков. Локальная память используется для хранения локальных данных потоков в случае нехватки регистров или объявления локальных массивов внутри ядра без ключевого слова __shared__.
Следующим типом памяти является разделяемая shared память. Эта память расположена непосредственно в мультипроцессоре, но она выделяется на уровне блоков – каждый блок получает в свое распоряжение одно и то же количество разделяемой памяти. Всего каждый мультипроцессор содержит 16Кбайт разделяемой памяти, и от того, сколько разделяемой памяти требуется блоку, зависит количество блоков, которое может быть запущено на одном мультипроцессоре.
Разделяемая память обладает очень большой латентностью, доступна всем тредам блока, как на чтение, так и на запись.
Глобальная память – это обычная DRAM-память, которая выделяется при помощи специальных функций на СPU. Все треды сетки могут читать и писать в глобальную память. Поскольку глобальная память расположена на центральном процессоре, а не на графическом, то она обладает высокой латентностью (от 400 до 600 тактов).
Ниже в таблице описаны вышеуказанные типы памяти с их характеристиками.
Таблица 2.
| Тип памяти | Расположение | Кеширование | Доступ | Уровень доступа | Время жизни | Возможные операции |
|
Регистры |
Мультипроцессор |
нет |
R/w |
Per-thread |
Тред |
чтение\запись |
|
Локальная |
DRAM |
нет |
R/w |
Per-thread |
Тред |
|
|
Разделяемая |
Мультипроцессор |
нет |
R/w |
Per-block (все треды блока) |
Блок |
чтение\запись |
|
Глобальная |
DRAM |
нет |
R/w |
Все треды и CPU (Per-grid) |
Выделяется CPU |
чтение\запись |
|
Константная |
DRAM |
да |
R/o |
Все треды и CPU (Per-grid) |
Выделяется CPUВыделяется CPU |
чтение |
|
Текстурная |
DRAM |
да |
R/o |
Все треды и CPU (Per-grid) |
Выделяется CPU |
чтение |
Работа с глобальной памятью в CUDA
Как уже говорилось выше, глобальная — это основная память в CUDA, однако она является при этом и самой медленной (она медленно обращается по случайному адресу, отсутствует кэш, и соответственно, каждый раз данные надо считывать заново).
Процесс работы с глобальной памятью очень похож на обычную работу с памятью на центральном процессоре:
• инициализация (выделение);
• заполнение (копирование);
• использование;
• обратное копирование;
• освобождение.
Стоит упомянуть базовые функции CUDA для работы с глобальной памятью, так как они являются неотъемлемой частью программ.
cudaError_t cudaMalloc (void **devPtr, size_t size) – выделение памяти на видеокарте. Здесь devPtr – указатель на участок памяти на устройстве, size – размер выделяемой памяти в байтах.
Возвращает:
cudaSuccess – при удачном выделении памяти;
cudaErrorMemoryAllocation – при ошибке выделения памяти;
cudaError_t cudaFree (void * devPtr) – освобождение памяти на устройстве, здесь devPtr также является указателем на участок памяти;
cudaError_t cudaMemcpy (void * dst, const void * src, size_t countm enum cudaMemcpyKind kind) – функция копирования данных между хостом (Host) и устройством (Device). Здесь dst и src – указатели на источник и приемник соответственно, count – размер копируемых данных в байтах, kind – направление копирования. kind имеет всего 4 значения:
cudaMemcpyHostToDevice – копирование данных с хоста на устройство,
cudaMemcpyDeviceToHost – копирвоание данных с устройства на хост,
cudaMemcpyHostToHost – копирование данных с хоста на хост,
cudaMemcpyDeviceToDevice – копирование данных с устройства на устройства.
Функция cudaError_t cudaMemcpy возвращает:
cudaSuccess – при удачном копировании
1. cudaErrorInvalidValue – неверные параметры аргумента (например, размер копирования отрицателен)
2. cudaErrorInvalidDevicePointer – неверный указатель памяти в видеокарте
3. cudaErrorInvalidMemcpyDirection – неверное направление (например, перепутан источник и место—назначение копирования)
Работа с разделяемой памятью в CUDA
Разделяемая память находится на каждом SM и используется совместно всеми тредами внутри блока. Разделяемая память, как правило, требует использования команды явной синхронизации __syncthreads() после операций копирования и изменения. Для повышения пропускной способности разделяемая память разбита на 16 ячеек, доступ к которым идет независимо. Банки состоят из 32-битовых слов, причем подряд идущие слова расположены в идущих подряд блоках. Если возникает ситуация, что несколько тредов обращаются к одной и той же ячейке (не элементу, а именно ячейке, то их обращения обрабатываются последовательно (происходит конфликт ячеек), однако если все треды обращаются к одному и тому же элементу, то конфликта не происходит. Например,
__shared__ float a [N]; float x = a [base + threadIdx.x];
конфликта не происходит, так как размер float равен 32 битам и последовательно идущие элементы типа float попадают в последовательно идущие ячейки. threadIdx.x используется для доступа к отдельной нити.
__shared__ char a [N]; char x = a [base + threadIdx.x];
происходят конфликты четвертого порядка, так как размер элемента типа char равен 8 битам и 4 последовательно идущих элемента типа char попадают в один и тот же банк.
Стандартная схема работы с разделяемой памятью выглядит следующим образом:
• загрузить необходимые данные в разделяемую память(из глобальной),
• __syncthreads(),
• выполнить вычисления с загруженными данными,
• __syncthreads(),
• записать результат в глобальную память.
Константная память используется тогда, когда в ядро необходимо передать много различных данных, которые будут одинаково использоваться всеми тредами ядра. __constant__ float constData[256] — объявление глобальной переменной с именем constData для использования в качестве константной памяти. cudaMemcpyToSymbol (constData, hostData, sizeof (data), 0, cudaMemcpyHostToDevice) — копирование данных с центрального процессора в константную память. Использование константной памяти внутри ядра ничем не отличается от использования обычной глобальной переменной на хосте.
