Технология Nvidia CUDA (Л/Р) 2009

Методические указания по выполнению лабораторной работы «Технология Nvidia CUDA» по курсу «Организация ЭВМ и систем»

Составители: доц. Андреев А.Е., Шаповалов О.В. Волгоградский государственный технический университет.

Содержание методички (без картинок):

Введение

CUDA (англ. Compute Unified Device Architecture) — технология GPGPU (англ. General-Purpose computing on Graphics Processing Units – вычисления общего назначения на графических процессорах), позволяющая программистам реализовывать на упрощённом языке программирования Си алгоритмы, выполнимые на графических процессорах видеокарт GeForce восьмого поколения и старше. Технология CUDA разработана компанией Nvidia. Фактически CUDA позволяет включать в текст Си программы специальные функции. Эти функции пишутся на упрощённом языке программирования Си и выполняются на графических процессорах Nvidia.

Чтобы понять, какие преимущества приносит перенос расчётов на видеокарты, приведём усреднённые цифры, полученные исследователями по всему миру. В среднем, при переносе вычислений на GPU, во многих задачах достигается ускорение в 5-30 раз, по сравнению с быстрыми универсальными процессорами.

Состав NVIDIA CUDA

CUDA включает два API: высокого уровня (CUDA Runtime API) и низкого (CUDA Driver API), хотя в одной программе одновременное использование обоих невозможно, нужно использовать или один или другой. Высокоуровневый работает «сверху» низкоуровневого, все вызовы runtime транслируются в простые инструкции, обрабатываемые низкоуровневым Driver API. Но даже «высокоуровневый» API предполагает знание основ об устройстве и работе видеокарт NVIDIA.

Есть и ещё один уровень, даже более высокий — две библиотеки:

CUBLAS — CUDA вариант BLAS (Basic Linear Algebra Subprograms), предназначенный для вычислений задач линейной алгебры и использующий прямой доступ к ресурсам GPU;

CUFFT — CUDA вариант библиотеки Fast Fourier Transform для расчёта быстрого преобразования Фурье, широко используемого при обработке сигналов. Поддерживаются следующие типы преобразований: complex-complex (C2C), real-complex (R2C) и complex-real (C2R).

Написание программ на CUDA

Для разработки собственных программ следует разбираться в базовых архитектурных особенностях видеочипов NVIDIA. Все инструкции в видекарте выполняются по принципу SIMD, когда одна инструкция применяется ко всем потокам в warp (в CUDA это группа из 32 потоков — минимальный объём данных, обрабатываемых мультипроцессорами). Этот способ выполнения назвали SIMT (single instruction multiple threads — одна инструкция и много потоков).

При исполнении программы, центральный процессор выполняет свои порции кода, а GPU выполняет CUDA код с наиболее тяжелыми параллельными вычислениями. Эта часть, предназначенная для GPU, называется ядром (kernel). В ядре определяются операции, которые будут исполнены над данными.

Таким образом, CUDA использует параллельную модель вычислений, когда каждый из SIMD процессоров выполняет ту же инструкцию над разными элементами данных параллельно. GPU является вычислительным устройством(device) для центрального процессора(host), обладающим собственной памятью и обрабатывающим параллельно большое количество потоков. Ядром (kernel) называется функция для GPU, исполняемая потоками. Видеочип отличается от CPU тем, что может обрабатывать одновременно сотни тысяч потоков, что обычно для графики, которая хорошо распараллеливается.

Модель программирования в CUDA предполагает группирование потоков. Потоки объединяются в блоки потоков (thread block) — одномерные или двумерные сетки потоков, взаимодействующих между собой при помощи разделяемой памяти и точек синхронизации. Программа (ядро, kernel) исполняется над сеткой (grid) блоков потоков (thread blocks), см. рисунок ниже. Одновременно исполняется одна сетка. Каждый блок может быть одно-, двух- или трехмерным по форме, и может состоять из 512 потоков на текущем аппаратном обеспечении.

Пример программы на CUDA

В качестве примера рассмотрим задачу умножения квадратной матрицы на вектор.

Исходные данные:
A[n][n] – матрица размерности n x n;
b[n] – вектор, состоящий из n элементов.

Результат:
c[n] – вектор из n элементов.
//Последовательный алгоритм умножения матрицы на вектор

for (i=0; i<n; i++)
{
        c[i]=0;
        for (j=0; j<m; j++)
        {
            c[i]=A[i][j]*b[j];
        }
}

Теперь рассмотрим решение этой задачи на видеокарте. Следующий код иллюстрирует пример вызова функции CUDA:

// инициализация CUDA
if(!InitCUDA())   { return 0; }

int Size = 1000;
// обычные массивы в оперативной памяти
float *h_a,*h_b,*h_c;
h_a = new float[Size*Size];
h_b = new float[Size];
h_c = new float[Size];

for (int i=0;i<Size;i++)     // инициализация массивов a и b
{
  for (int k=0;k<Size;k++)
    {
          h_a[i*Size+k]=1;
    }
  h_b[i]=2;
}

// указатели на массивы в видеопамяти
float *d_a,*d_b,*d_c;

// выделение видеопамяти
cudaMalloc((void **)&d_a, sizeof(float)*Size*Size);
cudaMalloc((void **)&d_b, sizeof(float)*Size);
cudaMalloc((void **)&d_c, sizeof(float)*Size);

// копирование из оперативной памяти в видеопамять
CUDA_SAFE_CALL(cudaMemcpy(d_a, h_a, sizeof(float)*Size*Size,
                              cudaMemcpyHostToDevice) );
CUDA_SAFE_CALL(cudaMemcpy(d_b, h_b, sizeof(float)*Size,
                              cudaMemcpyHostToDevice) );

// установка количества блоков
dim3 grid((Size+255)/256, 1, 1);
// установка количества потоков в блоке
dim3 threads(256, 1, 1);

// вызов функции
MatrVectMul<<< grid, threads >>> (d_c, d_a, d_b,Size);

// копирование из видеопамяти в оперативную память
CUDA_SAFE_CALL(cudaMemcpy(h_c, d_c, sizeof(float)*Size,
                              cudaMemcpyDeviceToHost) );

// освобождение памяти
CUDA_SAFE_CALL(cudaFree(d_a));
CUDA_SAFE_CALL(cudaFree(d_b));
CUDA_SAFE_CALL(cudaFree(d_c));



В принципе структуру любой программы с использованием CUDA можно представить аналогично рассмотренному выше примеру. Таким образом, можно предложить следующую последовательность действий:

1. инициализация CUDA
2. выделение видеопамяти для хранения данных программы
3. копирование необходимых для работы функции данных из оперативной памяти в видеопамять
4. вызов функции CUDA
5. копирование возвращаемых данных из видеопамяти в оперативную
6. освобождение видеопамяти

Пример функции, исполнимой на видеокарте

extern "C" __global__ void MatrVectMul(float *d_c, float *d_a,      float *d_b, int Size)
{

        int i = blockIdx.x*blockDim.x+threadIdx.x;
int k;
d_c[i]=0;
        for (k=0;k<Size;k++)
            {
            d_c[i]+=d_a[i*Size+k]*d_b[k];
          }
}

Здесь:

• threadIdx.x – идентификатор потока в блоке по координате x,
• blockIdx.x – идентификатор блока в гриде по координате x,
• blockDim.x – количество потоков в одном блоке.

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

Важно помнить, что функция, предназначенная для исполнения на видеокарте, не должна обращать к оперативной памяти. Подобное обращение приведет к ошибке. Если необходимо работать с каким-либо объектом в оперативной памяти, предварительно его надо скопировать в видеопамять, и обращаться из функции CUDA к этой копии.

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

Следует отметить следующее, что функция, работающая на видеокарте, должна выполняться не более 1 секунды. Иначе, функция будет не завершена, и программа завершится с ошибкой.

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

Замер времени в CUDA

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

unsigned int timer;

// создание таймера
cutCreateTimer(&timer);

// запуск таймера
cutStartTimer(timer);

… // код, время исполнения которого необходимо замерить

// остановка таймера
cutStopTimer( timer);

// получение времени
printf("time: %f (ms)\n", cutGetTimerValue(timer));

// удаление таймера
cutDeleteTimer( timer);

Создание проекта CUDA в VS 2005

После установки CUDA SDK, CUDA Toolkit и VS-интегратора в меню создания проекта появится новый пункт. Чтобы создать проект CUDA надо выбрать закладку Visual C++,  CUDA, CUDAWinApp. Далее можно выбрать тип проекта (по умолчанию создается консольное) и другие параметры. В итоге будет создан проект, который выводит на экран строку «Hello CUDA!». Чтобы приложение могло запуститься, в текущем каталоге должен присутствовать файл cutil32D.dll для работы в режиме Debug и файл cutil32.dll для работы в режиме Release (например, для проекта с названием lab4 необходимо положить эти файлы в директорию lab4/lab4). Причём, путь к файлу проекта не должен содержать русских букв.

Задания

Необходимо написать программу согласно варианту, при этом реализовать 2 функции: одну для выполнения на процессоре, вторую для выполнения на видеокарте. Затем сравнить результаты (возвращаемые значения) и скорость работы, задав большой размер матрицы (Size=1000). Массивы должны быть типа float.

Варианты заданий

1. Нахождение скалярного произведения векторов (протестировать для Size=1000000);
2. Нахождение суммы для каждой строки матрицы;
3. Вычисление среднего арифметического строк (столбцов) матрицы;
4. Умножение матрицы на множитель;
5. Сложение матриц;
6. Нахождение суммы квадратов элементов строк матрицы для всех строк матрицы;
7. Вычисление матричного выражения A=B+k*C, где A, B, C - матрицы, k - скалярный множитель;
8. Умножение матрицы на матрицу;
9. Вычитание матриц;
10. Транспонирование матрицы.

Контрольные вопросы

1. Что подразумевается под терминами kernel, host и device?
2. Объясните понятия grid, thread block, thread.
3. На каких видеокартах доступна технология CUDA?
4. Как вы считаете, будут ли перечисленные выше задания выполняться быстрее для Size=5 на видеокарте, чем на процессоре и почему?
5. Что общего в технологиях OpenMP и Nvidia CUDA и чем они различаются (с точки зрения программирования)