From 15800bc3cc44014f7f403fcb38e503c511431012 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Thu, 19 Sep 2019 09:15:36 +0300
Subject: [PATCH] Almost working attribute accum
---
components/renderers/cpp/src/splat_render.cpp | 40 +++++++++++-
components/renderers/cpp/src/splatter.cu | 63 ++++++++++++++-----
.../renderers/cpp/src/splatter_cuda.hpp | 15 +++++
.../include/ftl/rgbd/channels.hpp | 7 ++-
4 files changed, 104 insertions(+), 21 deletions(-)
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index db6fa402d..25dc21c01 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -33,7 +33,7 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
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::Contribution, 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));
@@ -81,6 +81,8 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
//LOG(INFO) << "Render ready: " << camera.width << "," << camera.height;
+ temp_.createTexture<int>(Channel::Depth);
+
// Render each camera into virtual view
for (size_t i=0; i<scene_->frames.size(); ++i) {
auto &f = scene_->frames[i];
@@ -104,13 +106,47 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
ftl::cuda::dibr_merge(
f.createTexture<float4>(Channel::Points),
- temp_.createTexture<int>(Channel::Depth),
+ temp_.getTexture<int>(Channel::Depth),
params, stream
);
//LOG(INFO) << "DIBR DONE";
}
+ temp_.createTexture<float4>(Channel::Colour);
+ temp_.createTexture<float>(Channel::Contribution);
+
+ // Accumulate attribute contributions for each pixel
+ for (auto &f : scene_->frames) {
+ // Convert colour from BGR to BGRA if needed
+ if (f.get<GpuMat>(Channel::Colour).type() == CV_8UC3) {
+ // Convert to 4 channel colour
+ auto &col = f.get<GpuMat>(Channel::Colour);
+ GpuMat tmp(col.size(), CV_8UC4);
+ cv::cuda::swap(col, tmp);
+ cv::cuda::cvtColor(tmp,col, cv::COLOR_BGR2BGRA);
+ }
+
+ ftl::cuda::dibr_attribute(
+ f.createTexture<uchar4>(Channel::Colour),
+ f.createTexture<float4>(Channel::Points),
+ temp_.getTexture<int>(Channel::Depth),
+ temp_.getTexture<float4>(Channel::Colour),
+ temp_.getTexture<float>(Channel::Contribution),
+ params, stream
+ );
+ }
+
+ // Normalise attribute contributions
+ //for (auto &f : scene_->frames) {
+ ftl::cuda::dibr_normalise(
+ temp_.createTexture<float4>(Channel::Colour),
+ out.createTexture<uchar4>(Channel::Colour),
+ temp_.createTexture<float>(Channel::Contribution),
+ stream
+ );
+ //}
+
//ftl::cuda::dibr(depth1_, colour1_, normal1_, depth2_, colour_tmp_, depth3_, scene_->cameraCount(), params, stream);
// Step 1: Put all points into virtual view to gather them
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index ad3a6e8a9..c1b46fc1d 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -51,6 +51,8 @@ void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<int> &de
cudaSafeCall( cudaGetLastError() );
}
+//==============================================================================
+
__device__ inline float4 make_float4(const uchar4 &c) {
return make_float4(c.x,c.y,c.z,c.w);
}
@@ -135,25 +137,54 @@ __global__ void dibr_attribute_contrib_kernel(
}
}
+void ftl::cuda::dibr_attribute(
+ TextureObject<uchar4> &colour_in, // Original colour image
+ TextureObject<float4> &points, // Original 3D points
+ TextureObject<int> &depth_in, // Virtual depth map
+ TextureObject<float4> &colour_out, // Accumulated output
+ //TextureObject<float4> normal_out,
+ TextureObject<float> &contrib_out,
+ SplatParams ¶ms, cudaStream_t stream) {
+ const dim3 gridSize((depth_in.width() + 2 - 1)/2, (depth_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(2*WARP_SIZE, T_PER_BLOCK);
+
+ dibr_attribute_contrib_kernel<<<gridSize, blockSize, 0, stream>>>(
+ colour_in,
+ points,
+ depth_in,
+ colour_out,
+ contrib_out,
+ params
+ );
+ cudaSafeCall( cudaGetLastError() );
+}
-
+//==============================================================================
__global__ void dibr_normalise_kernel(
- TextureObject<float4> colour_in,
- TextureObject<uchar4> colour_out,
- //TextureObject<float4> normals,
- TextureObject<float> contribs) {
-const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
-const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
-
-if (x < colour_in.width() && y < colour_in.height()) {
- const float4 colour = colour_in.tex2D((int)x,(int)y);
- //const float4 normal = normals.tex2D((int)x,(int)y);
- const float contrib = contribs.tex2D((int)x,(int)y);
-
- if (contrib > 0.0f) {
- colour_out(x,y) = make_uchar4(colour.x / contrib, colour.y / contrib, colour.z / contrib, 0);
- //normals(x,y) = normal / contrib;
+ TextureObject<float4> colour_in,
+ TextureObject<uchar4> colour_out,
+ //TextureObject<float4> normals,
+ TextureObject<float> contribs) {
+ const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x < colour_in.width() && y < colour_in.height()) {
+ const float4 colour = colour_in.tex2D((int)x,(int)y);
+ //const float4 normal = normals.tex2D((int)x,(int)y);
+ const float contrib = contribs.tex2D((int)x,(int)y);
+
+ if (contrib > 0.0f) {
+ colour_out(x,y) = make_uchar4(colour.x / contrib, colour.y / contrib, colour.z / contrib, 0);
+ //normals(x,y) = normal / contrib;
+ }
}
}
+
+void ftl::cuda::dibr_normalise(TextureObject<float4> &colour_in, TextureObject<uchar4> &colour_out, TextureObject<float> &contribs, cudaStream_t stream) {
+ const dim3 gridSize((colour_in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (colour_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ dibr_normalise_kernel<<<gridSize, blockSize, 0, stream>>>(colour_in, colour_out, contribs);
+ cudaSafeCall( cudaGetLastError() );
}
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index b1b9bbd3d..8f6557b7f 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -7,6 +7,21 @@
namespace ftl {
namespace cuda {
void dibr_merge(ftl::cuda::TextureObject<float4> &points, ftl::cuda::TextureObject<int> &depth, ftl::render::SplatParams params, cudaStream_t stream);
+
+ void dibr_attribute(
+ ftl::cuda::TextureObject<uchar4> &colour_in, // Original colour image
+ ftl::cuda::TextureObject<float4> &points, // Original 3D points
+ ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float4> &colour_out, // Accumulated output
+ //TextureObject<float4> normal_out,
+ ftl::cuda::TextureObject<float> &contrib_out,
+ ftl::render::SplatParams ¶ms, cudaStream_t stream);
+
+ void dibr_normalise(
+ ftl::cuda::TextureObject<float4> &colour_in,
+ ftl::cuda::TextureObject<uchar4> &colour_out,
+ ftl::cuda::TextureObject<float> &contribs,
+ cudaStream_t stream);
}
}
diff --git a/components/rgbd-sources/include/ftl/rgbd/channels.hpp b/components/rgbd-sources/include/ftl/rgbd/channels.hpp
index 00d26b937..a87d03e80 100644
--- a/components/rgbd-sources/include/ftl/rgbd/channels.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/channels.hpp
@@ -17,9 +17,10 @@ enum struct Channel : int {
Disparity = 3,
Depth2 = 3,
Deviation = 4,
- Normals, // 32FC4
- Points, // 32FC4
- Confidence, // 32F
+ Normals = 5, // 32FC4
+ Points = 6, // 32FC4
+ Confidence = 7, // 32F
+ Contribution = 7, // 32F
Flow, // 32F
Energy, // 32F
LeftGray,
--
GitLab