From fac2ab56ab738731cc1ddf15cbc93206370af984 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nicolas.pope@utu.fi>
Date: Wed, 13 Nov 2019 19:30:45 +0200
Subject: [PATCH] Implements #244 360 environment background
---
.../cpp/include/ftl/render/tri_render.hpp | 3 ++
components/renderers/cpp/src/reprojection.cu | 48 +++++++++++++++++++
.../renderers/cpp/src/splatter_cuda.hpp | 5 ++
components/renderers/cpp/src/tri_render.cpp | 24 +++++++++-
4 files changed, 79 insertions(+), 1 deletion(-)
diff --git a/components/renderers/cpp/include/ftl/render/tri_render.hpp b/components/renderers/cpp/include/ftl/render/tri_render.hpp
index 60d3dcbf6..3d9183643 100644
--- a/components/renderers/cpp/include/ftl/render/tri_render.hpp
+++ b/components/renderers/cpp/include/ftl/render/tri_render.hpp
@@ -43,6 +43,9 @@ class Triangular : public ftl::render::Renderer {
cudaStream_t stream_;
float3 light_pos_;
+ cv::cuda::GpuMat env_image_;
+ ftl::cuda::TextureObject<uchar4> env_tex_;
+
//ftl::Filters *filters_;
template <typename T>
diff --git a/components/renderers/cpp/src/reprojection.cu b/components/renderers/cpp/src/reprojection.cu
index f1a9022a5..8a075b653 100644
--- a/components/renderers/cpp/src/reprojection.cu
+++ b/components/renderers/cpp/src/reprojection.cu
@@ -274,3 +274,51 @@ template void ftl::cuda::reproject(
const ftl::render::SplatParams ¶ms,
const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream);
+
+
+// ===== Equirectangular Reprojection ==========================================
+
+__device__ inline float2 equirect_reprojection(int x_img, int y_img, double f, const float3x3 &rot, int w1, int h1, const ftl::rgbd::Camera &cam) {
+ float3 ray3d = cam.screenToCam(x_img, y_img, 1.0f);
+ ray3d /= length(ray3d);
+ ray3d = rot * ray3d;
+
+ //inverse formula for spherical projection, reference Szeliski book "Computer Vision: Algorithms and Applications" p439.
+ float theta = atan2(ray3d.y,sqrt(ray3d.x*ray3d.x+ray3d.z*ray3d.z));
+ float phi = atan2(ray3d.x, ray3d.z);
+
+ const float pi = 3.14f;
+
+ //get 2D point on equirectangular map
+ float x_sphere = (((phi*w1)/pi+w1)/2);
+ float y_sphere = (theta+ pi/2)*h1/pi;
+
+ return make_float2(x_sphere,y_sphere);
+};
+
+__global__ void equirectangular_kernel(
+ TextureObject<uchar4> image_in,
+ TextureObject<uchar4> image_out,
+ Camera camera, float3x3 pose) {
+
+ const int x = (blockIdx.x*blockDim.x + threadIdx.x);
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x >= 0 && y >= 0 && x < image_out.width() && y < image_out.height()) {
+ const float2 p = equirect_reprojection(x,y, camera.fx, pose, image_in.width(), image_in.height(), camera);
+ const float4 colour = image_in.tex2D(p.x, p.y);
+ image_out(x,y) = make_uchar4(colour.x, colour.y, colour.z, 0);
+ }
+}
+
+void ftl::cuda::equirectangular_reproject(
+ ftl::cuda::TextureObject<uchar4> &image_in,
+ ftl::cuda::TextureObject<uchar4> &image_out,
+ const ftl::rgbd::Camera &camera, const float3x3 &pose, cudaStream_t stream) {
+
+ const dim3 gridSize((image_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (image_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ equirectangular_kernel<<<gridSize, blockSize, 0, stream>>>(image_in, image_out, camera, pose);
+ cudaSafeCall( cudaGetLastError() );
+}
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 838eda409..010a63718 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -85,6 +85,11 @@ namespace cuda {
const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream);
+ void equirectangular_reproject(
+ ftl::cuda::TextureObject<uchar4> &image_in,
+ ftl::cuda::TextureObject<uchar4> &image_out,
+ const ftl::rgbd::Camera &camera, const float3x3 &pose, cudaStream_t stream);
+
template <typename A, typename B>
void dibr_normalise(
ftl::cuda::TextureObject<A> &in,
diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp
index f9e89ffb3..b3fba82a7 100644
--- a/components/renderers/cpp/src/tri_render.cpp
+++ b/components/renderers/cpp/src/tri_render.cpp
@@ -130,6 +130,19 @@ Triangular::Triangular(nlohmann::json &config, ftl::rgbd::FrameSet *fs) : ftl::r
on("light_y", [this](const ftl::config::Event &e) { light_pos_.y = value("light_y", 0.3f); });
on("light_z", [this](const ftl::config::Event &e) { light_pos_.z = value("light_z", 0.3f); });
+ // Load any environment texture
+ std::string envimage = value("environment", std::string(""));
+ if (envimage.size() > 0) {
+ cv::Mat envim = cv::imread(envimage);
+ if (!envim.empty()) {
+ if (envim.type() == CV_8UC3) {
+ cv::cvtColor(envim,envim, cv::COLOR_BGR2BGRA);
+ }
+ env_image_.upload(envim);
+ env_tex_ = std::move(ftl::cuda::TextureObject<uchar4>(env_image_, true));
+ }
+ }
+
cudaSafeCall(cudaStreamCreate(&stream_));
//filters_ = ftl::create<ftl::Filters>(this, "filters");
@@ -485,7 +498,16 @@ bool Triangular::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out) {
// Clear all channels to 0 or max depth
out.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(1000.0f), cvstream);
- out.get<GpuMat>(Channel::Colour).setTo(background_, cvstream);
+
+ if (env_image_.empty() || !value("environment_enabled", false)) {
+ out.get<GpuMat>(Channel::Colour).setTo(background_, cvstream);
+ } else {
+ auto pose = params.m_viewMatrixInverse.getFloat3x3();
+ ftl::cuda::equirectangular_reproject(
+ env_tex_,
+ out.createTexture<uchar4>(Channel::Colour, true),
+ camera, pose, stream_);
+ }
//LOG(INFO) << "Render ready: " << camera.width << "," << camera.height;
--
GitLab