From f9d94b8a9853c923025d00c27fcd48d3fb74d2d8 Mon Sep 17 00:00:00 2001 From: Nicolas Pope <nwpope@utu.fi> Date: Sun, 3 Nov 2019 12:20:42 +0200 Subject: [PATCH] Add radius option and correct coord bug --- components/operators/src/smoothchan.cu | 21 ++++++++++++++++----- components/operators/src/smoothing.cpp | 3 ++- 2 files changed, 18 insertions(+), 6 deletions(-) diff --git a/components/operators/src/smoothchan.cu b/components/operators/src/smoothchan.cu index 963fc85ff..edf53be23 100644 --- a/components/operators/src/smoothchan.cu +++ b/components/operators/src/smoothchan.cu @@ -18,8 +18,8 @@ __global__ void smooth_chan_kernel( const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; if (x < smoothing_out.width() && y < smoothing_out.height()) { - const int ix = (float)x / (float)smoothing_out.width() * (float)colour_in.width(); - const int iy = (float)x / (float)smoothing_out.height() * (float)colour_in.height(); + const int ix = int(((float)x / (float)smoothing_out.width()) * (float)colour_in.width()); + const int iy = int(((float)y / (float)smoothing_out.height()) * (float)colour_in.height()); // A distance has already been found if (smoothing_out(x,y) < 1.0f) return; @@ -27,13 +27,18 @@ __global__ void smooth_chan_kernel( float mindist = 1.0f; const uchar4 c0 = colour_in.tex2D(ix, iy); - const float3 pos = camera.screenToCam(ix, iy, depth_in.tex2D(ix, iy)); + const float d0 = depth_in.tex2D(ix, iy); + if (d0 < camera.minDepth || d0 > camera.maxDepth) return; + + const float3 pos = camera.screenToCam(ix, iy, d0); for (int v=-RADIUS; v<=RADIUS; ++v) { #pragma unroll for (int u=-RADIUS; u<=RADIUS; ++u) { const uchar4 c = colour_in.tex2D(ix+u, iy+v); - const float3 posN = camera.screenToCam(ix+u, iy+v, depth_in.tex2D(ix+u, iy+v)); + const float d = depth_in.tex2D(ix+u, iy+v); + if (d < camera.minDepth || d > camera.maxDepth) continue; + const float3 posN = camera.screenToCam(ix+u, iy+v, d); if (ftl::cuda::colourDistance(c, c0) >= alpha) mindist = min(mindist, length(pos - posN)); } @@ -56,7 +61,13 @@ void ftl::cuda::smooth_channel( const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); - smooth_chan_kernel<1><<<gridSize, blockSize, 0, stream>>>(colour_in, depth_in, smoothing_out, camera, alpha); + switch (radius) { + case 5: smooth_chan_kernel<5><<<gridSize, blockSize, 0, stream>>>(colour_in, depth_in, smoothing_out, camera, alpha); break; + case 4: smooth_chan_kernel<4><<<gridSize, blockSize, 0, stream>>>(colour_in, depth_in, smoothing_out, camera, alpha); break; + case 3: smooth_chan_kernel<3><<<gridSize, blockSize, 0, stream>>>(colour_in, depth_in, smoothing_out, camera, alpha); break; + case 2: smooth_chan_kernel<2><<<gridSize, blockSize, 0, stream>>>(colour_in, depth_in, smoothing_out, camera, alpha); break; + case 1: smooth_chan_kernel<1><<<gridSize, blockSize, 0, stream>>>(colour_in, depth_in, smoothing_out, camera, alpha); break; + } cudaSafeCall( cudaGetLastError() ); diff --git a/components/operators/src/smoothing.cpp b/components/operators/src/smoothing.cpp index 284658e86..01f7f7de6 100644 --- a/components/operators/src/smoothing.cpp +++ b/components/operators/src/smoothing.cpp @@ -8,6 +8,7 @@ using ftl::operators::SimpleMLS; using ftl::operators::ColourMLS; using ftl::operators::SmoothChannel; using ftl::codecs::Channel; +using ftl::rgbd::Format; using cv::cuda::GpuMat; HFSmoother::HFSmoother(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) { @@ -74,7 +75,7 @@ bool SmoothChannel::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd float threshold = config()->value("threshold", 30.0f); // Clear to max smoothing - out.get<GpuMat>(Channel::Smoothing).setTo(cv::Scalar(1.0f)); + out.create<GpuMat>(Channel::Smoothing, Format<float>(s->parameters().width, s->parameters().height)).setTo(cv::Scalar(1.0f)); // Reduce to nearest ftl::cuda::smooth_channel( -- GitLab