Лекция: Архитектура устройства, виды памяти.
Модель программирования в CUDA предполагает группирование потоков. Потоки объединяются в блоки потоков (thread block) — одномерные или двумерные сетки потоков, взаимодействующих между собой при помощи разделяемой памяти и точек синхронизации.
Программа (ядро, kernel) исполняется над сеткой (grid) блоков потоков (thread blocks). Верхний уровень ядра GPU состоит из блоков, которые группируются в сетку или грид (grid) размерностью N1 * N2 * N3. Это можно изобразить следующим образом:
Рис. 1. Вычислительное устройство GPU. Рис. 2. Устройство блока GPU.
Любой блок в свою очередь состоит из нитей (threads), которые являются непосредственными исполнителями вычислений. Нити в блоке сформированы в виде трехмерного массива (рис. 2).
GPU может читать из константной памяти и из текстурной памяти. CPU может читать и записывать в константную, текстурную память и глобальную память. Каждый из процессоров нутри блока может читать и писать в shared memory, в регистры, в локальную память и в глобальную память.
Основные ограничения CUDA:
· отсутствие поддержки рекурсии для выполняемых функций;
· минимальная ширина блока в 32 потока;
· закрытая архитектура CUDA, принадлежащая NVIDIA
Регистровая память - самая быстрая из всех видов памяти, компилятор старается разместить все локальные переменные в регистрах, если регистров не хватает, то располагает в локальный памяти.
Локальная память — достаточно быстрая память, но ее трудно контролировать и разработчики CUDA советуют ее не использовать.
Глобальная память — самый большой объём памяти, доступный для всех мультипроцессоров на видеочипе, размер составляет от 256 мегабайт до 1.5 гигабайт на текущих решениях (и до 4 Гбайт на Tesla). Обладает высокой пропускной способностью, более 100 гигабайт/с для топовых решений NVIDIA, но очень большими задержками в несколько сот тактов. Не кэшируется, поддерживает обобщённые инструкции load и store, и обычные указатели на память. Дает возможность читать из любой ячейки и писать в любую ячейку. Разработчики CUDA рекомендуют сохранять в глобальной памяти результаты вычислений, перед тем как их отправить на хост. То, что будет храниться в глобальной памяти имеют спецификатор:__global__. В глобальной памяти можно выделять память динамически с помощью функции cudaMalloc().
Константная память - кэшируется, в константную память с хоста можно записывать с помощью функции cudaMemcpyToSymbol(). Если константа хранится в глобальной памяти, то для нее существует спецификатор: __constant__. С device в константную память нельзя ничего писать.
Разделяемая память — не кэшируется, но работает очень быстро. На один микропроцессор доступно только 16кб разделяемой памяти. При размещении данных указывается спецификатор: __shared__.
Текстурная память - кэшируется, она сохранилась благодаря первоначальному назначению видеоадаптера, физически текстурная память не отделена от глобальной. С device в текстурную память писать нельзя.