256^3 rendering at 15FPS WITHOUT THE TRAVERSAL ALGORITHM! This thing's gonna FLY!

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
+			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];
 
 			// 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){

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<Map>(128, map.get());
+	octree = std::make_shared<Map>(256, map.get());
 	raycaster->assign_octree(octree);
 
 

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)) {

src/map/Octree.cpp

@@ -246,9 +246,6 @@ std::tuple<uint64_t, uint64_t> Octree::GenerationRecursion(char* data, sf::Vecto
 		}
 	}
 	
-	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 
 	// a far pointer for ever descriptor (size * 2)