diff --git a/components/filters/src/smoothing.cpp b/components/filters/src/smoothing.cpp index 19360840e65978a684a69edb289a21a430aa6b13..b288d3d5b7e9b248329d746974cdf069a3b9a04f 100644 --- a/components/filters/src/smoothing.cpp +++ b/components/filters/src/smoothing.cpp @@ -71,6 +71,7 @@ void MLSSmoother::filter(ftl::rgbd::Frame &f, ftl::rgbd::Source *s, cudaStream_t //bool do_smooth = value("mls_smooth", false); float thresh = value("mls_threshold", 0.04f); int iters = value("mls_iterations", 1); + int radius = value("mls_radius",2); //if (!do_smooth) return; @@ -89,6 +90,7 @@ void MLSSmoother::filter(ftl::rgbd::Frame &f, ftl::rgbd::Source *s, cudaStream_t f.createTexture<float>(Channel::Depth), f.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(f.get<cv::cuda::GpuMat>(Channel::Depth).size())), thresh, + radius, s->parameters(), 0 ); diff --git a/components/filters/src/smoothing.cu b/components/filters/src/smoothing.cu index b334d7c0e5de685c8e62363517a2be5398fa8673..903f90a199fb31d017afc1e7dcb5d7ce406848a7 100644 --- a/components/filters/src/smoothing.cu +++ b/components/filters/src/smoothing.cu @@ -77,13 +77,20 @@ void ftl::cuda::mls_smooth( ftl::cuda::TextureObject<float> &depth_in, ftl::cuda::TextureObject<float> &depth_out, float smoothing, + int radius, const ftl::rgbd::Camera &camera, cudaStream_t stream) { const dim3 gridSize((depth_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); - mls_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); + switch (radius) { + case 5: mls_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break; + case 4: mls_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break; + case 3: mls_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break; + case 2: mls_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break; + case 1: mls_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break; + } cudaSafeCall( cudaGetLastError() ); diff --git a/components/filters/src/smoothing_cuda.hpp b/components/filters/src/smoothing_cuda.hpp index d1ad1ce63859eb8bb983b6362221a419276ffd92..6ebe397e77d410a2559f4b18a44ca2b4127495ef 100644 --- a/components/filters/src/smoothing_cuda.hpp +++ b/components/filters/src/smoothing_cuda.hpp @@ -13,6 +13,7 @@ void mls_smooth( ftl::cuda::TextureObject<float> &depth_in, ftl::cuda::TextureObject<float> &depth_out, float smoothing, + int radius, const ftl::rgbd::Camera &camera, cudaStream_t stream);