cuda_course_task_filtering (Задания)

Лихогруд Николай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.