From a68d53fbca0fe5aa11591db6df7340423649be16 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Mon, 11 Nov 2019 20:27:11 +0200
Subject: [PATCH] Allow any radius
---
components/operators/src/mask.cu | 30 +++++++++++-------------------
1 file changed, 11 insertions(+), 19 deletions(-)
diff --git a/components/operators/src/mask.cu b/components/operators/src/mask.cu
index 397935c5d..b18b71e47 100644
--- a/components/operators/src/mask.cu
+++ b/components/operators/src/mask.cu
@@ -4,8 +4,7 @@
using ftl::cuda::Mask;
-template <int RADIUS>
-__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<uchar4> support, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float threshold) {
+__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<uchar4> support, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float threshold, int radius) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
@@ -28,20 +27,20 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
//}
const uchar4 sup = support.tex2D((int)x, (int)y);
- if (sup.x <= RADIUS) {
- float dS = depth.tex2D((int)x - sup.x - RADIUS, (int)y);
+ if (sup.x <= radius) {
+ float dS = depth.tex2D((int)x - sup.x - radius, (int)y);
if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
}
- if (sup.y <= RADIUS) {
- float dS = depth.tex2D((int)x + sup.y + RADIUS, (int)y);
+ if (sup.y <= radius) {
+ float dS = depth.tex2D((int)x + sup.y + radius, (int)y);
if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
}
- if (sup.z <= RADIUS) {
- float dS = depth.tex2D((int)x, (int)y - sup.z - RADIUS);
+ if (sup.z <= radius) {
+ float dS = depth.tex2D((int)x, (int)y - sup.z - radius);
if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
}
- if (sup.w <= RADIUS) {
- float dS = depth.tex2D((int)x, (int)y + sup.w + RADIUS);
+ if (sup.w <= radius) {
+ float dS = depth.tex2D((int)x, (int)y + sup.w + radius);
if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
}
}
@@ -54,14 +53,7 @@ void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda
const dim3 gridSize((params.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (params.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
- switch (discon) {
- case 5 : discontinuity_kernel<5><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
- case 4 : discontinuity_kernel<4><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
- case 3 : discontinuity_kernel<3><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
- case 2 : discontinuity_kernel<2><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
- case 1 : discontinuity_kernel<1><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
- default: break;
- }
+ discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh, discon);
cudaSafeCall( cudaGetLastError() );
#ifdef _DEBUG
@@ -77,7 +69,7 @@ __global__ void cull_discontinuity_kernel(ftl::cuda::TextureObject<int> mask, ft
if (x < depth.width() && y < depth.height()) {
Mask m(mask.tex2D((int)x,(int)y));
- if (m.isDiscontinuity()) depth(x,y) = 0.0f;
+ if (m.isDiscontinuity()) depth(x,y) = 1000.0f;
}
}
--
GitLab