ASSIGNMENT 3. 2011 Assignment 3.1 Растеризация на CUDA – Загрузить 3d модель (stanford bunny, dragon или buddha)stanford bunny, dragon или buddha – Преобразовать. - презентация

1 ASSIGNMENT 3

2 2011 Assignment 3.1 Растеризация на CUDA – Загрузить 3d модель (stanford bunny, dragon или buddha)stanford bunny, dragon или buddha – Преобразовать вершины – Растеризация в разделяемой памяти – Простое освещение

3 2011 Загрузка модели Используйте.obj формат – текстовый файл с простой разметкой v v v v v v v v vn vn vn vn vn f 1//1 3//3 4//4 f 4//4 2//2 1//1 f 5//5 6//6 8//8 f 8//8 7//7 5//5 f 1//1 2//2 6//6 f 6//6 5//5 1//1 f 2//2 4//4 8//8 f 8//8 6//6 2//2 f 4//4 3//3 7//7 f 7//7 8//8 4//4 f 3//3 1//1 5//5 f 5//5 7//7 3//3 V – префикс задает вершину Vn – префикс задает нормаль f – префикс задает грань в формате: индекс вершины/индекс текстуры/индекс нормали. Индекс текстуры для задания не нужен. Индексы начинаются с 1, в то время как массивы в С++ с 0 f – префикс задает грань в формате: индекс вершины/индекс текстуры/индекс нормали. Индекс текстуры для задания не нужен. Индексы начинаются с 1, в то время как массивы в С++ с 0

4 2011 Загрузка модели Загрузку можете написать самостоятельно Можете взять готовую в интернете – Например glmglm Можете спросить в группе

5 2011 Преобразование вершин Преобразование вершин бывает – Model или «модельное». Это преобразование вершин из координатного пространства модели в мировые координаты. Типичные преобразования – поворот, масштабирование и перенос. – View или «видовое». Это преобразование вершин из мировых координат в координаты камеры. Камера – это точка в пространстве, с направлением взгляда и вектором ориентации «вверх». В графических API камера распологается в точке (0,0,0), смотрит по вектору (0, 0, 1) (или (0, 0, -1)), и ориентирована по вектору вверх (0,1,0). – Projection или «проекции». Это преобразование вершин из координат камеры в координаты пространства отсечения. После умножения на матрицу проекции xyz координаты делятся на w. Полученные координаты лежат в пространстве отсечения (это куб [-1, 1] 2 x [0, 1]). Вершины которые оказались вне этого куба – не попадают в пирамиду видимости. Все преобразования вершин работают с векторами xyzw. – Для точек w == 1, для направлений w == 0. Все матрицы преобразований – матрицы 4х4 – Как построить матрицу – смотрим вызовы directx или спрашиваем в группе.

6 2011 Преобразование вершин Преобразование вершин : – Пусть N треугольников, треугольник это 3 вершины идущие подряд: A 1, B 1, C 1, … A N, B N, C N – Более оптимально разложить треугольники так: A 1, A 2, A 3,.. A N, B 1, B 2, B 3,..,B N, C 1, C 2, C 3,..,C N Входной список вершин (N*3) Входной список нормалей (N*3) Выходной список вершин (N*3) Выходной список нормалей (N*3) CUDA ядро (N*3) потоков: V * Model * View * Projection N * Model CUDA ядро (N*3) потоков: V * Model * View * Projection N * Model

7 2011 Глобальная память для буфера кадра Окно растеризации (framebuffer) WxH пикселей – два массива данных в глобальной памяти. Размер – параметр. Хорошие размеры для эксперимента – 1024х1024. Их необходимо проинициализовать начальными значениями Например – черный цвет и глубина равная 1.0 (бесконечно далеко) Depth (float) Color (uchar4)

8 2011 Растеризация Всё окно разбивается на тайлы (tile). Размер тайла – например 8x8, 16x8, 16x16, 32x2, 32x4, 32x8, 32x16 пикселей. Вероятнее всего, оптимальный размер – 32x8, но это не известно Для каждого тайла запускается CUDA блок, в каждом блоке кол-во потоков == кол- ву пикселей в тайле (это важно)

9 2011 Растеризация Рассмотрим работу отдельного тайла

10 2011 Растеризация CUDA ядро: пусть размер блока Bx*By – Прочитать из глобальной памяти в разделяемую исходные значения цвета и глубины (1 поток читает по одному значению) __shared __ uchar4 color[Bx*By]; __shared__ float depth[Bx*By]; – синхронизация

11 2011 Растеризация CUDA ядро: пусть размер блока Bx*By – Вычислить границы тайла в пространстве [-1, 1] 2 – Каждый поток читает треугольник из глобальной памяти и проверяет пересекает ли треугольник данный тайл Т.е. Если N треугольников, то каждый поток пересекает N/(Bx*By) с тайлом чем больше потоков, тем лучше. Выходной список вершин (N*3) Выходной список нормалей (N*3)

12 2011 Растеризация CUDA ядро: пусть размер блока Bx*By – Если треугольник пересекает тайл то его индекс в памяти записывается в массив разделяемой памяти __shared__ int smem_tri_idx[Bx*By] – Если треугольник не пересекает тайл то в разделяемую память записывается -1 – синхронизация

13 2011 Растеризация Пересекает поток : 0, 1, 2, 3, 4, 5, 6, 7 Результат пересечения: 1, 0, 0, 0, 1, 1, 1, 0 Поток записал в shared: 0, -1, -1, -1, 4, 5, 6, -1

14 2011 Растеризация Пересекает поток : 0, 1, 2, 3, 4, 5, 6, 7 Результат пересечения: 1, 0, 0, 0, 1, 1, 1, 0 Поток записал в shared: 0, -1, -1, -1, 4, 5, 6, -1 Не пересекающие тайл треугольники нас не интересуют!

15 2011 Растеризация CUDA ядро: – Синхронизация (индексы треугольников лежат в shared памяти) – Цикл (по размеру блока) Каждый поток берет индекс треугольника Читает треугольник из глобальной памяти – если индекс != -1 Проверяет пересекает ли данный треугольник тот пиксел, которому соответствует поток

16 2011 Растеризация: тайл крупным планом

17 2011 Растеризация: тайл крупным планом Как вычислить барицентрические коодинаты? – Разложить AP = u*AB + v*AC – 2 ур-ния (x, y) – 2 неизвестные (u, v) – Решение через определитель – Если u >= 0 и v >= 0 и u+v

18 2011 Растеризация: тайл крупным планом Нужно вычислить – Глубину точки Сравнить с глубиной в разделяемой памяти Если тест глубины пройден, тогда обновить глубину в разделяемой памяти и сделать расчет освещения – Расчет освещения и запись в разделяемую память Вычислить единичную нормаль в точке n Диффузное освещение L в точке P – Lp = max(0, dot(n, light_dir)) где light_dir – направление счета

19 2011 Растеризация CUDA ядро: – Продолжить процесс для следующих 64-ёх треугольников – Когда все треугольники закончились записать результат из разделяемой в глобальную память – Можно дальше растеризовать след. модель Разделяемая память Глобальная память

20 2011 Результат Резульатат растеризации записывайте в файл – Можете использовать devil course/wiki/Devil_Basics – Работает на нашем сервере

21 2011 Assignment 3.2 Тоже самое что задание 3.1 но дополнительно на выбор: – Растеризовать не весь список треугольников, а предварительно составить более короткие списки треугольников Используйте Thrust и операцию stream compaction – Более сложная модель освещения (например Phong) – Поддержка OpenGL (возможность полетать камерой) Тогда проверять будем на windows машине Быстрое создание gl окна course/wiki/Freeglut_GLEW_Basicshttp://code.google.com/p/msu-opengl- course/wiki/Freeglut_GLEW_Basics – Отложенное освещение (deferred shading)

22 2011 Глобальная память для буфера кадра Depth (float) Color (uchar4)

23 2011 Глобальная память для буфера кадра Depth (float) Color (uchar4) Q1 Q2 Q3 Q4 4 списка: Q1: Q2: Q3: Q4:

24 списка: Q11: Q12: Q13: Q14: Q2: Q3: Q4: Глобальная память для буфера кадра: можно иерархически Depth (float) Color (uchar4) Q11 Q2 Q3 Q4 Q12 Q13 Q14

25 2011 Assignment 3.2 Более сложная модель освещения: – Очень подробное описание моделей с кодом шейдеров glsl можно найти тут: – Спрашивайте в нашей группе

26 2011 Assignment 3.2 Deferred shading – Проблема в задании 3.1 в том, что когда часть потоков найдет пересечение пиксела с треугольником а часть нет, то возникнет простой в работе: Те потоки которые найдут пересечение с треугольником начнут вычислять освещение Если освещение сложное – то это много математических инструкций Остальные потоки (для которых пиксели не пересекают данный треугольник) ждут – Решение: отложить процесс освещения до самого конца Если пиксел пересекает треугольник, то запоминать индекс треугольника и барицентрические координаты u и v В разделяемой памяти записывать глубину, индекс треугольника, u, v После того как все треугольники проверены – вычислить освещение в тех пикселах для которых индекс треугольника != -1

27 2011 Общие правила по оформлению программ Если сдаете по – ДОЛЖЕН быть с темой CUDA ASSIGNEMENT 2011.N (N - номер задания, в данном случае 3) Задание сложное – И на него у вас 1 неделя – Не откладывайте с вопросами Недвусмысленный намек от кэпа

28 2011 Общие правила по оформлению прорамм – Программа должна делать проверки на ошибки: Наличие девайса? Выделилась память? И т.д. – Программа должна компилироваться с CUDA Toolkit 3.2 Если писали на windows то vcproj для VS2005 / VS2008 либо (makefile +.bat) Если писали на *nix то make

29 2011 Общие правила по оформлению программ Если вы используете любые другие инклюды кроме стандартных – не расчитывайте, что они прописаны на проверяющей машине. Пример того, чего не будет на машине: – cutil.h (требует установки CUDA SDK) Пример того, что будет на машине: – cudart.h (ставиться вместе с CUDA toolkit) – stdio.h (стандартная C библиотека) – thrust – cufft

30 2011 Вопросы