diff --git a/applications/gui/src/camera.cpp b/applications/gui/src/camera.cpp index 48858d57fb74273d712a0c2e0ef48d87197e6d31..5a0964ab6e44f3a74ed5b61fbf799ca3b4226fda 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/screen.cpp b/applications/gui/src/screen.cpp index 17f3291ef9d87f1ea5ac92307ab207739ee75e12..187ff7945d608f72ef45f4674101209bb3e2f7f7 100644 --- a/applications/gui/src/screen.cpp +++ b/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, diff --git a/applications/gui/src/src_window.cpp b/applications/gui/src/src_window.cpp index be7657c8c5c1d56ec2207aff2b02e7dcbbebee09..15e9e96f75f5b51c9e9a10f9004d3d597dd9974e 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 0000000000000000000000000000000000000000..3044342b0e33a47b7d4784b307a27b04f4d75cf0 --- /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 3affec4bdf172f61b88e7c24824996ac2e7fcd9b..6d0deb848509996105a550abc5197805b4939d8c 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" @@ -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_)); } diff --git a/components/renderers/cpp/src/colour_util.cu b/components/renderers/cpp/src/colour_util.cu index bc8c1603398da9e3f94bad6e55b7d0dc3e3a09a0..b40ae62d6f42ed699c1b78beb88121b72c68ccf7 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; @@ -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); diff --git a/components/renderers/cpp/src/dibr.cu b/components/renderers/cpp/src/dibr.cu index 7530f0a4717d99a9a9ec62944c814ac37e7a9bff..2aa5987371afc61348ffe94f4fe0fc41929e4cb6 100644 --- a/components/renderers/cpp/src/dibr.cu +++ b/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 diff --git a/components/renderers/cpp/src/overlay.cpp b/components/renderers/cpp/src/overlay.cpp index d1d83c9da5b19249029a7cf2d713890840ec4265..78415f5634f9ad90f695503adf1a5a3057f36856 100644 --- a/components/renderers/cpp/src/overlay.cpp +++ b/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)); diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp index da47b6fd3b2ef8ae0809728d27d876338207350f..e1021651ce1989cbf06b0a4b00fc04a1d457b844 100644 --- a/components/renderers/cpp/src/splatter_cuda.hpp +++ b/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>