Correct syncthreads bug

applications/reconstruct/src/ray_cast_sdf.cu

@@ -180,6 +180,7 @@ __global__ void nickRenderKernel(ftl::voxhash::HashData hashData, RayCastData ra
 	if (i == 0) blocks[0] = hashData.d_hashCompactified[blockIdx.x];
 	if (i >= 1 && i <= 7) blocks[i] = hashData.getHashEntryForSDFBlockPos(blockDelinear(blocks[0].pos, i));
 
+	// Make sure all hash entries are cached
 	__syncthreads();
 
 	const int3 pi_base = hashData.SDFBlockToVirtualVoxelPos(blocks[0].pos);
@@ -222,6 +223,8 @@ __global__ void nickRenderKernel(ftl::voxhash::HashData hashData, RayCastData ra
 		j+SDF_DX+SDF_DZ, j+SDF_DX+SDF_DY+SDF_DZ
 	};
 
+	__syncthreads();
+
 	// If any weight is 0, skip this voxel
 	const bool missweight = voxels[ix[0]].weight == 0 || voxels[ix[1]].weight == 0 || voxels[ix[2]].weight == 0 ||
 			voxels[ix[3]].weight == 0 || voxels[ix[4]].weight == 0 || voxels[ix[5]].weight == 0 ||
@@ -252,8 +255,6 @@ __global__ void nickRenderKernel(ftl::voxhash::HashData hashData, RayCastData ra
 
 	//return;*/
 
-	__syncthreads();
-
 	bool is_surface = false;
 	// Identify surfaces through sign change. Since we only check in one direction
 	// it is fine to check for any sign change?