diff --git a/include/map/Map.h b/include/map/Map.h index 41bb930..d7c364d 100644 --- a/include/map/Map.h +++ b/include/map/Map.h @@ -24,16 +24,16 @@ public: Map(uint32_t dimensions); - void generate_octree(); + void dump_logs(); void setVoxel(sf::Vector3i position, int val); bool getVoxelFromOctree(sf::Vector3i position); bool getVoxel(sf::Vector3i pos); - Octree a; + Octree octree; - void test_map(); + bool test(); private: @@ -42,6 +42,9 @@ private: std::stringstream output_stream; // ========================= + void generate_octree(unsigned int dimensions); + + // Generate children is the main recursive function uint64_t generate_children(sf::Vector3i pos, int dim); char* voxel_data; diff --git a/include/map/Octree.h b/include/map/Octree.h index 60d71ca..c2b01b2 100644 --- a/include/map/Octree.h +++ b/include/map/Octree.h @@ -10,13 +10,16 @@ public: Octree(); ~Octree() {}; - uint64_t *blob = new uint64_t[100000]; + uint64_t *trunk_buffer = new uint64_t[10000]; + uint64_t *descriptor_buffer = new uint64_t[100000]; + uint32_t *attachment_lookup = new uint32_t[100000]; + uint64_t *attachment_buffer = new uint64_t[100000]; uint64_t root_index = 0; uint64_t stack_pos = 0x8000; uint64_t global_pos = 0; - uint64_t copy_to_stack(std::vector children); + uint64_t copy_to_stack(std::vector children, unsigned int voxel_scale); // With a position and the head of the stack. Traverse down the voxel hierarchy to find // the IDX and stack position of the highest resolution (maybe set resolution?) oct @@ -26,22 +29,28 @@ public: private: + std::vector anchor_stack; + unsigned int octree_voxel_dimension = 32; + // (X, Y, Z) mask for the idx const uint8_t idx_set_x_mask = 0x1; const uint8_t idx_set_y_mask = 0x2; const uint8_t idx_set_z_mask = 0x4; - // Mask for + // Mask for checking if valid or leaf const uint8_t mask_8[8] = { 0x1, 0x2, 0x4, 0x8, 0x10, 0x20, 0x40, 0x80 }; + // Mask for counting the previous valid bits const uint8_t count_mask_8[8]{ 0x1, 0x3, 0x7, 0xF, 0x1F, 0x3F, 0x7F, 0xFF }; + + // uint64_t manipulation masks const uint64_t child_pointer_mask = 0x0000000000007fff; const uint64_t far_bit_mask = 0x8000; const uint64_t valid_mask = 0xFF0000; diff --git a/kernels/ray_caster_kernel.cl b/kernels/ray_caster_kernel.cl index 123bb6d..9a933de 100644 --- a/kernels/ray_caster_kernel.cl +++ b/kernels/ray_caster_kernel.cl @@ -385,6 +385,6 @@ __kernel void raycaster( } while (++dist < 700.0f); - write_imagef(image, pixel, white_light(mix(fog_color, (float4)(0.40, 0.00, 0.40, 0.2), 1.0 - max(dist / 700.0f, (float)0)), (float3)(lights[7], lights[8], lights[9]), face_mask)); + //write_imagef(image, pixel, white_light(mix(fog_color, (float4)(0.40, 0.00, 0.40, 0.2), 1.0 - max(dist / 700.0f, (float)0)), (float3)(lights[7], lights[8], lights[9]), face_mask)); return; } diff --git a/src/main.cpp b/src/main.cpp index 079030d..33201ac 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -92,9 +92,8 @@ int main() { // ============================= Map _map(32); - _map.generate_octree(); - _map.a.print_block(0); - _map.test_map(); + _map.test(); + _map.dump_logs(); //std::cin.get(); //return 0; // ============================= @@ -102,7 +101,7 @@ int main() { sf::RenderWindow window(sf::VideoMode(WINDOW_X, WINDOW_Y), "SFML"); window.setMouseCursorVisible(false); window.setKeyRepeatEnabled(false); - window.setFramerateLimit(120); + //window.setFramerateLimit(120); window.setVerticalSyncEnabled(false); ImGui::SFML::Init(window); diff --git a/src/map/Map.cpp b/src/map/Map.cpp index 5745e13..f5f0b7c 100644 --- a/src/map/Map.cpp +++ b/src/map/Map.cpp @@ -11,8 +11,16 @@ Map::Map(uint32_t dimensions) { if (rand() % 25 < 2) voxel_data[i] = 1; else - voxel_data[i] = 1; + voxel_data[i] = 0; } + + generate_octree(dimensions); +} + + +void Map::dump_logs() { + + octree.print_block(0); } uint64_t Map::generate_children(sf::Vector3i pos, int voxel_scale) { @@ -36,7 +44,6 @@ uint64_t Map::generate_children(sf::Vector3i pos, int voxel_scale) { // want to do chunking / loading of raw data I can edit the voxel access if (voxel_scale == 1) { - // uint64_t child_descriptor = 0; // Setting the individual valid mask bits @@ -49,8 +56,8 @@ uint64_t Map::generate_children(sf::Vector3i pos, int voxel_scale) { // We are querying leafs, so we need to fill the leaf mask child_descriptor |= 0xFF000000; - // This is where contours - // The CP will be left blank, contours will be added maybe + // The CP will be left blank, contour mask and ptr will need to + // be added here later return child_descriptor; } @@ -92,7 +99,7 @@ uint64_t Map::generate_children(sf::Vector3i pos, int voxel_scale) { // interlace them and allow the memory handler to work correctly. // Copy the children to the stack and set the child_descriptors pointer to the correct value - child_descriptor |= a.copy_to_stack(descriptor_array); + child_descriptor |= octree.copy_to_stack(descriptor_array, voxel_scale); // Free space may also be allocated here as well @@ -100,19 +107,18 @@ uint64_t Map::generate_children(sf::Vector3i pos, int voxel_scale) { return child_descriptor; } -void Map::generate_octree() { +void Map::generate_octree(unsigned int dimensions) { // Launch the recursive generator at (0,0,0) as the first point // and the octree dimension as the initial block size uint64_t root_node = generate_children(sf::Vector3i(0, 0, 0), OCT_DIM/2); - uint64_t tmp = 0; // ========= DEBUG ============== PrettyPrintUINT64(root_node, &output_stream); output_stream << " " << OCT_DIM << " " << counter++ << std::endl; // ============================== - a.root_index = a.copy_to_stack(std::vector{root_node}); + octree.root_index = octree.copy_to_stack(std::vector{root_node}, OCT_DIM); // Dump the debug log DumpLog(&output_stream, "raw_output.txt"); @@ -125,7 +131,7 @@ void Map::setVoxel(sf::Vector3i world_position, int val) { bool Map::getVoxelFromOctree(sf::Vector3i position) { - return a.get_voxel(position); + return octree.get_voxel(position); } bool Map::getVoxel(sf::Vector3i pos){ @@ -137,7 +143,7 @@ bool Map::getVoxel(sf::Vector3i pos){ } } -void Map::test_map() { +bool Map::test() { std::cout << "Validating map..." << std::endl; @@ -192,5 +198,7 @@ void Map::test_map() { std::cout << "Octree linear xyz access : "; std::cout << timer.restart().asMicroseconds() << " microseconds" << std::endl; + + return true; } diff --git a/src/map/Octree.cpp b/src/map/Octree.cpp index 873c4be..d53791e 100644 --- a/src/map/Octree.cpp +++ b/src/map/Octree.cpp @@ -5,24 +5,38 @@ Octree::Octree() { // initialize the first stack block for (int i = 0; i < 0x8000; i++) { - blob[i] = 0; + descriptor_buffer[i] = 0; } } -uint64_t Octree::copy_to_stack(std::vector children) { - // Check for the 15 bit boundry - if (stack_pos - children.size() > stack_pos) { - global_pos = stack_pos; - stack_pos = 0x8000; - } - else { - stack_pos -= children.size(); - } +// Copy to stack enables the hybrid depth-breadth first tree by taking +// a list of valid non-leaf child descriptors contained under a common parent. + +uint64_t Octree::copy_to_stack(std::vector children, unsigned int voxel_scale) { + + //// Check for the 15 bit boundry + //if (stack_pos - children.size() > stack_pos) { + // global_pos = stack_pos; + // stack_pos = 0x8000; + //} + //else { + // stack_pos -= children.size(); + //} + + // Copy to stack needs to keep track of an "anchor_stack" which will hopefully facilitate + // relative pointer generation for items being copied to the stack + + // We need to return the relative pointer to the child node list + // 16 bits, one far bit, one sign bit? 14 bits == +- 16384 + + // Worth halving the ptr reach to enable backwards ptrs? + // could increase packability allowing far ptrs and attachments to come before or after + - // Check for the far bit - memcpy(&blob[stack_pos + global_pos], children.data(), children.size() * sizeof(uint64_t)); + stack_pos -= children.size(); + memcpy(&descriptor_buffer[stack_pos + global_pos], children.data(), children.size() * sizeof(uint64_t)); // Return the bitmask encoding the index of that value // If we tripped the far bit, allocate a far index to the stack and place @@ -39,7 +53,7 @@ bool Octree::get_voxel(sf::Vector3i position) { oct_state state; // push the root node to the parent stack - uint64_t head = blob[root_index]; + uint64_t head = descriptor_buffer[root_index]; state.parent_stack[state.parent_stack_position] = head; // Set our initial dimension and the position at the corner of the oct to keep track of our position @@ -116,7 +130,7 @@ bool Octree::get_voxel(sf::Vector3i position) { // access the element at which head points to and then add the specified number of indices // to get to the correct child descriptor - head = blob[(head & child_pointer_mask) + count]; + head = descriptor_buffer[(head & child_pointer_mask) + count]; // Increment the parent stack position and put the new oct node as the parent state.parent_stack_position++; @@ -142,7 +156,7 @@ void Octree::print_block(int block_pos) { std::stringstream sss; for (int i = block_pos; i < (int)pow(2, 15); i++) { - PrettyPrintUINT64(blob[i], &sss); + PrettyPrintUINT64(descriptor_buffer[i], &sss); sss << "\n"; } DumpLog(&sss, "raw_data.txt"); diff --git a/src/raycaster/Hardware_Caster.cpp b/src/raycaster/Hardware_Caster.cpp index e30d80d..f09edf1 100644 --- a/src/raycaster/Hardware_Caster.cpp +++ b/src/raycaster/Hardware_Caster.cpp @@ -613,7 +613,7 @@ int Hardware_Caster::compile_kernel(std::string kernel_source, bool is_path, std // Try and build the program // "-cl-finite-math-only -cl-fast-relaxed-math -cl-unsafe-math-optimizations" - error = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); + error = clBuildProgram(program, 1, &device_id, "-cl-finite-math-only -cl-fast-relaxed-math -cl-unsafe-math-optimizations", NULL, NULL); // Check to see if it errored out if (vr_assert(error, "clBuildProgram")) {