From ca26f80f62a7c5cfbb9385c60cb805706c9ae7c8 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Sat, 31 Oct 2020 19:34:10 +0200
Subject: [PATCH] WIP correct render
---
components/renderers/cpp/src/CUDARender.cpp | 31 +++++++++++++
.../renderers/cpp/src/splatter_cuda.hpp | 9 ++++
.../renderers/cpp/src/triangle_render.cu | 44 ++++++++++++++++++-
3 files changed, 83 insertions(+), 1 deletion(-)
diff --git a/components/renderers/cpp/src/CUDARender.cpp b/components/renderers/cpp/src/CUDARender.cpp
index dab897d56..c274e136f 100644
--- a/components/renderers/cpp/src/CUDARender.cpp
+++ b/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
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 41088cc22..147ccf994 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/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,
diff --git a/components/renderers/cpp/src/triangle_render.cu b/components/renderers/cpp/src/triangle_render.cu
index dccdd276a..3b45d1315 100644
--- a/components/renderers/cpp/src/triangle_render.cu
+++ b/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 ========
/*
--
GitLab