2. Порядок выполнения работы

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

2. Написать второй вариант библиотеки при выбранном оптимальном размере блока потоков с использованием разделяемой памяти.

3. Засечь время выполнения алгоритмов.

4. Сравнить время выполнения двух алгоритмов (с разделяемой памятью и без нее).

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

Рассмотрим использование разделяемой памяти на примере перемножения двух квадратных матриц A и B размера N*N. В качестве размера блока возьмем16×16 и будем считать, что размер матриц N кратен 16. Необходимо вычислить произведение данных матриц, результатом которого является матрица С. Элемент С_i,j произведения двух матриц A и B вычисляется при помощи следующей формулы:

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

В этом примере будем использовать двухмерные блоки 16x16 и двухмерную сетку.

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

// Функция ядра

#define BLOCK_SIZE 16 // размерность блока 16x16

#define N 1024 // размер матрицы N*N

__global__ void matMult ( float * a, float * b, int n, float * c )

{

int bx = blockIdx.x; // индекс блока

int by = blockIdx.y;

int tx = threadIdx.x; // индекс нити

int ty = threadIdx.y;

float sum = 0.0f; // здесь накапливается результат

int ia = n * BLOCK_SIZE * by + n * ty; // смещение для a[i][0]

int ib = BLOCK_SIZE * bx + tx; // смещение для b[0][j]

// Перемноженние и суммирование в цикле;

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

sum += a [ia + k] * b [ib + k*n];

// смещение для записываемого элемента

int ic = n * BLOCK_SIZE * by + BLOCK_SIZE * bx;

// сохраняем результат в глобальную память;

c [ic + n * ty + tx] = sum;

}

void main ()

{

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

unsigned int hTimer;

cutilCheckError(cutCreateTimer(&hTimer));

// размер матриц в байтах

int numBytes = N * N * sizeof ( float );

// выделяем память для матриц на CPU

float * a = new float [N*N];

float * b = new float [N*N];

float * c = new float [N*N];

// заполняем исходные матрицы

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

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

{

int k = N*i + j;

a [k] = 0.0f;

b [k] = 1.0f;

}

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

cudaSetDevice(0);

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

float * adev = NULL;

float * bdev = NULL;

float * cdev = NULL;

cudaMalloc ( (void**)&adev, numBytes );

cudaMalloc ( (void**)&bdev, numBytes );

cudaMalloc ( (void**)&cdev, numBytes );

// устанавливаем размер блока и сетки

dim3 threads ( BLOCK_SIZE, BLOCK_SIZE );

dim3 blocks ( N / threads.x, N / threads.y);

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

cutilCheckError(cutResetTimer(hTimer));

cutilCheckError(cutStartTimer(hTimer));

// Копируем матрицы на GPU

cudaMemcpy ( adev, a, numBytes, cudaMemcpyHostToDevice );

cudaMemcpy ( bdev, b, numBytes, cudaMemcpyHostToDevice );

// запускаем ядро

matMult<<<blocks, threads>>> ( adev, bdev, N, cdev );

// Копируем матрицу обратно на CPU

cudaMemcpy ( c, cdev, numBytes, cudaMemcpyDeviceToHost );

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

cutilCheckError(cutStopTimer(hTimer));

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

// Очищаем ресурсы

cudaFree ( adev );

cudaFree ( bdev );

cudaFree ( cdev );

delete a;

delete b;

delete c;

}

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

Теперь рассмотрим другой подход решения этой задачи. Можно заметно повысить быстродействие программы за счет использования разделяемой памяти. Для этого разобьем результирующую матрицу на подматрицы 16×16, вычислением каждой такой подматрицы будет заниматься один блок (рис.3.14).

Рис. 3.14. Вычисления элементов C’

Стоит обратить внимание, что для вычисления такой подматрицы нужны только небольшие "полосы" матриц A и B. Также легко заметить, что результирующая матрица является суммой попарных произведений подматриц из этих двух полос:

За счет этого можно выполнить вычисление подматрицы всего за N/16 шагов. На каждом таком шаге в разделяемую память загружается по одной 16x16 подматрице A и одной подматрице B. Поскольку адресация разделяемой памяти в пределах одного блока одинакова для всех потоков, то на каждом шаге одна нить загружает по одному элементу из каждой подматрицы и вычисляет соответствующую им сумму членов. По окончании всех вычислений производится запись элемента в итоговую матрицу.

Ниже представлена функция ядра для данного случая.

__global__ void matMult ( float * a, float * b, int n, float * c )

{

int bx = blockIdx.x; // индекс блока

int by = blockIdx.y;

int tx = threadIdx.x; // индекс нити

int ty = threadIdx.y;

// Индекс начала первой подматрицы А

int aBegin = n * BLOCK_SIZE * by;

int aEnd = aBegin + n - 1;

// Шаг перебора подматрицы А

int aStep = BLOCK_SIZE;

// Индекс первой подматрицы В

int bBegin = BLOCK_SIZE * bx;

// Шаг перебора подматрицы В

int bStep = BLOCK_SIZE * n;

float sum = 0.0f; // Вычисляемый элемент

for (int ia = aBegin, ib = bBegin; ia <= aEnd; ia += aStep, ib += bStep)

{

// Разделяемая память для подматрицы А

__shared__ float as [BLOCK_SIZE][BLOCK_SIZE];

// Разделяемая память для подматрицы B

__shared__ float bs [BLOCK_SIZE][BLOCK_SIZE];

// Загружаем по одному элементу из матрицы А и В в разделяемую память

as [ty][tx] = a [ia + n * ty + tx];

bs [ty][tx] = b [ib + n * ty + tx];

__syncthreads(); // дожидаемся загрузки двух подматриц

// Вычисляем нужный элемент произведения загруженных подматриц;

for ( int k = 0; k < BLOCK_SIZE; k++ )

sum += as [ty][k] * bs [k][tx];

// Дожидаемся, пока все нити блока закончат

// вычислять свои элементы

__syncthreads();

}

// Записываем результат вычисления;

int ic = n * BLOCK_SIZE * by + BLOCK_SIZE * bx;

c [ic + n * ty + tx] = sum;

}

После копирования элементов подматриц А и B вызывается функция __syncthreads. Эта функция синхронизирует потоки в пределах блока, т.е. дожидается загрузки всех элементов подматриц. Её отличие от других способов синхронизации заключается в том, что она выполняется только на GPU. Точно так же синхронизация выполняется после вычисления произведения загруженных подматриц.

Теперь для вычисления одного элемента произведения матриц нужно всего 2*N/16 чтений из глобальной памяти. В табл. 3.7 приведено сравнение быстродействия двух подходов решения данной задачи (при N = 1024). По результатам видно, что за счет использования разделяемой памяти удается поднять быстродействие более чем в 2,5 раза. (Задача реализована на видеокарте NVIDIA GeForce GTS 450)

Таблица 3. 7.

Временные затраты на перемножение матриц различными методами

Способ перемножения матриц	Затраченное время (мс)
С использованием глобальной памяти	82.97
С использованием глобальной и разделяемой памяти	30.60

4. Защита лабораторной работы

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

9. Лабораторная работа № 4. Оптимизация приложений с глобальной памятью в технологии CUDA

   1. Цель работы

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

2. Порядок выполнения работы

1. Написать библиотеку методов сегментации изображения согласно индивидуальному варианту из практической работы №5. Вычисления производить на графическом процессоре при выбранном оптимальном размере блока потоков с использованием глобальной памяти.

2. Написать ту же библиотеку методов сегментации изображения, вычисления производить на центральном процессоре.

3. Засечь время выполнения алгоритмов.

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

5. Сделать выводы.