mitchellhansen/voxel-raycaster
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Merge pull request #2 from MitchellHansen/octree-develop

Browse Source 
Octree develop
This commit is contained in:
Mitchell
2017-10-12 00:17:22 -07:00
committed by GitHub
parent 2519532172 77e283a4ce
commit 618535fed7
6 changed files with 46 additions and 36 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
include/map/Octree.h
View File

@@ -22,7 +22,7 @@ struct OctState {
class Octree { class Octree {
public: public:
static const int buffer_size = 100000; static const int buffer_size = 300000;
Octree(); Octree();
~Octree() {}; ~Octree() {};

20
include/util.hpp
View File

@@ -150,20 +150,20 @@ inline std::string read_file(std::string file_name){
inline void PrettyPrintUINT64(uint64_t i, std::stringstream* ss) { inline void PrettyPrintUINT64(uint64_t i, std::stringstream* ss) {
*ss << "[" << std::bitset<15>(i) << "]"; //*ss << "[" << std::bitset<15>(i) << "]";
*ss << "[" << std::bitset<1>(i >> 15) << "]"; //*ss << "[" << std::bitset<1>(i >> 15) << "]";
*ss << "[" << std::bitset<8>(i >> 16) << "]"; //*ss << "[" << std::bitset<8>(i >> 16) << "]";
*ss << "[" << std::bitset<8>(i >> 24) << "]"; //*ss << "[" << std::bitset<8>(i >> 24) << "]";
*ss << "[" << std::bitset<32>(i >> 32) << "]\n"; //*ss << "[" << std::bitset<32>(i >> 32) << "]\n";
} }
inline void PrettyPrintUINT64(uint64_t i) { inline void PrettyPrintUINT64(uint64_t i) {
std::cout << "[" << std::bitset<15>(i) << "]"; //std::cout << "[" << std::bitset<15>(i) << "]";
std::cout << "[" << std::bitset<1>(i >> 15) << "]"; //std::cout << "[" << std::bitset<1>(i >> 15) << "]";
std::cout << "[" << std::bitset<8>(i >> 16) << "]"; //std::cout << "[" << std::bitset<8>(i >> 16) << "]";
std::cout << "[" << std::bitset<8>(i >> 24) << "]"; //std::cout << "[" << std::bitset<8>(i >> 24) << "]";
std::cout << "[" << std::bitset<32>(i >> 32) << "]" << std::endl; //std::cout << "[" << std::bitset<32>(i >> 32) << "]" << std::endl;
} }
inline void DumpLog(std::stringstream* ss, std::string file_name) { inline void DumpLog(std::stringstream* ss, std::string file_name) {

21
kernels/ray_caster_kernel.cl
View File

@@ -126,7 +126,7 @@ bool get_oct_vox(
parent_stack[parent_stack_position] = head; 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 // 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; int3 quad_position = zeroed_int3;
// While we are not at the required resolution // While we are not at the required resolution
@@ -192,16 +192,17 @@ bool get_oct_vox(
// Negate it by one as it counts itself // Negate it by one as it counts itself
int count = popcount((uchar)(head >> 16) & count_mask_8[mask_index]) - 1; int count = popcount((uchar)(head >> 16) & count_mask_8[mask_index]) - 1;
//bool jumping = false; // access the far point at which the head points too. Determine it's value, and add
//if (far_bit_mask & descriptor_buffer[current_index]) // a count of the valid bits to the index
// jumping = true; 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 // access the element at which head points to and then add the specified number of indices
// to get to the correct child descriptor // to get to the correct child descriptor
current_index = current_index + (head & child_pointer_mask) + count; else {
current_index = current_index + (head & child_pointer_mask) + count;
//if (jumping == true) }
// current_index = descriptor_buffer[current_index];
head = octree_descriptor_buffer[current_index]; head = octree_descriptor_buffer[current_index];
// Increment the parent stack position and put the new oct node as the parent // Increment the parent stack position and put the new oct node as the parent
@@ -323,7 +324,7 @@ __kernel void raycaster(
break; break;
} }
constant int vox_dim = 128; constant int vox_dim = OCTDIM;
// If we hit a voxel // If we hit a voxel
if (voxel.x < vox_dim && voxel.y < vox_dim && voxel.z < vox_dim){ if (voxel.x < vox_dim && voxel.y < vox_dim && voxel.z < vox_dim){

2
src/Application.cpp
View File

@@ -37,7 +37,7 @@ bool Application::init_clcaster() {
// Init the raycaster with a specified dimension and a pointer to the source // Init the raycaster with a specified dimension and a pointer to the source
// array style data // array style data
octree = std::make_shared<Map>(128, map.get()); octree = std::make_shared<Map>(256, map.get());
raycaster->assign_octree(octree); raycaster->assign_octree(octree);

6
src/CLCaster.cpp
View File

@@ -706,7 +706,11 @@ bool CLCaster::compile_kernel(std::string kernel_source, bool is_path, std::stri
// Try and build the program // Try and build the program
// "-cl-finite-math-only -cl-fast-relaxed-math -cl-unsafe-math-optimizations" // "-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 // Check to see if it error'd out
if (cl_assert(error)) { if (cl_assert(error)) {

31
src/map/Octree.cpp
View File

@@ -126,16 +126,17 @@ OctState Octree::GetVoxel(sf::Vector3i position) {
// Negate it by one as it counts itself // Negate it by one as it counts itself
int count = count_bits((uint8_t)(head >> 16) & count_mask_8[mask_index]) - 1; 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 // access the element at which head points to and then add the specified number of indices
// to get to the correct child descriptor // to get to the correct child descriptor
bool jumping = false; else {
if (far_bit_mask & descriptor_buffer[current_index]) current_index = current_index + (head & child_pointer_mask) + count;
jumping = true; }
current_index = current_index + (head & child_pointer_mask) + count;
if (jumping == true)
current_index = descriptor_buffer[current_index];
head = descriptor_buffer[current_index]; head = descriptor_buffer[current_index];
@@ -192,6 +193,7 @@ std::tuple<uint64_t, uint64_t> Octree::GenerationRecursion(char* data, sf::Vecto
// absolute position of it within the descriptor buffer // absolute position of it within the descriptor buffer
std::tuple<uint64_t, uint64_t> descriptor_and_position(0, 0); std::tuple<uint64_t, uint64_t> descriptor_and_position(0, 0);
// If we hit the 1th voxel scale then we need to query the 3D grid // 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 // 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 // want to do chunking / loading of raw data I can edit the voxel access
@@ -254,7 +256,7 @@ std::tuple<uint64_t, uint64_t> Octree::GenerationRecursion(char* data, sf::Vecto
if (page_header_counter - worst_case_insertion_size <= 0) { if (page_header_counter - worst_case_insertion_size <= 0) {
// Jump to the page headers position and reset the counter // 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; page_header_counter = 0x8000;
// Fill the space with blank // Fill the space with blank
@@ -264,11 +266,13 @@ std::tuple<uint64_t, uint64_t> Octree::GenerationRecursion(char* data, sf::Vecto
} }
for (int i = 0; i < descriptor_position_array.size(); i++) { //for (int i = 0; i < descriptor_position_array.size(); i++) {
std::get<1>(descriptor_position_array.at(i)) = descriptor_buffer_position - i; // std::get<1>(descriptor_position_array.at(i)) = descriptor_buffer_position - i;
} //}
unsigned int far_pointer_count = 0; 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; uint64_t far_pointer_block_position = descriptor_buffer_position;
// Count the far pointers we need to allocate // Count the far pointers we need to allocate
@@ -302,7 +306,8 @@ std::tuple<uint64_t, uint64_t> Octree::GenerationRecursion(char* data, sf::Vecto
if (relative_distance > 0x8000) { if (relative_distance > 0x8000) {
descriptor |= far_bit_mask; 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--; far_pointer_block_position--;