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/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 6665db9..0177b34 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 @@ -254,7 +256,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 +266,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 +306,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--;