Вызов функции ядра.

При запуске ядра используются специальные директивы для задания параметров ядра и передачи ему необходимых параметров. Объявление функции для ядра с параметром:

_global_ void Kernel_name(params);

Запуск ядра:

Kernel_name <<<grid, block, mem, stream>>> (params);

где dim3 grid – конфигурация сетки для запуска;

dim3 block – конфигурация блока при запуске;

size_t mem – количество разделяемой памяти на блок;

cudaStream_t stream – описание потока.

4. Методические рекомендации

Рассмотрим этапы создания CUDA-приложения на VisualStudio 2008 C++.

1. Создать новый проект и добавить к нему все используемые файлы CUDA (файлы с расширением - cu). Однако, поскольку среда «не знает», как их нужно компилировать, следует использо­вать специальный файл, задающий правила компиляции этих файлов (Custom Build Rules). Для этого необходимо скопировать с директории CUDA SDK все файлы с расширением - rules в директорию Visual Studio в папку, содержащую данные файлы.

2. Указать правила для компиляции CUDA файлов (рис. 3.9 и 3.10).

Рис. 3.9. Контекстное меню проекта

Рис. 3.10. Настройка правила компиляции

Далее необходимо подключить библиотеки к проекту. Для этого в свойстве проекта нужно открыть пункт Linker  General и в поле Additional Library Directories ввести «$(CUDA_LIB_PATH)» (рис. 3.11).

Рис. 3.11. Свойства проекта

4. В пункте Linker → Input в поле Additional Dependencies необходимо ввести используемые библиотеки CUDA (для работы с CUDA Runtime API будет cudart.lib) – рис. 3.12.

Рис. 3.12. Свойства проекта, пункт Linker → Input.

Пример программы:

#define BLOCK_SIZE 256 // Размер блока

#define N 1024 // Всего элементов для расчета

//Объявляем ядро

__global__ void kernel (int* data)

{

//Проводим вычисление абсолютной координаты в линейных блоке и сетке.

int idx = blockIdx.x * BLOCK_SIZE + threadIdx.x;

.}

//Основная программа CPU

int main ()

{

//Объявляем указатель на массив на ГА

int* data;

//Задаем конфигурацию ядра

dim3 block = dim3(BLOCK_SIZE,1);

dim3 grid = dim3(N / BLOCK_SIZE,1);

.

//Запускаем ядро с заданной конфигурацией и передаем ему параметры

kernel <<<grid, block>>> (data);

}

У функций ядра есть особенность – асинхронное исполнение. Это значит, что управление в основную программу может возвратиться до реального завершения выполнения расчетов на GPU. Для завершения работы заданной функции ядра необходимо использовать средства синхронизации, например cudaThreadSynchronize().

Рассмотрим пример задачи, очень хорошей для распараллеливания. Необходимо вычислить сумму двух векторов длиной N: А+ В = С.

(Задача реализована на видеокарте NVIDIA GeForce GTS 450)

Для начала необходимо разбить исходную последовательность из N элементов на блоки потоков. Потоком (нитью) в данной задаче является операция сложения каждого i-го элемента двух векторов А и В, результатом которой является i-й элемент вектора С. Блоки могут быть одно-, двух- и трёхмерными.

Количество потоков в блоке и количество блоков в сетке выбирается программистом исходя из загрузки ресурсов мультипроцессора и с учётом аппаратных ограничений.

При выборе стоит придерживаться следующих правил:

1. Число потоков в блоке должно быть не менее 32 для формирования полноценного warp и использования всех исполнительных элементов.

2. Желательно выбрать число потоков в блоке кратным 32, чтобы не было незаполненных варпов;

3. Размер блока выбирается таким, чтобы были задействованы все активные нити мультипроцессора;

4. Блок из максимального числа потоков неэффективен, так как обладает повышенными задержками;

5. Блоков должно быть не менее числа мультипроцессоров, чтобы SM не простаивал во время доступа к памяти.

Для оптимального выбора размера блока потоков и в целях достижения высокой производительности GPU можно воспользоваться CUDA Occupancy Calculator. Данный инструмент представляет собой таблицу MS Excel, входящую в состав CUDA SDK. Occupancy Calculator позволяет оценить отношение активных варпов к максимальному их количеству, поддерживаемому на мультипроцессоре GPU. Данное отношение показывает, насколько правильно выбран размер блока потоков исходя из количества регистров и объема разделяемой памяти мультипроцессора.

Максимальный размер блока на используемой видеокарте: 1024×1024×64 потока. Однако данное произведение не должно превышать 1024 – это максимальное количество потоков в блоке. Количество активных потоков на одном мультипроцессоре данной аппаратуры составляет 1536, а максимальное число одновременно запускаемых на нем блоков – 8. Таким образом, для использования всех нитей мультипроцессора размер блока должен быть неменьше 192 потока (192 = 1536/8). С другой стороны, блок из максимального количества потоков (1024) не эффективен, так как не все нити мультипроцессора будут задействованы. Исходя из этих данных и правил, описанных выше, оптимальный размер блока выбирается равным 192, 256 или 512 потокам.

Поскольку вектор одномерный, ограничимся использованием x-измерением нашего блока, т.е. задействуем только одну полосу нитей из блока (рис. 3.13).

Рис. 3.13. Разделение на блоки потоков

Ниже представлено описание и код программы:

#define BLOCK_SIZE 256

// Код вычислительного ядра для GPU: C = A + B.

__global__ void vectorAdd(const float* A, const float* B, float* C)

{

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

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

}

Функция vectorAdd является ядром, она и будет осуществлять сложение векторов.

Спецификатор __global__ показывает, что функция относится к ядру – ее вызов происходит на CPU, а выполняет GPU. Функция при помощи встроенных переменных вычисляет соответствующий данной нити глобальный индекс, для которого необходимо провести сложение соответствующих элементов и записать результат:

blockDim – число потоков в блоке;

blockIdx – номер блока внутри сетки;

threadIdx – номер потока внутри блока.

В нашем примере массивы одномерные, используется только координата x.

// Хостовый код для CPU.

void main()

{

const unsigned int N = 1024*1024; // Число элементов вектора.

dim3 blocks = dim3(N/ BLOCK_SIZE, 1); //Размер используемой сетки

dim3 threads = dim3(BLOCK_SIZE, 1); //Размер используемого блока

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

cudaSetDevice(0);

// Инициализируем таймер

unsigned int hTimer;

cutilCheckError(cutCreateTimer(&hTimer));

// Выделяем память для 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;

}

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

float *gpuA, *gpuB, *gpuC;

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

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

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

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

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

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

// Запуск таймера

cutilCheckError(cutResetTimer(hTimer));

cutilCheckError(cutStartTimer(hTimer));

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

{

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

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

cutilCheckMsg("kernel launch failure");

cutilSafeCall( cudaThreadSynchronize() );

}

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

cutilSafeCall( cudaMemcpy ((void *)cpuC, gpuC, N * sizeof(float), cudaMemcpyDeviceToHost) );

// Остановка таймера и вывод времени

cutilCheckError(cutStopTimer(hTimer));

printf("GPU time: %.3f (ms)\n", cutGetTimerValue(hTimer)/100);

// Проверяем результат.

int i;

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

if (fabs(cpuC[i] - (cpuA[i] + cpuB[i])) > 1e-7f) break;

printf("%s\n", (i == N) ? "PASSED" : "FAILED");

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

delete[] cpuA;

delete[] cpuB;

delete[] cpuC;

cudaFree(gpuA);

cudaFree(gpuB);

cudaFree(gpuC);

cutilCheckError(cutDeleteTimer(hTimer));

cudaThreadExit();

getchar();

}

Ниже представлена зависимость времени выполнения ядра от размера блока и длинны исходных векторов (табл. 6.3).

Таблица 3. 6.

Зависимость времени выполнения ядра от размера блока и исходных векторов

Длина вектора	Число нитей в блоке
	64	192	256	512	1024
	Время, мс
1024×64	0,037	0,026	0,025	0,025	0,032
1024×1024	0,394	0,256	0,256	0,256	0,375