4. Прикладные программные интерфейсы cuda

CUDA имеет два программных интерфейса: программный интерфейс драйвера (CUDA driver API) и интерфейс времени выполнения (C runtime for CUDA). Номера их версий позволяют разработчикам проверять наличие возможностей, связанных с этими API, и определять, требуется ли более новая версия по сравнению с установленной версией. Это важно, так как интерфейс драйвера обладает обратной совместимостью. Другими словами, приложения, скомпилированные под определенную версию интерфейса драйвера, будут работать на более новых версиях драйвера, но не обязательно на более старых. Это же касается и интерфейса времени выполнения CUDA, так как он реализован поверх интерфейса драйвера и взаимодействует с CUDA-устройствами только через него.

Указанные программные интерфейсы вместе образуют хостовый компонент времени выполнения программной платформы CUDA, т. е. их можно использовать только в функциях, выполняющихся на хосте. Библиотеки предоставляют следующую функциональность:

• управление устройствами;

• управление контекстом;

• управление памятью;

• управление модулями кода;

• управление выполнением;

• управление ссылками на текстуры;

• взаимодействие с графическими API: OpenGL или Direct3D. Эти два API взаимно исключают друг друга: программа может использовать только один из них.

Интерфейс времени выполнения CUDA, который используется чаще, упрощает код управления устройствами за счет предоставления неявной инициализации, управления контекстом и модулями кода. Хостовый код, генерируемый компилятором nvcc, основан на интерфейсе времени выполнения CUDA, поэтому приложения, ссылающиеся на этот код, должны использовать именно этот программный интерфейс.

Напротив, интерфейс драйвера требует больше кода, который труднее отлаживать, но дает возможность более гибкого управления. В частности, конфигурация и запуск CUDA-функции с использованием интерфейса драйвера более сложны, так как конфигурация CUDA-функции и передача параметров должны быть произведены при помощи прямых вызовов соответствующих функций вместо простого синтаксиса вида <<<…>>>. В отличие от кода, выполняемого на хосте, код CUDA-функции для CUDA-устройства одинаков для обоих API. Определить, какой из интерфейсов используется, можно так: интерфейс драйвера доступен через динамическую библиотеку nvcuda.dll, все его точки входа имеют префикс “cu”, а интерфейс времени выполнения – через библиотеку cudart.dll с префиксом точек входа “cuda”.

Интерфейс времени выполнения CUDA осуществляет контроль над загрузкой CUDA-функции, передачей ей параметров и конфигурацией ее выполнения. При этом API берет на себя инициализацию кода, управление контекстом и модулями, конфигурацию CUDA-функции и передачу параметров. Он включает две главные части:

• cuda_runtime_api.h – низкоуровневые функции с интерфейсом в стиле “C”;

• cuda_runtime.h – высокоуровневые функции с интерфейсом в стиле “C++”, которые основаны на низкоуровневых функциях.

Интерфейс драйвера CUDA более низкоуровневый и имеет по сравнению с интерфейсом времени выполнения следующие преимущества:

• отсутствие зависимости от библиотеки времени выполнения;

• больший контроль над устройствами (к примеру, только интерфейс драйвера позволяет с одного потока CPU управлять несколькими GPU);

• нет расширений языка C в хостовом коде, поэтому он может быть скомпилирован без участия компилятора nvcc.

Важно отметить, что для каждой функции интерфейса времени выполнения имеется эквивалентная функция интерфейса драйвера. Однако последний включает дополнительные функции, такие как функция для миграции контекста от одного потока хоста к другому. Дополнительными возможностями интерфейса драйвера, отсутствующими у интерфейса времени выполнения, являются:

• управление контекстом;

• поддержка текстур на базе 16-битных чисел с плавающей запятой;

• доступ к библиотеке обработки изображений MCL.

Иногда именно эти особенности заставляют разработчиков выбирать определенный API из двух. В остальных случаях рекомендуется использовать интерфейс времени выполнения, так как его проще использовать. Кроме того, так как все его функции имеют эквиваленты в интерфейсе драйвера, то позже при необходимости легко можно осуществить переход к нему.

Для иллюстрации различий между двумя имеющимися API, приведены листинги 1.1 и 1.2, реализующие сложение векторов. Первый состоит из 27 строк кода, а второй – из 50 и задает более низкоуровневые операции, чем при использовании интерфейса времени выполнения.

Листинг 1.1. Хостовый код для сложения двух векторов с использо-ванием интерфейса времени выполнения (C runtime for CUDA).

const unsigned int block_size = 512;

const unsigned int blocks = 3;

const unsigned int N = blocks * block_size;

// Выделяем память для 3-х векторов.

float *cpuA = new float[N];

float *cpuB = new float[N];

float *cpuC = new float[N];

// Инициализируем два вектора случайными числами.

srand(123);

for (int i = 0; i < N; i++)

{

cpuA[i] = rand() / (float)RAND_MAX;

cpuB[i] = rand() / (float)RAND_MAX;

}

// Создаем CUDA-устройство и контекст.

cudaSetDevice(0); // Выбираем первое устройство.

// Выделяем память на GPU.

float *gpuA, *gpuB, *gpuC;

cudaMalloc((void **)&gpuA, N * sizeof(float));

cudaMalloc((void **)&gpuB, N * sizeof(float));

cudaMalloc((void **)&gpuC, N * sizeof(float));

// Копируем два вектора из памяти CPU на GPU.

cudaMemcpy(gpuA, cpuA, N*sizeof(float), cudaMemcpyHostToDevice);

cudaMemcpy(gpuB, cpuB, N*sizeof(float), cudaMemcpyHostToDevice);

// Запускаем CUDA-функцию на GPU.

vectorAdd<<<blocks, block_size>>> (gpuA, gpuB, gpuC);

// Копируем результат из памяти GPU на CPU.

cudaMemcpy ((void *) cpuC, gpuC, N * sizeof(float),

cudaMemcpyDeviceToHost);

// Освобождаем выделенные ресурсы.

delete[] cpuA; delete[] cpuB; delete[] cpuC;

cudaFree(gpuA); cudaFree(gpuB); cudaFree(gpuC);

Листинг 1.2. Хостовый код для сложения двух векторов с использованием интерфейса драйвера (CUDA driver API).

const unsigned int block_size = 512;

const unsigned int blocks = 3;

const unsigned int N = blocks * block_size;

// Выделяем память для 3-х векторов.

float *cpuA = new float[N];

float *cpuB = new float[N];

float *cpuC = new float[N];

// Инициализируем два вектора случайными числами.

srand(123);

for (int i = 0; i < N; i++)

{

cpuA[i] = rand() / (float)RAND_MAX;

cpuB[i] = rand() / (float)RAND_MAX;

}

CUdevice hDevice;

CUcontext hContext;

CUmodule hModule;

CUfunction hFunction;

// Создаем CUDA-устройство и контекст.

cuInit(0);

cuDeviceGet(&hContext, 0); // Выбираем первое устройство.

cuCtxCreate(&hContext, 0, hDevice));

cuModuleLoad(&hModule, “vectorAdd.cubin”);

cuModuleGetFunction(&hFunction, hModule, "vectorAdd");

// Выделяем память на GPU.

CUdeviceptr gpuA, gpuB, gpuC;

cuMemAlloc(&gpuA, N * sizeof(float));

cuMemAlloc(&gpuB, N * sizeof(float));

cuMemAlloc(&gpuC, N * sizeof(float));

// Копируем два вектора из памяти CPU на GPU.

cuMemcpyHtoD(gpuA, cpuA, N * sizeof(float));

cuMemcpyHtoD(gpuB, cpuB, N * sizeof(float));

// Задаем параметры, передаваемые CUDA-функции, и

// конфигурацию выполнения CUDA-функции.

#define ALIGN_UP(offset, alignment) \

(offset) = ((offset) + (alignment) – 1) & ~((alignment) – 1)

int offset = 0;

void* ptr;

ptr = (void*)(size_t)gpuA;

ALIGN_UP(offset, __alignof(ptr));

cuParamSetv(cuFunction, offset, &ptr, sizeof(ptr));

offset += sizeof(ptr);

ptr = (void*)(size_t)gpuB;

ALIGN_UP(offset, __alignof(ptr));

cuParamSetv(cuFunction, offset, &ptr, sizeof(ptr));

offset += sizeof(ptr);

ptr = (void*)(size_t)gpuC;

ALIGN_UP(offset, __alignof(ptr));

cuParamSetv(cuFunction, offset, &ptr, sizeof(ptr));

offset += sizeof(ptr);

// Запускаем CUDA-функцию на GPU.

cuParamSetSize(cuFunction, offset);

cuFuncSetBlockShape(cuFunction, block_size, 1, 1);

cuLaunchGrid(cuFunction, blocks, 1);

// Копируем результат из памяти GPU на CPU.

cuMemcpyDtoH(cpuC, gpuC, N * sizeof(float));

// Освобождаем выделенные ресурсы.

delete[] cpuA; delete[] cpuB; delete[] cpuC;

cuMemFree(gpuA); cuMemFree(gpuB); cuMemFree(gpuC);

В интерфейсе времени выполнения не существует функции для явной инициализации. Она автоматически происходит при вызове любой функции этого интерфейса (точнее, за исключением функций из разделов Device Management и Version Management справочного руководства CUDA). Это необходимо помнить при замерах времени выполнения функций и при интерпретации кодов ошибок после первого обращения к интерфейсу времени выполнения.

После инициализации интерфейса времени выполнения в каком-то потоке на хосте использование любого ресурса (память, поток, событие и т.п.), выделенного посредством вызова функции этого интерфейса, допустимо только в контексте этого же потока. Следовательно, только вызовы из этого потока (копирование участков памяти, запуски CUDA-функций и др.) способны работать с этими ресурсами. Это происходит из-за того, что инициализация включает создание контекста CUDA (context) и привязку его к данному потоку на хосте, после чего данный контекст не может быть привязан ни к какому другому потоку.