From 927914ca9ac9a4f0f9db0aa6a6e77b10ba96b0fb Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Wed, 1 Apr 2020 08:56:23 +0300
Subject: [PATCH] Allow composite on original sources
---
applications/gui/src/camera.cpp | 3 ++
.../{src => include/ftl/cuda}/colour_cuda.hpp | 5 +++
components/renderers/cpp/src/CUDARender.cpp | 2 +-
components/renderers/cpp/src/colour_util.cu | 45 ++++++++++++++++++-
components/renderers/cpp/src/colouriser.cpp | 2 +-
5 files changed, 54 insertions(+), 3 deletions(-)
rename components/renderers/cpp/{src => include/ftl/cuda}/colour_cuda.hpp (83%)
diff --git a/applications/gui/src/camera.cpp b/applications/gui/src/camera.cpp
index 24fb0aedb..1ebe9cb4f 100644
--- a/applications/gui/src/camera.cpp
+++ b/applications/gui/src/camera.cpp
@@ -13,6 +13,7 @@
#include <ftl/render/colouriser.hpp>
#include <ftl/cuda/transform.hpp>
#include <ftl/operators/gt_analysis.hpp>
+#include <ftl/cuda/colour_cuda.hpp>
#include <ftl/render/overlay.hpp>
#include "statsimage.hpp"
@@ -160,6 +161,8 @@ void ftl::gui::Camera::draw(std::vector<ftl::rgbd::FrameSet*> &fss) {
if (!frame->hasChannel(channel_)) return;
auto &buf = colouriser_->colourise(*frame, channel_, 0);
+ auto &buf2 = frame->getTexture<uchar4>(Channel::Colour);
+ ftl::cuda::compositeInverse(buf2, buf, 0);
// For non-virtual cameras, copy the CUDA texture into the opengl
// texture device-to-device.
diff --git a/components/renderers/cpp/src/colour_cuda.hpp b/components/renderers/cpp/include/ftl/cuda/colour_cuda.hpp
similarity index 83%
rename from components/renderers/cpp/src/colour_cuda.hpp
rename to components/renderers/cpp/include/ftl/cuda/colour_cuda.hpp
index 1e44bdc93..492b441ec 100644
--- a/components/renderers/cpp/src/colour_cuda.hpp
+++ b/components/renderers/cpp/include/ftl/cuda/colour_cuda.hpp
@@ -22,6 +22,11 @@ void composite(
ftl::cuda::TextureObject<uchar4> &out,
cudaStream_t stream);
+void compositeInverse(
+ ftl::cuda::TextureObject<uchar4> &in,
+ ftl::cuda::TextureObject<uchar4> &out,
+ cudaStream_t stream);
+
}
}
diff --git a/components/renderers/cpp/src/CUDARender.cpp b/components/renderers/cpp/src/CUDARender.cpp
index b00efb036..5a6665d68 100644
--- a/components/renderers/cpp/src/CUDARender.cpp
+++ b/components/renderers/cpp/src/CUDARender.cpp
@@ -7,7 +7,7 @@
#include <ftl/render/colouriser.hpp>
#include <ftl/cuda/transform.hpp>
-#include "colour_cuda.hpp"
+#include <ftl/cuda/colour_cuda.hpp>
#define LOGURU_REPLACE_GLOG 1
#include <loguru.hpp>
diff --git a/components/renderers/cpp/src/colour_util.cu b/components/renderers/cpp/src/colour_util.cu
index b40ae62d6..625261ffd 100644
--- a/components/renderers/cpp/src/colour_util.cu
+++ b/components/renderers/cpp/src/colour_util.cu
@@ -1,4 +1,4 @@
-#include "colour_cuda.hpp"
+#include <ftl/cuda/colour_cuda.hpp>
#include <ftl/cuda/transform.hpp>
using ftl::cuda::TextureObject;
@@ -158,6 +158,49 @@ void ftl::cuda::composite(
cudaSafeCall( cudaGetLastError() );
}
+__global__ void composite_inverse_kernel(
+ const uchar4* __restrict__ in,
+ int in_pitch,
+ uchar4* __restrict__ out,
+ int out_pitch,
+ int width, int height) {
+
+ for (STRIDE_Y(y, height)) {
+ for (STRIDE_X(x, width)) {
+ const uchar4 c1 = in[x+y*in_pitch];
+ const uchar4 c2 = out[x+y*out_pitch];
+
+ const float b = (float(c2.w)/255.0f);
+ const float a = 1.0f - b;
+
+ out[x+y*out_pitch] = make_uchar4(
+ clamp(float(c1.x)*a + float(c2.x)*b, 255.0f),
+ clamp(float(c1.y)*a + float(c2.y)*b, 255.0f),
+ clamp(float(c1.z)*a + float(c2.z)*b, 255.0f),
+ 255.0f
+ );
+ }
+ }
+}
+
+void ftl::cuda::compositeInverse(
+ TextureObject<uchar4> &in,
+ TextureObject<uchar4> &out,
+ cudaStream_t stream) {
+
+ static constexpr int THREADS_X = 32;
+ static constexpr int THREADS_Y = 8;
+
+ const dim3 gridSize(6,64);
+ const dim3 blockSize(THREADS_X, THREADS_Y);
+
+ composite_inverse_kernel<<<gridSize, blockSize, 0, stream>>>(
+ in.devicePtr(), in.pixelPitch(),
+ out.devicePtr(), out.pixelPitch(),
+ out.width(), out.height());
+ cudaSafeCall( cudaGetLastError() );
+}
+
// ==== Flipping ===============================================================
diff --git a/components/renderers/cpp/src/colouriser.cpp b/components/renderers/cpp/src/colouriser.cpp
index be92ed854..6341f1a86 100644
--- a/components/renderers/cpp/src/colouriser.cpp
+++ b/components/renderers/cpp/src/colouriser.cpp
@@ -1,6 +1,6 @@
#include <ftl/render/colouriser.hpp>
#include "splatter_cuda.hpp"
-#include "colour_cuda.hpp"
+#include <ftl/cuda/colour_cuda.hpp>
#include <ftl/cuda/normals.hpp>
#include <opencv2/cudaarithm.hpp>
--
GitLab