GPU Ray Tracing Vladimir Frolov (Nvidia, MSU, Keldysh Institute of Applied Math)

2011 План Трассировка лучей и растеризация – Определение видимости – Global illumination Ускоряющие структуры Особенности трассировки лучей на GPU Узкие места Работа мультипроцессора на низком уровне – CPU vs GPU Оптимизация и отладка неоднородного кода

2011 Растеризация vs трассировка

2011 Растеризация vs трассировка Растеризация – Обработка по 1 треугольнику – Трудно параллелится – Эквивалентна задаче построения дерева или сортировке

2011 Трассировка лучей Алгоритм

2011 Трассировка лучей Представление 3D объектов – Аналитическое – Меши из треугольников

2011 Трассировка лучей Поверхность задана как массив треугольников Узкое место – поиск пересечения луча с поверхностью – треугольников – лучей – => операций – log(N) k * 10 6 (k~[1..2])

2011 Пересечение луча и треугольника Простой вариант – Ax + By + Cz + D = 0 – Найти t x = p.x + v.x*t y = p.y + v.y*t z = p.z + v.z*t

2011 Пересечение луча и треугольника Простой вариант – t известно z = p + v*t S = cross(v1-v0, v2-v0) u = cross(v1-z, v0-z) v = cross(v1-z, v2-z) t1 = cross(v2-z, v0-z) – |u + v + t1 – S | < ε

2011 Пересечение луча и треугольника Оптимизированный вариант – Барицентрические координаты u := u/S, v := v/S, t1 := t1/S t1 = 1 – u - v

2011 Пересечение луча и треугольника Оптимизированный вариант

2011 Пересечение луча и треугольника Unit test float3 o = mul3x4(m, origin); float3 d = mul3x3(m, dir); float t = -o.z/d.z; float u = o.x + t*d.x; float v = o.y + t*d.y; Где m это T

2011 Пересечение луча и треугольника Где m это T float3 o = mul3x4(m, origin); float3 d = mul3x3(m, dir); float t = -o.z/d.z; float u = o.x + t*d.x; float v = o.y + t*d.y;

2011 Пересечение луча и треугольника Простой вариант – Операции (* : 39, +/- : 53, / : 1) Оптимизированный вариант – Операции (* : 23, +/- : 24, / : 1) Юнит тест – Операции (*: 20, + : 20, / : 1) – Экономит регистры GPU

2011 Регистры и локальные переменные float3 RayTrace (const Ray& ray) { for (int depth=0 ;depth

2011 Занятость мультипроцессора Регистры – GF100: maxreg = 32768, maxwarps = 48 – block size = 256 рег

2011 GPU vs CPU – Локальные переменные CPU : стэк GPU : регистры – Задержки в конвейере (Pipeline hazards) CPU : Out Of Order GPU : Hyperthreading – Латентность памяти CPU: огромный кэш GPU: сокрытие латентности за счет параллелизма

2011 Проблеммы конвейерных ВУ add R1,R2 mul R0,R1 add R1,R2 nop mul R0,R1 add R1,R2 mul R0,R1 Регистровый файл Read After Write (RAW) Hazard

2011 Почему GPU быстрее? add R1,R2 mul R0,R1 add R1,R2 * mul R0,R1 * add R1,R2 ** mul R0,R1 ** add R1,R2 add R1,R2 * add R1,R2 ** mul R0,R1 mul R0,R1 * mul R0,R1 **

2011 GPU vs CPU Латентность памяти

2011 Занятость мультипроцессора Регистры

2011 Занятость мультипроцессора Регистры – G80: maxreg = 8192, maxwarps = 24 – GT200: maxreg = 16384, maxwarps = 32 – GF1**: maxreg = 32768, maxwarps = 48

2011 Занятость мультипроцессора Регистры – GF100: maxreg = 32768, maxwarps = 48 – block size = 256 рег

2011 Занятость мультипроцессора Разделяемая память – GF100: size per SM = bytes, maxwarps = 48 – block size = 256, 8 warps – байт на блок (64 байта на тред) – байт на блок (48 байт на тред) – 8192 байта на блок (32 байта на тред)

2011 Занятость мультипроцессора CUDA occupancy calculator

2011 Занятость мультипроцессора Текстурный кэш

2011 Пересечение луча и треугольника Оптимизированный вариант – регистра Unit test – регистра И это только пересечение луча и треугольника!

2011 Регистры и локальные переменные float3 RayTrace (const Ray& ray) { for (int depth=0 ;depth

2011 Экономим регистры Два подхода: – Uber kernel – Separate kernel

2011 Uber kernel vs Separate kernel Uber kernel zРекурсия (+) zЛокальность данных (+) zНепрерывное выполнение (+) zТрудно профилировать (-) zТрудно оптимизировать (-) zСамая сложная часть лимитирует все ядро (-) Separate kernel z Нет рекурсии(-) z Обращения к DRAM (-) z Латентность остановки/запуска (-) z Легче отлаживать (+) z Легко профилировать (+) z Легко оптимизировать (+) z Гибкость (+)

2011 Организация рей трейсера Ядро пересечений – регистра Но это только пересечения Нужно разбить алгоритм трассировки на несколько ядер

2011 Ускоряющие структуры regular grid kd-tree BVH

2011 Регулярная сетка

2011 Регулярная сетка traversal if (tMaxX

2011 Регулярная сетка Преимущества – Просто и быстро строится – Простой алгоритм траверса Недостатки – Плохо справляется с пустым пространством – Требует много памяти – Много повторных пересечений – отвратительно разбивает геометрию Только для небольших сцен (1-50K)

2011 Регулярная сетка Почему сетка плохо разбивает геометрию? zПеребрали 15 вокселей z7 раз посчитали пересечение с одним и тем же треугольником!

2011 Иерархическая сетка Небольшое число вокселей Рекурсивно разбиваем воксели в местах с плотной геометрией

2011 Иерархическая сетка + Решает проблему чайника на стадионе - Переход между узлами вычислительно сложен - 12 доп. регистров как минимум - Нужно устранять рекурсию

2011 kd-деревья

2011 kd-деревья L R R L

2011 kd-деревья R L LLLRRLRR LL LRRL RR

2011 kd-деревья R L LLLRRLRR LLLLLR LRL LRR RL RRL RRR

2011 kd-деревья R L LLLRRLRR RL RRR RRL LRR LLRLLLR LLLL LRLL LRLR

2011 kd-деревья Алгоритм траверса Регистры – 13 min: – луч - 6 – t, tmin, tmax – 3 – node – 2 – tid, stack_top – 2 – На практике удалось уложиться в 16! – Стек: локальная память

2011 kd-деревья RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: Текущий узел: Алгоритм траверса

2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: R Текущий узел: L

2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: R Текущий узел: LL

2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: LLR, R Текущий узел: LLL

2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: LLR, R Текущий узел: LLLR

2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: R Текущий узел: LLRМожно было бы остановиться!

2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: Текущий узел: R

2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: RR Текущий узел: RL

2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: Текущий узел: RR

2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: Текущий узел: RRRКонец, результат: LLR, RRR

2011 BVH деревья Bounding Volume Hierarchy

2011 BVH деревья Bounding Volume Hierarchy

2011 BVH деревья Bounding Volume Hierarchy L R RL RR LR RLRR

2011 BVH деревья Траверс на CPU L R RL RR LR RLRR

2011 BVH деревья Траверс на CPU L R RL RR LR RLRR Стек: L

2011 BVH деревья Траверс на CPU L R RL RR LR RLRR Стек: L

2011 BVH деревья Траверс на CPU L R RL RR LR RLRR Стек: L

2011 BVH деревья Траверс на GPU

2011 BVH деревья Безстековый траверс на GPU

2011 BVH деревья Траверс на GPU L R RL RR LR RLRR

2011 BVH деревья Траверс на GPU L R RL RR LR RLRR

2011 BVH деревья Траверс на GPU L R RL RR LR RLRR

2011 BVH деревья Траверс на GPU L R RL RR LR RLRR

2011 BVH деревья Траверс на GPU L R RL RR LR RLRR

2011 BVH деревья Как делать на CUDA? – Луч – 6 регистров – Бокс – 6 регистров – t_near, t_far, - 2 – nodeOffset, leftOffset, tid – 3 Пересечение луча с боксом – Минимум по всем 6 плоскостям (vmin[0] + rInv.pos[0]) * rInv.dir[0]; vmin vmax t_near t_far vmax

2011 BVH деревья Как делать на CUDA? – Стек на локальной памяти – Локальная память это не так медленно, как может показаться 1. Бесстековый алгоритм – Перебираем массив всегда строго слева направо

2011 Улучшаем производительность Persistent threads __global__ void my_kernel() { for (int warp = 0; int warp < 16; warp++) { int tid = blockDim.x*blockIdx.x + threadIdx.x + warp*step; … // тут код вашего кёрнела }

K сфер, 3 отражения, 1920х1200

2011 Path tracing

2011 Path tracing

2011 Задание 3: растеризация! Программная растеризация на CUDA

2011 Задание 3: растеризация! Программная растеризация на CUDA 3 kernel-а Vertex Shader Rasterizer Pixel shader

2011 Задание 3: растеризация! Простой вариант: – Растеризация 2D треугольников Rasterizer

2011 Задание 3: растеризация! Блок 16x16 пикселей = 256 threads – Или 32x8 triangles in global memory Threads (256) shared memory (256)

2011 Задание 3: растеризация! Блок 16x16 пикселей = 256 threads – Или 32x8 for (int i=0;i

2011 Задание 3: растеризация! Требуется: – Реализация базового алгоритма растеризации – Z-буффер – Вывод картинки куда-нибудь, хотя бы в файл – Замеры времени работы

2011 Задание 3: растеризация! Приветствуются: – Всевозможные оптимизации – Сложные пиксельные и вершинные шейдеры Анимация Освещение по фонгу, кук-торрансу Текстурирование И.т.д. – Вывод с использованим OpenGL