Resolves #336 Flipped VR images

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();

applications/gui/src/screen.cpp

@@ -38,7 +38,7 @@ namespace {
 		in vec2 vertex;
 		out vec2 uv;
 		void main() {
-			uv = vec2(vertex.x, 1.0 - vertex.y);
+			uv = vec2(vertex.x, vertex.y);
 			vec2 scaledVertex = (vertex * scaleFactor) + position;
 			gl_Position  = vec4(2.0*scaledVertex.x - 1.0,
 								2.0*scaledVertex.y - 1.0,

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;
 		}
 

components/renderers/cpp/include/ftl/cuda/transform.hpp

+#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 

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"
 
@@ -538,6 +539,7 @@ void CUDARender::_endSubmit() {
 		accum_,
 		out_->getTexture<uchar4>(out_chan_),
 		contrib_,
+		false,  // Flip
 		stream_
 	);
 }
@@ -545,6 +547,10 @@ void CUDARender::_endSubmit() {
 void CUDARender::_end() {
 	_postprocessColours(*out_);
 
+	// Final OpenGL flip
+	ftl::cuda::flip(out_->getTexture<uchar4>(out_chan_), stream_);
+	ftl::cuda::flip(out_->getTexture<float>(_getDepthChannel()), stream_);
+
 	cudaSafeCall(cudaStreamSynchronize(stream_));
 }
 

components/renderers/cpp/src/colour_util.cu

 #include "colour_cuda.hpp"
+#include <ftl/cuda/transform.hpp>
 
 using ftl::cuda::TextureObject;
 
@@ -156,3 +157,63 @@ void ftl::cuda::composite(
 		out.width(), out.height());
 	cudaSafeCall( cudaGetLastError() );
 }
+
+
+// ==== Flipping ===============================================================
+
+template <typename T>
+__global__ void flip_kernel(
+		T* __restrict__ img,
+		int pitch, 
+		int width, int height) {
+
+	for (STRIDE_Y(y, height/2)) {
+	for (STRIDE_X(x, width)) {
+		const T c1 = img[x+y*pitch];
+		const T c2 = img[x+(height-y-2)*pitch];
+
+		img[x+y*pitch] = c2;
+		img[x+(height-y-2)*pitch] = c1;
+	}
+	}
+}
+
+template <typename T>
+void ftl::cuda::flip(
+		TextureObject<T> &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>>>(
+		img.devicePtr(), img.pixelPitch(),
+		img.width(), img.height());
+	cudaSafeCall( cudaGetLastError() );
+}
+
+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);

components/renderers/cpp/src/dibr.cu

@@ -81,7 +81,7 @@ void ftl::cuda::dibr_merge(TextureObject<int> &depth_out, const float4x4 &transf
 
 // ==== Normalize ==============================================================
 
-template <typename A, typename B>
+template <typename A, typename B, bool FLIPY>
 __global__ void dibr_normalise_kernel(
         TextureObject<A> in,
         TextureObject<B> out,
@@ -97,24 +97,29 @@ __global__ void dibr_normalise_kernel(
 		//out(x,y) = (contrib == 0.0f) ? make<B>(a) : make<B>(a / contrib);
 
         if (contrib > 0.0f) {
-            out(x,y) = make<B>(a / contrib);
+            if (FLIPY) out(x,out.height()-y-1) = make<B>(a / contrib);
+            else out(x,y) = make<B>(a / contrib);
             //normals(x,y) = normal / contrib;
         }
     }
 }
 
 template <typename A, typename B>
-void ftl::cuda::dibr_normalise(TextureObject<A> &in, TextureObject<B> &out, TextureObject<int> &contribs, cudaStream_t stream) {
+void ftl::cuda::dibr_normalise(TextureObject<A> &in, TextureObject<B> &out, TextureObject<int> &contribs, bool flip, cudaStream_t stream) {
     const dim3 gridSize((in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (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>>>(in, out, contribs);
+    if (flip) {
+        dibr_normalise_kernel<A,B,true><<<gridSize, blockSize, 0, stream>>>(in, out, contribs);
+    } else {
+        dibr_normalise_kernel<A,B,false><<<gridSize, blockSize, 0, stream>>>(in, out, contribs);
+    }
     cudaSafeCall( cudaGetLastError() );
 }
 
-template void ftl::cuda::dibr_normalise<float4,uchar4>(TextureObject<float4> &in, TextureObject<uchar4> &out, TextureObject<int> &contribs, cudaStream_t stream);
-template void ftl::cuda::dibr_normalise<float,float>(TextureObject<float> &in, TextureObject<float> &out, TextureObject<int> &contribs, cudaStream_t stream);
-template void ftl::cuda::dibr_normalise<float4,float4>(TextureObject<float4> &in, TextureObject<float4> &out, TextureObject<int> &contribs, cudaStream_t stream);
+template void ftl::cuda::dibr_normalise<float4,uchar4>(TextureObject<float4> &in, TextureObject<uchar4> &out, TextureObject<int> &contribs, bool, cudaStream_t stream);
+template void ftl::cuda::dibr_normalise<float,float>(TextureObject<float> &in, TextureObject<float> &out, TextureObject<int> &contribs, bool, cudaStream_t stream);
+template void ftl::cuda::dibr_normalise<float4,float4>(TextureObject<float4> &in, TextureObject<float4> &out, TextureObject<int> &contribs, bool, cudaStream_t stream);
 
 // Float version
 

components/renderers/cpp/src/overlay.cpp

@@ -288,7 +288,7 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const
 
 	if (value("show_poses", false)) {
 		for (size_t i=0; i<fs.frames.size(); ++i) {
-			auto pose = fs.frames[i].getPose().inverse(); //.inverse() * state.getPose();
+			auto pose = fs.frames[i].getPose(); //.inverse() * state.getPose();
 
 			auto name = fs.frames[i].get<std::string>("name");
             _drawOutlinedShape(Shape::CAMERA, state.getPose().inverse() * pose, Eigen::Vector3f(0.2f,0.2f,0.2f), make_uchar4(255,0,0,80), make_uchar4(255,0,0,255));

components/renderers/cpp/src/splatter_cuda.hpp

@@ -141,6 +141,7 @@ namespace cuda {
 		ftl::cuda::TextureObject<A> &in,
 		ftl::cuda::TextureObject<B> &out,
 		ftl::cuda::TextureObject<int> &contribs,
+		bool flipy,
 		cudaStream_t stream);
 
 	template <typename A, typename B>