ТРАССИРОВКА ЛУЧЕЙ НА CUDA Докладчик: Фролов В.А. (ВМиК МГУ, G&M Lab; Nvidia)Фролов В.А. (ВМиК МГУ, G&M Lab; Nvidia) Научный руководитель: Игнатенко А.В. (ВМиК МГУ, G&M Lab) Игнатенко А.В. (ВМиК МГУ, G&M Lab) Лекторы: Боресков А.В. (ВМиК МГУ) Харламов А.А. (ВМиК МГУ, Nvidia)

План z OpenGL interoperability z RT – что, зачем и как? z Пересечение луча и треугольника z Ускоряющие структуры z Организация трассировки лучей z Ray marching z Задание

OpenGL interop z OpenGL 2.0 y VBO, PBO z OpenGL 3.0 y VBO, PBO y texture y renderbuffer z OpenGL 4.0 y ?

OpenGL 2.0 interop z Pack y glReadPixels (…) y glGetTexImage (…) z Unpack y glDrawPixels (…) y glTexImage2D (…) y glTexSubImage2D (…)

OpenGL 2.0 interop Pointer (CUDA) uint* gpu_values; PBO (CUDA) GLuint buffer_id; register, unregister map, unmap my_kernel >>(gpu_values);

OpenGL 2.0 interop glGenBuffers(1, &pixelBuffer); glBindBuffer (GL_PIXEL_UNPACK_BUFFER, pixelBuffer); glBufferData(GL_PIXEL_UNPACK_BUFFER, w*h*sizeof(int), screen_buffer, GL_STATIC_DRAW); cudaGLRegisterBufferObject (pixelBuffer); cudaGLMapBufferObject ((void**)&my_pointer, pixelBuffer); my_kernel >>(my_pointer); cudaGLUnmapBufferObject(pixelBuffer); cudaGLUnregisterBufferObject(pixelBuffer); glBindTexture(GL_TEXTURE_2D, tex_id); glTexImage2D (…, 0, …, w, h, 0, …,..., (GLvoid*)0);

OpenGL 3.0 interop z Старый API y сохраняется z Новый API Pointer (CUDA) uint* gpu_values; CUDA Array; cudaArray *arr; cudaGraphicsResource FBO (renderbuffer)

OpenGL 3.0 interop GLuint tex_id; glGenTextures(1, & tex_id); … struct cudaGraphicsResource *tex_resource; cudaGraphicsGLRegisterImage (&tex_resource, tex_id, GL_TEXTURE_2D, cudaGraphicsMapFlagsReadOnly); cudaGraphicsMapResources (1, & tex_resource, 0); cudaArray *array_ptr; cudaGraphicsSubResourceGetMappedArray (&array_ptr, tex_resource, 0, 0); cudaMemcpyToArray(array_ptr,0, 0,dest,size,cudaMemcpyDeviceToDevice); cudaGraphicsUnmapResources(1, & tex_resource, 0);

Ray Tracing zФотореалистичный синтез изображений zPOV-Ray

Ray Tracing zФотореалистичный синтез изображений zPOV-Ray

Real Time Ray Tracing z Скорость в ущерб качеству

Ray Tracing z Точность y Path tracing y Фотонные карты y Распределенная трассировка лучей (стохастическая) z Скорость y Обратная трассировка лучей y Растеризация + обратная трассировка лучей

Обратная трассировка лучей z Алгоритм y Первичные, теневые, отраженные лучи

Ray Tracing z Представление 3D объектов y Аналитическое y Меши из треугольников

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

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

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

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

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

Пересечение луча и треугольника zПростой вариант y Операции (* : 39, +/- : 53, / : 1) x тактов zОптимизированный вариант y Операции (* : 23, +/- : 24, / : 1) x такта z Как считали нижнюю оценку? y использование mad вместо mul и add y 4*(N_mul + |N_add - N_mul|)

Другие примитивы z Бокс – это 6 плоскостей y(vmin.x - r.pos.x) / r.dir.x; y(vmin.x + rInv.pos.x) * rInv.dir.x; y 6 add и 6 mul == 12 mad, 48 тактов z Сфера y ~13 mad + sqrtf == = 84 такта y меньше ветвлений y Иерархия из сфер не лучше иерархии из боксов vmin vmax t_near t_far vmax

Occupancy zОчень быстрое переключение zВлияет: yРегистры yshared память yРазмер блока

Occupancy z Регистры y 8192 регистра на SM y Блоки по 8x8 нитей y 128 регистров на нить x nvcc не дает столько регистров, почему? x рег

Пересечение луча и треугольника z Регистры y 6 регистров на луч y 9 регистров на вершины y 3 регистра на (t, u, v) y 1 регистр на triNum y 1 на счетчик в цикле y 1 как минимум на tid y 2 на min_t и min_id z 23 уже занято!

Пересечение луча и треугольника zUnit 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;

Пересечение луча и треугольника zUnit 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

Пересечение луча и треугольника zUnit test float4 row2 = tex1Dfetch(tex, triAddress+0); (oz,dz) = … float t = -oz/dz; if (t = tTriangle) continue; float4 row0 = tex1Dfetch(tex, triAddress+1); (ox, dx) = … float u = ox + t*dx; if (u < 0.f) continue; float4 row1 = tex1Dfetch(tex, triAddress+2); (oy dy) = … float v = oy + t*dy; if(v >= 0 && (u+v)

Экономия регистров zДва подхода к экономии регистров y uber kernel y separate kernel

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

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

Архитектура zКак хранить геометрию? struct Triangle { uint v[3]; uint index; }; struct Triangle { float4 v[3]; // strore undex in v[2].w }; struct Vertex { float3 pos[3]; float3 norm[3]; float2 texCoord; uint materialIndex; }; Vertex array: Index array:

Архитектура zКак хранить геометрию? kd tree Triangle list indices vertices intersection

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

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

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

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

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

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

BVH деревья zBounding Volume Hierarchy

BVH деревья zBounding Volume Hierarchy

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

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

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

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

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

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

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

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

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

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

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

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

BVH деревья z Как делать на CUDA? y 24 mad-а покрывают латентность текстурной памяти 1. Стек на локальной памяти y Локальная память это не так медленно, как может показаться 2. Бесстековый алгоритм y Перебираем массив всегда строго слева направо

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

kd-деревья struct KdTreeNode { float split; uint leftOffset: 29; uint splitAxis: 2; uint leaf: 1; };

kd-деревья L R R L struct KdTreeNode { float split; uint leftOffset: 29; uint splitAxis: 2; uint leaf: 1; };

kd-деревья R L LLLRRLRR LL LRRL RR struct KdTreeNode { float split; uint leftOffset: 29; uint splitAxis: 2; uint leaf: 1; };

kd-деревья R L LLLRRLRR LLLLLR LRL LRR RL RRL RRR struct KdTreeNode { float split; uint leftOffset: 29; uint splitAxis: 2; uint leaf: 1; };

kd-деревья R L LLLRRLRR RL RRR RRL LRR LLRLLLR LLLL LRLL LRLR struct KdTreeNode { float split; uint leftOffset: 29; uint splitAxis: 2; uint leaf: 1; };

kd-деревья R L LLLRRLRR RL RRR RRL LRR LLR LLLL LLLR LRLL LRLRLLRLRR struct KdTreeNode { float split; uint leftOffset: 29; uint splitAxis: 2; uint leaf: 1; };

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

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

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

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

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

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

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

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

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

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

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

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

Производителность z Conference Room (281K треугольников) z~40M лучей в секунду на GTX260

Производителность z Стулья (1.5M треугольников)

Производительность

kd-tree vs BVH на CUDA z BVH со стеком на локальной памяти y Покрывается латентность текстурной памяти y Меньше глубина y Лишние плоскости z kd-tree y Экономит регистры y Можно эффективнее задействовать кэш? y 1 mad и пара ветвлений на одну tex1Dfetch

Резюме zOpenGL interop y текстуры, PBO, VBO, FBO zOccupancy y Оптимизации по локальным переменным y Uber kernel y Separate kernel zPersistent threads

Ray marching