Skip to content
GitLab
Commit ffe4df92 authored by 's avatar
Browse files

Added LS generation with center raycasts through a base bvh

parent 8038a3bd
Hide whitespace changes
Inline Side-by-side
src/executables/FullExecutable/Application.cpp
View file @ ffe4df92
......@@ -64,12 +64,12 @@ namespace glare
{
m_collector->collect();
const auto &mesh_collector = m_collector->meshCollectors().begin()->second;
math::UniformBoundsSubdivision subdivision(mesh_collector.bounds(), 20);
/*math::UniformBoundsSubdivision subdivision(mesh_collector.bounds(), 20);
size_t lines = advanced::lsmath::calculateLineCount(subdivision);
Log_Info << subdivision.subdivisions.x << "; " << subdivision.subdivisions.y << "; " << subdivision.subdivisions.z;
Log_Info << subdivision.subdivisions.x << "; " << subdivision.subdivisions.y << "; " << subdivision.subdivisions.z;*/
advanced::lsmath::generateLines(mesh_collector, 20);
quitPromptDefault(0);
}
......
src/libraries/advanced/BVH.cpp deleted 100644 → 0
View file @ 8038a3bd
#include "BVH.h"
#include <engine/rendering/Shader.h>
#include <engine/EngineState.h>
#include <engine/Math.h>
namespace glare
{
namespace advanced
{
BVH::BVH()
{
}
void BVH::initialize(const std::shared_ptr<SceneCollector> &collector)
{
m_scene_collector = collector;
m_mesh_index_offset = 0;
build();
buildBuffers();
}
BVH::~BVH()
{
}
std::shared_ptr<core::Buffer<gl::BufferType::eShaderStorage>> BVH::getTrianglesBuffer() const
{
return m_triangles_buffer;
}
std::shared_ptr<core::Buffer<gl::BufferType::eShaderStorage>> BVH::getBVHNodesBuffer() const
{
return m_bvh_nodes_buffer;
}
void BVH::buildBuffers()
{
if (!m_triangles_buffer)
m_triangles_buffer = std::make_shared<core::Buffer<gl::BufferType::eShaderStorage>>();
if (!m_bvh_nodes_buffer)
m_bvh_nodes_buffer = std::make_shared<core::Buffer<gl::BufferType::eShaderStorage>>();
m_triangles_buffer->setData(m_sorted_triangles, gl::BufferUsage::eDynamicRead);
m_bvh_nodes_buffer->setData(m_nodes, gl::BufferUsage::eDynamicRead);
}
void BVH::build()
{
m_sorted_triangles = std::vector<Triangle>(m_scene_collector->triangles());
m_nodes.resize(2 * m_sorted_triangles.size());
m_triangle_centroids.resize(m_sorted_triangles.size());
m_centroid_boxes.resize(m_sorted_triangles.size());
m_safe_centroid_boxes.resize(m_sorted_triangles.size());
m_ranges.clear();
m_node_counter = 0;
math::Bounds centroid_box;
#pragma omp parallel for schedule(static) //static, because workload is the same for every iteration
for (int i = 0; i < int(m_sorted_triangles.size()); ++i)
{
// [] is not necessarily range-checked, whereas at() has to be range checked by standard
const Triangle& triangle = m_sorted_triangles[i];
//calculate triangle centroid
glm::vec4 centroid = (m_scene_collector->vertices()[triangle.indices[0]].position
+ m_scene_collector->vertices()[triangle.indices[1]].position
+ m_scene_collector->vertices()[triangle.indices[2]].position)
/ 3.f;
m_triangle_centroids[i] = centroid;
centroid_box.expand(centroid);
}
m_centroid_boxes[0] = centroid_box;
TriangleIndexRange range;
range.start = 0;
range.end = int(m_sorted_triangles.size() - 1);
range.parent = -1;
m_ranges.push_back(range);
m_ranges.resize(m_sorted_triangles.size());
m_safe_ranges.resize(m_sorted_triangles.size());
m_kernel_counter = 1;
m_bvh_depth = 0;
while (m_kernel_counter > 0)
{
m_bvh_depth++;
m_buffer_position = 0;
const int kernel_counter = m_kernel_counter;
#pragma omp parallel for schedule(dynamic) //dynamic, because construct takes different time dynamically
for (int global_id = 0; global_id < kernel_counter; global_id++)
{
split(global_id);
}
m_node_counter += kernel_counter;
//secure reads
if (m_kernel_counter > 0)
{
#pragma omp parallel for schedule(static) //static, because workload is the same for every iteration
for (int global_id = 0; global_id < m_kernel_counter; global_id++)
{
std::swap(m_ranges[global_id], m_safe_ranges[global_id]);
std::swap(m_centroid_boxes[global_id], m_safe_centroid_boxes[global_id]);
}
}
}
m_centroid_boxes.clear();
m_safe_centroid_boxes.clear();
m_ranges.clear();
m_safe_ranges.clear();
m_triangle_centroids.clear();
}
bool BVH::shouldSplit(const math::Bounds &centroid)
{
glm::vec4 diff = centroid[1] - centroid[0];
return diff.x > 0.f && diff.y > 0.f && diff.z > 0.f;
}
void BVH::addAxisAlignedBoundingBox(BVHNode& node, int start, int end) const
{
glm::vec4 best_min = glm::vec4(std::numeric_limits<float>::max(), std::numeric_limits<float>::max(), std::numeric_limits<float>::max(), 1.0f);
glm::vec4 best_max = glm::vec4(std::numeric_limits<float>::lowest(), std::numeric_limits<float>::lowest(), std::numeric_limits<float>::lowest(), 1.0f);
for (int id = start; id <= end; id++)
{
Triangle t = m_sorted_triangles.at(id);
glm::vec4 a = m_scene_collector->vertices().at(t.indices[0]).position;
glm::vec4 b = m_scene_collector->vertices().at(t.indices[1]).position;
glm::vec4 c = m_scene_collector->vertices().at(t.indices[2]).position;
glm::vec4 current_min = glm::min(glm::min(a, b), c);
glm::vec4 current_max = glm::max(glm::max(a, b), c);
best_min = glm::min(current_min, best_min);
best_max = glm::max(current_max, best_max);
}
node.bounds.min = best_min;
node.bounds.max = best_max;
}
void BVH::split(uint32_t global_id)
{
static const float flt_max = std::numeric_limits<float>::max();
static const float flt_min = std::numeric_limits<float>::lowest();
int start = m_ranges[global_id].start;
int end = m_ranges[global_id].end;
int parent = m_ranges[global_id].parent;
//centboxes cover the centroids of the triangles, not the whole objects
math::Bounds cebo = m_centroid_boxes[global_id];
bool splitting = parent < 0 || shouldSplit(cebo);
// since the thread will now be processed the value of the kernel_counter is decreased
--m_kernel_counter;
int leaf_threshold = 1;
BVHNode node;
// current_node_id is the position in the node buffer in which the new node is stored
int current_node_id = m_node_counter + global_id;
int diff = end - start;
if ((diff > leaf_threshold) && splitting)
{ //inner node
//pos is the position in which the new ranges will be stored
int pos = m_buffer_position.fetch_add(1);
TriangleIndexRange left_range, right_range;
math::Bounds left_centbox, right_centbox;
// since the range is divided in two parts more threads can process the construction.
// therefore the kernel_counter is increased by 2
m_kernel_counter.fetch_add(2);
// 1. the AABB which contains the range's triangle information is added for later processing (traversal)
addAxisAlignedBoundingBox(node, start, end);
// 0 is the flag for inner node in the node buffer
node.type = 0;
// 2. the split plane which causes the least cost is determined by iterating over the
// 3 axes (since it's axis aligned)
float best = flt_max;
int best_axis = 0;
int best_plane = 0;
AxisConstants axis_constants[3];
Bin bins[3][16];
for (int axis = 0; axis < 3; ++axis)
{
axis_constants[axis].cbmin = cebo[0][axis];
axis_constants[axis].cbmax = cebo[1][axis];
axis_constants[axis].cbdiff = axis_constants[axis].cbmax - axis_constants[axis].cbmin;
axis_constants[axis].k = (m_bin_count * (1 - m_bin_epsilon)) / axis_constants[axis].cbdiff;
//Initialize Bins with default values.
for (int32_t i = 0; i < m_bin_count; ++i)
{
bins[axis][i].bounds.max = glm::vec4(flt_min);
bins[axis][i].bounds.min = glm::vec4(flt_max);
bins[axis][i].object_count = 0;
}
}
// This is the thing that takes too damn long
for (int id = start; id <= end; ++id)
{
glm::vec4 centroidAxis = m_triangle_centroids[id];
for (int axis = 0; axis < 3; ++axis)
{
int bin_id = int(axis_constants[axis].k * (centroidAxis[axis] - axis_constants[axis].cbmin));
bins[axis][bin_id].bounds.expand(centroidAxis);
++(bins[axis][bin_id].object_count);
}
}
for (int axis = 0; axis < 3; ++axis)
{
Bin left_bounds[15];
//Again, initialize all bounds.
for (int i = 0; i < m_plane_count; ++i)
{
left_bounds[i].bounds.max = glm::vec4(flt_min);
left_bounds[i].bounds.min = glm::vec4(flt_max);
left_bounds[i].object_count = 0;
}
left_bounds[0].bounds.expand(bins[axis][0].bounds);
left_bounds[0].object_count = bins[axis][0].object_count;
for (int plane = 1; plane < m_plane_count; ++plane)
{
left_bounds[plane].bounds.expand(left_bounds[plane - 1].bounds);
left_bounds[plane].bounds.expand(bins[axis][plane].bounds);
left_bounds[plane].object_count = left_bounds[plane - 1].object_count + bins[axis][plane].object_count;
}
Bin right_bounds[15];
//Again and again, initialize bounds.
for (int i = 0; i < m_plane_count; ++i)
{
right_bounds[i].bounds.max = glm::vec4(flt_min);
right_bounds[i].bounds.min = glm::vec4(flt_max);
right_bounds[i].object_count = 0;
}
for (int plane = m_plane_count - 1; plane >= 0; --plane)
{
right_bounds[plane].bounds.expand(bins[axis][plane + 1].bounds);
right_bounds[plane].object_count = bins[axis][plane + 1].object_count;
if (plane != m_plane_count - 1)
{
right_bounds[plane].bounds.expand(right_bounds[plane + 1].bounds);
right_bounds[plane].object_count += right_bounds[plane + 1].object_count;
}
float surface_area_left = left_bounds[plane].bounds.surface();
float surface_area_right = right_bounds[plane].bounds.surface();
float cost = surface_area_left * left_bounds[plane].object_count + surface_area_right * right_bounds[plane].object_count;
if (cost < best)
{
best = cost;
best_axis = axis;
best_plane = plane;
left_centbox[0] = left_bounds[plane].bounds.min;
left_centbox[1] = left_bounds[plane].bounds.max;
right_centbox[0] = right_bounds[plane].bounds.min;
right_centbox[1] = right_bounds[plane].bounds.max;
}
}
}
//3. sorting the triangle ids to their correct ranges
const float cbmin = axis_constants[best_axis].cbmin;
const float k = axis_constants[best_axis].k;
int left = start;
int right = end;
bool left_stopped = false;
bool right_stopped = false;
while (left < right)
{
if (!left_stopped)
{
float centroid_best_axis_left = m_triangle_centroids[left][best_axis];
int bin_id_left = int(k * (centroid_best_axis_left - cbmin));
if (bin_id_left > best_plane) { left_stopped = true; }
else { ++left; }
}
if (!right_stopped)
{
float centroid_best_axis_right = m_triangle_centroids[right][best_axis];
int bin_id_right = int(k * (centroid_best_axis_right - cbmin));
if (bin_id_right <= best_plane) { right_stopped = true; }
else { --right; }
}
if (left_stopped && right_stopped)
{
//swap triangles and centroids
std::swap(m_sorted_triangles[left], m_sorted_triangles[right]);
std::swap(m_triangle_centroids[left], m_triangle_centroids[right]);
left_stopped = false;
right_stopped = false;
++left;
--right;
}
}
//determing left_range, right_range for the following threads
left_range.start = start;
right_range.end = end;
if (left > right)
{
left_range.end = right;
right_range.start = left;
}
else
{
if (left_stopped)
{
left_range.end = left - 1;
right_range.start = left;
}
else if (right_stopped)
{
left_range.end = right;
right_range.start = right + 1;
}
else
{
float centroid_best_axis_left = m_triangle_centroids[left][best_axis];
int bin_id_left = int(k * (centroid_best_axis_left - cbmin));
if (bin_id_left > best_plane)
{
left_range.end = left - 1;
right_range.start = left;
}
else
{
left_range.end = left;
right_range.start = left + 1;
}
}
}
left_range.parent = current_node_id;
right_range.parent = current_node_id;
//writing the id_ranges and centboxes to the buffers for processing in following threads
m_safe_ranges[pos * 2] = left_range;
m_safe_ranges[pos * 2 + 1] = right_range;
m_safe_centroid_boxes[pos * 2] = left_centbox;
m_safe_centroid_boxes[pos * 2 + 1] = right_centbox;
node.parent = parent;
//5. writing to the node_buffer for traversal in another kernel
m_nodes[current_node_id] = node;
if (parent > -1)
{
if (current_node_id % 2 == 0) { m_nodes[parent].right_idx = current_node_id; }
else { m_nodes[parent].left_idx = current_node_id; }
}
}
else //leaf node
{
addAxisAlignedBoundingBox(node, start, end);
node.type = 1;
node.left_idx = start;
node.right_idx = end;
node.parent = parent;
m_nodes[current_node_id] = node;
if (parent > -1)
{
if (current_node_id % 2 == 0) { m_nodes[parent].right_idx = current_node_id; }
else { m_nodes[parent].left_idx = current_node_id; }
}
++m_leaf_counter;
}
}
}
}
\ No newline at end of file
src/libraries/advanced/BVH.h deleted 100644 → 0
View file @ 8038a3bd
#ifndef __RAYTRACER_BVH_H
#define __RAYTRACER_BVH_H
#include <vector>
#include <engine/rendering/GraphNode.h>
#include <engine/rendering/MeshRenderer.h>
#include <atomic>
#include <engine/Math.h>
#include "SceneCollector.h"
namespace glare
{
namespace advanced
{
struct BVHNode
{
math::Bounds bounds;
uint32_t type;
uint32_t left_idx;
uint32_t right_idx;
uint32_t parent;
};
struct TriangleIndexRange
{
uint32_t start;
uint32_t end;
uint32_t parent;
uint32_t padding;
};
struct AxisConstants
{
float cbmin;
float cbmax;
float cbdiff;
float k;
};
struct Bin
{
math::Bounds bounds;
int object_count;
};
class BVH {
public:
BVH();
~BVH();
void initialize(const std::shared_ptr<SceneCollector> &collector);
std::shared_ptr<core::Buffer<gl::BufferType::eShaderStorage>> getTrianglesBuffer() const;
std::shared_ptr<core::Buffer<gl::BufferType::eShaderStorage>> getBVHNodesBuffer() const;
private:
void build();
void buildBuffers();
void split(uint32_t global_id);
static bool shouldSplit(const math::Bounds &centroid);
void addAxisAlignedBoundingBox(BVHNode& node, int start, int end) const;
const int32_t m_bin_count = 16;
const int32_t m_plane_count = m_bin_count - 1;
const float m_bin_epsilon = 0.1f;
std::shared_ptr<SceneCollector> m_scene_collector;
std::vector<Triangle> m_sorted_triangles;
std::vector<BVHNode> m_nodes;
std::shared_ptr<core::Buffer<gl::BufferType::eShaderStorage>> m_triangles_buffer;
std::shared_ptr<core::Buffer<gl::BufferType::eShaderStorage>> m_bvh_nodes_buffer;
std::vector<math::Bounds> m_centroid_boxes;
std::vector<math::Bounds> m_safe_centroid_boxes;
std::vector<TriangleIndexRange> m_ranges;
std::vector<TriangleIndexRange> m_safe_ranges;
std::vector<glm::vec4> m_triangle_centroids;
std::atomic_int m_buffer_position;
std::atomic_int m_leaf_counter;
std::atomic_int m_kernel_counter;
std::atomic_int m_mesh_index_offset;
uint32_t m_bvh_depth = 0;
uint32_t m_node_counter = 0;
};
}
}
#endif // !__RAYTRACER_BVH_H
\ No newline at end of file
src/libraries/advanced/FlatVoxelizer.cpp deleted 100644 → 0
View file @ 8038a3bd
#include "FlatVoxelizer.h"
#include <glm/glm.hpp>
#include <engine/Time.h>
namespace glare
{
namespace advanced
{
FlatVoxelizer::FlatVoxelizer()
{
m_size_buffer = std::make_unique<core::Buffer<gl::BufferType::eShaderStorage>>();
m_range_offset_buffer = std::make_unique<core::Buffer<gl::BufferType::eShaderStorage>>();
m_vertex_buffer = std::make_unique<core::Buffer<gl::BufferType::eShaderStorage>>();
m_triangle_buffer = std::make_unique<core::Buffer<gl::BufferType::eShaderStorage>>();
m_primitive_ids_buffer = std::make_unique<core::Buffer<gl::BufferType::eShaderStorage>>();
m_voxel_buffer = std::make_unique<core::Buffer<gl::BufferType::eShaderStorage>>();
m_find_overlap_count_shader = core::ShaderProgram::makeUnique("/voxelizer/flat/voxelpass_find_overlap_count.compute");
m_find_voxel_ranges_shader = core::ShaderProgram::makeUnique("/voxelizer/flat/voxelpass_find_voxel_ranges.compute");
m_push_primitives_shader = core::ShaderProgram::makeUnique("/voxelizer/flat/voxelpass_push_primitives.compute");
}
void FlatVoxelizer::voxelize(const MeshCollector& collector, unsigned subdiv_longest)
{
core::ClockNS clock;
// ######### SETUP #########
// calculate voxel dimensions and spans considering the
// subdiv_longest parameter. For that, calculate, which
// is the longest side of the mesh bounds and subdivide
// it by subdiv_longest. After that, fill up as many
// voxels as it needs to cover the whole bounding box.
m_bounds = collector.bounds();
m_subdivisions = math::UniformBoundsSubdivision(m_bounds, subdiv_longest);
int largest_axis = collector.bounds().largest();
glm::vec3 size = collector.bounds().size().xyz + glm::vec3(1e-5f);
float largest_axis_value = glm::compMax(size);
m_voxels.resize(m_subdivisions.childCount());
Log_Info << "Size calculation took " << clock.restart() << "ns";
Log_Info << "Having a resolution of " << m_subdivisions.subdivisions.x << "x" << m_subdivisions.subdivisions.y << "x" << m_subdivisions.subdivisions.z;
Log_Info << "Having a bounds size of " << size.x << "x" << size.y << "x" << size.z;
Log_Info << "Having a voxel size of " << m_subdivisions.child_size;
Log_Info << "Having " << collector.triangles().size() << " triangles";
// Setup part
// Initialize buffers
// m_primitive_ids_buffer will be resized, after computing it's extend.
m_size_buffer->setData(sizeof(uint32_t), gl::BufferUsage::eDynamicCopy, 0);
m_range_offset_buffer->setData(sizeof(uint32_t), gl::BufferUsage::eDynamicCopy, 0);
m_vertex_buffer->setData(collector.vertices(), gl::BufferUsage::eDynamicRead);
m_triangle_buffer->setData(collector.triangles(), gl::BufferUsage::eDynamicRead);
m_voxel_buffer->setData(m_voxels, gl::BufferUsage::eDynamicCopy);
Log_Info << "Size calculation took " << clock.restart() << "ns";
//Assign buffers.
m_find_overlap_count_shader->updateStorageBuffer("atomic_counter_buffer", m_size_buffer);
m_find_overlap_count_shader->updateStorageBuffer("triangle_buffer", m_triangle_buffer);
m_find_overlap_count_shader->updateStorageBuffer("vertex_buffer", m_vertex_buffer);
m_find_overlap_count_shader->updateStorageBuffer("flat_voxel_buffer", m_voxel_buffer);
m_find_overlap_count_shader->updateUniformFloat("u_flat_voxelizer.voxel_size", m_subdivisions.child_size);
m_find_overlap_count_shader->updateUniformVec3("u_flat_voxelizer.resolution", m_subdivisions.subdivisions);
m_find_overlap_count_shader->updateUniformVec4("u_scene_bounds.min", m_bounds.min);
m_find_overlap_count_shader->updateUniformVec4("u_scene_bounds.max", m_bounds.max);
// ######### PASS 01 #########
// Determine, how many triangles each voxel overlaps.
// Here, the full buffer size will be calculated too.
m_find_overlap_count_shader->dispatchCompute(uint32_t(collector.triangles().size()), 256);
m_buffer_size = m_size_buffer->read<uint32_t>(1)[0];
m_primitive_ids_buffer->setData(sizeof(uint32_t) * m_buffer_size, gl::BufferUsage::eDynamicCopy);
Log_Info << "Pass 01 took " << clock.restart() << "ns";
Log_Info << "Having a size of " << m_buffer_size;
//Assign buffers.
m_find_voxel_ranges_shader->updateStorageBuffer("atomic_range_buffer", m_range_offset_buffer);
m_find_voxel_ranges_shader->updateStorageBuffer("flat_voxel_buffer", m_voxel_buffer);
// ######### PASS 02 #########
// Compute the real triangle-id-buffer sub-ranges for each voxel
m_find_voxel_ranges_shader->dispatchCompute(uint32_t(m_voxels.size()), 256);
Log_Info << "Pass 02 took " << clock.restart() << "ns";
//Assign buffers.
m_push_primitives_shader->updateStorageBuffer("triangles_buffer", m_triangle_buffer);
m_push_primitives_shader->updateStorageBuffer("vertex_buffer", m_vertex_buffer);
m_push_primitives_shader->updateStorageBuffer("flat_voxel_buffer", m_voxel_buffer);
m_push_primitives_shader->updateStorageBuffer("primitive_ids_buffer", m_primitive_ids_buffer);
m_push_primitives_shader->updateUniformFloat("u_flat_voxelizer.voxel_size", m_subdivisions.child_size);
m_push_primitives_shader->updateUniformVec3("u_flat_voxelizer.resolution", m_subdivisions.subdivisions);
m_push_primitives_shader->updateUniformVec4("u_scene_bounds.min", m_bounds.min);
m_push_primitives_shader->updateUniformVec4("u_scene_bounds.max", m_bounds.max);
// ######### PASS 03 #########
// For each triangle, check again which voxel it overlaps and push it onto it's sub-range.
m_push_primitives_shader->dispatchCompute(uint32_t(collector.triangles().size()), 256);
Log_Info << "Pass 03 took " << clock.restart() << "ns";
/* auto bdata = m_voxel_buffer->read<FlatVoxel>(uint32_t(m_voxels.size()));
int i = 0;
for(auto &&a : bdata)
{
Log_Info << i++ << ": " << a.primitive_count << ", " << a.primitive_range_start << ", " << a.atomic_push_pointer;
if (a.primitive_count > 0);
}*/