Add unimpl depth fxaa

components/operators/src/antialiasing.cpp

@@ -13,10 +13,18 @@ FXAA::~FXAA() {
 }
 
 bool FXAA::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t stream) {
+	if (in.hasChannel(Channel::Depth)) {
 		ftl::cuda::fxaa(
 			in.getTexture<uchar4>(Channel::Colour),
+			in.getTexture<float>(Channel::Depth),
 			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,41 @@ __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) {
+
+    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
+        fxaa2(x, y, data);
+    }
+}
+
+void ftl::cuda::fxaa(ftl::cuda::TextureObject<uchar4> &colour, ftl::cuda::TextureObject<float> &depth, 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);
+	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, cudaStream_t stream);
 
 }
 }