CUDA поддерживает два программных интерфейса, CUDA Runtime API и CUDA Driver API, эти интерфейсы взаимоисключающие, т.е. если вы используете runtime, то driver вы воспользоваться не сможете и наоборот. Driver API более низкоуровневый интерфейс, предлагающий написание большего количества кода, но более гибкое управление и написание внешних модулей (что-то типа шейдеров) на языке близком к ассемблеру (PTX), Runtime попроще. В функциях Driver API используется префикс cu, в Runtime API cuda.

Kernel могут работать только с памятью устройства, значит перед вызовом kernel надо научиться ее резервировать, освобождать и копировать. Освобождаемая память устройства может быть двух видов линейная и CUDA arrays (специально для работы с текстурами). Линейная использует 32 битное адресное пространство. Самые распространенные функции для работы с линейной памятью это cudaMalloc () – для резервирования, cudaFree() – для освобождения и cudaMemcpy() – для ее копирования между host и device(GPU не имеет непосредственного доступа к памяти CPU и наоборот). Остановимся подробней на каждой.

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

const char * cudaGetErrorString (cudaError_t error)

Так же можно получить код последней ошибки с помощью функции

cudaError_t cudaGetLastError (void)

cudaError_t cudaMalloc ( void ** devPtr, size_t size ) – выделяет size байт памяти (в пространстве GPU) и возвращает указатель на нее.

devPtr – указатель на выделенную память.

size – размер запрашиваемой память в байтах.

Если память не выделена возвращается код ошибки cudaErrorMemoryAllocation иначе cudaSuccess.

Пусть нам например надо выделить память на массив из N элементов типа float:

float* d_A;

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

cudaError_t cudaFree (void *devPtr) - освобождение выделенной памяти, например

cudaFree(d_A);

cudaError_t cudaMemcpy (void * dst, const void * src, size_t count, enum cudaMemcpyKind kind) копирует память из host в device и наоборот.

dst – куда копируем данные

src – откуда копируем

size – объем копируемой памяти в байтах

kind – перечисление, указывающие направление копирования, имеет значения

• cudaMemcpyHostToHost Host -> Host.
• cudaMemcpyHostToDevice Host -> Device.  
• cudaMemcpyDeviceToHost Device -> Host.  
• cudaMemcpyDeviceToDevice Device -> Device

Фу, вроде минимум знаний для написания простейшей программы готов, давайте напишем простейшую программу сложения N мерных векторов(так скажем без всяких понтов, типа проверок на ошибки, event’ов и использования шаровой памяти). Так, вот ее код.

#include <stdio.h>

#include <stdlib.h>

#include <cuda.h>

__global__ void VecAdd(float *A,float *B,float *C)

{

int i=blockDim.x*blockIdx.x+threadIdx.x;

C[i]=A[i]+B[i];

}

int main(int argc, char** argv)

{

int n=1024*1024;

int nsz=n*sizeof(float);

float *a=NULL;

float *b=NULL;

float *c=NULL;

float *da;

float *db;

float *dc;

a=(float *)malloc(nsz);

b=(float *)malloc(nsz);

c=(float *)malloc(nsz);

cudaMalloc((void**)&da,nsz);

cudaMalloc((void**)&db,nsz);

cudaMalloc((void**)&dc,nsz);

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

{

a[i]=i;

b[i]=(n-i);

c[i]=0.0f;

}

cudaMemcpy(da,a,nsz,cudaMemcpyHostToDevice);

cudaMemcpy(db,b,nsz,cudaMemcpyHostToDevice);

dim3 threads=dim3(512,1,1);

dim3 blocks=dim3(n/512,1,1);

VecAdd<<<blocks,threads>>>(da,db,dc);

cudaMemcpy(c,dc,nsz,cudaMemcpyDeviceToHost);

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

{

printf("%f+%f=%f\n",a[i],b[i],c[i]);

}

cudaFree(da);

cudaFree(db);

cudaFree(dc);

free(a);

free(b);

free(c);

system("pause");

return 0;

}

Начнем рассматривать функцию main. Как видно размер массива равен 1048576 . В переменную nsz помещаем размер массива в байтах (4194304). Затем выделяем память на host и device (в DRAM CPU и DRAM GPU), инициализируем массив в памяти host, переменными, копируем массив данных из памяти host в память device. Определяем размер и размерность grid (здесь 2048Х1Х1), затем размер и размерность block(здесь 512Х1Х1) и вызываем kernel. Копируем результат работы kernel в память host, выводим на терминал. Освобождаем память зарезервированную на host и device.

верстаем