Скачать презентацию
Идет загрузка презентации. Пожалуйста, подождите
Презентация была опубликована 10 лет назад пользователемГерман Прозоров
1 GPU Ray Tracing Vladimir Frolov (Nvidia, MSU, Keldysh Institute of Applied Math)
2 2011 План Трассировка лучей и растеризация – Определение видимости – Global illumination Ускоряющие структуры Особенности трассировки лучей на GPU Узкие места Работа мультипроцессора на низком уровне – CPU vs GPU Оптимизация и отладка неоднородного кода
3 2011 Растеризация vs трассировка
4 2011 Растеризация vs трассировка Растеризация – Обработка по 1 треугольнику – Трудно параллелится – Эквивалентна задаче построения дерева или сортировке
5 2011 Трассировка лучей Алгоритм
6 2011 Трассировка лучей Представление 3D объектов – Аналитическое – Меши из треугольников
7 2011 Трассировка лучей Поверхность задана как массив треугольников Узкое место – поиск пересечения луча с поверхностью – треугольников – лучей – => операций – log(N) k * 10 6 (k~[1..2])
8 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
9 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 | < ε
10 2011 Пересечение луча и треугольника Оптимизированный вариант – Барицентрические координаты u := u/S, v := v/S, t1 := t1/S t1 = 1 – u - v
11 2011 Пересечение луча и треугольника Оптимизированный вариант
12 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
13 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;
14 2011 Пересечение луча и треугольника Простой вариант – Операции (* : 39, +/- : 53, / : 1) Оптимизированный вариант – Операции (* : 23, +/- : 24, / : 1) Юнит тест – Операции (*: 20, + : 20, / : 1) – Экономит регистры GPU
15 2011 Регистры и локальные переменные float3 RayTrace (const Ray& ray) { for (int depth=0 ;depth
16 2011 Занятость мультипроцессора Регистры – GF100: maxreg = 32768, maxwarps = 48 – block size = 256 рег
17 2011 GPU vs CPU – Локальные переменные CPU : стэк GPU : регистры – Задержки в конвейере (Pipeline hazards) CPU : Out Of Order GPU : Hyperthreading – Латентность памяти CPU: огромный кэш GPU: сокрытие латентности за счет параллелизма
18 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
19 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 **
20 2011 GPU vs CPU Латентность памяти
21 2011 Занятость мультипроцессора Регистры
22 2011 Занятость мультипроцессора Регистры – G80: maxreg = 8192, maxwarps = 24 – GT200: maxreg = 16384, maxwarps = 32 – GF1**: maxreg = 32768, maxwarps = 48
23 2011 Занятость мультипроцессора Регистры – GF100: maxreg = 32768, maxwarps = 48 – block size = 256 рег
24 2011 Занятость мультипроцессора Разделяемая память – GF100: size per SM = bytes, maxwarps = 48 – block size = 256, 8 warps – байт на блок (64 байта на тред) – байт на блок (48 байт на тред) – 8192 байта на блок (32 байта на тред)
25 2011 Занятость мультипроцессора CUDA occupancy calculator
26 2011 Занятость мультипроцессора Текстурный кэш
27 2011 Пересечение луча и треугольника Оптимизированный вариант – регистра Unit test – регистра И это только пересечение луча и треугольника!
28 2011 Регистры и локальные переменные float3 RayTrace (const Ray& ray) { for (int depth=0 ;depth
29 2011 Экономим регистры Два подхода: – Uber kernel – Separate kernel
30 2011 Uber kernel vs Separate kernel Uber kernel zРекурсия (+) zЛокальность данных (+) zНепрерывное выполнение (+) zТрудно профилировать (-) zТрудно оптимизировать (-) zСамая сложная часть лимитирует все ядро (-) Separate kernel z Нет рекурсии(-) z Обращения к DRAM (-) z Латентность остановки/запуска (-) z Легче отлаживать (+) z Легко профилировать (+) z Легко оптимизировать (+) z Гибкость (+)
31 2011 Организация рей трейсера Ядро пересечений – регистра Но это только пересечения Нужно разбить алгоритм трассировки на несколько ядер
32 2011 Ускоряющие структуры regular grid kd-tree BVH
33 2011 Регулярная сетка
34 2011 Регулярная сетка traversal if (tMaxX
35 2011 Регулярная сетка Преимущества – Просто и быстро строится – Простой алгоритм траверса Недостатки – Плохо справляется с пустым пространством – Требует много памяти – Много повторных пересечений – отвратительно разбивает геометрию Только для небольших сцен (1-50K)
36 2011 Регулярная сетка Почему сетка плохо разбивает геометрию? zПеребрали 15 вокселей z7 раз посчитали пересечение с одним и тем же треугольником!
37 2011 Иерархическая сетка Небольшое число вокселей Рекурсивно разбиваем воксели в местах с плотной геометрией
38 2011 Иерархическая сетка + Решает проблему чайника на стадионе - Переход между узлами вычислительно сложен - 12 доп. регистров как минимум - Нужно устранять рекурсию
39 2011 kd-деревья
40 2011 kd-деревья L R R L
41 2011 kd-деревья R L LLLRRLRR LL LRRL RR
42 2011 kd-деревья R L LLLRRLRR LLLLLR LRL LRR RL RRL RRR
43 2011 kd-деревья R L LLLRRLRR RL RRR RRL LRR LLRLLLR LLLL LRLL LRLR
44 2011 kd-деревья Алгоритм траверса Регистры – 13 min: – луч - 6 – t, tmin, tmax – 3 – node – 2 – tid, stack_top – 2 – На практике удалось уложиться в 16! – Стек: локальная память
45 2011 kd-деревья RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: Текущий узел: Алгоритм траверса
46 2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: R Текущий узел: L
47 2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: R Текущий узел: LL
48 2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: LLR, R Текущий узел: LLL
49 2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: LLR, R Текущий узел: LLLR
50 2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: R Текущий узел: LLRМожно было бы остановиться!
51 2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: Текущий узел: R
52 2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: RR Текущий узел: RL
53 2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: Текущий узел: RR
54 2011 kd-деревья Алгоритм траверса RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR Стек: Текущий узел: RRRКонец, результат: LLR, RRR
55 2011 BVH деревья Bounding Volume Hierarchy
56 2011 BVH деревья Bounding Volume Hierarchy
57 2011 BVH деревья Bounding Volume Hierarchy L R RL RR LR RLRR
58 2011 BVH деревья Траверс на CPU L R RL RR LR RLRR
59 2011 BVH деревья Траверс на CPU L R RL RR LR RLRR Стек: L
60 2011 BVH деревья Траверс на CPU L R RL RR LR RLRR Стек: L
61 2011 BVH деревья Траверс на CPU L R RL RR LR RLRR Стек: L
62 2011 BVH деревья Траверс на GPU
63 2011 BVH деревья Безстековый траверс на GPU
64 2011 BVH деревья Траверс на GPU L R RL RR LR RLRR
65 2011 BVH деревья Траверс на GPU L R RL RR LR RLRR
66 2011 BVH деревья Траверс на GPU L R RL RR LR RLRR
67 2011 BVH деревья Траверс на GPU L R RL RR LR RLRR
68 2011 BVH деревья Траверс на GPU L R RL RR LR RLRR
69 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
70 2011 BVH деревья Как делать на CUDA? – Стек на локальной памяти – Локальная память это не так медленно, как может показаться 1. Бесстековый алгоритм – Перебираем массив всегда строго слева направо
71 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; … // тут код вашего кёрнела }
72 K сфер, 3 отражения, 1920х1200
73 2011 Path tracing
74 2011 Path tracing
75 2011 Задание 3: растеризация! Программная растеризация на CUDA
76 2011 Задание 3: растеризация! Программная растеризация на CUDA 3 kernel-а Vertex Shader Rasterizer Pixel shader
77 2011 Задание 3: растеризация! Простой вариант: – Растеризация 2D треугольников Rasterizer
78 2011 Задание 3: растеризация! Блок 16x16 пикселей = 256 threads – Или 32x8 triangles in global memory Threads (256) shared memory (256)
79 2011 Задание 3: растеризация! Блок 16x16 пикселей = 256 threads – Или 32x8 for (int i=0;i
80 2011 Задание 3: растеризация! Требуется: – Реализация базового алгоритма растеризации – Z-буффер – Вывод картинки куда-нибудь, хотя бы в файл – Замеры времени работы
81 2011 Задание 3: растеризация! Приветствуются: – Всевозможные оптимизации – Сложные пиксельные и вершинные шейдеры Анимация Освещение по фонгу, кук-торрансу Текстурирование И.т.д. – Вывод с использованим OpenGL
Еще похожие презентации в нашем архиве:
© 2024 MyShared Inc.
All rights reserved.