From af7e0bf00be56e0ebb52bb9cdc4295de7a3f0cbd Mon Sep 17 00:00:00 2001 From: mitchellhansen Date: Thu, 12 Apr 2018 22:06:12 -0700 Subject: [PATCH] going to need to either add method of getting sub voxel position from the idx mask, or keep track of both vox and sub vox position in the DFS. Converted some logic steps in the DFS over to branchless selects --- kernels/ray_caster_kernel.cl | 37 +++++++++++------------------------- 1 file changed, 11 insertions(+), 26 deletions(-) diff --git a/kernels/ray_caster_kernel.cl b/kernels/ray_caster_kernel.cl index 761dd7a..1eab2ea 100644 --- a/kernels/ray_caster_kernel.cl +++ b/kernels/ray_caster_kernel.cl @@ -178,26 +178,14 @@ struct TraversalState get_oct_vox( ts.idx_stack[ts.scale] = 0; // Do the logic steps to find which sub oct we step down into - if (position.x >= (dimension / 2) + ts.oct_pos.x) { + uchar3 thing = select((uchar3)(0, 0, 0), + (uchar3)(idx_set_x_mask, idx_set_y_mask, idx_set_z_mask), + convert_char3(position >= (int3)(dimension/2) + ts.oct_pos)); - // Set our voxel position to the (0,0) of the correct oct - // ts.oct_pos.x += (dimension / 2); + ts.idx_stack[ts.scale] = thing.x | thing.y | thing.z; - // Set the idx to represent the move - ts.idx_stack[ts.scale] |= idx_set_x_mask; - - } - if (position.y >= (dimension / 2) + ts.oct_pos.y) { - - // ts.oct_pos.y += (dimension / 2); - ts.idx_stack[ts.scale] |= idx_set_y_mask; - - } - if (position.z >= (dimension / 2) + ts.oct_pos.z) { - - // ts.oct_pos.z += (dimension / 2); - ts.idx_stack[ts.scale] |= idx_set_z_mask; - } + // Set our voxel position to the (0,0) of the correct oct + ts.oct_pos += select((int3)(0), (int3)(dimension/2), position >= (int3)(dimension/2) + ts.oct_pos); int mask_index = ts.idx_stack[ts.scale]; @@ -217,11 +205,9 @@ struct TraversalState get_oct_vox( // If all went well and we found a valid non-leaf oct then we will traverse further down the hierarchy ts.scale++; ts.parent_stack_position++; - ts.oct_pos += select((int3)(0), (int3)(dimension/2), position >= (int3)(dimension/2) + ts.oct_pos); dimension /= 2; ts.resolution /= 2; - // Count the number of valid octs that come before and add it to the index to get the position // Negate it by one as it counts itself int count = popcount((uchar)(ts.current_descriptor >> 16) & count_mask_8[mask_index]) - 1; @@ -251,7 +237,6 @@ struct TraversalState get_oct_vox( // Parent stack is only populated up to the current descriptors parent. // So that would be the current voxels grandparent - ts.oct_pos += select((int3)(0), (int3)(dimension/2), position >= (int3)(dimension/2) + ts.oct_pos); ts.found = 0; return ts; } @@ -333,7 +318,7 @@ __kernel void raycaster( intersection_t += delta_t * -1 * convert_float3(isless(intersection_t, 0)); int distance_traveled = 0; - int max_distance = 100; + int max_distance = 10; uint bounce_count = 0; int3 face_mask = { 0, 0, 0 }; int voxel_data = 0; @@ -390,8 +375,8 @@ __kernel void raycaster( prev_voxel = voxel; // not working, wish I would have commented!!! - //voxel.xyz += voxel_step.xyz * face_mask.xyz * convert_int3((traversal_state.oct_pos - voxel.xyz) + traversal_state.resolution); - voxel.xyz += voxel_step.xyz * face_mask.xyz * traversal_state.resolution; + voxel.xyz += voxel_step.xyz * face_mask.xyz * convert_int3((traversal_state.oct_pos - voxel.xyz) + traversal_state.resolution); + //voxel.xyz += voxel_step.xyz * face_mask.xyz * traversal_state.resolution; // Test for out of bounds contions, add fog if (any(voxel >= *map_dim) || any(voxel < 0)){ @@ -465,7 +450,7 @@ __kernel void raycaster( mask_index = traversal_state.idx_stack[traversal_state.scale]; is_valid = (traversal_state.parent_stack[traversal_state.parent_stack_position] >> 16) & mask_8[mask_index]; failsafe++; - if (failsafe > 50) + if (failsafe > 10) break; } @@ -540,7 +525,7 @@ __kernel void raycaster( is_valid = (traversal_state.parent_stack[traversal_state.parent_stack_position] >> 16) & mask_8[mask_index]; failsafe++; - if (failsafe > 50) + if (failsafe > 10) break; }