From 481a87086a730fe520beaddb5ed92b927ab893a3 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nicolas.pope@utu.fi>
Date: Thu, 26 Sep 2019 15:37:34 +0300
Subject: [PATCH] Allow for float attributes
---
components/renderers/cpp/src/splat_render.cpp | 66 +++++++---
components/renderers/cpp/src/splatter.cu | 114 ++++++++++++++++++
.../renderers/cpp/src/splatter_cuda.hpp | 23 +++-
3 files changed, 182 insertions(+), 21 deletions(-)
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index 9c507fc43..1b39ccebf 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -28,6 +28,8 @@ void Splatter::renderChannel(
temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
+
+ bool is_float = ftl::rgbd::isFloatChannel(channel);
// Render each camera into virtual view
for (size_t i=0; i < scene_->frames.size(); ++i) {
@@ -75,23 +77,53 @@ void Splatter::renderChannel(
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
- );
+ if (is_float) {
+ ftl::cuda::dibr_attribute(
+ f.createTexture<float>(channel),
+ f.createTexture<float4>(Channel::Points),
+ temp_.getTexture<int>(Channel::Depth),
+ temp_.getTexture<float4>(Channel::Colour),
+ temp_.getTexture<float>(Channel::Contribution),
+ params, stream
+ );
+ } else if (channel == Channel::Colour || channel == Channel::Right) {
+ 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
+ );
+ } else {
+ ftl::cuda::dibr_attribute(
+ f.createTexture<uchar4>(channel),
+ 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
- ftl::cuda::dibr_normalise(
- temp_.createTexture<float4>(Channel::Colour),
- out.createTexture<uchar4>(channel),
- temp_.createTexture<float>(Channel::Contribution),
- stream
- );
+ if (is_float) {
+ // Normalise attribute contributions
+ ftl::cuda::dibr_normalise(
+ temp_.createTexture<float4>(Channel::Colour),
+ out.createTexture<float>(channel),
+ temp_.createTexture<float>(Channel::Contribution),
+ stream
+ );
+ } else {
+ // Normalise attribute contributions
+ ftl::cuda::dibr_normalise(
+ temp_.createTexture<float4>(Channel::Colour),
+ out.createTexture<uchar4>(channel),
+ temp_.createTexture<float>(Channel::Contribution),
+ stream
+ );
+ }
}
bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cudaStream_t stream) {
@@ -165,9 +197,9 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
{
temp_.get<GpuMat>(Channel::Depth).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
}
- else if (chan == Channel::Energy)
+ else if (chan == Channel::Contribution)
{
- cv::cuda::swap(temp_.get<GpuMat>(Channel::Energy), out.create<GpuMat>(Channel::Energy));
+ cv::cuda::swap(temp_.get<GpuMat>(Channel::Contribution), out.create<GpuMat>(Channel::Contribution));
}
else if (chan == Channel::Right)
{
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index c1b46fc1d..3b1ae4b47 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -137,6 +137,70 @@ __global__ void dibr_attribute_contrib_kernel(
}
}
+__global__ void dibr_attribute_contrib_kernel(
+ TextureObject<float> 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 params) {
+
+ //const ftl::voxhash::DepthCameraCUDA &camera = c_cameras[cam];
+
+ const int tid = (threadIdx.x + threadIdx.y * blockDim.x);
+ //const int warp = tid / WARP_SIZE;
+ 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(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;
+
+ const float3 camPos = params.m_viewMatrix * worldPos;
+ if (camPos.z < params.camera.minDepth) return;
+ if (camPos.z > params.camera.maxDepth) return;
+ const uint2 screenPos = params.camera.camToScreen<uint2>(camPos);
+
+ const int upsample = 8; //min(UPSAMPLE_MAX, int((5.0f*r) * params.camera.fx / camPos.z));
+
+ // Not on screen so stop now...
+ if (screenPos.x >= depth_in.width() || screenPos.y >= depth_in.height()) return;
+
+ // Is this point near the actual surface and therefore a contributor?
+ const float d = ((float)depth_in.tex2D((int)screenPos.x, (int)screenPos.y)/1000.0f);
+ //if (abs(d - camPos.z) > DEPTH_THRESHOLD) return;
+
+ // TODO:(Nick) Should just one thread load these to shared mem?
+ const float colour = (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.
+ const int lane = tid % WARP_SIZE;
+ for (int i=lane; i<upsample*upsample; i+=WARP_SIZE) {
+ const float u = (i % upsample) - (upsample / 2);
+ const float v = (i / upsample) - (upsample / 2);
+
+ // Use the depth buffer to determine this pixels 3D position in camera space
+ const float d = ((float)depth_in.tex2D(screenPos.x+u, screenPos.y+v)/1000.0f);
+ const float3 nearest = params.camera.screenToCam((int)(screenPos.x+u),(int)(screenPos.y+v),d);
+
+ // What is contribution of our current point at this pixel?
+ 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 float wcolour = colour * weight;
+ //const float4 wnormal = normal * weight;
+
+ //printf("Z %f\n", d);
+
+ // Add this points contribution to the pixel buffer
+ atomicAdd((float*)&colour_out(screenPos.x+u, screenPos.y+v), wcolour);
+ atomicAdd(&contrib_out(screenPos.x+u, screenPos.y+v), weight);
+ }
+ }
+}
+
void ftl::cuda::dibr_attribute(
TextureObject<uchar4> &colour_in, // Original colour image
TextureObject<float4> &points, // Original 3D points
@@ -159,6 +223,28 @@ void ftl::cuda::dibr_attribute(
cudaSafeCall( cudaGetLastError() );
}
+void ftl::cuda::dibr_attribute(
+ TextureObject<float> &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(
@@ -181,6 +267,26 @@ __global__ void dibr_normalise_kernel(
}
}
+__global__ void dibr_normalise_kernel(
+ TextureObject<float4> colour_in,
+ TextureObject<float> 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) = colour.x / contrib;
+ //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);
@@ -188,3 +294,11 @@ void ftl::cuda::dibr_normalise(TextureObject<float4> &colour_in, TextureObject<u
dibr_normalise_kernel<<<gridSize, blockSize, 0, stream>>>(colour_in, colour_out, contribs);
cudaSafeCall( cudaGetLastError() );
}
+
+void ftl::cuda::dibr_normalise(TextureObject<float4> &colour_in, TextureObject<float> &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 3cd22a144..8c57d5848 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -13,17 +13,32 @@ namespace cuda {
cudaStream_t stream);
void dibr_attribute(
- ftl::cuda::TextureObject<uchar4> &colour_in, // Original colour image
+ ftl::cuda::TextureObject<uchar4> &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
+ ftl::cuda::TextureObject<float4> &out, // Accumulated output
//TextureObject<float4> normal_out,
ftl::cuda::TextureObject<float> &contrib_out,
ftl::render::SplatParams ¶ms, cudaStream_t stream);
+ void dibr_attribute(
+ ftl::cuda::TextureObject<float> &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> &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> &in,
+ ftl::cuda::TextureObject<uchar4> &out,
+ ftl::cuda::TextureObject<float> &contribs,
+ cudaStream_t stream);
+
void dibr_normalise(
- ftl::cuda::TextureObject<float4> &colour_in,
- ftl::cuda::TextureObject<uchar4> &colour_out,
+ ftl::cuda::TextureObject<float4> &in,
+ ftl::cuda::TextureObject<float> &out,
ftl::cuda::TextureObject<float> &contribs,
cudaStream_t stream);
}
--
GitLab