МОСКОВСКИЙ АВИАЦИОННЫЙ ИНСТИТУТ

(НАЦИОНАЛЬНЫЙ ИССЛЕДОВАТЕЛЬСКИЙ УНИВЕРСИТЕТ)

Кафедра вычислительной математики и программирования

спецкурс «Параллельное программирование»

ОТЧЕТ

Лабораторная работа № 4

«Работа с текстурной памятью»

Выполнила:

Группа: 8Д-301БК, вариант 8

Преподаватель:

Москва, 2015

Оглавление

1. Постановка задачи 3

2. Описание решения 3

3. Основные моменты кода 3

4. Результат работы программы 4

5. Сравнение скорости выполнения на CPU и GPU 5

6. Выводы 5

Приложения 6

1. Постановка задачи

Определение границы.

2. Описание решения

Для определения границ объектов на изображении применяется оператор Собеля. Для простоты изображение, с которым работает программа, находится в формате PGM.

3. Основные моменты кода

__device__ unsigned char

ComputeSobel(unsigned char ul, // upper left

unsigned char um, // upper middle

unsigned char ur, // upper right

unsigned char ml, // middle left

unsigned char mm, // middle (unused)

unsigned char mr, // middle right

unsigned char ll, // lower left

unsigned char lm, // lower middle

unsigned char lr, // lower right

float fScale)

{

short Horz = ur + 2*mr + lr - ul - 2*ml - ll;

short Vert = ul + 2*um + ur - ll - 2*lm - lr;

НЕ нашли? Не то? Что вы ищете?

short Sum = (short)(fScale*(abs((int)Horz)+abs((int)Vert)));

if (Sum < 0)

{

return 0;

}

else if (Sum > 255)

{

return 255;

}

return (unsigned char) Sum;

}

__global__ void SobelTex(int w, int h, float fScale, unsigned char *dev)

{

if(threadIdx. x > 0 && threadIdx. x < w-1 && blockIdx. x > 0 && blockIdx. x < h-1)

{

unsigned char pix00 = tex2D(tex, (float) threadIdx. x-1, (float) blockIdx. x-1);

unsigned char pix01 = tex2D(tex, (float) threadIdx. x+0, (float) blockIdx. x-1);

unsigned char pix02 = tex2D(tex, (float) threadIdx. x+1, (float) blockIdx. x-1);

unsigned char pix10 = tex2D(tex, (float) threadIdx. x-1, (float) blockIdx. x+0);

unsigned char pix11 = tex2D(tex, (float) threadIdx. x+0, (float) blockIdx. x+0);

unsigned char pix12 = tex2D(tex, (float) threadIdx. x+1, (float) blockIdx. x+0);

unsigned char pix20 = tex2D(tex, (float) threadIdx. x-1, (float) blockIdx. x+1);

unsigned char pix21 = tex2D(tex, (float) threadIdx. x+0, (float) blockIdx. x+1);

unsigned char pix22 = tex2D(tex, (float) threadIdx. x+1, (float) blockIdx. x+1);

dev[(threadIdx. x-1)+(blockIdx. x-1)*(blockDim. x-2)] = ComputeSobel(pix00, pix01, pix02, pix10, pix11, pix12, pix20, pix21, pix22, fScale);

}

4. Результат работы программы

Рис. 1. Окно вывода консольного приложения

Рис. 2. Исходное изображение

Рис. 3. Результат расчета на GPU

Рис. 4. Результат расчета на CPU

5. Сравнение скорости выполнения на CPU и GPU

При достаточно большом размере изображения чистое время обработки с использованием GPU значительно меньше, чем использованием CPU (в 32 раза при размере исходного изображения 1640х1024). С изменением размера изображения время обработки на GPU меняется не так сильно (с 2.85 мс до 2.7 мс при уменьшении размера изображения до 1024х1024), как на CPU (c 85 мс до 51 мс при таком же изменении размера изображения).

6. Выводы

Несмотря на неизбежные накладные расходы, связанные с необходимостью разбиения задачи и пересылки данных, выигрыш в производительности при достаточно больших размерах данных довольно существенен при использовании в программе NVIDIA CUDA в сравнении с традиционным алгоритмом решения на CPU.

Приложения

Код программы

Файл main. cpp

#include <cuda_runtime. h>

#include <stdlib. h>

#include <stdio. h>

#include <string. h>

#include<iostream>

#include<cmath>

#include "SobelGPU. h"

#include <helper_functions. h>

#include <helper_cuda. h>

#include <helper_timer. h>

static int imWidth = 0;

static int imHeight = 0;

float imageScale = 1.f;

void display(unsigned char *pixels)

{

sobelFilter(pixels, imWidth, imHeight, imageScale);

}

unsigned char

ComputeSobelCPU(unsigned char ul, // upper left

unsigned char um, // upper middle

unsigned char ur, // upper right

unsigned char ml, // middle left

unsigned char mm, // middle (unused)

unsigned char mr, // middle right

unsigned char ll, // lower left

unsigned char lm, // lower middle

unsigned char lr, // lower right

float fScale)

{

short Horz = ur + 2*mr + lr - ul - 2*ml - ll;

short Vert = ul + 2*um + ur - ll - 2*lm - lr;

short Sum = (short)(fScale*(abs((int)Horz)+abs((int)Vert)));

if (Sum < 0)

{

return 0;

}

else if (Sum > 255)

{

return 255;

}

return (unsigned char) Sum;

}

void SobelCPU(unsigned char *pSobel, int w, int h, float fScale, unsigned char *pixels)

{

for(int i=1; i<w-1; i++)

{

for(int j=1; j<h-1; j++)

{

pSobel[(i-1)+(w-2)*(j-1)] = ComputeSobelCPU(pixels[(i-1)+w*(j-1)], pixels[(i-1)+w*(j-0)],pixels[(i-1)+w*(j+1)],

pixels[(i-0)+w*(j-1)],pixels[(i-0)+w*(j-0)],pixels[(i-0)+w*(j+1)],

pixels[(i+1)+w*(j-1)],pixels[(i+1)+w*(j-0)],pixels[(i+1)+w*(j+1)], fScale);

}

void initializeData(char *file, unsigned char *pixels)

{

unsigned int w, h;

size_t file_length= strlen(file);

printf("%s\n", file);

if (!strcmp(&file[file_length-3], "pgm"))

{

if (sdkLoadPGM<unsigned char>(file, &pixels, &w, &h) != true)

{

printf("Failed to load PGM image file: %s\n", file);

exit(EXIT_FAILURE);

}

else

{

cudaDeviceReset();

exit(EXIT_FAILURE);

}

imWidth = (int)w;

imHeight = (int)h;

setupTexture(imWidth, imHeight, pixels);

unsigned char * sobelcpu = (unsigned char*)calloc((imWidth-2) * (imHeight-2), sizeof(unsigned char));

StopWatchInterface *timer = NULL;

sdkCreateTimer(&timer);

sdkStartTimer(&timer);

SobelCPU(sobelcpu, imWidth, imHeight, imageScale, pixels);

sdkStopTimer(&timer);

float elapsedTimeInMs = sdkGetTimerValue(&timer);

printf("CPU: %3.7f ms\n", elapsedTimeInMs);

sdkSavePGM<unsigned char>("./data/2.pgm", sobelcpu, imWidth-2, imHeight-2);

}

void loadDefaultImage(char *image_path, unsigned char *pixels)

{

printf("Reading image: car. pgm\n");

const char *image_filename = "car. pgm";

image_path = sdkFindFilePath(image_filename, NULL);

if (image_path == NULL)

{

printf("Failed to read image file: <%s>\n", image_filename);

exit(EXIT_FAILURE);

}

initializeData(image_path, pixels);

free(image_path);

}

int main(int argc, char **argv)

{

int deviceCount = 0;

cudaError_t error_id = cudaGetDeviceCount(&deviceCount);

if (error_id!= cudaSuccess)

{

printf("cudaGetDeviceCount returned %d\n-> %s\n", (int)error_id, cudaGetErrorString(error_id));

printf("Result = FAIL\n");

exit(EXIT_FAILURE);

}

// This function call returns 0 if there are no CUDA capable devices.

if (deviceCount == 0)

{

printf("There are no available device(s) that support CUDA\n");

exit(EXIT_FAILURE);

}

else

{

printf("Detected %d CUDA Capable device(s)\n", deviceCount);

}

unsigned char *pixels = NULL;

char *image_path = NULL;

loadDefaultImage(image_path, pixels);

display(pixels);

}

Файл SobelGPU. cu

#include <stdio. h>

#include <stdlib. h>

#include <cuda_runtime. h>

#include <helper_string. h>

#include <helper_image. h>

#include <helper_timer. h>

#include "SobelGPU. h"

texture<unsigned char, 2> tex;

static cudaArray *array = NULL;

#define checkCudaErrors(err) __checkCudaErrors (err, __FILE__, __LINE__)

inline void __checkCudaErrors(cudaError err, const char *file, const int line)

{

if (cudaSuccess!= err)

{

fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n",

file, line, (int)err, cudaGetErrorString(err));