4.1 Первая программа

Нить — непосредственный исполнитель вычислений. Каким же тогда образом происходит распараллеливание вычислений между нитями? Рассмотрим работу отдельно взятого блока.

Задача. Требуется вычислить сумму двух векторов размерностью N элементов.

Нам известны максимальные размеры нашего блока: 512*512*64 нитей. Так как вектор у нас одномерный, то пока ограничимся использованием x-измерения нашего блока, то есть задействуем только одну полосу нитей из блока (рис. 4.1).

PIC

Рис. 4.1 Полоса нитей в блоке.

Заметим, что x-размерность блока 512, то есть, мы можем сложить за один раз векторы, длина которых N <= 512 элементов. В прочем, при более массивных вычислениях, можно использовать большее число блоков и многомерные массивы. Так же есть одна интересная особенность, возможно подумать, что в одном блоке можно задействовать 512*512*64 = 16777216 нитей, это не так, в целом, это произведение не может превышать 512.

В самой программе необходимо выполнить следующие этапы:

1.
Получить данные для расчетов.
2.
Скопировать эти данные в GPU память.
3.
Произвести вычисление в GPU через функцию ядра.
4.
Скопировать вычисленные данные из GPU памяти в CPU.
5.
Посмотреть результаты.
6.
Высвободить используемые ресурсы.

Переходим непосредственно к написанию кода:

Первым делом напишем функцию ядра, которая и будет осуществлять сложение векторов:

// Функция сложения двух векторов 
__global__ void addVector(float* left, float* right, float* result) 
{ 
  // Получаем id текущей нити. 
  int idx = threadIdx.x; 
  // Расчитываем результат. 
  result[idx] = left[idx] + right[idx]; 
}

Таким образом, распараллеливание будет выполнено автоматически при запуске ядра. В этой функции так же используется встроенная переменная threadIdx и её поле x, которая позволяет задать соответствие между расчетом элемента вектора и нитью в блоке. Делаем расчет каждого элемента вектора в отдельной нити.

Пишем код, которые отвечает за 1 и 2 пункт в программе:

#define SIZE 512 
__host__ int main() 
 
{ 
  //Выделяем память под вектора 
  float* vec1 = new float[SIZE]; 
  float* vec2 = new float[SIZE]; 
  float* vec3 = new float[SIZE]; 
  //Инициализируем значения векторов 
  for (int i = 0; i < SIZE; i++) 
  { 
    vec1[i] = i; 
    vec2[i] = i; 
  } 
  //Указатели на память видеокарте 
  float* devVec1; 
  float* devVec2; 
  float* devVec3; 
  //Выделяем память для векторов на видеокарте 
  cudaMalloc((void**)&devVec1, sizeof(float) * SIZE); 
  cudaMalloc((void**)&devVec2, sizeof(float) * SIZE); 
  cudaMalloc((void**)&devVec3, sizeof(float) * SIZE); 
  //Копируем данные в память видеокарты 
  cudaMemcpy(devVec1, vec1, sizeof(float) * SIZE, cudaMemcpyHostToDevice); 
  cudaMemcpy(devVec2, vec2, sizeof(float) * SIZE, cudaMemcpyHostToDevice); 
}

Для выделения памяти на видеокарте используется функция cudaMalloc, которая имеет следующий прототип:

cudaError_t cudaMalloc( void** devPtr, size_t count ), где

1.
devPtr — указатель, в который записывается адрес выделенной памяти,
2.
count — размер выделяемой памяти в байтах.

Возвращает:

1.
cudaSuccess — при удачном выделении памяти
2.
cudaErrorMemoryAllocation — при ошибке выделения памяти

Для копирования данных в память видеокарты используется cudaMemcpy, которая имеет следующий прототип:

cudaError_t cudaMemcpy(void* dst, const void* src ,size_t count, enum cudaMemcpyKind kind), где

1.
dst — указатель, содержащий адрес места-назначения копирования,
2.
src — указатель, содержащий адрес источника копирования,
3.
count — размер копируемого ресурса в байтах,
4.
cudaMemcpyKind — перечисление, указывающее направление копирования (может быть cudaMemcpyHostToDevice, cudaMemcpyDevice ToHost, cudaMemcpyHostToHost, cudaMemcpyDeviceToDevice).

Возвращает:

1.
cudaSuccess – при удачном копировании
2.
cudaErrorInvalidValue – неверные параметры аргумента (например, размер копирования отрицателен)
3.
cudaErrorInvalidDevicePointer – неверный указатель памяти в видеокарте
4.
cudaErrorInvalidMemcpyDirection – неверное направление (например, перепутан источник и место-назначение копирования)

Теперь переходим к непосредственному вызову ядра для вычисления на GPU.

... 
dim3 gridSize = dim3(1, 1, 1);    //Размер используемого грида 
dim3 blockSize = dim3(SIZE, 1, 1); //Размер используемого блока 
//Выполняем вызов функции ядра 
addVector<<<gridSize,blockSize>>>(devVec1, devVec2, devVec3); 
...

В этом случае определять размер грида и блока необязательно, так как используем всего один блок и одно измерение в блоке, поэтому код выше можно записать:

addVector<<<1, SIZE>>>(devVec1, devVec2, devVec3);

Остается скопировать результат расчета из видеопамяти в память хоста. Но у функций ядра при этом есть особенность – асинхронное исполнение, то есть, если после вызова ядра начал работать следующий участок кода, то это ещё не значит, что GPU выполнил расчеты. Для завершения работы заданной функции ядра необходимо использовать средства синхронизации, например event’ы. Поэтому, перед копированием результатов на хост выполняем синхронизацию нитей GPU через event.

Код после вызова ядра:

  //Выполняем вызов функции ядра 
  addVector<<<blocks, threads>>>(devVec1, devVec2, devVec3); 
  //Хендл event’а 
  cudaEvent_t syncEvent; 
  cudaEventCreate(&syncEvent);    //Создаем event 
  cudaEventRecord(syncEvent, 0);  //Записываем event 
  cudaEventSynchronize(syncEvent);  //Синхронизируем event 
  //Только теперь получаем результат расчета 
  cudaMemcpy(vec3, devVec3, sizeof(float) * SIZE, cudaMemcpyDeviceToHost);

Рассмотрим более подробно функции из Event Managment API.

Event создается с помощью функции cudaEventCreate, прототип которой имеет вид:

cudaError_t cudaEventCreate( cudaEvent_t* event ), где *event – указатель для записи хендла event’а.

Возвращает:

1.
cudaSuccess – в случае успеха
2.
cudaErrorInitializationError – ошибка инициализации
3.
cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
4.
cudaErrorInvalidValue – неверное значение
5.
cudaErrorMemoryAllocation – ошибка выделения памяти

Запись event’а осуществляется с помощью функции cudaEvent Record, прототип которой имеет вид:

cudaError_t cudaEventRecord( cudaEvent_t event, CUstream stream ), где

1.
event – хендл хаписываемого event’а,
2.
stream – номер потока, в котором записываем (в нашем случае это основной нулевой по-ток).

Возвращает:

1.
cudaSuccess — в случае успеха
2.
cudaErrorInvalidValue — неверное значение
3.
cudaErrorInitializationError — ошибка инициализации
4.
cudaErrorPriorLaunchFailure — ошибка при предыдущем асинхронном запуске функции
5.
cudaErrorInvalidResourceHandle — неверный хендл event’а

Синхронизация event’а выполняется функцией cudaEvent Synchronize. Данная функция ожидает окончание работы всех нитей GPU и прохождение заданного event’а и только потом отдает управление вызывающей программе. Прототип функции имеет вид:
cudaError_t cudaEventSynchronize(cudaEvent_t event), где event – хендл event’а, прохождение которого ожидается.

Возвращает:

1.
cudaSuccess — в случае успеха
2.
cudaErrorInitializationError — ошибка инициализации
3.
cudaErrorPriorLaunchFailure — ошибка при предыдущем асинхронном запуске функции
4.
cudaErrorInvalidValue — неверное значение
5.
cudaErrorInvalidResourceHandle — неверный хендл event’а
  //Результаты расчета 
  for (int i = 0; i < SIZE; i++) 
  { 
    printf("Element #%i: %.1f\n", i, vec3[i]); 
  } 
  // Высвобождаем ресурсы 
  cudaEventDestroy(syncEvent); 
  cudaFree(devVec1); 
  cudaFree(devVec2); 
  cudaFree(devVec3); 
  delete[] vec1; vec1 = 0; 
  delete[] vec2; vec2 = 0; 
  delete[] vec3; vec3 = 0;

Для освождения памяти видеокарты используется функция cudaFree, которая имеет прототип:

_t cudaFree(void *devPtr)

Необходимым моментом работы с CUDA, на которое следует обратить внимание, является то, что многие функции API – асинхронные, то есть управление возвращается еще до реального завершения требуемой операции. К числу асинхронных операций относятся: