3.2 Типы памяти

В CUDA для GPU существуют 6 видов памяти. Это регистры, локальная, глобальная, разделяемая, константная и текстурная память. Каждая из этих типов памяти имеет определенное назначение, которое обуславливается её техническими параметрами.






Тип памятиДоступУровень выделенияСкорость работы




регистры R/W per-thread высокая




local R/W per-thread низкая




global R/W per-grid низкая




shared R/W per-block высокая




constant R/O per-grid высокая




texture R/O per-grid высокая





Таблица 3.1: Типы памяти

При этом CPU имеет R/W доступ только к глобальной, константной и текстурной памяти (находящейся в DRAM GPU) и только через функции копирования памяти между CPU и GPU (предоставляемые CUDA API).

1.
Регистры. Компилятор размещает все локальные переменные функций в регистрах. Доступ к таким переменным осуществляется с максимальной скоростью. Для того чтобы рассчитать количество регистров, доступных одной нити GPU, необходимо разделить общее число регистров на произведение количества нитей в блоке и количества блоков в гриде. Все регистры GPU 32 разрядные. Компилятор старается максимизировать число одновременно работающих блоков. Для получения большей эффективности надо занимать меньше чем 32 регистра. В этом случае теоретически будет запущено 4 блока на одном мультипроцессоре. Однако здесь следует учитывать объем разделяемой памяти, занимаемой потоками, так как если один блок занимает все разделяемую память, два таких блока не могут выполняться на мультипроцессоре одновременно.
2.
Локальная память. По скоростным характеристикам локальная память работает намного медленнее, чем регистровая. В какой-либо функции она может быть использована компилятором при большем количестве локальных переменных. Ее рекомендуется использовать при необходимости. Физически локальная память является аналогом глобальной памяти.
3.
Глобальная память. Одним из достижений технологии является возможность произвольность адресации глобальной памяти, но одновременно с тем же является самым медленным типом памяти. Глобальная память не кэшируется. В основном она необходима для хранения результатов работы программы перед их отправкой в хост. В алгоритмах, где требуется высокая производительность, количество операций с глобальной памятью необходимо минимизировать. С помощью спецификатора __global__, можно выделить глобальные переменные, а так же динамически, с помощью функций из семейства cudMallocXXX. Из устройства эту функцию вызывать нельзя. Таким образом, распределением памяти должна заниматься программа-хост, работающая на CPU.
4.
Разделяемая память. Относиться к быстрому типу памяти. Разделяемую память рекомендуется использовать для минимизации обращение к глобальной памяти, а так же для хранения локальных переменных функций. Адресация разделяемой памяти между нитями потока одинакова в пределах одного блока, что может быть использовано для обмена данными между потоками в пределах одного блока. Для размещения данных в разделяемой памяти используется спецификатор __shared__. Следует отметить, что разделяемая память для одного блока. Поэтому если нужно использовать ее как управляемый кэш, нужно обращаться к разным элементам массива, например: float x = mem\_shared[threadIdx.x]; Где threadIdx.x — индекс x потока внутри блока.
5.
Константная память. Является достаточно быстрым типом памяти. Она кэшируется. Кэш существует в единственном экземпляре для одного мультипроцессора, следовательно, общий для всех задач внутри блока. В константной памяти можно размещать данные любого типа и читать их при помощи простого присваивания. Для размещения данных в константной памяти предусмотрен спецификатор __constant__. Динамическое выделение в отличие от глобальной памяти в константной не поддерживается, если необходимо использовать массив в константной памяти, то его размер необходимо указать заранее.
6.
Текстурная память. Предназначена главным образом для работы с текстурами. Она оптимизирована под выборку 2D данных и имеет следующие возможности:

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

Синхронизация

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

__syncthreads() позволяет устанавливать точку синхронизации как барьер, и до того момента как все нити блока не закончат выполнение всех предыдущих инструкций, ни один последующий код не будет выполняться. Ее не рекомендуется использовать внутри условий.

Также можно использовать event’ы для синхронизации на CPU.

Для отслеживания выполнения кода GPU в CUDA используются event’ы. Для работы с ними служат следующие функции:

cudaError_t cudaEventCreate ( cudaEvent_t * event );

Функция служит для создания event’а.

cudaError_t cudaEventRecord ( cudaEvent_t event, CUstream stream );

Функция cudaEventRecord используется для задания места, прохождение которого должен сигнализировать данный event. Если параметр stream не равен нулю, то отслеживается только завершение выполнения всех операций в данном потоке. Надо заметить, что этот запрос является асинхронным - он фактически только обозначает место в потоке команд, прохождение которого потом будет запрашиваться.

cudaError_t cudaEventQuery ( cudaEvent_t event );

Функция cudaEventQuery выполняет мгновенную проверку "прохождения"данного eventа - управление из нее сразу же возвращается. В случае, если все операции, предшествующие данному event’у были закончены, то возвращается cudaSuccess, иначе возвращается значение
cudaErrorNotReady.

cudaError_t cudaEventSynchronize ( cudaEvent_t event );

Явная синхронизация, т.е. ожидание пока все операции для данного event’а не будут завершены, обеспечивается этой командой.

cudaError_t cudaEventElapsedTime ( float * time, cudaEvent_t startEvent, cudaEvent_t stopEvent );

При помощи функции cudaEventElapsedTime можно узнать время в миллисекундах (с точностью до половины микросекунды), прошедшее между данными event’ами (между моментами, когда каждый из этих event’ов был "записан").

cudaError_t cudaEventDestroy ( cudaEvent_t event );

Функция выполняет для удаления event’а.