From 2135cc2043f0cbc61ed169223f5044562ca25c22 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Fri, 28 Feb 2020 13:17:25 +0200
Subject: [PATCH] Flip original sources also

---
 applications/gui/src/camera.cpp               |  3 +++
 applications/gui/src/src_window.cpp           |  4 +++-
 .../cpp/include/ftl/cuda/transform.hpp        | 23 +++++++++++++++++++
 components/renderers/cpp/src/CUDARender.cpp   |  1 +
 components/renderers/cpp/src/colour_cuda.hpp  |  5 ----
 components/renderers/cpp/src/colour_util.cu   | 21 +++++++++++++++++
 6 files changed, 51 insertions(+), 6 deletions(-)
 create mode 100644 components/renderers/cpp/include/ftl/cuda/transform.hpp

diff --git a/applications/gui/src/camera.cpp b/applications/gui/src/camera.cpp
index 48858d57f..5a0964ab6 100644
--- a/applications/gui/src/camera.cpp
+++ b/applications/gui/src/camera.cpp
@@ -11,6 +11,7 @@
 #include <ftl/operators/antialiasing.hpp>
 #include <ftl/cuda/normals.hpp>
 #include <ftl/render/colouriser.hpp>
+#include <ftl/cuda/transform.hpp>
 
 #include <ftl/codecs/faces.hpp>
 
@@ -134,6 +135,7 @@ void ftl::gui::Camera::draw(std::vector<ftl::rgbd::FrameSet*> &fss) {
 	if (fid_ != 255) {
 		for (auto *fs : fss) {
 			if (!usesFrameset(fs->id)) continue;
+			UNIQUE_LOCK(fs->mtx, lk);
 
 			ftl::rgbd::Frame *frame = nullptr;
 
@@ -147,6 +149,7 @@ void ftl::gui::Camera::draw(std::vector<ftl::rgbd::FrameSet*> &fss) {
 			texture1_.make(buf.width(), buf.height());
 			auto dst1 = texture1_.map(0);
 			cudaMemcpy2D(dst1.data, dst1.step1(), buf.devicePtr(), buf.pitch(), buf.width()*4, buf.height(), cudaMemcpyDeviceToDevice);
+			ftl::cuda::flip<uchar4>(dst1, 0);
 			texture1_.unmap(0);
 
 			width_ = texture1_.width();
diff --git a/applications/gui/src/src_window.cpp b/applications/gui/src/src_window.cpp
index be7657c8c..15e9e96f7 100644
--- a/applications/gui/src/src_window.cpp
+++ b/applications/gui/src/src_window.cpp
@@ -387,11 +387,13 @@ void SourceWindow::draw(NVGcontext *ctx) {
 			if (!camera.second.thumbview) camera.second.thumbview = new ftl::gui::ThumbView(ipanel_, screen_, cam);
 			camera.second.thumbview->setFixedSize(nanogui::Vector2i(320,180));
 
+			auto *iv = dynamic_cast<nanogui::ImageView*>(camera.second.thumbview);
+
 			/*if ((size_t)ipanel_->childCount() < i+1) {
 				new ftl::gui::ThumbView(ipanel_, screen_, cam);
 			}*/
 			//if (thumbs_[i].isValid()) dynamic_cast<nanogui::ImageView*>(camera.second.thumbview)->bindImage(thumbs_[i].texture());
-			if (cam->getLeft().isValid()) dynamic_cast<nanogui::ImageView*>(camera.second.thumbview)->bindImage(cam->getLeft().texture());
+			if (cam->getLeft().isValid()) iv->bindImage(cam->getLeft().texture());
 			++i;
 		}
 
diff --git a/components/renderers/cpp/include/ftl/cuda/transform.hpp b/components/renderers/cpp/include/ftl/cuda/transform.hpp
new file mode 100644
index 000000000..3044342b0
--- /dev/null
+++ b/components/renderers/cpp/include/ftl/cuda/transform.hpp
@@ -0,0 +1,23 @@
+#ifndef _FTL_CUDA_TRANSFORM_HPP_
+#define _FTL_CUDA_TRANSFORM_HPP_
+
+#include <ftl/cuda_common.hpp>
+
+namespace ftl {
+namespace cuda {
+
+    template <typename T>
+	void flip(
+		ftl::cuda::TextureObject<T> &out,
+		cudaStream_t stream);
+
+
+    template <typename T>
+	void flip(
+		cv::cuda::GpuMat &out,
+		cudaStream_t stream);
+
+}
+}
+
+#endif 
diff --git a/components/renderers/cpp/src/CUDARender.cpp b/components/renderers/cpp/src/CUDARender.cpp
index a48d754ca..6d0deb848 100644
--- a/components/renderers/cpp/src/CUDARender.cpp
+++ b/components/renderers/cpp/src/CUDARender.cpp
@@ -5,6 +5,7 @@
 #include <ftl/cuda/normals.hpp>
 #include <ftl/operators/mask_cuda.hpp>
 #include <ftl/render/colouriser.hpp>
+#include <ftl/cuda/transform.hpp>
 
 #include "colour_cuda.hpp"
 
diff --git a/components/renderers/cpp/src/colour_cuda.hpp b/components/renderers/cpp/src/colour_cuda.hpp
index 40ce5cce4..1e44bdc93 100644
--- a/components/renderers/cpp/src/colour_cuda.hpp
+++ b/components/renderers/cpp/src/colour_cuda.hpp
@@ -22,11 +22,6 @@ void composite(
 		ftl::cuda::TextureObject<uchar4> &out,
 		cudaStream_t stream);
 
-template <typename T>
-	void flip(
-		ftl::cuda::TextureObject<T> &out,
-		cudaStream_t stream);
-
 }
 }
 
diff --git a/components/renderers/cpp/src/colour_util.cu b/components/renderers/cpp/src/colour_util.cu
index b9c589384..b40ae62d6 100644
--- a/components/renderers/cpp/src/colour_util.cu
+++ b/components/renderers/cpp/src/colour_util.cu
@@ -1,4 +1,5 @@
 #include "colour_cuda.hpp"
+#include <ftl/cuda/transform.hpp>
 
 using ftl::cuda::TextureObject;
 
@@ -196,3 +197,23 @@ void ftl::cuda::flip(
 
 template void ftl::cuda::flip<float>(TextureObject<float> &,cudaStream_t stream);
 template void ftl::cuda::flip<uchar4>(TextureObject<uchar4> &,cudaStream_t stream);
+
+template <typename T>
+void ftl::cuda::flip(
+		cv::cuda::GpuMat &img,
+		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);
+
+	flip_kernel<T><<<gridSize, blockSize, 0, stream>>>(
+		(T*)img.data, img.step/sizeof(T),
+		img.cols, img.rows);
+	cudaSafeCall( cudaGetLastError() );
+}
+
+template void ftl::cuda::flip<float>(cv::cuda::GpuMat &,cudaStream_t stream);
+template void ftl::cuda::flip<uchar4>(cv::cuda::GpuMat &,cudaStream_t stream);
-- 
GitLab