В данном уроке будет изложено описание программной модели 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 включают в себя:
-
спецификаторы функций, которые показывают, как и откуда будут выполняться функции;
-
спецификаторы переменных, которые служат для указания типа используемой памяти GPU;
-
спецификаторы запуска ядра ГПУ;
-
встроенные переменные для идентификации тредов и блоков;
-
дополнительные типы переменных.
Итак, разберем каждую ключевую особенность отдельно. Сперва нужно пояснить, что же такое спецификаторы в CUDA и какие они бывают.
Спецификаторы определяют вызов функции, а всего их 3:
__global__ — исполнение программы происходит на GPU, а вызов производится с CPU;
__device__ — исполнение и вызов с GPU;
__host__ — исполнение и вызов с CPU.
Существуют дополнительные ограничения на выполняемые на видеокарте функции:
-
нельзя брать адрес от функции (за исключением __global__);
-
нет поддержки рекурсии;
-
внутри функций нельзя применять спецификатор static к переменным;
-
не поддерживается переменное число аргументов в функциях.
Спецификаторы переменных, которые служат для указания типа используемой памяти GPU, приведены в таблице 1.
Таблица 1.
Спецификаторы переменных в CUDA
| Спецификатор | Находится в | Доступна для | Вид доступа |
| __device__ | устройстве | устройства | только чтение |
| __constant__ | устройстве | устройства и хоста | чтение/запись |
| __shared__ | устройстве | блока потоков | чтение/запись |
На их использование также накладывается ряд ограничений:
— спецификатор __device__ является аналогом const на центральном процессоре.
— спецификатор __shared__ применяется для задания разделяемой памяти и не может быть инициализирован при объявлении. Он требует явной синхронизации потоков внутри одного блока после любой операции с такой памятью.
— запись в __constant__ память осуществляется только через специальные функции с CPU,
— никакие спецификаторы нельзя применять к полям структур.
Директивой для вызова ядра является конструкция kernelName<<Dg, Db, Ns, S>>.
kernelName – это имя функции ядра спецификатора __global__.
Dg или gridSize – это размерность сетки для блоков dim3, выделенная для расчетов.
Db или blockSize – размер блока dim3, выделенного для расчетов.
Ns или sharedMemSize — размер дополнительной памяти, выделяемой при запуске ядра.
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) функция, в которой есть следующая строка:
kernel <<<1,1>>>();
Тройные угловые скобки указывают на то, что происходит вызов от хоста к девайсу, цифры внутри скобок обозначают количество блоков и количество тредов (потоков в блоке) соответственно. Например, мы хотим использовать 1 блок и указываем N потоков в этом блоке, тогда строчка имеет следующий вид:
kernel <<<1,N>>>();
Обращаться к отдельному блоку можно через индекс массива блоков blockIdx.x. Подробное описание переменной будет дано в следующей главе в листинге программ.
Точно также можно менять число тредов, а блок будет один, либо использовать конструкцию, где задействовано и несколько блоков и несколько тредов одновременно.
Продолжая тему встроенных переменных, следует отметить, что в CUDA добавлено несколько векторных типов для удобства копирования и доступа к данным. Типы (u)char, (u)int, (u)short, (u)long, float могут быть 1-, 2-, 3-, 4-мерными векторами, а longlong, double — только 1- и 2-мерными.
Для создания переменных таких типов требуется применять функции вида make_(тип)(размерность), например:
int2 a = make_int2 ( 1, 2 ); float4 b = make_float4 ( 1, 2, 3, 4 );
Для всех типов в CUDA Toolkit версии > 3.0 определены покомпонентные операции. Также существует специальный тип dim3, основанный на типе uint3, имеющий нормальный конструктор. Этот конструктор позволяет задавать не все компоненты вектора (недостающие компоненты инициализируются единицами). Данный тип используется для задания параметров запуска ядра. dim3 block = dim3 ( 64, 2 ).
К встроенным переменным-идентификаторам тредов и блоков относят:
dim3 gridDim – размерность грида. Содержит в себе информацию о конфигурации сетки.
uint3 blockIdx – индекс текущего блока в вычислении на GPU. Иными словами, позволяет узнать координаты текущего блока внутри сетки.
dim3 blockDim – размерность блока при запуске ядра. Позволяет узнать размерность блока, выделенного при текущем вызове ядра.
uint3 threadIdx – индекс текущего треда в вычислении на GPU, или другими словами, координаты текущего потока внутри блока.
