From 5efd2b7c35209673445fc704cda95db10d4710b5 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Wed, 18 Sep 2019 10:07:35 +0300
Subject: [PATCH] WIP fixing compile problems in refactor
---
applications/reconstruct/CMakeLists.txt | 2 +-
.../cpp/include/ftl/cuda/weighting.hpp | 23 ++++++++++++
.../cpp/include/ftl/render/splat_render.hpp | 2 +-
.../renderers/cpp}/src/mls_cuda.hpp | 0
components/renderers/cpp/src/splat_render.cpp | 36 ++++++++++---------
components/renderers/cpp/src/splatter.cu | 10 +++++-
.../rgbd-sources/include/ftl/rgbd/source.hpp | 2 +-
.../rgbd-sources/include/ftl/rgbd/virtual.hpp | 1 +
8 files changed, 56 insertions(+), 20 deletions(-)
create mode 100644 components/renderers/cpp/include/ftl/cuda/weighting.hpp
rename {applications/reconstruct => components/renderers/cpp}/src/mls_cuda.hpp (100%)
diff --git a/applications/reconstruct/CMakeLists.txt b/applications/reconstruct/CMakeLists.txt
index 584e1e7a8..de3aca69c 100644
--- a/applications/reconstruct/CMakeLists.txt
+++ b/applications/reconstruct/CMakeLists.txt
@@ -32,6 +32,6 @@ set_property(TARGET ftl-reconstruct PROPERTY CUDA_SEPARABLE_COMPILATION ON)
endif()
#target_include_directories(cv-node PUBLIC ${PROJECT_SOURCE_DIR}/include)
-target_link_libraries(ftl-reconstruct ftlcommon ftlrgbd Threads::Threads ${OpenCV_LIBS} ftlctrl ftlnet)
+target_link_libraries(ftl-reconstruct ftlcommon ftlrgbd Threads::Threads ${OpenCV_LIBS} ftlctrl ftlnet ftlrender)
diff --git a/components/renderers/cpp/include/ftl/cuda/weighting.hpp b/components/renderers/cpp/include/ftl/cuda/weighting.hpp
new file mode 100644
index 000000000..9498d0508
--- /dev/null
+++ b/components/renderers/cpp/include/ftl/cuda/weighting.hpp
@@ -0,0 +1,23 @@
+#ifndef _FTL_CUDA_WEIGHTING_HPP_
+#define _FTL_CUDA_WEIGHTING_HPP_
+
+namespace ftl {
+namespace cuda {
+
+/*
+ * Guennebaud, G.; Gross, M. Algebraic point set surfaces. ACMTransactions on Graphics Vol. 26, No. 3, Article No. 23, 2007.
+ * Used in: FusionMLS: Highly dynamic 3D reconstruction with consumer-grade RGB-D cameras
+ * r = distance between points
+ * h = smoothing parameter in meters (default 4cm)
+ */
+__device__ inline float spatialWeighting(float r, float h) {
+ if (r >= h) return 0.0f;
+ float rh = r / h;
+ rh = 1.0f - rh*rh;
+ return rh*rh*rh*rh;
+}
+
+}
+}
+
+#endif // _FTL_CUDA_WEIGHTING_HPP_
diff --git a/components/renderers/cpp/include/ftl/render/splat_render.hpp b/components/renderers/cpp/include/ftl/render/splat_render.hpp
index 5357bf5c3..89934aa5b 100644
--- a/components/renderers/cpp/include/ftl/render/splat_render.hpp
+++ b/components/renderers/cpp/include/ftl/render/splat_render.hpp
@@ -21,7 +21,7 @@ class Splatter : public ftl::render::Renderer {
explicit Splatter(nlohmann::json &config, const ftl::rgbd::FrameSet &fs);
~Splatter();
- void render(ftl::rgbd::VirtualSource *src, cudaStream_t stream=0);
+ bool render(ftl::rgbd::VirtualSource *src, cudaStream_t stream=0) override;
//void setOutputDevice(int);
diff --git a/applications/reconstruct/src/mls_cuda.hpp b/components/renderers/cpp/src/mls_cuda.hpp
similarity index 100%
rename from applications/reconstruct/src/mls_cuda.hpp
rename to components/renderers/cpp/src/mls_cuda.hpp
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index ebd2cb487..dc07c126e 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -14,22 +14,24 @@ Splatter::~Splatter() {
}
-void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
+bool Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
if (!src->isReady()) return;
const auto &camera = src->parameters();
//cudaSafeCall(cudaSetDevice(scene_->getCUDADevice()));
- output_.create<cv::cuda::GpuMat>(Channel::Depth, Format<float>(camera.width, camera.height));
- output_.create<cv::cuda::GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height));
+ output_.create<GpuMat>(Channel::Depth, Format<float>(camera.width, camera.height));
+ output_.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height));
- temp_.create<cv::cuda::GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
- temp_.create<cv::cuda::GpuMat>(Channel::Colour2, Format<uchar4>(camera.width, camera.height));
- temp_.create<cv::cuda::GpuMat>(Channel::Confidence, Format<float>(camera.width, camera.height));
- temp_.create<cv::cuda::GpuMat>(Channel::Depth, Format<int>(camera.width, camera.height));
- temp_.create<cv::cuda::GpuMat>(Channel::Depth2, Format<int>(camera.width, camera.height));
- temp_.create<cv::cuda::GpuMat>(Channel::Normals, Format<float4>(camera.width, camera.height));
+ temp_.create<GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
+ temp_.create<GpuMat>(Channel::Colour2, Format<uchar4>(camera.width, camera.height));
+ temp_.create<GpuMat>(Channel::Confidence, Format<float>(camera.width, camera.height));
+ temp_.create<GpuMat>(Channel::Depth, Format<int>(camera.width, camera.height));
+ temp_.create<GpuMat>(Channel::Depth2, Format<int>(camera.width, camera.height));
+ temp_.create<GpuMat>(Channel::Normals, Format<float4>(camera.width, camera.height));
+
+ cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
// Create buffers if they don't exist
/*if ((unsigned int)depth1_.width() != camera.width || (unsigned int)depth1_.height() != camera.height) {
@@ -90,10 +92,10 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
//ftl::cuda::clear_depth(depth2_, stream);
//ftl::cuda::clear_colour(colour2_, stream);
- temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0), stream);
- temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0), stream);
- output_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0.0f), stream);
- output_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0,0,0), stream);
+ temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0), cvstream);
+ temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0), cvstream);
+ output_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0.0f), cvstream);
+ output_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0,0,0), cvstream);
// Step 1: Put all points into virtual view to gather them
//ftl::cuda::dibr_raw(depth1_, scene_->cameraCount(), params, stream);
@@ -107,11 +109,11 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
//ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream);
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
- temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, stream);
+ temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
//src->writeFrames(ts, colour1_, depth2_, stream);
src->write(scene_.timestamp, output_, stream);
} else {
- temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, stream);
+ temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
//src->writeFrames(ts, colour1_, depth2_, stream);
src->write(scene_.timestamp, output_, stream);
@@ -145,7 +147,7 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
src->write(scene_.timestamp, output_, stream);
} else {
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
- temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, stream);
+ temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
//src->writeFrames(ts, colour1_, depth2_, stream);
src->write(scene_.timestamp, output_, stream);
}
@@ -156,6 +158,8 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
//ftl::cuda::splat_points(depth1_, depth2_, params, stream);
// TODO: Second pass
+
+ return true;
}
void Splatter::setOutputDevice(int device) {
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index 570a2b076..3c7d94f39 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -2,6 +2,8 @@
#include <ftl/rgbd/camera.hpp>
#include <ftl/cuda_common.hpp>
+#include <ftl/cuda/weighting.hpp>
+
#define T_PER_BLOCK 8
#define UPSAMPLE_FACTOR 1.8f
#define WARP_SIZE 32
@@ -44,6 +46,12 @@ __device__ inline float4 make_float4(const uchar4 &c) {
return make_float4(c.x,c.y,c.z,c.w);
}
+
+#define ENERGY_THRESHOLD 0.1f
+#define SMOOTHING_MULTIPLIER_A 10.0f // For surface search
+#define SMOOTHING_MULTIPLIER_B 4.0f // For z contribution
+#define SMOOTHING_MULTIPLIER_C 4.0f // For colour contribution
+
/*
* Pass 2: Accumulate attribute contributions if the points pass a visibility test.
*/
@@ -100,7 +108,7 @@ __global__ void dibr_attribute_contrib_kernel(
const float weight = ftl::cuda::spatialWeighting(length(nearest - camPos), SMOOTHING_MULTIPLIER_C*(nearest.z/params.camera.fx));
if (screenPos.x+u < colour_out.width() && screenPos.y+v < colour_out.height() && weight > 0.0f) { // TODO: Use confidence threshold here
const float4 wcolour = colour * weight;
- const float4 wnormal = normal * weight;
+ //const float4 wnormal = normal * weight;
//printf("Z %f\n", d);
diff --git a/components/rgbd-sources/include/ftl/rgbd/source.hpp b/components/rgbd-sources/include/ftl/rgbd/source.hpp
index e14950e63..5d296e699 100644
--- a/components/rgbd-sources/include/ftl/rgbd/source.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/source.hpp
@@ -55,7 +55,7 @@ class Source : public ftl::Configurable {
explicit Source(ftl::config::json_t &cfg);
Source(ftl::config::json_t &cfg, ftl::rgbd::SnapshotReader *);
Source(ftl::config::json_t &cfg, ftl::net::Universe *net);
- ~Source();
+ virtual ~Source();
public:
/**
diff --git a/components/rgbd-sources/include/ftl/rgbd/virtual.hpp b/components/rgbd-sources/include/ftl/rgbd/virtual.hpp
index e8e8edd40..55dfb9e89 100644
--- a/components/rgbd-sources/include/ftl/rgbd/virtual.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/virtual.hpp
@@ -9,6 +9,7 @@ namespace rgbd {
class VirtualSource : public ftl::rgbd::Source {
public:
explicit VirtualSource(ftl::config::json_t &cfg);
+ ~VirtualSource();
/**
* Write frames into source buffers from an external renderer. Virtual
--
GitLab