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

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

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

3 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 ?

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

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

6 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);

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

8 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);

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

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

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

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

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

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

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

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

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

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

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

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

21 Другие примитивы 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

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

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

24 Пересечение луча и треугольника 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 уже занято!

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

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

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

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

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

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

31 Архитектура 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:

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

33

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

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

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

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

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

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

40 BVH деревья zBounding Volume Hierarchy

41 BVH деревья zBounding Volume Hierarchy

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

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

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

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

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

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

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

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

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

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

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

53 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

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

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

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

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

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

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

60 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; };

61 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; };

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

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

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

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

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

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

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

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

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

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

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

73 Производителность 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; … // тут код вашего кёрнела }

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

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

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

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

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

79 Ray marching

80

81