diff --git a/components/operators/src/antialiasing.cpp b/components/operators/src/antialiasing.cpp index 25986e63ea3196993a7c385ee763e93a8fd22a9f..1e6ef8f6c2f2f33de83fb7d113f2bd60b6f1d1f8 100644 --- a/components/operators/src/antialiasing.cpp +++ b/components/operators/src/antialiasing.cpp @@ -17,6 +17,7 @@ bool FXAA::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t strea ftl::cuda::fxaa( in.getTexture<uchar4>(Channel::Colour), in.getTexture<float>(Channel::Depth), + config()->value("threshold", 0.1f), stream ); } else { diff --git a/components/operators/src/antialiasing.cu b/components/operators/src/antialiasing.cu index 5d991d8d6c7c1e91746f88809c1beba7a693616a..2ad573db0b206b304e43ff4e5ab86bbbf7c5e780 100644 --- a/components/operators/src/antialiasing.cu +++ b/components/operators/src/antialiasing.cu @@ -66,7 +66,7 @@ __device__ void fxaa2(int x, int y, ftl::cuda::TextureObject<uchar4> &data) { data(x,y) = out_color; } -__global__ void filter_fxaa2(ftl::cuda::TextureObject<uchar4> data, ftl::cuda::TextureObject<float> depth) { +__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; @@ -74,15 +74,25 @@ __global__ void filter_fxaa2(ftl::cuda::TextureObject<uchar4> data, ftl::cuda::T if(x >= 0 && x < data.width() && y >= 0 && y < data.height()) { // Do a depth discon test - fxaa2(x, y, data); + 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, cudaStream_t stream) { +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); + filter_fxaa2<<<gridSize, blockSize, 0, stream>>>(colour, depth, threshold); cudaSafeCall( cudaGetLastError() ); #ifdef _DEBUG diff --git a/components/operators/src/antialiasing_cuda.hpp b/components/operators/src/antialiasing_cuda.hpp index 708b41de12e437499519431c9d04563ca55ef328..36249ccc7f686680819704dc0fe8e1adcce533ab 100644 --- a/components/operators/src/antialiasing_cuda.hpp +++ b/components/operators/src/antialiasing_cuda.hpp @@ -7,7 +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); +void fxaa(ftl::cuda::TextureObject<uchar4> &colour, ftl::cuda::TextureObject<float> &depth, float threshold, cudaStream_t stream); } }