Лихогруд Николай 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]; - презентация

1 Лихогруд Николай Задание

2

3 Свертка с квадратным ядром 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 ; Квадратное ядро

4 Фильтрация изображений

5 Базовое задание Выбрать любые два фильтра из статьи Разобраться с какой-либо библиотекой для работы с изображениям (предлагаю DevIL ) Реализовать загрузку, модификацию и сохранение изображений в формате png/jpeg с палитрой RGBA Реализовать добавление к изображению рамки с шириной, равной радиусу фильтра Рамка заполняется крайними пикселами изображения

6 Базовое задание Написать ядро, принимающее на вход Фильтр (через константную память) Исходную матрицу с рамкой Память под матрицу-результат (тоже с рамкой) Каждая нить вычисляет один элемент матрицы-результата Ядро фильтрации передается через константную память #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) {}

7 Работа с пикселями Ядро работает с массивами пикселов, каждая нить рассчитывает пиксел выходного изображения – компоненты R, G, B Удобно использовать класс Pixel с полями uint8_t r,g,b,a ; и FloatPixel с полями float r,g,b,a Через Pixel читаем/пишем память В FloatPixel накапливаем результат Для считывания/записи используем get_pixel – аналог get_elem, но приводит указатель на начало строки к типу Pixel, а не float

8 Pixel #include struct Pixel { uint8_t r,g,b,a; __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) }; Оператор присваивания и конструктор копирования копирует сразу 4 байта, а не побайтовое

9 Замечания к классу Pixel __host__ __device__ означают, что функция будет скомпилирована и для CPU, и для GPU При работе с памятью важно считывать/записывать пиксел одной 4- байтной инструкцией, вместо 4 однобайтовых С этой целью переопределены конструктор копирования и оператор присваивания Pixel *pixels; // массив пикселов Pixel pixel2 = pixels[2]; // Конструктор копирования Pixels[0] = pixel1 + pixel2; // Оператор присваивания Проверьте через $ cuobjdump -sass, чтобы в ассемблере не было инструкций ST/LD c суффиксом U8

10 FloatPixel #include struct FloatPixel { float r,g,b,a; __host__ __device__ FloatPixel(float r, float g, float b, float a); __host__ __device__ FloatPixel operator+(const FloatPixel &otherPixel); __host__ __device__ Pixel getPixel(); }; Оператор сложения с другим FloatPixel для накопления результата

11 Про эрозию и наращивание В эрозии и наращивании нужно найти в окрестности элемент с максимальной/минимальной яркостью Окрестность задается матрицей из нулей и единиц Яркость пиксела можно записать в четвертую компоненту, все равно мы её не используем Тогда можно добавить оператор Pixel &operator>(Pixel &otherPixel), возвращающий пиксел с большей четвертой компонентой Яркость всех пикселов нужно посчитать один раз, на хосте или GPU, по формуле 0.3*R *G *B (или 0.21R G B)

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

13 Требования к базовой части Входные параметры: Радиус фильтра Размеры синтетической матрицы или файл с входным изображением Программа: Выводися время применения фильтров Если на входе было изображение – сохраняет на диск результат применения двух фильтров

14

15 Добавляем потоки Программа дополнительно принимает число потоков Разрежем матрицу на горизонтальные полосы по числу потоков Каждый поток рассчитывает одну из полос

16 Добавляем потоки Для старта вычислений в нулевом потоке достаточно скопировать первую полосу вместе с её нижней рамкой

17 Добавляем потоки Для старта вычислений в следующем потоке часть данных уже скопирована Что уже есть на GPU

18 Добавляем потоки Для следующего потока так же часть данных уже будет на GPU Что уже есть на GPU

19 Ускорение При помощи профилировщика подобрать радиус фильтра и размеры синтетической матрицы так, чтобы копирования совпадали со временем выполнения Максимальное ускорение

20

21 Используем multi-GPU Разрезать матрицу на горизонтальные полосы по числу видеокарт Разослать по видеокартам полосы + их рамки Каждая видеокарта считает независимо свою часть С использованием потоков Использовать OpenMP

22 3 GPU

23 Второй фильтр & multi GPU Для применения второго фильтра понадобятся рамки полос, вычисленные на соседних устройствах Их нужно переслать после применения первого фильтра через cudaMemCpyAsync ()/ cudaMemcpyPeerAsync () Включить прямой доступ (peerAccess) там, где это возможно

24

25 end