Merge branch 'feature/283/fxaaenable' into 'master'

Resolves #283 re-enable fxaa approx

Closes #283

See merge request nicolas.pope/ftl!227

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

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

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(

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

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