From 81a2a6fb4ee8a97ff3fead8f1a3d8dfa0bc018e1 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Thu, 31 Oct 2019 10:44:37 +0200
Subject: [PATCH] Simple mesh blend
---
.../renderers/cpp/src/splatter_cuda.hpp | 7 ++++
components/renderers/cpp/src/tri_render.cpp | 31 +++++++++++++--
.../renderers/cpp/src/triangle_render.cu | 38 +++++++++++++++++++
3 files changed, 73 insertions(+), 3 deletions(-)
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 46ef5012d..038f5a3e5 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -21,6 +21,13 @@ namespace cuda {
ftl::cuda::TextureObject<short2> &screen,
const ftl::render::SplatParams ¶ms,
cudaStream_t stream);
+
+ void mesh_blender(
+ ftl::cuda::TextureObject<int> &depth_in,
+ ftl::cuda::TextureObject<int> &depth_out,
+ const ftl::rgbd::Camera &camera,
+ float alpha,
+ cudaStream_t stream);
void dibr_merge(
ftl::cuda::TextureObject<float4> &points,
diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp
index c579a04bd..be84e5a64 100644
--- a/components/renderers/cpp/src/tri_render.cpp
+++ b/components/renderers/cpp/src/tri_render.cpp
@@ -292,8 +292,17 @@ void Triangular::_dibr(cudaStream_t stream) {
void Triangular::_mesh(cudaStream_t stream) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
- temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
+ bool do_blend = value("mesh_blend", true);
+ float blend_alpha = value("blend_alpha", 0.02f);
+ if (do_blend) {
+ temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
+ temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
+ } else {
+ temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
+ }
+
+ // For each source depth map
for (size_t i=0; i < scene_->frames.size(); ++i) {
auto &f = scene_->frames[i];
auto *s = scene_->sources[i];
@@ -305,6 +314,7 @@ void Triangular::_mesh(cudaStream_t stream) {
auto pose = MatrixConversion::toCUDA(s->getPose().cast<float>());
+ // Calculate and save virtual view screen position of each source pixel
ftl::cuda::screen_coord(
f.createTexture<float>(Channel::Depth),
f.createTexture<float>(Channel::Depth2, Format<float>(f.get<GpuMat>(Channel::Depth).size())),
@@ -312,14 +322,29 @@ void Triangular::_mesh(cudaStream_t stream) {
params_, pose, s->parameters(), stream
);
+ // Must reset depth channel if blending
+ if (do_blend) {
+ temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
+ }
+
+ // Decide on and render triangles around each point
ftl::cuda::triangle_render1(
f.getTexture<float>(Channel::Depth2),
- temp_.createTexture<int>(Channel::Depth2),
+ temp_.createTexture<int>((do_blend) ? Channel::Depth : Channel::Depth2),
f.getTexture<short2>(Channel::Screen),
params_, stream
);
- //LOG(INFO) << "DIBR DONE";
+ if (do_blend) {
+ // Blend this sources mesh with previous meshes
+ ftl::cuda::mesh_blender(
+ temp_.getTexture<int>(Channel::Depth),
+ temp_.createTexture<int>(Channel::Depth2),
+ params_.camera,
+ blend_alpha,
+ stream
+ );
+ }
}
}
diff --git a/components/renderers/cpp/src/triangle_render.cu b/components/renderers/cpp/src/triangle_render.cu
index 77694ed1d..cdf42fff9 100644
--- a/components/renderers/cpp/src/triangle_render.cu
+++ b/components/renderers/cpp/src/triangle_render.cu
@@ -3,6 +3,8 @@
#include <ftl/rgbd/camera.hpp>
#include <ftl/cuda_common.hpp>
+#include <ftl/cuda/weighting.hpp>
+
using ftl::rgbd::Camera;
using ftl::cuda::TextureObject;
using ftl::render::SplatParams;
@@ -162,3 +164,39 @@ void ftl::cuda::triangle_render1(TextureObject<float> &depth_in, TextureObject<i
triangle_render_1_kernel<-1,-1><<<gridSize, blockSize, 0, stream>>>(depth_in, depth_out, screen, params);
cudaSafeCall( cudaGetLastError() );
}
+
+// ==== BLENDER ========
+
+/*
+ * Merge two depth maps together
+ */
+ __global__ void mesh_blender_kernel(
+ TextureObject<int> depth_in,
+ TextureObject<int> depth_out,
+ ftl::rgbd::Camera camera,
+ float alpha) {
+ 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;
+
+ int a = depth_in.tex2D(x,y);
+ int b = depth_out.tex2D(x,y);
+
+ float mindepth = (float)min(a,b) / 100000.0f;
+ float maxdepth = (float)max(a,b) / 100000.0f;
+ float weight = ftl::cuda::weighting(maxdepth-mindepth, alpha);
+
+ //depth_out(x,y) = (int)(((float)mindepth + (float)maxdepth*weight) / (1.0f + weight) * 100000.0f);
+
+ float depth = (mindepth + maxdepth*weight) / (1.0f + weight);
+ depth_out(x,y) = (int)(depth * 100000.0f);
+}
+
+void ftl::cuda::mesh_blender(TextureObject<int> &depth_in, TextureObject<int> &depth_out, const ftl::rgbd::Camera &camera, float alpha, 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);
+
+ mesh_blender_kernel<<<gridSize, blockSize, 0, stream>>>(depth_in, depth_out, camera, alpha);
+ cudaSafeCall( cudaGetLastError() );
+}
--
GitLab