From 1909d7d4e93ff0926057b6fa8f9c01c8a955a4c3 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Wed, 26 Feb 2020 12:09:18 +0200
Subject: [PATCH] Fixes for colourising and blending
---
applications/gui/src/camera.cpp | 6 +++---
.../cpp/include/ftl/render/CUDARender.hpp | 2 +-
.../renderers/cpp/include/ftl/render/renderer.hpp | 2 +-
components/renderers/cpp/src/CUDARender.cpp | 10 +++++-----
components/renderers/cpp/src/colour_cuda.hpp | 2 +-
components/renderers/cpp/src/colour_util.cu | 14 +++++++-------
components/renderers/cpp/src/colouriser.cpp | 12 ++++++++++--
components/renderers/cpp/src/normals.cu | 2 +-
8 files changed, 29 insertions(+), 21 deletions(-)
diff --git a/applications/gui/src/camera.cpp b/applications/gui/src/camera.cpp
index ab94a7848..db1ee404d 100644
--- a/applications/gui/src/camera.cpp
+++ b/applications/gui/src/camera.cpp
@@ -333,14 +333,14 @@ void ftl::gui::Camera::_draw(std::vector<ftl::rgbd::FrameSet*> &fss) {
}
if (channel_ != Channel::Left && channel_ != Channel::Right && channel_ != Channel::None) {
- renderer_->blend(0.5f, channel_);
+ renderer_->blend(channel_);
if (isStereo()) {
- renderer2_->blend(0.5f, mapToSecondChannel(channel_));
+ renderer2_->blend(mapToSecondChannel(channel_));
}
}
if (enable_overlay) {
- renderer_->blend(overlayer_->value("alpha", 0.8f), Channel::Overlay);
+ renderer_->blend(Channel::Overlay);
}
renderer_->end();
diff --git a/components/renderers/cpp/include/ftl/render/CUDARender.hpp b/components/renderers/cpp/include/ftl/render/CUDARender.hpp
index 7a9d46e5b..25f0ffaa6 100644
--- a/components/renderers/cpp/include/ftl/render/CUDARender.hpp
+++ b/components/renderers/cpp/include/ftl/render/CUDARender.hpp
@@ -28,7 +28,7 @@ class CUDARender : public ftl::render::Renderer {
bool submit(ftl::rgbd::FrameSet *in, ftl::codecs::Channels<0>, const Eigen::Matrix4d &t) override;
//void setOutputDevice(int);
- void blend(float alpha, ftl::codecs::Channel) override;
+ void blend(ftl::codecs::Channel) override;
void setViewPort(ftl::render::ViewPortMode mode, const ftl::render::ViewPort &vp) {
params_.viewport = vp;
diff --git a/components/renderers/cpp/include/ftl/render/renderer.hpp b/components/renderers/cpp/include/ftl/render/renderer.hpp
index 5d9d55f13..2273220cf 100644
--- a/components/renderers/cpp/include/ftl/render/renderer.hpp
+++ b/components/renderers/cpp/include/ftl/render/renderer.hpp
@@ -56,7 +56,7 @@ class Renderer : public ftl::Configurable {
*/
virtual bool submit(ftl::rgbd::FrameSet *, ftl::codecs::Channels<0>, const Eigen::Matrix4d &)=0;
- virtual void blend(float, ftl::codecs::Channel)=0;
+ virtual void blend(ftl::codecs::Channel)=0;
protected:
Stage stage_;
diff --git a/components/renderers/cpp/src/CUDARender.cpp b/components/renderers/cpp/src/CUDARender.cpp
index 557392d1c..9a62807cc 100644
--- a/components/renderers/cpp/src/CUDARender.cpp
+++ b/components/renderers/cpp/src/CUDARender.cpp
@@ -490,7 +490,7 @@ void CUDARender::begin(ftl::rgbd::Frame &out, ftl::codecs::Channel chan) {
stage_ = Stage::ReadySubmit;
}
-void CUDARender::blend(float alpha, Channel c) {
+void CUDARender::blend(Channel c) {
if (stage_ == Stage::Finished) {
throw FTL_Error("Cannot call blend at this time");
} else if (stage_ == Stage::ReadySubmit) {
@@ -504,11 +504,11 @@ void CUDARender::blend(float alpha, Channel c) {
//cv::cuda::addWeighted(buf.to_gpumat(), alpha, out_->get<GpuMat>(out_chan_), 1.0f-alpha, 0.0f,
// out_->get<GpuMat>(out_chan_), -1, cvstream);
- if (alpha < 0.0f) {
+ //if (alpha < 0.0f) {
ftl::cuda::composite(buf, out_->getTexture<uchar4>(out_chan_), stream_);
- } else {
- ftl::cuda::blend_alpha(buf, out_->getTexture<uchar4>(out_chan_), alpha, 1.0f-alpha, stream_);
- }
+ //} else {
+ // ftl::cuda::blend_alpha(buf, out_->getTexture<uchar4>(out_chan_), alpha, 1.0f-alpha, stream_);
+ //}
}
void CUDARender::end() {
diff --git a/components/renderers/cpp/src/colour_cuda.hpp b/components/renderers/cpp/src/colour_cuda.hpp
index bad912c96..1e44bdc93 100644
--- a/components/renderers/cpp/src/colour_cuda.hpp
+++ b/components/renderers/cpp/src/colour_cuda.hpp
@@ -9,7 +9,7 @@ namespace cuda {
template <typename T>
void lut(ftl::cuda::TextureObject<T> &in, ftl::cuda::TextureObject<uchar4> &out,
const cv::cuda::PtrStepSz<uchar3> &lut, float minval, float maxval,
- bool invert, cudaStream_t stream);
+ float alpha, bool invert, cudaStream_t stream);
void blend_alpha(
ftl::cuda::TextureObject<uchar4> &in,
diff --git a/components/renderers/cpp/src/colour_util.cu b/components/renderers/cpp/src/colour_util.cu
index 01939ebd3..bc8c16033 100644
--- a/components/renderers/cpp/src/colour_util.cu
+++ b/components/renderers/cpp/src/colour_util.cu
@@ -14,12 +14,12 @@ template <typename T, bool INVERT>
int out_pitch,
int width, int height,
const uchar3* __restrict__ lut,
- float minval, float maxval) {
+ float minval, float maxval, float alpha) {
__shared__ uchar4 table[256];
int id = threadIdx.x + blockDim.x*threadIdx.y;
- table[id] = make_uchar4(lut[id], (id == 0 || id == 255) ? 0 : 255);
+ table[id] = make_uchar4(lut[id], (id == 0 || id == 255) ? 0 : alpha);
__syncthreads();
@@ -36,7 +36,7 @@ template <typename T, bool INVERT>
template <typename T>
void ftl::cuda::lut(TextureObject<T> &in, TextureObject<uchar4> &out,
const cv::cuda::PtrStepSz<uchar3> &lut, float minval, float maxval,
- bool invert,
+ float alpha, bool invert,
cudaStream_t stream) {
static constexpr int THREADS_X = 64; // Must total 256
@@ -46,19 +46,19 @@ void ftl::cuda::lut(TextureObject<T> &in, TextureObject<uchar4> &out,
const dim3 blockSize(THREADS_X, THREADS_Y);
if (invert) {
- lut_kernel<T,true><<<gridSize, blockSize, 0, stream>>>(in.devicePtr(), in.pixelPitch(), out.devicePtr(), out.pixelPitch(), out.width(), out.height(), lut.data, minval, maxval);
+ lut_kernel<T,true><<<gridSize, blockSize, 0, stream>>>(in.devicePtr(), in.pixelPitch(), out.devicePtr(), out.pixelPitch(), out.width(), out.height(), lut.data, minval, maxval, alpha);
} else {
- lut_kernel<T,false><<<gridSize, blockSize, 0, stream>>>(in.devicePtr(), in.pixelPitch(), out.devicePtr(), out.pixelPitch(), out.width(), out.height(), lut.data, minval, maxval);
+ lut_kernel<T,false><<<gridSize, blockSize, 0, stream>>>(in.devicePtr(), in.pixelPitch(), out.devicePtr(), out.pixelPitch(), out.width(), out.height(), lut.data, minval, maxval, alpha);
}
cudaSafeCall( cudaGetLastError() );
}
template void ftl::cuda::lut<float>(TextureObject<float> &in, TextureObject<uchar4> &out,
- const cv::cuda::PtrStepSz<uchar3> &lut, float minval, float maxval, bool invert,
+ const cv::cuda::PtrStepSz<uchar3> &lut, float minval, float maxval, float, bool invert,
cudaStream_t stream);
template void ftl::cuda::lut<short>(TextureObject<short> &in, TextureObject<uchar4> &out,
- const cv::cuda::PtrStepSz<uchar3> &lut, float minval, float maxval, bool invert,
+ const cv::cuda::PtrStepSz<uchar3> &lut, float minval, float maxval, float, bool invert,
cudaStream_t stream);
// ==== Blending ===============================================================
diff --git a/components/renderers/cpp/src/colouriser.cpp b/components/renderers/cpp/src/colouriser.cpp
index 95e85c471..bcaaf28d9 100644
--- a/components/renderers/cpp/src/colouriser.cpp
+++ b/components/renderers/cpp/src/colouriser.cpp
@@ -133,6 +133,10 @@ TextureObject<uchar4> &Colouriser::_processNormals(ftl::rgbd::Frame &f, Channel
auto light_diffuse = parseCUDAColour(value("diffuse", std::string("#e0e0e0")));
auto light_ambient = parseCUDAColour(value("ambient", std::string("#0e0e0e")));
+
+ light_diffuse.w = value("alpha", 0.5f)*255.0f;
+ light_ambient.w = light_diffuse.w;
+
auto light_pos = make_float3(value("light_x", 0.3f), value("light_y", 0.2f), value("light_z", 1.0f));
ftl::cuda::normal_visualise(f.createTexture<half4>(c), buf,
@@ -152,7 +156,9 @@ TextureObject<uchar4> &Colouriser::_processSingle(ftl::rgbd::Frame &f, Channel c
depth_lut.upload(lut);
}
- ftl::cuda::lut(f.createTexture<T>(c), buf, depth_lut, 0, std::numeric_limits<T>::max(), true, stream);
+ float alpha = value("alpha",0.5f)*255.0f;
+
+ ftl::cuda::lut(f.createTexture<T>(c), buf, depth_lut, 0, std::numeric_limits<T>::max(), alpha, true, stream);
return buf;
}
@@ -165,7 +171,9 @@ TextureObject<uchar4> &Colouriser::_processFloat(ftl::rgbd::Frame &f, Channel c,
depth_lut.upload(lut);
}
- ftl::cuda::lut(f.createTexture<float>(c), buf, depth_lut, minval, maxval, false, stream);
+ float alpha = value("alpha",0.5f)*255.0f;
+
+ ftl::cuda::lut(f.createTexture<float>(c), buf, depth_lut, minval, maxval, alpha, false, stream);
return buf;
}
diff --git a/components/renderers/cpp/src/normals.cu b/components/renderers/cpp/src/normals.cu
index 1375dc035..ddeb6294e 100644
--- a/components/renderers/cpp/src/normals.cu
+++ b/components/renderers/cpp/src/normals.cu
@@ -430,7 +430,7 @@ __global__ void vis_normals_kernel(ftl::cuda::TextureObject<half4> norm,
output(x,y) = make_uchar4(
min(255.0f, diffuse.x*d + ambient.x),
min(255.0f, diffuse.y*d + ambient.y),
- min(255.0f, diffuse.z*d + ambient.z), 255);
+ min(255.0f, diffuse.z*d + ambient.z), ambient.w);
}
}
}
--
GitLab