Fixes #66 virtual compactify

applications/reconstruct/include/ftl/ray_cast_sdf.hpp

@@ -42,7 +42,7 @@ public:
 		return params;
 	}
 
-	void render(const ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams, const DepthCameraParams& cameraParams, const Eigen::Matrix4f& lastRigidTransform);
+	void render(ftl::voxhash::HashData& hashData, ftl::voxhash::HashParams& hashParams, const DepthCameraParams& cameraParams, const Eigen::Matrix4f& lastRigidTransform);
 
 	const RayCastData& getRayCastData(void) {
 		return m_data;
@@ -60,6 +60,8 @@ private:
 	void create(const RayCastParams& params);
 	void destroy(void);
 
+	void compactifyHashEntries(ftl::voxhash::HashData& hashData, ftl::voxhash::HashParams& hashParams);
+
 	void rayIntervalSplatting(const ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams, const Eigen::Matrix4f& lastRigidTransform); // rasterize
 
 	RayCastParams m_params;

applications/reconstruct/include/ftl/scene_rep_hash_sdf.hpp

@@ -98,7 +98,7 @@ class SceneRep : public ftl::Configurable {
 		alloc(depthCameraData, depthCameraParams, d_bitMask);
 
 		//generate a linear hash array with only occupied entries
-		compactifyHashEntries(depthCameraData);
+		compactifyHashEntries();
 
 		//volumetrically integrate the depth data into the depth SDFBlocks
 		integrateDepthMap(depthCameraData, depthCameraParams);
@@ -115,7 +115,7 @@ class SceneRep : public ftl::Configurable {
 
 	void setLastRigidTransformAndCompactify(const Eigen::Matrix4f& lastRigidTransform, const DepthCameraData& depthCameraData) {
 		setLastRigidTransform(lastRigidTransform);
-		compactifyHashEntries(depthCameraData);
+		compactifyHashEntries();
 	}
 
 
@@ -145,7 +145,7 @@ class SceneRep : public ftl::Configurable {
 		return m_hashData;
 	} 
 
-	const HashParams& getHashParams() const {
+	HashParams& getHashParams() {
 		return m_hashParams;
 	}
 
@@ -312,7 +312,7 @@ private:
 	}
 
 
-	void compactifyHashEntries(const DepthCameraData& depthCameraData) {
+	void compactifyHashEntries() { //const DepthCameraData& depthCameraData) {
 		//Start Timing
 		//if(GlobalAppState::get().s_timingsDetailledEnabled) { cutilSafeCall(cudaDeviceSynchronize()); m_timer.start(); }
 

applications/reconstruct/src/ray_cast_sdf.cpp

@@ -36,7 +36,17 @@ void CUDARayCastSDF::destroy(void)
 	//m_rayIntervalSplatting.OnD3D11DestroyDevice();
 }
 
-void CUDARayCastSDF::render(const ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams, const DepthCameraParams& cameraParams, const Eigen::Matrix4f& lastRigidTransform)
+extern "C" unsigned int compactifyHashAllInOneCUDA(ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams);
+
+
+void CUDARayCastSDF::compactifyHashEntries(ftl::voxhash::HashData& hashData, ftl::voxhash::HashParams& hashParams) { //const DepthCameraData& depthCameraData) {
+		
+	hashParams.m_numOccupiedBlocks = compactifyHashAllInOneCUDA(hashData, hashParams);		//this version uses atomics over prefix sums, which has a much better performance
+	std::cout << "Ray blocks = " << hashParams.m_numOccupiedBlocks << std::endl;
+	hashData.updateParams(hashParams);	//make sure numOccupiedBlocks is updated on the GPU
+}
+
+void CUDARayCastSDF::render(ftl::voxhash::HashData& hashData, ftl::voxhash::HashParams& hashParams, const DepthCameraParams& cameraParams, const Eigen::Matrix4f& lastRigidTransform)
 {
 	updateConstantDepthCameraParams(cameraParams);
 	//rayIntervalSplatting(hashData, hashParams, lastRigidTransform);
@@ -48,6 +58,8 @@ void CUDARayCastSDF::render(const ftl::voxhash::HashData& hashData, const ftl::v
 	m_params.m_viewMatrixInverse = MatrixConversion::toCUDA(lastRigidTransform);
 	m_data.updateParams(m_params);
 
+	compactifyHashEntries(hashData, hashParams);
+
 	if (hash_render_) nickRenderCUDA(hashData, hashParams, m_data, m_params);
 	else renderCS(hashData, m_data, m_params);
 

applications/reconstruct/src/ray_cast_sdf.cu

@@ -298,6 +298,8 @@ __global__ void nickRenderKernel(ftl::voxhash::HashData hashData, RayCastData ra
 	const float2 screenPosf = DepthCameraData::cameraToKinectScreenFloat(camPos);
 	const uint2 screenPos = make_uint2(make_int2(screenPosf)); //  + make_float2(0.5f, 0.5f)
 
+	if (camPos.z < 0.0f) return;
+
 	/*if (screenPos.x < params.m_width && screenPos.y < params.m_height && 
 			rayCastData.d_depth[(screenPos.y)*params.m_width+screenPos.x] > camPos.z) {
 		rayCastData.d_depth[(screenPos.y)*params.m_width+screenPos.x] = camPos.z;