diff --git a/applications/gui/src/camera.cpp b/applications/gui/src/camera.cpp index 12c7165e0c377ef29d84839d560c84624c03313a..ab94a7848ae05b7c1715530a9909bbf158975110 100644 --- a/applications/gui/src/camera.cpp +++ b/applications/gui/src/camera.cpp @@ -173,7 +173,7 @@ void ftl::gui::Camera::draw(std::vector<ftl::rgbd::FrameSet*> &fss) { auto &d = *data.rbegin(); cv::Mat over_depth; - over_depth.create(im1_.size(), CV_32F); + over_depth.create(overlay_.size(), CV_32F); auto cam = ftl::rgbd::Camera::from(intrinsics_); @@ -225,7 +225,7 @@ void ftl::gui::Camera::draw(std::vector<ftl::rgbd::FrameSet*> &fss) { squre_pts.push_back(cv::Point2f(cam.width,cam.height)); cv::Mat transmtx = cv::getPerspectiveTransform(quad_pts,squre_pts); - cv::Mat transformed = cv::Mat::zeros(im1_.rows, im1_.cols, CV_8UC4); + cv::Mat transformed = cv::Mat::zeros(overlay_.rows, overlay_.cols, CV_8UC4); //cv::warpPerspective(im1_, im1_, transmtx, im1_.size()); ftl::render::ViewPort vp; @@ -301,6 +301,12 @@ void ftl::gui::Camera::_draw(std::vector<ftl::rgbd::FrameSet*> &fss) { frame_.create<cv::cuda::GpuMat>(Channel::Colour) = texture1_.map(renderer_->getCUDAStream()); if (isStereo()) frame_.create<cv::cuda::GpuMat>(Channel::Colour2) = texture2_.map((renderer2_) ? renderer2_->getCUDAStream() : 0); + overlay_.create(state_.getLeft().height, state_.getLeft().width, CV_8UC4); + frame_.create<cv::Mat>(Channel::Overlay) = overlay_; + + overlay_.setTo(cv::Scalar(0,0,0,0)); + bool enable_overlay = overlayer_->value("enabled", false); + { FTL_Profile("Render",0.034); renderer_->begin(frame_, Channel::Colour); @@ -318,6 +324,12 @@ void ftl::gui::Camera::_draw(std::vector<ftl::rgbd::FrameSet*> &fss) { fs->mtx.lock(); renderer_->submit(fs, ftl::codecs::Channels<0>(Channel::Colour), transforms_[fs->id]); if (isStereo()) renderer2_->submit(fs, ftl::codecs::Channels<0>(Channel::Colour), transforms_[fs->id]); + + if (enable_overlay) { + // Generate and upload an overlay image. + overlayer_->apply(*fs, overlay_, state_); + frame_.upload(Channel::Overlay, renderer_->getCUDAStream()); + } } if (channel_ != Channel::Left && channel_ != Channel::Right && channel_ != Channel::None) { @@ -327,6 +339,10 @@ void ftl::gui::Camera::_draw(std::vector<ftl::rgbd::FrameSet*> &fss) { } } + if (enable_overlay) { + renderer_->blend(overlayer_->value("alpha", 0.8f), Channel::Overlay); + } + renderer_->end(); if (isStereo()) renderer2_->end(); } catch(std::exception &e) { @@ -355,12 +371,6 @@ void ftl::gui::Camera::_draw(std::vector<ftl::rgbd::FrameSet*> &fss) { width_ = texture1_.width(); height_ = texture1_.height(); - // Apply any overlays to the image - for (auto *fs : fss) { - if (!usesFrameset(fs->id)) continue; - overlayer_->apply(*fs, im1_, state_); - } - if (record_stream_ && record_stream_->active()) { // TODO: Allow custom channel selection ftl::rgbd::FrameSet fs2; @@ -561,16 +571,6 @@ void ftl::gui::Camera::setChannel(Channel c) { cv::addWeighted(edges, weight, out, 1.0, 0.0, out, CV_8UC4); }*/ -bool ftl::gui::Camera::thumbnail(cv::Mat &thumb) { - { - UNIQUE_LOCK(mutex_, lk); - if (im1_.empty()) return false; - // FIXME: Use correct aspect ratio? - cv::resize(im1_, thumb, cv::Size(320,180)); - } - cv::flip(thumb, thumb, 0); - return true; -} void ftl::gui::Camera::active(bool a) { if (a) { @@ -686,16 +686,6 @@ const void ftl::gui::Camera::captureFrame() { } if (framesets_) draw(*framesets_); - - { - UNIQUE_LOCK(mutex_, lk); - if (im1_.rows != 0) { - //texture1_.update(im1_); - } - if (isStereo() && im2_.rows != 0) { - //texture2_.update(im2_); - } - } } //return texture1_; @@ -706,7 +696,7 @@ void ftl::gui::Camera::snapshot(const std::string &filename) { { UNIQUE_LOCK(mutex_, lk); - cv::flip(im1_, flipped, 0); + //cv::flip(im1_, flipped, 0); } cv::cvtColor(flipped, flipped, cv::COLOR_BGRA2BGR); cv::imwrite(filename, flipped); diff --git a/applications/gui/src/camera.hpp b/applications/gui/src/camera.hpp index 95939a8f71200cc99e708fdd2d9e45a30a66cc41..c64bdfcaa12596460b63c00d303f7df1cab8ada9 100644 --- a/applications/gui/src/camera.hpp +++ b/applications/gui/src/camera.hpp @@ -87,8 +87,6 @@ class Camera { const GLTexture &getLeft() const { return texture1_; } const GLTexture &getRight() const { return texture2_; } - bool thumbnail(cv::Mat &thumb); - void snapshot(const std::string &filename); void startVideoRecording(const std::string &filename); @@ -135,8 +133,8 @@ class Camera { bool pause_; ftl::codecs::Channel channel_; ftl::codecs::Channels<0> channels_; - cv::Mat im1_; // first channel (left) - cv::Mat im2_; // second channel ("right") + cv::Mat overlay_; // first channel (left) + //cv::Mat im2_; // second channel ("right") bool stereo_; std::atomic_flag stale_frame_; int rx_; diff --git a/components/codecs/include/ftl/codecs/channels.hpp b/components/codecs/include/ftl/codecs/channels.hpp index 44096367899572965d1c08eb1d2459275d72575a..f2f8a833bfb2f5515a407dbcaf135929ebce4e63 100644 --- a/components/codecs/include/ftl/codecs/channels.hpp +++ b/components/codecs/include/ftl/codecs/channels.hpp @@ -36,6 +36,7 @@ enum struct Channel : int { Smoothing = 19, // 32F RightHighRes = 20, // 8UC3 or 8UC4 Colour2HighRes = 20, + Overlay = 21, // 8UC4 Audio = 32, AudioMono = 32, diff --git a/components/renderers/cpp/src/CUDARender.cpp b/components/renderers/cpp/src/CUDARender.cpp index 9dec33bcbf46369f18ce5989f02bd5bd4bcc7ee5..557392d1ced004f232f7bcf5fef3dd6262468766 100644 --- a/components/renderers/cpp/src/CUDARender.cpp +++ b/components/renderers/cpp/src/CUDARender.cpp @@ -504,7 +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); - ftl::cuda::blend_alpha(buf, out_->getTexture<uchar4>(out_chan_), alpha, 1.0f-alpha, stream_); + 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_); + } } void CUDARender::end() { diff --git a/components/renderers/cpp/src/colour_cuda.hpp b/components/renderers/cpp/src/colour_cuda.hpp index fe499af68e956c4b8c9d2fd7bab2b87264cd2df0..bad912c96f99e1d36091076f23da436dde3f8c26 100644 --- a/components/renderers/cpp/src/colour_cuda.hpp +++ b/components/renderers/cpp/src/colour_cuda.hpp @@ -17,6 +17,11 @@ void blend_alpha( float alpha, float beta, cudaStream_t stream); +void composite( + ftl::cuda::TextureObject<uchar4> &in, + ftl::cuda::TextureObject<uchar4> &out, + cudaStream_t stream); + } } diff --git a/components/renderers/cpp/src/colour_util.cu b/components/renderers/cpp/src/colour_util.cu index ae3d72f8cde48af1a615addd8782076a09e6ec3d..01939ebd37609c690ab4ec8f4cf873a357be3192 100644 --- a/components/renderers/cpp/src/colour_util.cu +++ b/components/renderers/cpp/src/colour_util.cu @@ -63,6 +63,10 @@ template void ftl::cuda::lut<short>(TextureObject<short> &in, TextureObject<ucha // ==== Blending =============================================================== +__device__ inline float clamp(float a, float c) { + return (a > c) ? c : a; +} + __global__ void blend_alpha_kernel( const uchar4* __restrict__ in, int in_pitch, @@ -77,12 +81,12 @@ __global__ void blend_alpha_kernel( const uchar4 c2 = out[x+y*out_pitch]; const float a = alpha*(float(c1.w)/255.0f); - const float b = 1.0f - a; + const float b = 1.0f - (float(c1.w)/255.0f); out[x+y*out_pitch] = make_uchar4( - float(c1.x)*a + float(c2.x)*b, - float(c1.y)*a + float(c2.y)*b, - float(c1.z)*a + float(c2.z)*b, + 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 ); } @@ -107,3 +111,48 @@ void ftl::cuda::blend_alpha( out.width(), out.height(), alpha); cudaSafeCall( cudaGetLastError() ); } + +// ==== Composite ============================================================== + +__global__ void composite_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 a = (float(c1.w)/255.0f); + const float b = 1.0f - a; + + 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::composite( + 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_kernel<<<gridSize, blockSize, 0, stream>>>( + in.devicePtr(), in.pixelPitch(), + out.devicePtr(), out.pixelPitch(), + out.width(), out.height()); + cudaSafeCall( cudaGetLastError() ); +} diff --git a/components/renderers/cpp/src/colouriser.cpp b/components/renderers/cpp/src/colouriser.cpp index 70521ba4104d6d457ab1e502170fbb22b740a60a..95e85c47116857799276e52c18b8812033b735c6 100644 --- a/components/renderers/cpp/src/colouriser.cpp +++ b/components/renderers/cpp/src/colouriser.cpp @@ -114,6 +114,7 @@ Colouriser::~Colouriser() { TextureObject<uchar4> &Colouriser::colourise(ftl::rgbd::Frame &f, Channel c, cudaStream_t stream) { switch (c) { + case Channel::Overlay : return f.createTexture<uchar4>(c); case Channel::ColourHighRes : case Channel::Colour : case Channel::Colour2 : return _processColour(f,c,stream); diff --git a/components/renderers/cpp/src/overlay.cpp b/components/renderers/cpp/src/overlay.cpp index 3607dc01521976303928e13b55cc44d60abb078e..a539c6724c3abebd140ac9c218f83943181a3c74 100644 --- a/components/renderers/cpp/src/overlay.cpp +++ b/components/renderers/cpp/src/overlay.cpp @@ -64,6 +64,8 @@ void Overlay::apply(ftl::rgbd::FrameSet &fs, cv::Mat &out, ftl::rgbd::FrameState } } } + + cv::flip(out, out, 0); } void ftl::overlay::draw3DLine(