From 3fc5561212f4f2dd43251754b0255a3257b1f3d8 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Wed, 18 Sep 2019 08:55:50 +0300
Subject: [PATCH] WIP Refactor of render
---
applications/reconstruct/CMakeLists.txt | 6 +--
applications/reconstruct/src/ilw.cpp | 2 +
applications/reconstruct/src/ilw.hpp | 2 +-
applications/reconstruct/src/main.cpp | 12 ++++-
components/renderers/cpp/src/splat_render.cpp | 46 +++++++++++++------
components/renderers/cpp/src/splatter.cu | 14 ++++--
6 files changed, 58 insertions(+), 24 deletions(-)
diff --git a/applications/reconstruct/CMakeLists.txt b/applications/reconstruct/CMakeLists.txt
index 9d9065d39..584e1e7a8 100644
--- a/applications/reconstruct/CMakeLists.txt
+++ b/applications/reconstruct/CMakeLists.txt
@@ -12,9 +12,9 @@ set(REPSRC
#src/virtual_source.cpp
#src/splat_render.cpp
#src/dibr.cu
- src/mls.cu
- src/depth_camera.cu
- src/depth_camera.cpp
+ #src/mls.cu
+ #src/depth_camera.cu
+ #src/depth_camera.cpp
src/ilw.cpp
)
diff --git a/applications/reconstruct/src/ilw.cpp b/applications/reconstruct/src/ilw.cpp
index bef1f806a..ae6b6ae4e 100644
--- a/applications/reconstruct/src/ilw.cpp
+++ b/applications/reconstruct/src/ilw.cpp
@@ -12,6 +12,8 @@ ILW::~ILW() {
}
bool ILW::process(ftl::rgbd::FrameSet &fs) {
+ return true;
+
_phase0(fs);
for (int i=0; i<2; ++i) {
diff --git a/applications/reconstruct/src/ilw.hpp b/applications/reconstruct/src/ilw.hpp
index 81d400629..1d269b83a 100644
--- a/applications/reconstruct/src/ilw.hpp
+++ b/applications/reconstruct/src/ilw.hpp
@@ -40,7 +40,7 @@ class ILW : public ftl::Configurable {
/**
* Take a frameset and perform the iterative lattice warping.
*/
- bool process(ftl::rgbd::FrameSet &in, ftl::rgbd::FrameSet &out);
+ bool process(ftl::rgbd::FrameSet &fs);
private:
/*
diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 36d4862a0..2af70b290 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/applications/reconstruct/src/main.cpp
@@ -16,7 +16,7 @@
#include <ftl/rgbd/group.hpp>
#include "ilw.hpp"
-#include "splat_render.hpp"
+#include <ftl/render/splat_render.hpp>
#include <string>
#include <vector>
@@ -100,7 +100,7 @@ static void run(ftl::Configurable *root) {
ftl::rgbd::VirtualSource *virt = ftl::create<ftl::rgbd::VirtualSource>(root, "virtual", net);
ftl::render::Splatter *splat = ftl::create<ftl::render::Splatter>(root, "renderer", scene);
ftl::rgbd::Group group;
- ftl::ILW align;
+ ftl::ILW *align = ftl::create<ftl::ILW>(root, "merge");
// Generate virtual camera render when requested by streamer
virt->onRender([splat,&scene_B](ftl::rgbd::Frame &out) {
@@ -149,6 +149,14 @@ static void run(ftl::Configurable *root) {
});
return true;
});
+
+ ftl::timer::stop();
+ net->shutdown();
+ delete align;
+ delete splat;
+ delete virt;
+ delete stream;
+ delete net;
}
int main(int argc, char **argv) {
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index fc92d5410..ebd2cb487 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -4,6 +4,7 @@
using ftl::render::Splatter;
using ftl::rgbd::Channel;
using ftl::rgbd::Format;
+using cv::cuda::GpuMat;
Splatter::Splatter(nlohmann::json &config, const ftl::rgbd::FrameSet &fs) : ftl::render::Renderer(config), scene_(fs) {
@@ -83,11 +84,16 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
//ftl::cuda::dibr(depth2_, colour1_, scene_->cameraCount(), params, stream);
//src->writeFrames(ts, colour1_, depth2_, stream);
//} else {
- ftl::cuda::clear_depth(depth1_, stream);
- ftl::cuda::clear_depth(depth3_, stream);
- ftl::cuda::clear_depth(depth2_, stream);
- ftl::cuda::clear_colour(colour2_, stream);
- ftl::cuda::dibr(depth1_, colour1_, normal1_, depth2_, colour_tmp_, depth3_, scene_->cameraCount(), params, stream);
+
+ //ftl::cuda::clear_depth(depth1_, stream);
+ //ftl::cuda::clear_depth(depth3_, 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);
// Step 1: Put all points into virtual view to gather them
//ftl::cuda::dibr_raw(depth1_, scene_->cameraCount(), params, stream);
@@ -99,18 +105,24 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
if (value("splatting", false)) {
//ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream);
- ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
- src->writeFrames(ts, colour1_, depth2_, 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);
+ //src->writeFrames(ts, colour1_, depth2_, stream);
+ src->write(scene_.timestamp, output_, stream);
} else {
- ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
- src->writeFrames(ts, colour1_, depth2_, stream);
+ temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, stream);
+ //ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
+ //src->writeFrames(ts, colour1_, depth2_, stream);
+ src->write(scene_.timestamp, output_, stream);
}
} else if (src->getChannel() == Channel::Energy) {
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
//if (src->value("splatting", false)) {
//ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream);
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
- src->writeFrames(ts, colour1_, depth2_, stream);
+ //src->writeFrames(ts, colour1_, depth2_, stream);
+ src->write(scene_.timestamp, output_, stream);
//} else {
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
// src->writeFrames(colour1_, depth2_, stream);
@@ -122,16 +134,20 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
params.m_viewMatrix = MatrixConversion::toCUDA(matrix.inverse());
params.m_viewMatrixInverse = MatrixConversion::toCUDA(matrix);
- ftl::cuda::clear_depth(depth1_, stream);
+ //ftl::cuda::clear_depth(depth1_, stream);
ftl::cuda::dibr(depth1_, colour1_, normal1_, depth2_, colour_tmp_, depth3_, scene_->cameraCount(), params, stream);
- src->writeFrames(ts, colour1_, colour2_, stream);
+ //src->writeFrames(ts, colour1_, colour2_, stream);
+ src->write(scene_.timestamp, output_, stream);
} else {
if (value("splatting", false)) {
//ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream);
- src->writeFrames(ts, colour1_, depth2_, stream);
+ //src->writeFrames(ts, colour1_, depth2_, stream);
+ src->write(scene_.timestamp, output_, stream);
} else {
- ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
- src->writeFrames(ts, colour1_, depth2_, 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);
+ //src->writeFrames(ts, colour1_, depth2_, stream);
+ src->write(scene_.timestamp, output_, stream);
}
}
//}
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index 586eb40fd..570a2b076 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -2,6 +2,14 @@
#include <ftl/rgbd/camera.hpp>
#include <ftl/cuda_common.hpp>
+#define T_PER_BLOCK 8
+#define UPSAMPLE_FACTOR 1.8f
+#define WARP_SIZE 32
+#define DEPTH_THRESHOLD 0.05f
+#define UPSAMPLE_MAX 60
+#define MAX_ITERATIONS 32 // Note: Must be multiple of 32
+#define SPATIAL_SMOOTHING 0.005f
+
using ftl::cuda::TextureObject;
using ftl::render::SplatParams;
@@ -13,7 +21,7 @@ using ftl::render::SplatParams;
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
- const float3 worldPos = make_float3(tex2D<float4>(points, x, y));
+ const float3 worldPos = make_float3(points.tex2D(x, y));
if (worldPos.x == MINF) return;
// Find the virtual screen position of current point
@@ -55,7 +63,7 @@ __global__ void dibr_attribute_contrib_kernel(
const int x = (blockIdx.x*blockDim.x + threadIdx.x) / WARP_SIZE;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
- const float3 worldPos = make_float3(tex2D<float4>(camera.points, x, y));
+ const float3 worldPos = make_float3(points.tex2D(x, y));
//const float3 normal = make_float3(tex2D<float4>(camera.normal, x, y));
if (worldPos.x == MINF) return;
//const float r = (camera.poseInverse * worldPos).z / camera.params.fx;
@@ -75,7 +83,7 @@ __global__ void dibr_attribute_contrib_kernel(
//if (abs(d - camPos.z) > DEPTH_THRESHOLD) return;
// TODO:(Nick) Should just one thread load these to shared mem?
- const float4 colour = make_float4(colour_in.tex2D<uchar4>(x, y));
+ const float4 colour = make_float4(colour_in.tex2D(x, y));
//const float4 normal = tex2D<float4>(camera.normal, x, y);
// Each thread in warp takes an upsample point and updates corresponding depth buffer.
--
GitLab