From b9c1bef7bcf7de4dfa788e4034b14b74b7d74c07 Mon Sep 17 00:00:00 2001 From: MitchellHansen Date: Thu, 12 Oct 2017 00:04:06 -0700 Subject: [PATCH 1/2] Fixed a small bug that was breaking the far ptr's --- include/map/Octree.h | 2 +- include/util.hpp | 20 ++++++++++---------- src/map/Octree.cpp | 34 +++++++++++++++++++++------------- 3 files changed, 32 insertions(+), 24 deletions(-) diff --git a/include/map/Octree.h b/include/map/Octree.h index d3e7a7a..c96688e 100644 --- a/include/map/Octree.h +++ b/include/map/Octree.h @@ -22,7 +22,7 @@ struct OctState { class Octree { public: - static const int buffer_size = 100000; + static const int buffer_size = 300000; Octree(); ~Octree() {}; diff --git a/include/util.hpp b/include/util.hpp index af2f239..0ef49cb 100644 --- a/include/util.hpp +++ b/include/util.hpp @@ -150,20 +150,20 @@ inline std::string read_file(std::string file_name){ inline void PrettyPrintUINT64(uint64_t i, std::stringstream* ss) { - *ss << "[" << std::bitset<15>(i) << "]"; - *ss << "[" << std::bitset<1>(i >> 15) << "]"; - *ss << "[" << std::bitset<8>(i >> 16) << "]"; - *ss << "[" << std::bitset<8>(i >> 24) << "]"; - *ss << "[" << std::bitset<32>(i >> 32) << "]\n"; + //*ss << "[" << std::bitset<15>(i) << "]"; + //*ss << "[" << std::bitset<1>(i >> 15) << "]"; + //*ss << "[" << std::bitset<8>(i >> 16) << "]"; + //*ss << "[" << std::bitset<8>(i >> 24) << "]"; + //*ss << "[" << std::bitset<32>(i >> 32) << "]\n"; } inline void PrettyPrintUINT64(uint64_t i) { - std::cout << "[" << std::bitset<15>(i) << "]"; - std::cout << "[" << std::bitset<1>(i >> 15) << "]"; - std::cout << "[" << std::bitset<8>(i >> 16) << "]"; - std::cout << "[" << std::bitset<8>(i >> 24) << "]"; - std::cout << "[" << std::bitset<32>(i >> 32) << "]" << std::endl; + //std::cout << "[" << std::bitset<15>(i) << "]"; + //std::cout << "[" << std::bitset<1>(i >> 15) << "]"; + //std::cout << "[" << std::bitset<8>(i >> 16) << "]"; + //std::cout << "[" << std::bitset<8>(i >> 24) << "]"; + //std::cout << "[" << std::bitset<32>(i >> 32) << "]" << std::endl; } inline void DumpLog(std::stringstream* ss, std::string file_name) { diff --git a/src/map/Octree.cpp b/src/map/Octree.cpp index 6665db9..6d977f2 100644 --- a/src/map/Octree.cpp +++ b/src/map/Octree.cpp @@ -126,16 +126,17 @@ OctState Octree::GetVoxel(sf::Vector3i position) { // Negate it by one as it counts itself int count = count_bits((uint8_t)(head >> 16) & count_mask_8[mask_index]) - 1; + // access the far point at which the head points too. Determine it's value, and add + // a count of the valid bits to the index + if (far_bit_mask & descriptor_buffer[current_index]) { + int far_pointer_index = current_index + (head & child_pointer_mask); + current_index = descriptor_buffer[far_pointer_index] + count; + } // access the element at which head points to and then add the specified number of indices // to get to the correct child descriptor - bool jumping = false; - if (far_bit_mask & descriptor_buffer[current_index]) - jumping = true; - - current_index = current_index + (head & child_pointer_mask) + count; - - if (jumping == true) - current_index = descriptor_buffer[current_index]; + else { + current_index = current_index + (head & child_pointer_mask) + count; + } head = descriptor_buffer[current_index]; @@ -192,6 +193,7 @@ std::tuple Octree::GenerationRecursion(char* data, sf::Vecto // absolute position of it within the descriptor buffer std::tuple descriptor_and_position(0, 0); + // If we hit the 1th voxel scale then we need to query the 3D grid // and get the voxel at that position. I assume in the future when I // want to do chunking / loading of raw data I can edit the voxel access @@ -243,6 +245,9 @@ std::tuple Octree::GenerationRecursion(char* data, sf::Vecto descriptor_position_array.push_back(child); } } + + if (voxel_scale == 64) + std::cout << "WHoA"; // We are working bottom up so we need to subtract from the stack position // the amount of elements we want to use. In the worst case this will be @@ -254,7 +259,7 @@ std::tuple Octree::GenerationRecursion(char* data, sf::Vecto if (page_header_counter - worst_case_insertion_size <= 0) { // Jump to the page headers position and reset the counter - descriptor_buffer_position -= 0x8000 - page_header_counter; + descriptor_buffer_position -= page_header_counter; page_header_counter = 0x8000; // Fill the space with blank @@ -264,11 +269,13 @@ std::tuple Octree::GenerationRecursion(char* data, sf::Vecto } - for (int i = 0; i < descriptor_position_array.size(); i++) { - std::get<1>(descriptor_position_array.at(i)) = descriptor_buffer_position - i; - } + //for (int i = 0; i < descriptor_position_array.size(); i++) { + // std::get<1>(descriptor_position_array.at(i)) = descriptor_buffer_position - i; + //} unsigned int far_pointer_count = 0; + + // If looking "up" to zero, the far ptr is entered first before the cp block. Get it's position uint64_t far_pointer_block_position = descriptor_buffer_position; // Count the far pointers we need to allocate @@ -302,7 +309,8 @@ std::tuple Octree::GenerationRecursion(char* data, sf::Vecto if (relative_distance > 0x8000) { descriptor |= far_bit_mask; - descriptor |= far_pointer_block_position; + // The distance from this cp to the far ptr + descriptor |= far_pointer_block_position - descriptor_buffer_position; far_pointer_block_position--; From 77e283a4cef69ca1807aac40e9073b4fcdac30e4 Mon Sep 17 00:00:00 2001 From: MitchellHansen Date: Thu, 12 Oct 2017 00:15:56 -0700 Subject: [PATCH 2/2] 256^3 rendering at 15FPS WITHOUT THE TRAVERSAL ALGORITHM! This thing's gonna FLY! --- kernels/ray_caster_kernel.cl | 21 +++++++++++---------- src/Application.cpp | 2 +- src/CLCaster.cpp | 6 +++++- src/map/Octree.cpp | 3 --- 4 files changed, 17 insertions(+), 15 deletions(-) diff --git a/kernels/ray_caster_kernel.cl b/kernels/ray_caster_kernel.cl index dce496f..8ba1e2d 100644 --- a/kernels/ray_caster_kernel.cl +++ b/kernels/ray_caster_kernel.cl @@ -126,7 +126,7 @@ bool get_oct_vox( parent_stack[parent_stack_position] = head; // Set our initial dimension and the position at the corner of the oct to keep track of our position - int dimension = 128; + int dimension = OCTDIM; int3 quad_position = zeroed_int3; // While we are not at the required resolution @@ -192,16 +192,17 @@ bool get_oct_vox( // Negate it by one as it counts itself int count = popcount((uchar)(head >> 16) & count_mask_8[mask_index]) - 1; - //bool jumping = false; - //if (far_bit_mask & descriptor_buffer[current_index]) - // jumping = true; + // access the far point at which the head points too. Determine it's value, and add + // a count of the valid bits to the index + if (far_bit_mask & octree_descriptor_buffer[current_index]) { + int far_pointer_index = current_index + (head & child_pointer_mask); + current_index = octree_descriptor_buffer[far_pointer_index] + count; + } // access the element at which head points to and then add the specified number of indices // to get to the correct child descriptor - current_index = current_index + (head & child_pointer_mask) + count; - - //if (jumping == true) - // current_index = descriptor_buffer[current_index]; - + else { + current_index = current_index + (head & child_pointer_mask) + count; + } head = octree_descriptor_buffer[current_index]; // Increment the parent stack position and put the new oct node as the parent @@ -323,7 +324,7 @@ __kernel void raycaster( break; } - constant int vox_dim = 128; + constant int vox_dim = OCTDIM; // If we hit a voxel if (voxel.x < vox_dim && voxel.y < vox_dim && voxel.z < vox_dim){ diff --git a/src/Application.cpp b/src/Application.cpp index a45a36e..86ea2e9 100644 --- a/src/Application.cpp +++ b/src/Application.cpp @@ -37,7 +37,7 @@ bool Application::init_clcaster() { // Init the raycaster with a specified dimension and a pointer to the source // array style data - octree = std::make_shared(128, map.get()); + octree = std::make_shared(256, map.get()); raycaster->assign_octree(octree); diff --git a/src/CLCaster.cpp b/src/CLCaster.cpp index b9384bb..acb394c 100644 --- a/src/CLCaster.cpp +++ b/src/CLCaster.cpp @@ -706,7 +706,11 @@ bool CLCaster::compile_kernel(std::string kernel_source, bool is_path, std::stri // Try and build the program // "-cl-finite-math-only -cl-fast-relaxed-math -cl-unsafe-math-optimizations" - error = clBuildProgram(program, 1, &device_id, "-cl-finite-math-only -cl-fast-relaxed-math -cl-unsafe-math-optimizations", NULL, NULL); + + // need a ref to the oct dimensions + //std::string oct_dimensions = std::to_string(map->getDimensions().x); + std::string build_string = "-DOCTDIM=256 -cl-finite-math-only -cl-fast-relaxed-math -cl-unsafe-math-optimizations"; + error = clBuildProgram(program, 1, &device_id, build_string.c_str(), NULL, NULL); // Check to see if it error'd out if (cl_assert(error)) { diff --git a/src/map/Octree.cpp b/src/map/Octree.cpp index 6d977f2..0177b34 100644 --- a/src/map/Octree.cpp +++ b/src/map/Octree.cpp @@ -245,9 +245,6 @@ std::tuple Octree::GenerationRecursion(char* data, sf::Vecto descriptor_position_array.push_back(child); } } - - if (voxel_scale == 64) - std::cout << "WHoA"; // We are working bottom up so we need to subtract from the stack position // the amount of elements we want to use. In the worst case this will be