WIP correct render

components/renderers/cpp/src/CUDARender.cpp

@@ -363,6 +363,37 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
 	// Now merge new render to any existing frameset render, detecting collisions
 	ftl::cuda::touch_merge(depth_out_, out.createTexture<float>(_getDepthChannel()), collisions_, 1024, touch_dist_, stream_);
 
+	// For each source depth map, verify results
+	for (size_t i=0; i < scene_->frames.size(); ++i) {
+		//if (!scene_->hasFrame(i)) continue;
+		auto &f = scene_->frames[i].cast<ftl::rgbd::Frame>();
+		//auto *s = scene_->sources[i];
+
+		if (!f.has(Channel::Colour)) {
+			//LOG(ERROR) << "Missing required channel";
+			continue;
+		}
+
+		// We have the needed depth data?
+		if (use_depth && !f.hasOwn(Channel::Depth) && !f.hasOwn(Channel::GroundTruth)) {
+			continue;
+		}
+
+		//auto pose = MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
+		auto transformR = MatrixConversion::toCUDA(f.getPose().cast<float>().inverse() * t.cast<float>().inverse()) * poseInverse_;
+		auto transform = pose_ * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
+
+		ftl::cuda::reverse_verify(
+			out.getTexture<float>(_getDepthChannel()),
+			f.getTexture<float>(Channel::Depth),
+			transformR,
+			transform,
+			params_.camera,
+			f.getLeft(),
+			stream_
+		);
+	}
+
 	//filters_->filter(out, src, stream);
 
 	// Generate normals for final virtual image

components/renderers/cpp/src/splatter_cuda.hpp

@@ -170,6 +170,15 @@ namespace cuda {
 		ftl::cuda::TextureObject<float> &d2,
         float factor, cudaStream_t stream);
 
+	void reverse_verify(
+		ftl::cuda::TextureObject<float> &depth_in,
+		ftl::cuda::TextureObject<float> &depth_original,
+		const float4x4 &transformR,
+		const float4x4 &transform,
+		const ftl::rgbd::Camera &vintrin,
+		const ftl::rgbd::Camera &ointrin,
+		cudaStream_t stream);
+
 	void fix_bad_colour(
 		ftl::cuda::TextureObject<float> &depth,
 		ftl::cuda::TextureObject<uchar4> &out,

components/renderers/cpp/src/triangle_render.cu

@@ -202,7 +202,8 @@ __global__ void merge_convert_kernel(
 
 	if (x < 0 || x >= depth_in.width() || y < 0 || y >= depth_in.height()) return;
 
-	float a = float(depth_in.tex2D(x,y))*alpha;
+	int d = depth_in.tex2D(x,y);
+	float a = float(d)*alpha;
 	float b = depth_out.tex2D(x,y);
 	depth_out(x,y) = min(a,b);
 }
@@ -215,6 +216,47 @@ void ftl::cuda::merge_convert_depth(TextureObject<int> &depth_in, TextureObject<
 	cudaSafeCall( cudaGetLastError() );
 }
 
+// ==== Reverse Verify Result ==================================================
+
+// ==== Merge convert ===========
+
+__global__ void reverse_check_kernel(
+	TextureObject<float> depth_in,
+	TextureObject<float> depth_original,
+	float4x4 transformR,
+	float4x4 transform,
+	ftl::rgbd::Camera vintrin,
+	ftl::rgbd::Camera ointrin
+) {
+	const int x = blockIdx.x*blockDim.x + threadIdx.x;
+	const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+	if (x < 0 || x >= depth_in.width() || y < 0 || y >= depth_in.height()) return;
+
+	float d = depth_in.tex2D(x,y);
+	float3 campos = transformR * vintrin.screenToCam(x,y,d);
+	int2 spos = ointrin.camToScreen<int2>(campos);
+	int ox = spos.x;
+	int oy = spos.y;
+
+	if (campos.z > 0.0f && ox >= 0 && ox < ointrin.width && oy >= 0 && oy < ointrin.height) {
+		float d2 = depth_original.tex2D(ox,oy);
+
+		if (d2 < ointrin.maxDepth && d2 - campos.z > 0.002f) {
+			//printf("Original %f, %f\n", d2, campos.z);
+			depth_in(x,y) = 0.0f; //(transform * ointrin.screenToCam(ox,oy,d2)).z;
+		}
+	}
+}
+
+void ftl::cuda::reverse_verify(TextureObject<float> &depth_in, TextureObject<float> &depth_original, const float4x4 &transformR, const float4x4 &transform, const ftl::rgbd::Camera &vintrin, const ftl::rgbd::Camera &ointrin, cudaStream_t stream) {
+	const dim3 gridSize((depth_in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+	const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+	reverse_check_kernel<<<gridSize, blockSize, 0, stream>>>(depth_in, depth_original, transformR, transform, vintrin, ointrin);
+	cudaSafeCall( cudaGetLastError() );
+}
+
 // ==== BLENDER ========
 
 /*