Иерархия памяти CUDA. Текстуры в CUDA. Цифровая обработка сигналов презентация

Содержание

Слайд 1Лекторы:
Боресков А.В. (ВМиК МГУ)
Харламов А.А. (Харламов А.А. (NVidiaХарламов А.А. (NVidia)
Иерархия памяти

CUDA. Текстуры в CUDA. Цифровая обработка сигналов

Слайд 2План
Работа с Текстурами
Свертка
Шумоподавление
Масштабирование изображений
Сжатие изображений

Слайд 3РАБОТА С ТЕКСТУРАМИ

Слайд 4Типы памяти в CUDA

Слайд 5Архитектура Tesla 10

Interconnection Network

Слайд 6Архитектура Tesla Мультипроцессор Tesla 10

Слайд 7Texture в 3D
В CUDA есть доступ к fixed-function HW: Texture

Unit



Слайд 8Texture HW
Нормализация координат:
Обращение по координатам, которые лежат в диапазоне [0,1]

Слайд 9Texture HW
Преобразование координат:
Координаты, которые не лежат в диапазоне [0,1] (или [w,

h])

Clamp:
Координата «обрубается» по допустимым границам
Wrap
- Координата «заворачивается» в допустимый диапозон


[1.1, 0.7]



[1.1, 0.3]


[1.0, 0.7]

[0.1, 0.3]

Слайд 10Texture HW
Фильтрация:
Если вы используете float координаты, что должно произойти если координата

не попадает точно в texel?

Point:
Берется ближайший texel
Linear:
- Билинейная фильтрация

Приблизимся вот в эту точку



Слайд 11Texture HW
Фильтрация
Point:
Берется ближайший texel
[0.1, 0.3]






Слайд 12Texture HW
Фильтрация
Linear:
Билинейная фильтрация

(P1*(1-α)+P2*(α)) * (1-β) +
(P3*(1-α)+P4*(α)) * (β)


[0.1, 0.3]










P1
P2
P3
P4
α
β

Слайд 13Texture в CUDA
Преобразование данных:
cudaReadModeNormalizedFloat :
Исходный массив содержит данные в integer, возвращаемое

значение во floating point представлении (доступный диапазон значений отображается в интервал [0, 1] или [-1,1])
cudaReadModeElementType
Возвращаемое значение то же, что и во внутреннем представлении

Слайд 14cudaArray
Особый контейнер памяти
Черный ящик для приложения
Позволяет организовывать данные в 1D/ 2D/3D

массивы данных вида:
1/2/4 компонентные векторы
8/16/32 bit signed/unsigned integers
32 bit float
16 bit float (driver API)
На CC 2.x возможна запись из ядра

Слайд 15Texture в CUDA (cudaArray)
Только чтение
Обращение к 1D/2D/3D массивам данных по:
Целочисленным

индексам
Нормализованным координатам
Преобразование адресов на границах
Clamp, Wrap
Фильтрация данных
Point, Linear
Преобразование данных
Данные могут храниться в формате uchar4
Возвращаемое значение – float4

Слайд 16Surface в CUDA (cudaArray)
Чтение и запись
cudaArraySurfaceLoadStore   
Обращение к 1D/2D массивам данных

по:
Целочисленным индексам
Byte-адресация по x
Преобразование адресов на границах
Clamp, Zero, Trap

Слайд 17Texture в CUDA (linear)
Можно использовать обычную линейную память
Ограничения:
1D / 2D
Нет фильтрации
Доступ

по целочисленным координатам
Обращение по адресу вне допустимого диапазона возвращает ноль

Слайд 18Texture в CUDA (linear)
texture g_TexRef;

__global__ void kernel1 ( float

* data )
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;

data [idx] = tex1Dfetch(g_TexRef, idx);
}

int main(int argc, char ** argv)
{
float *phA = NULL, *phB = NULL, *pdA = NULL, *pdB = NULL;

for (int idx = 0; idx < nThreads * nBlocks; idx++)
phA[idx] = sinf(idx * 2.0f * PI / (nThreads * nBlocks) );

CUDA_SAFE_CALL( cudaMemcpy ( pdA, phA, nMemSizeInBytes, cudaMemcpyHostToDevice ) );

CUDA_SAFE_CALL( cudaBindTexture(0, g_TexRef, pdA, nMemSizeInBytes) );

dim3 threads = dim3( nThreads );
dim3 blocks = dim3( nBlocks );

kernel1 <<>> ( pdB );
CUDA_SAFE_CALL( cudaThreadSynchronize() );

CUDA_SAFE_CALL( cudaMemcpy ( phB, pdB, nMemSizeInBytes, cudaMemcpyDeviceToHost ) );

return 0; }

Слайд 19Texture в CUDA (cudaArray)
texture g_TexRef;
__global__ void kernel ( float

* data )
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
data [idx + blockIdx.y * gridDim.x * blockDim.x] = tex2D(g_TexRef, idx, blockIdx.y);
}
int main ( int argc, char * argv [] )
{
float *phA = NULL, *phB = NULL, *pdA = NULL, *pdB = NULL;// linear memory pointers
cudaArray * paA = NULL; // device cudaArray pointer

cudaChannelFormatDesc cfD=cudaCreateChannelDesc(32,0,0,0, cudaChannelFormatKindFloat);
CUDA_SAFE_CALL( cudaMallocArray(&paA, &cfD, nBlocksX * nThreads, nBlocksY) );

for (int idx = 0; idx < nThreads * nBlocksX; idx++) {
phA[idx] = sinf(idx * 2.0f * PI / (nThreads * nBlocksX) );
phA[idx + nThreads * nBlocksX] = cosf(idx * 2.0f * PI / (nThreads * nBlocksX) );
}

CUDA_SAFE_CALL(cudaMemcpyToArray(paA,0,0,phA,nMemSizeInBytes,cudaMemcpyHostToDevice) );
CUDA_SAFE_CALL( cudaBindTextureToArray(g_TexRef, paA) );

dim3 threads = dim3( nThreads );
dim3 blocks = dim3( nBlocksX, nBlocksY );
kernel2<<>> ( pdB );
CUDA_SAFE_CALL( cudaThreadSynchronize() );
CUDA_SAFE_CALL( cudaMemcpy ( phB, pdB, nMemSizeInBytes, cudaMemcpyDeviceToHost ) );

return 0;
}

Слайд 20Texture HW
Латентность больше, чем у прямого обращения в память
Дополнительные стадии в

конвеере:
Преобразование адресов
Фильтрация
Преобразование данных
Есть кэш
Разумно использовать, если:
Объем данных не влезает в shared память
Шаблон доступа хаотичный
Данные переиспользуются разными потоками

Слайд 21СВЕРТКА

Слайд 22Свертка
Определение свертки:

В Дискретном случае:

В 2D для изображений:

Слайд 23i
Свертка


31
∑

j

Слайд 24Свертка
Вычислительная сложность:
W x H x N x K – умножений


Сепарабельные фильтры


Размер

входного сигнала

Размер ядра



Слайд 25Примеры
Gaussian Blur
Edge Detection

Слайд 26Gaussian Blur
Свертка с ядром:

Слайд 27Gaussian Blur
#define SQR(x) ((x) * (x))
texture g_TexRef;

__global__ void GaussBlur(

float * pFilteredImage, int W, int H, float r)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;

float wSum = 0.0f;
float rResult = 0.0f;
for (int ix = -r; ix <= r; ix++){
for (int iy = -r; iy <= r; iy++)
{
float w = exp( -(SQR(ix) + SQR(iy)) / SQR(r) );
rResult += w * tex2D(g_TexRef, idx + ix, idy + iy);
wSum += w;
}
}
rResult = rResult / wSum;

pFilteredImage[idx + idy * W] = rResult;
}

Слайд 28Свертка Оптимизации
Использовать сепарабельные фильтры
Существенно меньше алгоритмическая сложность
Использовать shared память

Слайд 29Свертка Оптимизации
Исходное изображение

Слайд 30Эффективно ли такое разбиение изображения

Слайд 31Свертка Smem Оптимизации

Слайд 32Свертка Smem Оптимизации

Слайд 33Свертка Smem Оптимизации

Слайд 34Свертка Smem Оптимизации

Слайд 35EDGE DETECTION

Слайд 36Edge Detection
Обнаружение границ – поиск разрывов в яркости изображения
Идеальная граница
«реальная» граница
«шумная»

граница

Слайд 37Edge Detection
Градиент функции f(x,y)
Это вектор который показывает направление роста
Определяется как


G

Gy

Gx


θ(x,y) = atan(Gy/Gx)

Слайд 38Edge Detection
Разностная производная:



Свертка с ядром:

Слайд 39Edge Detection
Разностная производная:



Свертка с ядром:

Слайд 40Edge Detection
Prewitt mask:



Слайд 41Edge Detection
Sobel mask:

Слайд 42
Edge Detection
Оператор Лапласа:

Слайд 43ШУМОПОДАВЛЕНИЕ

Слайд 44Discrete Cosine Transform
Является основой современных алгоритмов сжатия данных с потерями (JPEG,

MPEG)


JPEG, 2/10

Слайд 45Discrete Cosine Transform
Представитель семейства пространственно-частотных 1D преобразований, задается формулами:
Прямое:
Обратное:
Нормировочные коэффициенты:

Слайд 46Discrete Cosine Transform
8-точечный случай: u=0

Слайд 47Discrete Cosine Transform
8-точечный случай: u=1

Слайд 48Discrete Cosine Transform
8-точечный случай: u=2

Слайд 49Discrete Cosine Transform
8-точечный случай: u=3

Слайд 50Discrete Cosine Transform
8-точечный случай: u=4

Слайд 51Discrete Cosine Transform
8-точечный случай: u=5

Слайд 52Discrete Cosine Transform
8-точечный случай: u=6

Слайд 53Discrete Cosine Transform
8-точечный случай: u=7

Слайд 54Discrete Cosine Transform
N-мерное преобразование обладает свойством сепарабельности


Коэффициенты A[8x8] преобразования вычисляются один

раз

Слайд 55Discrete Cosine Transform
Наивный: 64 нити на блок (8х8)
Загрузка одного пикселя из

текстуры
Барьер
Поток вычисляет один коэффициент
Барьер
Запись коэффициента в глобальную память

Слайд 56Насколько это эффективно?

Слайд 57Discrete Cosine Transform
Блок потоков обрабатывает несколько блоков 8х8
Один поток обрабатывает вектор

8х1 (1x8)

Слайд 58Image Denoising
Шумы в изображении
Импульсный
Salt & pepper
Аддитивный
Uniform
Gaussian

Слайд 59Ранговые фильтры
Алгоритм Р.Ф. ранга N:
Для каждого отсчета сигнала i
Выбор окрестности вокруг

отсчета i
Сортируем по значению
Выбираем N-ое значение как результат

Слайд 60Медиана
Ранговый фильтр N=0.5
Сортировка не обязательна для 8bit значений
Строим гистограмму
for(int i

i++) h[ signal[i] ]++;
Сканируем Гистограмму:
int sum = 0;
int targetSum = N*rank;
for(int i<0; i<256; i++)
{
sum += h[i]; if (sum > targetSum) return i;
}

Слайд 61Медиана Построение гистограммы


Сигнал + фартук в Smem





Слайд 62Что с банк-конфликтами?

Слайд 63Фильтрация (Аддитивный шум)
Размытие – это low-pass фильтр
Каким должен быть фильтр?
Подавлять шум?
Сохранять

детальность?











noise

E(noise) = 0

Слайд 64Gaussian Blur
Blur (размытие) изображение
Свертка с ядром:

Слайд 65Gaussian Blur
Blur (размытие) изображение
Свертка с ядром:

Слайд 66Адаптивное размытие
Свертка с ядром:

ClrSpaceDist – это фотометрическая близость

Слайд 67Bilateral

































































Слайд 68Bilateral
































































Слайд 69Bilateral Kernel
#define SQR(x) ((x) * (x))
texture g_TexRef;

__global__ void BilateralBlur(

float * pFilteredImage, int W, int H, float r)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;

float wSum = 0.0f;
float rResult = 0.0f;
float c = tex2D(g_TexRef, idx, idy);
for (int ix = -r; ix <= r; ix++)
for (int iy = -r; iy <= r; iy++)
{
float clr = tex2D(g_TexRef, idx + ix, idy + iy);
float w = exp( -(SQR(ix) + SQR(iy)) / SQR(r) – SQR(clr-c)/SQR(h));
rResult += w * clr;
wSum += w;
}
rResult = rResult / wSum;

pFilteredImage[idx + idy * W] = rResult;
}

Слайд 70Bilateral Оптимизации
Bilateral не сепарабельный фильтр
Но можно его разделить
Смешивать исходное изображение с фильтрованным
Если

в блоке много ненулевых коэф., то с большой вероятностью в этом блоке шум был подавлен успешно
Если в блоке много нулевых коэф., то с большой вероятностью в блоке много деталей (границы, текстура и т.д.)

Слайд 71Non Local Means
ClrSpaceDist – оценивать по блокам пикселей



































































Слайд 72Non Local Means
На вычисление одного веса:
NbxNb вычислений, N размер блока
На

фильтрацую одного пиксела:
NbxNb x RxR, R размер окна

Слайд 73Сравнение

Слайд 74Сравнение Bilateral

Слайд 75Сравнение NLM

Слайд 76МАСШТАБИРОВАНИЕ ИЗОБРАЖЕНИЙ

Слайд 77Артифакты
Алиасинг
При увелечении – ступенчатость
При уменьшении - муар
Ringing
Потеря четкости
Субпиксельный сдвиг
Влияет на формальные

метрики

Слайд 78Простые методы
Билинейная интерполяция
P1
P7
P3
P9
P4
P2
P5
P6
P8

Слайд 79Простые методы
Билинейная интерполяция
Сепарабельная
Очень быстрая
Поддерживается в HW
Точность фильтрации

Слайд 80Простые методы
Бикубическая интерполяция
Сепарабельная
Лучше качество

t=0
t=1
t=0.5

Слайд 81Сравнение

Слайд 82Lanczos


Слайд 83Gradient interpolation

P1
P7
P3
P9
P7
P2
P5
P6
P8
Dxd = abs(P3 – P5)
Dyd = abs(P1 – P9)
If (Dxd

> Dyd)
//граница P1P5P9
P5 = (P1 + P9) * 0.5f;

If (Dyd > Dxd)
//граница P3P5P7
P5 = (P1 + P9) * 0.5f;

If (Dyd ~= Dxd)
//граница не определена
P5 = (P1 + P3 + P7 + P9) * 0.25f;

Слайд 84NEDI (New Edge-Directed Interpolation)















































Слайд 85NEDI (New Edge-Directed Interpolation)















































Слайд 86NEDI (New Edge-Directed Interpolation)

Слайд 87NEDI (New Edge-Directed Interpolation)

Слайд 88NEDI
2i
2j
X = F(2i+1,2j+1) = α1*F(2i,2j)+

α2*F(2i+2,2j)+
α3*F(2i,2j+2)+
α4*F(2i+2,2j+2);

α = {α1, α2, α3, α4} ?

Слайд 89NEDI

0 1 2 3 …
0

1 2 3…




Слайд 90NEDI

0 1 2 3 …
0

1 2 3…




Xi,j = α1*F(i-2,j-2) + α2*F(i+2,j-2) +
α3*F(i-2,j+2) + α4*F(i+2,j+2);

For i = 2,4,6,8 | For j = 2,4,6,8

Ei,j = Pj,j – Xi,j – Approximation Error

SSE = ΣΣ sqr( Ei,j );

α = Arg min(SSE);
α
P = {Pij}
Cα = P

∂(SSE) / ∂α = 0.0

α = (CTC)-1 CTP

Слайд 91NEDI (improvement)
…
…

Слайд 92NEDI (improvement)
…
…





Слайд 93NEDI: Pros and Cons
Pros: NEDI
Четкие тонкие края
Cons: Очень медленно на CPU
Умножение

матриц 4х16
Обращение матрицы
Рингинг
CUDA:
Большой объем данных на тред
Много регистров
Сложно использовать Smem
Много ветвлений

Слайд 94Сравнение

Слайд 95Сравнение

Слайд 96ФРАКТАЛЬНОЕ СЖАТИЕ АЛГОРИТМА

Слайд 97Мат.часть
Сжимающее отображение f, X?X:

Если f сжимающее, то

Слайд 98Фрактальное сжатие
Ранг
Домен

Слайд 99Фрактальное сжатие
Ранг
Домен

Слайд 100Фрактальное сжатие
Ранг
Домен

Слайд 101Фрактальное сжатие
Ранг
Домен

Слайд 102...
Фрактальное сжатие
Ранг
Домен



Слайд 103Фрактальное сжатие
Ранг
Домен

Слайд 104Фрактальное сжатие
Ранг
Домен

Слайд 105Фрактальное сжатие

Слайд 106Фрактальное сжатие
Блок пикселей заменяется преобразованием
Сдвиг, поворот / отражение
s,t – преобразование яркости
float2

x rgb
s,t – сжимают во время компресси
Например s:5,t:3
в интервале [0, 1]

Слайд 107Декомпрессия
Сжимающее отображение f, X?X:

Если f сжимающее, то

Берем любое изображение
Строим Доменное

изображение
Применим полученный набор преобразований к блокам

Слайд 108Сколько раз необходимо применять преобразования?

Слайд 109Декомпрессия

Слайд 110Артифакты

Слайд 111Артефакты

Слайд 112Вопросы

Слайд 113Задание 1
Вычислить число pi методом монтекарло
Генерация точки (x, y) в [-1,

1] x [-1, 1]
Расчет сколько точек попало в ед. круг

Построение Kd дерева для точек
Задние на thrust

Слайд 114Задание 2
Изменение времени звуковой дорожки без изменения тона
Задание на CUFFT
Обучение фильтра
Разбить

изображение на блоки
Классифицировать каждый блок
Обучить фильтр имея пару (R, T)
R – «правильное» изображение
T – «входной» сигнал

Слайд 115Ресурсы нашего курса
Steps3d.Narod.Ru
Google Site CUDA.CS.MSU.SU
Google Group CUDA.CS.MSU.SU
Google Mail CS.MSU.SU
Google SVN
Tesla.Parallel.Ru
Twirpx.Com
Nvidia.Ru



Слайд 116Преобразование Фурье
Линейный оператор вида:


Обратный оператор:

Слайд 117Преобразование Фурье
Условие существования



Конечное число устранимых разрывов

Слайд 118Преобразование Фурье Пример 1D

Слайд 119Преобразование Фурье Пример 2D

Слайд 120Преобразование Фурье Свойства







Слайд 121Преобразование Фурье Свойства









Слайд 122Преобразование Фурье Свойства









Слайд 123Фильтры
Низкочастотные (low-pass)
Высокочастотные (high-pass)

Слайд 124Фильтры

Похожие презентации

ЗАБЫТЫЙ ХУДОЖНИК РОССИИ 316
Водный режим почв, дренажная система (Лекция 2) 272
Машины переменного тока 226
5 класс Урок 3 Ввод информации в память компьютера. Клавиатура. Группы клавиш.Практическая работа №1. Знакомство с клавиатурой. Учитель: Фетисова С.А. Г. Киселевск Кемеровской области СОШ №33 270
Если хочешь быть здоров 181
DragAndDrop 235

Обратная связь

Если не удалось найти и скачать презентацию, Вы можете заказать его на нашем сайте. Мы постараемся найти нужный Вам материал и отправим по электронной почте. Не стесняйтесь обращаться к нам, если у вас возникли вопросы или пожелания:

Email: Нажмите что бы посмотреть 

Что такое ThePresentation.ru?

Это сайт презентаций, докладов, проектов, шаблонов в формате PowerPoint. Мы помогаем школьникам, студентам, учителям, преподавателям хранить и обмениваться учебными материалами с другими пользователями.

Для правообладателей

Яндекс.Метрика