diff --git a/applications/gui/src/camera.cpp b/applications/gui/src/camera.cpp index 610be4ad8a7deacebedc16b4c55d914421642d9f..585d99a4ce9c3ed5383fddab8a698892ce453339 100644 --- a/applications/gui/src/camera.cpp +++ b/applications/gui/src/camera.cpp @@ -3,6 +3,8 @@ #include "screen.hpp" #include <nanogui/glutil.h> +#include <ftl/operators/antialiasing.hpp> + #include <fstream> #ifdef HAVE_OPENVR @@ -158,6 +160,7 @@ ftl::gui::Camera::Camera(ftl::gui::Screen *screen, int fsid, int fid, ftl::codec //posewin_->setVisible(false); posewin_ = nullptr; renderer_ = nullptr; + post_pipe_ = nullptr; record_stream_ = nullptr; /*src->setCallback([this](int64_t ts, ftl::rgbd::Frame &frame) { @@ -231,6 +234,13 @@ void ftl::gui::Camera::_draw(ftl::rgbd::FrameSet &fs) { // TODO: Insert post-render pipeline. // FXAA + Bad colour removal + if (!post_pipe_) { + post_pipe_ = ftl::config::create<ftl::operators::Graph>(screen_->root(), "post_filters"); + post_pipe_->append<ftl::operators::FXAA>("fxaa"); + } + + post_pipe_->apply(frame_, frame_, 0); + _downloadFrames(&frame_); if (record_stream_ && record_stream_->active()) { diff --git a/applications/gui/src/camera.hpp b/applications/gui/src/camera.hpp index 3c68015ee0b2f181f66412e1770fabe1b984fa36..c556961086d0b3c8336ee30574ba12dd570c2324 100644 --- a/applications/gui/src/camera.hpp +++ b/applications/gui/src/camera.hpp @@ -118,6 +118,7 @@ class Camera { cv::Mat im2_; // second channel ("right") ftl::render::Triangular *renderer_; + ftl::operators::Graph *post_pipe_; ftl::rgbd::Frame frame_; ftl::rgbd::FrameState state_; ftl::stream::File *record_stream_; diff --git a/components/operators/src/antialiasing.cpp b/components/operators/src/antialiasing.cpp index 00bfe3890983f060fe490e003e6c849ab0c47357..1e6ef8f6c2f2f33de83fb7d113f2bd60b6f1d1f8 100644 --- a/components/operators/src/antialiasing.cpp +++ b/components/operators/src/antialiasing.cpp @@ -13,10 +13,19 @@ FXAA::~FXAA() { } bool FXAA::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t stream) { - ftl::cuda::fxaa( - in.getTexture<uchar4>(Channel::Colour), - stream - ); + if (in.hasChannel(Channel::Depth)) { + ftl::cuda::fxaa( + in.getTexture<uchar4>(Channel::Colour), + in.getTexture<float>(Channel::Depth), + config()->value("threshold", 0.1f), + stream + ); + } else { + ftl::cuda::fxaa( + in.getTexture<uchar4>(Channel::Colour), + stream + ); + } if (in.hasChannel(Channel::Right)) { ftl::cuda::fxaa( diff --git a/components/operators/src/antialiasing.cu b/components/operators/src/antialiasing.cu index 7ad851a2af3508b6a0728b0081fc00ec879ac088..2ad573db0b206b304e43ff4e5ab86bbbf7c5e780 100644 --- a/components/operators/src/antialiasing.cu +++ b/components/operators/src/antialiasing.cu @@ -6,16 +6,7 @@ __device__ inline uchar4 toChar(const float4 rgba) { return make_uchar4(rgba.x*255.0f, rgba.y*255.0f, rgba.z*255.0f, 255); } -__global__ void filter_fxaa2(ftl::cuda::TextureObject<uchar4> data) { - - int x = blockIdx.x*blockDim.x + threadIdx.x; - int y = blockIdx.y*blockDim.y + threadIdx.y; - - if(x >= data.width() || y >= data.height()) - { - return; - } - +__device__ void fxaa2(int x, int y, ftl::cuda::TextureObject<uchar4> &data) { uchar4 out_color; cudaTextureObject_t texRef = data.cudaTexture(); @@ -75,6 +66,51 @@ __global__ void filter_fxaa2(ftl::cuda::TextureObject<uchar4> data) { data(x,y) = out_color; } +__global__ void filter_fxaa2(ftl::cuda::TextureObject<uchar4> data, ftl::cuda::TextureObject<float> depth, float threshold) { + + int x = blockIdx.x*blockDim.x + threadIdx.x; + int y = blockIdx.y*blockDim.y + threadIdx.y; + + if(x >= 0 && x < data.width() && y >= 0 && y < data.height()) + { + // Do a depth discon test + bool discon = false; + float d = depth.tex2D(x,y); + #pragma unroll + for (int u=-1; u<=1; ++u) { + #pragma unroll + for (int v=-1; v<=1; ++v) { + discon |= fabsf(d-depth.tex2D(x+u,y+v)) > threshold; + } + } + + if (discon) fxaa2(x, y, data); + } +} + +void ftl::cuda::fxaa(ftl::cuda::TextureObject<uchar4> &colour, ftl::cuda::TextureObject<float> &depth, float threshold, cudaStream_t stream) { + const dim3 gridSize((colour.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (colour.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); + const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); + + filter_fxaa2<<<gridSize, blockSize, 0, stream>>>(colour, depth, threshold); + cudaSafeCall( cudaGetLastError() ); + +#ifdef _DEBUG + cudaSafeCall(cudaDeviceSynchronize()); +#endif +} + +__global__ void filter_fxaa2(ftl::cuda::TextureObject<uchar4> data) { + + int x = blockIdx.x*blockDim.x + threadIdx.x; + int y = blockIdx.y*blockDim.y + threadIdx.y; + + if(x >= 0 && x < data.width() && y >= 0 && y < data.height()) + { + fxaa2(x, y, data); + } +} + void ftl::cuda::fxaa(ftl::cuda::TextureObject<uchar4> &colour, cudaStream_t stream) { const dim3 gridSize((colour.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (colour.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); diff --git a/components/operators/src/antialiasing_cuda.hpp b/components/operators/src/antialiasing_cuda.hpp index afe2c5246f43f1027c2cecfe25b368bbb4cf3e20..36249ccc7f686680819704dc0fe8e1adcce533ab 100644 --- a/components/operators/src/antialiasing_cuda.hpp +++ b/components/operators/src/antialiasing_cuda.hpp @@ -7,6 +7,7 @@ namespace ftl { namespace cuda { void fxaa(ftl::cuda::TextureObject<uchar4> &colour, cudaStream_t stream); +void fxaa(ftl::cuda::TextureObject<uchar4> &colour, ftl::cuda::TextureObject<float> &depth, float threshold, cudaStream_t stream); } }