Занятие 1. Программная модель CUDA - Лаборатория суперкомпьютерных и квантовых вычислений

В данном уроке будет изложено описание программной модели CUDA, основные термины и понятия, а также коротко будет описана концепция модели памяти GPU.

Итак, ниже приведены основные термины и понятия, которые будут необходимы для работы с CUDA:

Хост (Host) — центральный процессор, управляющий выполнением программы.

Устройство (Device) — видеоадаптер, выступающий в роли сопроцессора центрального процессора.

Грид (Grid) — объединение блоков, которые выполняются на одном устройстве.

Блок (Block) — объединение тредов, которое выполняется целиком на одном SM. Имеет свой уникальный идентификатор внутри грида.

Тред (Thread, поток) — единица выполнения программы. Имеет свой уникальный идентификатор внутри блока.

Варп (Warp) — 32 последовательно идущих треда, выполняется физически одновременно.

Ядро (Kernel) — параллельная часть алгоритма, выполняется на гриде.

Потоковые мультипроцессоры (Streaming Multiprocessors) — это система, состоящая из нескольких планировщиков варпов (Warps), 32-х процессорных ядра CUDA, использующие совместно память и кэш, то есть они видимые для всех процессоров.

Планировщик варпов (Warp Scheduler) — это система, создающая варпы из 32 потоков и распределяющая их по разным SM-процессорам. Планирование варпов осуществляется таким образом, чтобы максимально скрыть время ожидания данных из глобальной памяти: пока одни группы потоков ожидают данные, над другими исполняются инструкции.

Для реализации программы под GPU компания NVIDIA выпустила свои расширения для языка С, компилятор NVCC для сборки таких программ, ввела в обиход новое расширение *.cu для файлов, которые содержат CUDA вызовы.

Проще говоря, с программной точки зрения архитектура CUDA есть нечто иное, как расширение языков программирования С/С++, но с некоторыми аннотациями для их отличия от кода хоста, а также аннотации для различения типов памяти данных, которые существуют на графическом процессоре. Такие функции могут иметь параметры, и их можно вызвать с помощью синтаксиса, который очень похож на обычный вызов функции C, но немного расширен для возможности указать матрицу потоков GPU, которые должны выполнять вызываемую функцию. В течение своего жизненного цикла хост-процесс может отправлять много параллельных задач графического процессора.

Исходные файлы для приложений CUDA состоят из комбинации обычного хост-кода C ++ и функций устройства GPU. Трассировка CUDA-компиляции отделяет функции устройства от главного кода, компилирует функции устройства с использованием собственных компиляторов NVIDIA (nvcc.exe), компилирует код хоста с использованием компилятора хоста C ++.

Ключевые особенности программной стороны CUDA включают в себя:

дополнительные типы переменных.

Итак, разберем каждую ключевую особенность отдельно. Сперва нужно пояснить, что же такое спецификаторы в CUDA и какие они бывают.

Спецификаторы определяют вызов функции, а всего их 3:

Существуют дополнительные ограничения на выполняемые на видеокарте функции:

не поддерживается переменное число аргументов в функциях.

Спецификаторы переменных, которые служат для указания типа используемой памяти GPU, приведены в таблице 1.

Таблица 1.

Спецификаторы переменных в CUDA

Спецификатор	Находится в	Доступна для	Вид доступа
__device__	устройстве	устройства	только чтение
__constant__	устройстве	устройства и хоста	чтение/запись
__shared__	устройстве	блока потоков	чтение/запись

На их использование также накладывается ряд ограничений:

__device__ является аналогом const на центральном процессоре.

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

__constant__ память осуществляется только через специальные функции с CPU,

kernelName – это имя функции ядра спецификатора __global__.

Dg или gridSize – это размерность сетки для блоков dim3, выделенная для расчетов.

Db или blockSize – размер блока dim3, выделенного для расчетов.

S или cudaStream – переменная cudaStream_t, задающая поток, в котором будет произведен вызов. В дальнейшем будут приведены примеры, где это еще раз будет отмечено.

Рассмотрим следующую конструкцию:

__global__void sum_kernel (int *a, int *a, int *c, int N) {
int main ( void ) {
kernel <<<1,1>>>();
}
...
//код для GPU
}

Атрибут __global__, как уже было сказано, запускает программу с хоста, а исполняет ее на девайсе.

А void говорит нам о том, что возвращаемое значение отсутствует. Далее идет название самой функции name_kernel, которая обладает аргументами, определенными в круглых скобках (указатели на переменные). Дальше открывается основная (main) функция, в которой есть следующая строка:

Тройные угловые скобки указывают на то, что происходит вызов от хоста к девайсу, цифры внутри скобок обозначают количество блоков и количество тредов (потоков в блоке) соответственно. Например, мы хотим использовать 1 блок и указываем N потоков в этом блоке, тогда строчка имеет следующий вид:

Обращаться к отдельному блоку можно через индекс массива блоков blockIdx.x. Подробное описание переменной будет дано в следующей главе в листинге программ.

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

Продолжая тему встроенных переменных, следует отметить, что в CUDA добавлено несколько векторных типов для удобства копирования и доступа к данным. Типы (u)char, (u)int, (u)short, (u)long, float могут быть 1-, 2-, 3-, 4-мерными векторами, а longlong, double — только 1- и 2-мерными.

Для создания переменных таких типов требуется применять функции вида make_(тип)(размерность), например:

dim3, основанный на типе uint3, имеющий нормальный конструктор. Этот конструктор позволяет задавать не все компоненты вектора (недостающие компоненты инициализируются единицами). Данный тип используется для задания параметров запуска ядра. dim3 block = dim3 ( 64, 2 ).

К встроенным переменным-идентификаторам тредов и блоков относят:

dim3 gridDim – размерность грида. Содержит в себе информацию о конфигурации сетки.

uint3 blockIdx – индекс текущего блока в вычислении на GPU. Иными словами, позволяет узнать координаты текущего блока внутри сетки.

dim3 blockDim – размерность блока при запуске ядра. Позволяет узнать размерность блока, выделенного при текущем вызове ядра.

uint3 threadIdx – индекс текущего треда в вычислении на GPU, или другими словами, координаты текущего потока внутри блока.