Task08 Vladimir Evmenov CSSpace

CMakeLists.txt

@@ -64,7 +64,14 @@ set(LIBRARIES
 set(OPENCL_SOURCES
         src/kernels/cl/aplusb.cl
         src/kernels/cl/ray_tracing_render_brute_force.cl
-        src/kernels/cl/ray_tracing_render_using_lbvh.cl)
+        src/kernels/cl/ray_tracing_render_using_lbvh.cl
+        src/kernels/cl/merge_sort.cl
+        src/kernels/cl/morton_code_generation.cl
+        src/kernels/cl/lbvh_construction.cl
+        src/kernels/cl/lbvh_aabb_generation.cl
+        src/kernels/cl/zeros.cl
+        src/kernels/cl/bigbox_calc.cl
+        )
 
 set(OPENCL_INCLUDES src/kernels/defines.h src/kernels/cl/helpers/rassert.cl src/kernels/cl/camera_helpers.cl src/kernels/cl/geometry_helpers.cl src/kernels/cl/random_helpers.cl src/kernels/shared_structs/aabb_gpu_shared.h src/kernels/shared_structs/bvh_node_gpu_shared.h src/kernels/shared_structs/camera_gpu_shared.h src/kernels/shared_structs/morton_code_gpu_shared.h src/kernels/shared_structs/struct_helpers.h)
 set(OPENCL_DEFINES)

src/debug/debug_bvh.h

@@ -8,7 +8,7 @@
 #include <fstream>
 #include <limits>
 
-#include "scene_reader.h"
+// #include "scene_reader.h"
 #include "../kernels/shared_structs/bvh_node_gpu_shared.h"
 
 namespace debug {

src/kernels/cl/bigbox_calc.cl

@@ -0,0 +1,56 @@
+#ifdef __CLION_IDE__
+#include <libgpu/opencl/cl/clion_defines.cl>
+#endif
+
+#include "helpers/rassert.cl"
+#include "../defines.h"
+
+#include "../shared_structs/camera_gpu_shared.h"
+#include "../shared_structs/bvh_node_gpu_shared.h"
+#include "../shared_structs/aabb_gpu_shared.h"
+#include "../shared_structs/morton_code_gpu_shared.h"
+
+#include "camera_helpers.cl"
+#include "geometry_helpers.cl"
+#include "random_helpers.cl"
+
+// BVH traversal: closest hit along ray
+__kernel void bigbox_calc(    
+    uint                      nfaces,
+    __global const float*     vertices,
+    __global const uint*      faces,
+    __global AABBGPU*         bigbox)
+{
+    uint i = get_global_id(0);
+
+    if (i == 0) {
+        bigbox[0].min_x = +INFINITY;
+        bigbox[0].min_y = +INFINITY;
+        bigbox[0].min_z = +INFINITY;
+        bigbox[0].max_x = -INFINITY;
+        bigbox[0].max_y = -INFINITY;
+        bigbox[0].max_z = -INFINITY;
+
+        for (int j = 0; j < nfaces; j++) {
+            uint3  f  = loadFace(faces, j);
+            float3 v0 = loadVertex(vertices, f.x);
+            float3 v1 = loadVertex(vertices, f.y);
+            float3 v2 = loadVertex(vertices, f.z);
+
+            bigbox[0].min_x = min(bigbox[0].min_x, min(v0.x, min(v1.x, v2.x)));
+            bigbox[0].min_y = min(bigbox[0].min_y, min(v0.y, min(v1.y, v2.y)));
+            bigbox[0].min_z = min(bigbox[0].min_z, min(v0.z, min(v1.z, v2.z)));
+            bigbox[0].max_x = max(bigbox[0].max_x, min(v0.x, min(v1.x, v2.x)));
+            bigbox[0].max_y = max(bigbox[0].max_y, min(v0.y, min(v1.y, v2.y)));
+            bigbox[0].max_z = max(bigbox[0].max_z, min(v0.z, min(v1.z, v2.z)));
+        }
+
+
+        printf("Biggest box : (%f %f %f) -- (%f %f %f)\n", bigbox[0].min_x,
+                                             bigbox[0].min_y,
+                                             bigbox[0].min_z,
+                                             bigbox[0].max_x,
+                                             bigbox[0].max_y,
+                                             bigbox[0].max_z);
+    }
+}

src/kernels/cl/helpers/rassert.cl

@@ -1,6 +1,6 @@
 #include "../../defines.h"
 
-#if RASSERT_ENABLED
+#if 1
     #define rassert(condition, error_code)					                    \
             do {									            \
                     if (!(condition)) {						                    \

src/kernels/cl/lbvh_aabb_generation.cl

@@ -0,0 +1,84 @@
+#ifdef __CLION_IDE__
+#include <libgpu/opencl/cl/clion_defines.cl>
+#endif
+
+#include "helpers/rassert.cl"
+#include "../defines.h"
+
+#include "../shared_structs/camera_gpu_shared.h"
+#include "../shared_structs/bvh_node_gpu_shared.h"
+#include "../shared_structs/aabb_gpu_shared.h"
+
+#include "camera_helpers.cl"
+#include "geometry_helpers.cl"
+#include "random_helpers.cl"
+
+#define INVALID 0xffffffff
+
+void process_tri(
+    __global const uint*       faces,
+    __global const float*      vertices,
+    uint                       tri_id,
+    __global BVHNodeGPU*       node
+) {
+    uint3  f  = loadFace(faces, tri_id);
+    float3 v0 = loadVertex(vertices, f.x);
+    float3 v1 = loadVertex(vertices, f.y);
+    float3 v2 = loadVertex(vertices, f.z);
+
+    node->aabb.min_x = min(v0.x, min(v1.x, v2.x));
+    node->aabb.min_y = min(v0.y, min(v1.y, v2.y));
+    node->aabb.min_z = min(v0.z, min(v1.z, v2.z));
+
+    node->aabb.max_x = max(v0.x, max(v1.x, v2.x));
+    node->aabb.max_y = max(v0.y, max(v1.y, v2.y));
+    node->aabb.max_z = max(v0.z, max(v1.z, v2.z));
+
+    node->leftChildIndex = INVALID;
+    node->rightChildIndex = INVALID;
+}
+
+__kernel void lbvh_aabb_generation(
+    __global const uint*       morton_codes,
+    __global uint*             face_indexes,
+    __global const uint*       faces,
+    __global const float*      vertices,
+    __global BVHNodeGPU*       nodes,
+    __global int*              parent,
+    __global int*              counter,
+    uint                       nfaces,
+    int                        calc_leafs
+   )
+{
+    int i = get_global_id(0);
+    const int leafStart = (int)nfaces - 1; 
+
+    if (calc_leafs) { 
+        if (i >= leafStart && i < nfaces + leafStart) {
+            int face_id = face_indexes[i - leafStart];
+            process_tri(faces, vertices, face_id - leafStart, &nodes[i]);
+            // printf("%d <-- %d\n", i, parent[i]);
+            atomic_add(&counter[parent[i]], 1);
+            face_indexes[i - leafStart] -= leafStart;
+        }
+    }
+    else {
+        if (i < leafStart && counter[i] >= 2) {
+            BVHNodeGPU l = nodes[nodes[i].leftChildIndex];
+            BVHNodeGPU r = nodes[nodes[i].rightChildIndex];
+
+            nodes[i].aabb.min_x = min(l.aabb.min_x, r.aabb.min_x);
+            nodes[i].aabb.min_y = min(l.aabb.min_y, r.aabb.min_y);
+            nodes[i].aabb.min_z = min(l.aabb.min_z, r.aabb.min_z);
+
+            nodes[i].aabb.max_x = max(l.aabb.max_x, r.aabb.max_x);
+            nodes[i].aabb.max_y = max(l.aabb.max_y, r.aabb.max_y);
+            nodes[i].aabb.max_z = max(l.aabb.max_z, r.aabb.max_z);
+
+            if (parent[i] != -1)
+                atomic_add(&counter[parent[i]], 1);
+        }
+    }
+}
+
+

src/kernels/cl/lbvh_construction.cl

@@ -0,0 +1,128 @@
+#ifdef __CLION_IDE__
+#include <libgpu/opencl/cl/clion_defines.cl>
+#endif
+
+#include "helpers/rassert.cl"
+#include "../defines.h"
+
+#include "../shared_structs/camera_gpu_shared.h"
+#include "../shared_structs/bvh_node_gpu_shared.h"
+#include "../shared_structs/aabb_gpu_shared.h"
+
+#include "camera_helpers.cl"
+#include "geometry_helpers.cl"
+#include "random_helpers.cl"
+
+inline int common_pref_len(
+    __global const uint*       morton_codes,
+    int l,
+    int r, 
+    int nfaces
+) 
+{
+    if (l > r) {
+        int tmp = r;
+        r = l;
+        l = tmp;
+    }
+
+    rassert(r == l + 1, 4358904);
+
+    if (l < 0)
+        return -1;
+
+    if (r >= nfaces)
+        return -1;
+
+    int bit = 31;
+    while (bit >= 0 && (((morton_codes[l] + l) >> bit) & 1) == (((morton_codes[r] + r) >> bit) & 1))
+        bit--;
+
+    return 32 - bit - 1;
+}
+
+__kernel void lbvh_construction(
+    __global const uint*       morton_codes,
+    __global uint*             face_indexes,
+    __global const uint*       faces,
+    __global const float*      vertices,
+    __global BVHNodeGPU*       nodes,
+    __global int*              parent,
+    uint                       nfaces
+   )
+{
+    int i = get_global_id(0);
+    const int leafStart = (int)nfaces - 1; 
+
+    int d, pref_len, j, min_pref, min_pref_ind;
+    if (i < leafStart) {
+        d = 0, pref_len = -1;
+
+        if (i == 0)
+            d = +1;
+        else
+            d = (common_pref_len(morton_codes, i, i + 1, nfaces) > common_pref_len(morton_codes, i - 1, i, nfaces) ? +1 : -1);
+
+        pref_len = common_pref_len(morton_codes, i, i - d, nfaces);
+
+        j = i;
+        min_pref = 100, min_pref_ind = -1;
+        while (j + d < nfaces &&
+               j + d >= 0     &&
+               pref_len < common_pref_len(morton_codes, j, j + d, nfaces)) { 
+            if (min_pref > common_pref_len(morton_codes, j, j + d, nfaces)) {
+                min_pref = common_pref_len(morton_codes, j, j + d, nfaces);
+                min_pref_ind = j;
+            }
+            j += d;
+        }
+        // min_pref_ind, min_pref_ind + d - max on segment
+
+        if (i == 0) {
+            parent[0] = -1;
+        }
+        if (j == i + d) {
+            nodes[i].leftChildIndex  = face_indexes[min(i, j)];
+            nodes[i].rightChildIndex = face_indexes[max(i, j)];
+        }
+        else if (i == min_pref_ind) {
+            if (d == +1) {
+                nodes[i].leftChildIndex  = face_indexes[i];
+                nodes[i].rightChildIndex = i + d;
+            }
+            if (d == -1) {
+                nodes[i].leftChildIndex  = i + d;
+                nodes[i].rightChildIndex = face_indexes[i];
+            }
+        }
+        else if (min_pref_ind + d == j) { 
+            if (d == +1) {
+                nodes[i].leftChildIndex  = j - d;
+                nodes[i].rightChildIndex = face_indexes[j];
+            }
+            if (d == -1) {
+                nodes[i].leftChildIndex  = face_indexes[j];
+                nodes[i].rightChildIndex = j - d;
+            }
+        }
+        else {
+            nodes[i].leftChildIndex  = min(min_pref_ind, min_pref_ind + d);
+            nodes[i].rightChildIndex = max(min_pref_ind, min_pref_ind + d);
+        }
+    }
+
+    barrier(CLK_GLOBAL_MEM_FENCE);
+
+    if (i < leafStart) { 
+        parent[nodes[i].leftChildIndex] = i;
+        parent[nodes[i].rightChildIndex] = i;
+
+        // atomic_add(&counter[nodes[i].leftChildIndex], 1);
+        // atomic_add(&counter[nodes[i].rightChildIndex], 1);
+    }
+
+    barrier(CLK_GLOBAL_MEM_FENCE);
+
+    if (i == 0)
+        printf("tree build done\n");
+}