diff --git a/components/operators/src/smoothchan.cu b/components/operators/src/smoothchan.cu index edf53be232aa951ac716890e823495b452effa3c..08fc3cb43a0f48dcace0bac7dd679144016d5a1c 100644 --- a/components/operators/src/smoothchan.cu +++ b/components/operators/src/smoothchan.cu @@ -9,10 +9,10 @@ using ftl::cuda::TextureObject; template <int RADIUS> __global__ void smooth_chan_kernel( ftl::cuda::TextureObject<uchar4> colour_in, - ftl::cuda::TextureObject<float> depth_in, + //ftl::cuda::TextureObject<float> depth_in, ftl::cuda::TextureObject<float> smoothing_out, ftl::rgbd::Camera camera, - float alpha) { + float alpha, float scale) { const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; @@ -24,27 +24,29 @@ __global__ void smooth_chan_kernel( // A distance has already been found if (smoothing_out(x,y) < 1.0f) return; - float mindist = 1.0f; + float mindist = 50.0f; const uchar4 c0 = colour_in.tex2D(ix, iy); - const float d0 = depth_in.tex2D(ix, iy); - if (d0 < camera.minDepth || d0 > camera.maxDepth) return; + //const float d0 = depth_in.tex2D(ix, iy); + //if (d0 < camera.minDepth || d0 > camera.maxDepth) return; - const float3 pos = camera.screenToCam(ix, iy, d0); + //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 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); + //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)); + float d = sqrt((float)u*(float)u + (float)v*(float)v) * scale; + + if (ftl::cuda::colourDistance(c, c0) >= alpha) mindist = min(mindist, (float)d); } } - smoothing_out(x,y) = min(mindist, 1.0f); + smoothing_out(x,y) = min(mindist / 50.0f, 1.0f); } } @@ -54,6 +56,7 @@ void ftl::cuda::smooth_channel( ftl::cuda::TextureObject<float> &smoothing_out, const ftl::rgbd::Camera &camera, float alpha, + float scale, int radius, cudaStream_t stream) { @@ -62,11 +65,11 @@ void ftl::cuda::smooth_channel( 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; + case 5: smooth_chan_kernel<5><<<gridSize, blockSize, 0, stream>>>(colour_in, smoothing_out, camera, alpha, scale); break; + case 4: smooth_chan_kernel<4><<<gridSize, blockSize, 0, stream>>>(colour_in, smoothing_out, camera, alpha, scale); break; + case 3: smooth_chan_kernel<3><<<gridSize, blockSize, 0, stream>>>(colour_in, smoothing_out, camera, alpha, scale); break; + case 2: smooth_chan_kernel<2><<<gridSize, blockSize, 0, stream>>>(colour_in, smoothing_out, camera, alpha, scale); break; + case 1: smooth_chan_kernel<1><<<gridSize, blockSize, 0, stream>>>(colour_in, smoothing_out, camera, alpha, scale); break; } diff --git a/components/operators/src/smoothing.cpp b/components/operators/src/smoothing.cpp index 61dd9f8ce24fa26dea6854d271c56eb284a4e630..bb3b70677393ac7f15ccc6cc3fd533c65fdc19b7 100644 --- a/components/operators/src/smoothing.cpp +++ b/components/operators/src/smoothing.cpp @@ -77,6 +77,7 @@ bool SmoothChannel::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd int width = s->parameters().width; int height = s->parameters().height; + float scale = 1.0f; // Clear to max smoothing out.create<GpuMat>(Channel::Smoothing, Format<float>(width, height)).setTo(cv::Scalar(1.0f)); @@ -88,6 +89,7 @@ bool SmoothChannel::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd out.createTexture<float>(Channel::Smoothing), s->parameters(), threshold, + scale, radius, stream ); @@ -95,6 +97,7 @@ bool SmoothChannel::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd for (int i=0; i<iters; ++i) { width /= 2; height /= 2; + scale *= 2.0f; ftl::rgbd::Camera scaledCam = s->parameters().scaled(width, height); @@ -108,6 +111,7 @@ bool SmoothChannel::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd out.getTexture<float>(Channel::Smoothing), scaledCam, threshold, + scale, radius, stream ); diff --git a/components/operators/src/smoothing_cuda.hpp b/components/operators/src/smoothing_cuda.hpp index dae25e221b1121bbfd02bab732b5c6ba5f465aff..1b5079d1c94b5164f024aacde9d3e4897fc87f11 100644 --- a/components/operators/src/smoothing_cuda.hpp +++ b/components/operators/src/smoothing_cuda.hpp @@ -35,6 +35,7 @@ void smooth_channel( ftl::cuda::TextureObject<float> &smoothing_out, const ftl::rgbd::Camera &camera, float alpha, + float scale, int radius, cudaStream_t stream);