cuda_course_task_filtering (1265197)

Просмтор этого файла доступен только зарегистрированным пользователям. Но у нас супер быстрая регистрация: достаточно только электронной почты!

Текст из файла

Лихогруд Николайn.lihogrud@gmail.comЗаданиеСвертка с квадратным ядромtmp = 0;for(ik=-r..r)for(jk=-r..r)tmp += matrix[i+ik][j+jk]*filter[ik+r][jk+r];result[i][j]=tmp;КвадратноеядроФильтрация изображенийБазовое задание Выбрать любые два фильтра из статьиhttp://habrahabr.ru/post/142818/ Разобраться с какой-либо библиотекой для работы сизображениям (предлагаю DevIL ) Реализовать загрузку, модификацию и сохранение изображенийв формате png/jpeg с палитрой RGBA Реализовать добавление к изображению рамки с шириной,равной радиусу фильтра Рамка заполняется крайними пикселами изображенияБазовое задание Написать ядро, принимающее на вход Фильтр (через константную память) Исходную матрицу с рамкой Память под матрицу-результат (тоже с рамкой) Каждая нить вычисляет один элемент матрицы-результата Ядро фильтрации передается через константную память#define MAX_KERNEL_WIDTH (2 * MAX_KERNEL_RADIUS + 1)__constant__ float kernel[MAX_KERNEL_WIDTH * MAX_KERNEL_WIDTH]__global__ void filter(cudaPitchedPtr inMatrix,cudaPitchedPtr outMatrix, int kernelRadius) {}Работа с пикселями Ядро работает с массивами пикселов, каждая нить рассчитываетпиксел выходного изображения – компоненты R, G, B Удобно использовать класс Pixel с полями uint8_t r,g,b,a; иFloatPixel с полями float r,g,b,a Через Pixel читаем/пишем память В FloatPixel накапливаем результат Для считывания/записи используем get_pixel – аналогget_elem, но приводит указатель на начало строки к типу Pixel, ане floatPixel#include <stdint.h>struct Pixel {uint8_t r,g,b,a;Оператор присваивания и конструктор копированиякопирует сразу 4 байта, а не побайтово__host__ __device__ Pixel &operator=(const Pixel &otherPixel) {*((uint32_t *)this) = (*(uint32_t *)&otherPixel);return *this;}__host__ __device__ Pixel(const Pixel &otherPixel) {*((uint32_t *)this) = (*(uint32_t *)&otherPixel);}__host__ __device__ Pixel(uint8_t r, uint8_t g, uint8_t b, uint8_t a);__host__ __device__ FloatPixel operator*(float coefficient)};Замечания к классу Pixel __host__ __device__ означают, что функция будет скомпилирована идля CPU, и для GPU При работе с памятью важно считывать/записывать пиксел одной 4байтной инструкцией, вместо 4 однобайтовых С этой целью переопределены конструктор копирования и операторприсваиванияPixel *pixels; // массив пикселовPixel pixel2 = pixels[2]; // Конструктор копированияPixels[0] = pixel1 + pixel2; // Оператор присваивания Проверьте через $cuobjdump -sass, чтобы в ассемблере не былоинструкций ST/LD c суффиксом U8FloatPixel#include <stdint.h>struct FloatPixel {float r,g,b,a;__host__ __device__ FloatPixel(float r, float g, float b, float a);Оператор сложения с другим FloatPixelдля накопления результата__host__ __device__ FloatPixel operator+(const FloatPixel &otherPixel);__host__ __device__ Pixel getPixel();};Про эрозию и наращивание В эрозии и наращивании нужно найти в окрестности элемент смаксимальной/минимальной яркостью Окрестность задается матрицей из нулей и единиц Яркость пиксела можно записать в четвертую компоненту, все равно мыеё не используем Тогда можно добавить оператор Pixel &operator>(Pixel&otherPixel), возвращающий пиксел с большей четвертойкомпонентой Яркость всех пикселов нужно посчитать один раз, на хосте или GPU, поформуле 0.3*R + 0.59*G + 0.11*B (или 0.21R + 0.72G + 0.07B)Синтетические данные Для экспериментов с большими матрицами: Добавить возможность генерации случайных матриц изчисел (не пикселов) Добавить отдельное ядро, работающее с матрицами извещественных чисел Каждая нить считает не пиксел а просто суммупроизведений элементов окрестности на коэффициентыядраТребования к базовой части Входные параметры: Радиус фильтра Размеры синтетической матрицы или файл с входнымизображением Программа: Выводися время применения фильтров Если на входе было изображение – сохраняет на дискрезультат применения двух фильтровДобавляем потоки Программа дополнительно принимает число потоков Разрежем матрицу на горизонтальные полосы почислу потоков Каждый поток рассчитывает одну из полосДобавляем потоки Для старта вычислений в нулевом потоке достаточно скопироватьпервую полосу вместе с её нижней рамкойДобавляем потоки Для старта вычислений в следующем потоке часть данных ужеЧто уже есть на GPUскопированаДобавляем потокиЧто уже есть на GPU Для следующего потока так же часть данных уже будет на GPUУскорение При помощи профилировщика подобрать радиусфильтра и размеры синтетической матрицы так, чтобыкопирования совпадали со временем выполнения Максимальное ускорениеИспользуем multi-GPU Разрезать матрицу на горизонтальные полосы почислу видеокарт Разослать по видеокартам полосы + их рамки Каждая видеокарта считает независимо свою часть С использованием потоков Использовать OpenMP3 GPUВторой фильтр & multi GPU Для применения второго фильтра понадобятся рамкиполос, вычисленные на соседних устройствах Их нужно переслать после применения первого фильтрачерез cudaMemCpyAsync()/ cudaMemcpyPeerAsync() Включить прямой доступ (peerAccess) там, где этовозможноend.

Характеристики

Тип файла PDF

PDF-формат наиболее широко используется для просмотра любого типа файлов на любом устройстве. В него можно сохранить документ, таблицы, презентацию, текст, чертежи, вычисления, графики и всё остальное, что можно показать на экране любого устройства. Именно его лучше всего использовать для печати.

Например, если Вам нужно распечатать чертёж из автокада, Вы сохраните чертёж на флешку, но будет ли автокад в пункте печати? А если будет, то нужная версия с нужными библиотеками? Именно для этого и нужен формат PDF - в нём точно будет показано верно вне зависимости от того, в какой программе создали PDF-файл и есть ли нужная программа для его просмотра.

Список файлов учебной работы