Fix for odd no camera bug

applications/reconstruct/include/ftl/depth_camera.hpp

@@ -46,7 +46,7 @@ struct DepthCamera {
 
 	__host__ void free();
 
-	__host__ void _computeNormals(cudaStream_t stream);
+	__host__ void computeNormals(cudaStream_t stream);
 
 	ftl::cuda::TextureObject<float> *depth_tex_;
 	ftl::cuda::TextureObject<int> *depth2_tex_;

applications/reconstruct/src/depth_camera.cpp

@@ -45,11 +45,11 @@ void DepthCamera::updateData(const cv::Mat &depth, const cv::Mat &rgb, cv::cuda:
 	depth_tex_->upload(depth, cv::cuda::StreamAccessor::getStream(stream));
 	colour_tex_->upload(rgb, cv::cuda::StreamAccessor::getStream(stream));
 	//if (normal_mat_) {
-		_computeNormals(cv::cuda::StreamAccessor::getStream(stream));
+		//computeNormals(cv::cuda::StreamAccessor::getStream(stream));
 	//}
 }
 
-void DepthCamera::_computeNormals(cudaStream_t stream) {
-	ftl::cuda::point_cloud(*points_tex_, data, stream);
-	ftl::cuda::compute_normals(*points_tex_, *normal_tex_, stream);
+void DepthCamera::computeNormals(cudaStream_t stream) {
+	//ftl::cuda::point_cloud(*points_tex_, data, stream);
+	ftl::cuda::compute_normals(*depth_tex_, *normal_tex_, data, stream);
 }

applications/reconstruct/src/depth_camera.cu

@@ -604,7 +604,7 @@ void ftl::cuda::point_cloud(ftl::cuda::TextureObject<float4> &output, const Dept
 /// ===== NORMALS =====
 
 
-__global__ void compute_normals_kernel(const ftl::cuda::TextureObject<float4> input, ftl::cuda::TextureObject<float4> output)
+__global__ void compute_normals_kernel(const ftl::cuda::TextureObject<float> input, ftl::cuda::TextureObject<float4> output, DepthCameraCUDA camera)
 {
 	const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
@@ -617,11 +617,11 @@ __global__ void compute_normals_kernel(const ftl::cuda::TextureObject<float4> in
 
 	if(x > 0 && x < output.width()-1 && y > 0 && y < output.height()-1)
 	{
-		const float3 CC = make_float3(input(x,y)); //input[(y+0)*width+(x+0)];
-		const float3 PC = make_float3(input(x,y+1)); //input[(y+1)*width+(x+0)];
-		const float3 CP = make_float3(input(x+1,y)); //input[(y+0)*width+(x+1)];
-		const float3 MC = make_float3(input(x,y-1)); //input[(y-1)*width+(x+0)];
-		const float3 CM = make_float3(input(x-1,y)); //input[(y+0)*width+(x-1)];
+		const float3 CC = camera.pose * camera.params.kinectDepthToSkeleton(x,y,input(x,y)); //input[(y+0)*width+(x+0)];
+		const float3 PC = camera.pose * camera.params.kinectDepthToSkeleton(x,y,input(x,y+1)); //input[(y+1)*width+(x+0)];
+		const float3 CP = camera.pose * camera.params.kinectDepthToSkeleton(x,y,input(x+1,y)); //input[(y+0)*width+(x+1)];
+		const float3 MC = camera.pose * camera.params.kinectDepthToSkeleton(x,y,input(x,y-1)); //input[(y-1)*width+(x+0)];
+		const float3 CM = camera.pose * camera.params.kinectDepthToSkeleton(x,y,input(x-1,y)); //input[(y+0)*width+(x-1)];
 
 		if(CC.x != MINF && PC.x != MINF && CP.x != MINF && MC.x != MINF && CM.x != MINF)
 		{
@@ -637,11 +637,11 @@ __global__ void compute_normals_kernel(const ftl::cuda::TextureObject<float4> in
 	}
 }
 
-void ftl::cuda::compute_normals(const ftl::cuda::TextureObject<float4> &input, ftl::cuda::TextureObject<float4> &output, cudaStream_t stream) {
+void ftl::cuda::compute_normals(const ftl::cuda::TextureObject<float> &input, ftl::cuda::TextureObject<float4> &output, const DepthCameraCUDA &camera, 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);
 
-	compute_normals_kernel<<<gridSize, blockSize, 0, stream>>>(input, output);
+	compute_normals_kernel<<<gridSize, blockSize, 0, stream>>>(input, output, camera);
 
 #ifdef _DEBUG
 	cudaSafeCall(cudaDeviceSynchronize());

applications/reconstruct/src/depth_camera_cuda.hpp

@@ -27,7 +27,7 @@ void hole_fill(const TextureObject<int> &depth_in, const TextureObject<float> &d
 
 void point_cloud(ftl::cuda::TextureObject<float4> &output, const ftl::voxhash::DepthCameraCUDA &depthCameraData, cudaStream_t stream);
 
-void compute_normals(const ftl::cuda::TextureObject<float4> &points, ftl::cuda::TextureObject<float4> &normals, cudaStream_t stream);
+void compute_normals(const ftl::cuda::TextureObject<float> &depth, ftl::cuda::TextureObject<float4> &normals, const ftl::voxhash::DepthCameraCUDA &camera, cudaStream_t stream);
 
 }
 }

applications/reconstruct/src/voxel_scene.cpp

@@ -368,6 +368,7 @@ void SceneRep::_integrateDepthMaps() {
 	for (size_t i=0; i<cameras_.size(); ++i) {
 		if (!cameras_[i].source->isReady()) continue;
 		//ftl::cuda::clear_depth(*(cameras_[i].gpu.depth2_tex_), integ_stream_);
+		cameras_[i].gpu.computeNormals(integ_stream_);
 		ftl::cuda::clear_points(*(cameras_[i].gpu.points_tex_), integ_stream_);
 		ftl::cuda::mls_smooth(*(cameras_[i].gpu.points_tex_), m_hashParams, cameras_.size(), i, integ_stream_);
 		//ftl::cuda::int_to_float(*(cameras_[i].gpu.depth2_tex_), *(cameras_[i].gpu.depth_tex_), 1.0f / 1000.0f, integ_stream_);