Revert "Switch to storing point cloud from MLS step"

This reverts commit 55795d67.

applications/reconstruct/include/ftl/depth_camera.hpp

@@ -24,7 +24,6 @@ namespace voxhash {
 struct __align__(16) DepthCameraCUDA {
 	cudaTextureObject_t depth;
 	cudaTextureObject_t depth2;
-	cudaTextureObject_t point;
 	cudaTextureObject_t colour;
 	cudaTextureObject_t normal;
 	DepthCameraParams params;
@@ -55,7 +54,6 @@ struct DepthCamera {
 	cv::cuda::GpuMat *normal_mat_;
 	ftl::cuda::TextureObject<float> *depth_tex_;
 	ftl::cuda::TextureObject<float> *depth2_tex_;
-	ftl::cuda::TextureObject<float4> *point_tex_;
 	ftl::cuda::TextureObject<uchar4> *colour_tex_;
 	ftl::cuda::TextureObject<float4> *normal_tex_;
 

applications/reconstruct/src/depth_camera.cpp

@@ -15,26 +15,22 @@ DepthCamera::DepthCamera() {
 	depth2_tex_ = nullptr;
 	colour_tex_ = nullptr;
 	normal_tex_ = nullptr;
-	point_tex_ = nullptr;
 }
 
 void DepthCamera::alloc(const DepthCameraParams& params, bool withNormals) { //! todo resizing???
 	depth_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_32FC1);
 	depth2_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_32FC1);
 	colour_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_8UC4);
-	point_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_32FC4);
 	depth_tex_ = new ftl::cuda::TextureObject<float>((cv::cuda::PtrStepSz<float>)*depth_mat_);
 	depth2_tex_ = new ftl::cuda::TextureObject<float>((cv::cuda::PtrStepSz<float>)*depth2_mat_);
-	point_tex_ = new ftl::cuda::TextureObject<float4>((cv::cuda::PtrStepSz<float4>)*point_mat_);
 	colour_tex_ = new ftl::cuda::TextureObject<uchar4>((cv::cuda::PtrStepSz<uchar4>)*colour_mat_);
 	data.depth = depth_tex_->cudaTexture();
 	data.depth2 = depth2_tex_->cudaTexture();
 	data.colour = colour_tex_->cudaTexture();
-	data.point = point_tex_->cudaTexture();
 	data.params = params;
 
 	if (withNormals) {
-		
+		point_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_32FC3);
 		normal_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_32FC4);
 		normal_tex_ = new ftl::cuda::TextureObject<float4>((cv::cuda::PtrStepSz<float4>)*normal_mat_);
 		data.normal = normal_tex_->cudaTexture();
@@ -52,7 +48,6 @@ void DepthCamera::free() {
 	delete depth_tex_;
 	delete colour_tex_;
 	if (normal_tex_) delete normal_tex_;
-	if (point_tex_) delete normal_tex_;
 }
 
 void DepthCamera::updateData(const cv::Mat &depth, const cv::Mat &rgb, cv::cuda::Stream &stream) {
@@ -64,6 +59,6 @@ void DepthCamera::updateData(const cv::Mat &depth, const cv::Mat &rgb, cv::cuda:
 }
 
 void DepthCamera::_computeNormals(cudaStream_t stream) {
-	//ftl::cuda::point_cloud((float3*)point_mat_->data, data, stream);
-	//ftl::cuda::compute_normals((float3*)point_mat_->data, normal_tex_, stream);
+	ftl::cuda::point_cloud((float3*)point_mat_->data, data, stream);
+	ftl::cuda::compute_normals((float3*)point_mat_->data, normal_tex_, stream);
 }

applications/reconstruct/src/depth_camera.cu

@@ -12,17 +12,17 @@ using ftl::voxhash::HashParams;
 
 extern __constant__ ftl::voxhash::DepthCameraCUDA c_cameras[MAX_CAMERAS];
 
-__global__ void clear_depth_kernel(ftl::cuda::TextureObject<float4> depth) {
+__global__ void clear_depth_kernel(ftl::cuda::TextureObject<float> depth) {
 	const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
 
 	if (x < depth.width() && y < depth.height()) {
-		depth(x,y) = make_float4(MINF, MINF, MINF, MINF); //PINF;
+		depth(x,y) = 1000.0f; //PINF;
 		//colour(x,y) = make_uchar4(76,76,82,0);
 	}
 }
 
-void ftl::cuda::clear_depth(const ftl::cuda::TextureObject<float4> &depth, cudaStream_t stream) {
+void ftl::cuda::clear_depth(const ftl::cuda::TextureObject<float> &depth, cudaStream_t stream) {
 	const dim3 clear_gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
 	const dim3 clear_blockSize(T_PER_BLOCK, T_PER_BLOCK);
 	clear_depth_kernel<<<clear_gridSize, clear_blockSize, 0, stream>>>(depth);
@@ -66,7 +66,7 @@ __device__ float colordiffFloat2(const uchar4 &pa, const uchar4 &pb) {
 
 #define WINDOW_RADIUS 5
 
-__global__ void mls_smooth_kernel(ftl::cuda::TextureObject<float4> output, HashData hashData, HashParams hashParams, int numcams, int cam) {
+__global__ void mls_smooth_kernel(ftl::cuda::TextureObject<float> output, HashData hashData, HashParams hashParams, int numcams, int cam) {
 	const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
 
@@ -120,18 +120,18 @@ __global__ void mls_smooth_kernel(ftl::cuda::TextureObject<float4> output, HashD
 			//float sdf = wnorm.x * (mPos.x - wpos.x) + wnorm.y * (mPos.y - wpos.y) + wnorm.z * (mPos.z - wpos.z);
 
 			//wpos = wpos - (wnorm * (mPos - wpos));
-			//wpos = mainCamera.poseInverse * wpos;
-			//const uint2 screenPos = make_uint2(mainCamera.params.cameraToKinectScreenInt(wpos));
-			//if (screenPos.x < output.width() && screenPos.y < output.height()) {
-				//output(screenPos.x,screenPos.y) = (weights >= hashParams.m_confidenceThresh) ? wpos.z : 1000.0f;
-				output(x,y) = (weights >= hashParams.m_confidenceThresh) ? make_float4(wpos, 0.0f) : make_float4(0.0f, 0.0f, 1000.0f, 0.0f);
-			//}
+			wpos = mainCamera.poseInverse * wpos;
+			const uint2 screenPos = make_uint2(mainCamera.params.cameraToKinectScreenInt(wpos));
+			if (screenPos.x < output.width() && screenPos.y < output.height()) {
+				output(screenPos.x,screenPos.y) = (weights >= hashParams.m_confidenceThresh) ? wpos.z : 1000.0f;
+				//output(x,y) = (weights >= hashParams.m_confidenceThresh) ? wpos.z : 1000.0f;
+			}
 
 		}
 	}
 }
 
-void ftl::cuda::mls_smooth(TextureObject<float4> &output, const HashData &hashData, const HashParams &hashParams, int numcams, int cam, cudaStream_t stream) {
+void ftl::cuda::mls_smooth(TextureObject<float> &output, const HashData &hashData, const HashParams &hashParams, int numcams, int cam, cudaStream_t stream) {
 	const dim3 gridSize((output.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (output.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
 	const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
 

applications/reconstruct/src/depth_camera_cuda.hpp

@@ -7,9 +7,9 @@
 namespace ftl {
 namespace cuda {
 
-void clear_depth(const TextureObject<float4> &depth, cudaStream_t stream);
+void clear_depth(const TextureObject<float> &depth, cudaStream_t stream);
 
-void mls_smooth(TextureObject<float4> &output, const ftl::voxhash::HashData &hashData, const ftl::voxhash::HashParams &hashParams, int numcams, int cam, cudaStream_t stream);
+void mls_smooth(TextureObject<float> &output, const ftl::voxhash::HashData &hashData, const ftl::voxhash::HashParams &hashParams, int numcams, int cam, cudaStream_t stream);
 
 void point_cloud(float3* output, const ftl::voxhash::DepthCameraCUDA &depthCameraData, cudaStream_t stream);
 

applications/reconstruct/src/dibr.cu

@@ -94,7 +94,7 @@ __global__ void dibr_kernel_rev(
         const unsigned int cy = screenPos.y;
 
         if (cx < camera.params.m_imageWidth && cy < camera.params.m_imageHeight) {
-			float d = tex2D<float>(camera.depth2, (int)cx, (int)cy);
+			float d = tex2D<float>(camera.depth, (int)cx, (int)cy);
 			float camdiff = fabs(camPos.z-d);
             if (camdiff < 0.1f) {
             	colour_out(x,y) = tex2D<uchar4>(camera.colour,cx,cy);

applications/reconstruct/src/integrators.cu

@@ -246,10 +246,10 @@ __global__ void integrateMLSKernel(HashData hashData, HashParams hashParams, int
 		for (int v=-WINDOW_RADIUS; v<=WINDOW_RADIUS; ++v) {
 			for (int u=-WINDOW_RADIUS; u<=WINDOW_RADIUS; ++u) {
 				if (screenPos.x+u < width && screenPos.y+v < height) {	//on screen
-					float3 camPos = make_float3(tex2D<float4>(camera.point, screenPos.x+u, screenPos.y+v));
+					float depth = tex2D<float>(camera.depth, screenPos.x+u, screenPos.y+v);
 
 					//float4 normal = tex2D<float4>(camera.normal, screenPos.x+u, screenPos.y+v);
-					const float3 worldPos = camPos; //camera.pose * camPos; //camera.params.kinectDepthToSkeleton(screenPos.x+u, screenPos.y+v, depth);
+					const float3 worldPos = camera.pose * camera.params.kinectDepthToSkeleton(screenPos.x+u, screenPos.y+v, depth);
 					const float weight = spatialWeighting(length(pfb - worldPos));
 
 					wpos += weight*worldPos;

applications/reconstruct/src/scene_rep_hash_sdf.cu

@@ -182,7 +182,7 @@ __global__ void allocKernel(HashData hashData, HashParams hashParams, int camnum
 
 	if (x < cameraParams.m_imageWidth && y < cameraParams.m_imageHeight)
 	{
-		float d = tex2D<float>(camera.depth2, x, y);
+		float d = tex2D<float>(camera.depth, x, y);
 		//if (d == MINF || d < cameraParams.m_sensorDepthWorldMin || d > cameraParams.m_sensorDepthWorldMax)	return;
 		if (d == MINF || d == 0.0f)	return;
 

applications/reconstruct/src/voxel_scene.cpp

@@ -141,7 +141,7 @@ int SceneRep::upload() {
 				cam.params.m_imageHeight = in->parameters().height;
 				cam.params.m_sensorDepthWorldMax = in->parameters().maxDepth;
 				cam.params.m_sensorDepthWorldMin = in->parameters().minDepth;
-				cam.gpu.alloc(cam.params);
+				cam.gpu.alloc(cam.params, true);
 			}
 		}
 
@@ -357,8 +357,8 @@ void SceneRep::_compactifyAllocated() {
 
 void SceneRep::_integrateDepthMaps() {
 	for (size_t i=0; i<cameras_.size(); ++i) {
-		ftl::cuda::clear_depth(*(cameras_[i].gpu.point_tex_), integ_stream_);
-		ftl::cuda::mls_smooth(*(cameras_[i].gpu.point_tex_), m_hashData, m_hashParams, cameras_.size(), i, integ_stream_);
+		ftl::cuda::clear_depth(*(cameras_[i].gpu.depth_tex_), integ_stream_);
+		ftl::cuda::mls_smooth(*(cameras_[i].gpu.depth_tex_), m_hashData, m_hashParams, cameras_.size(), i, integ_stream_);
 	}
 	ftl::cuda::integrateDepthMaps(m_hashData, m_hashParams, cameras_.size(), integ_stream_);
 }