From 98e6cbe7f73fec86839b433459fc85c70b5094bc Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nicolas.pope@utu.fi>
Date: Sun, 17 Nov 2019 14:37:18 +0200
Subject: [PATCH] Implements #247 symmetric supports
---
components/operators/src/segmentation.cpp | 6 +++--
components/operators/src/segmentation.cu | 22 ++++++++++++++-----
.../operators/src/segmentation_cuda.hpp | 4 ++--
3 files changed, 22 insertions(+), 10 deletions(-)
diff --git a/components/operators/src/segmentation.cpp b/components/operators/src/segmentation.cpp
index 08b4de467..fb015b6b1 100644
--- a/components/operators/src/segmentation.cpp
+++ b/components/operators/src/segmentation.cpp
@@ -22,7 +22,8 @@ bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd:
out.createTexture<uchar4>(Channel::Support2, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())),
config()->value("depth_tau", 0.04f),
config()->value("v_max", 5),
- config()->value("h_max", 5), stream
+ config()->value("h_max", 5),
+ config()->value("symmetric", true), stream
);
} //else {
ftl::cuda::support_region(
@@ -30,7 +31,8 @@ bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd:
out.createTexture<uchar4>(Channel::Support1, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())),
config()->value("tau", 5.0f),
config()->value("v_max", 5),
- config()->value("h_max", 5), stream
+ config()->value("h_max", 5),
+ config()->value("symmetric", true), stream
);
//}
diff --git a/components/operators/src/segmentation.cu b/components/operators/src/segmentation.cu
index 3bfcbc0f5..c16e64793 100644
--- a/components/operators/src/segmentation.cu
+++ b/components/operators/src/segmentation.cu
@@ -22,7 +22,7 @@ __device__ inline float cross<float>(float p1, float p2) {
return fabs(p1-p2);
}
-template <typename T>
+template <typename T, bool SYM>
__device__ uchar4 calculate_support_region(const TextureObject<T> &img, int x, int y, float tau, int v_max, int h_max) {
int x_min = max(0, x - h_max);
int x_max = min(img.width()-1, x + h_max);
@@ -79,17 +79,24 @@ __device__ uchar4 calculate_support_region(const TextureObject<T> &img, int x, i
}
if (v > y_max) result.w = y_max - y;
+ // Make symetric left/right and up/down
+ if (SYM) {
+ result.x = min(result.x, result.y);
+ result.y = result.x;
+ result.z = min(result.z, result.w);
+ result.w = result.z;
+ }
return result;
}
-template <typename T>
+template <typename T, bool SYM>
__global__ void support_region_kernel(TextureObject<T> img, TextureObject<uchar4> region, float tau, int v_max, int h_max) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x < 0 || y < 0 || x >= img.width() || y >= img.height()) return;
- region(x,y) = calculate_support_region(img, x, y, tau, v_max, h_max);
+ region(x,y) = calculate_support_region<T,SYM>(img, x, y, tau, v_max, h_max);
}
void ftl::cuda::support_region(
@@ -97,13 +104,15 @@ void ftl::cuda::support_region(
ftl::cuda::TextureObject<uchar4> ®ion,
float tau,
int v_max,
- int h_max,
+ int h_max,
+ bool sym,
cudaStream_t stream) {
const dim3 gridSize((region.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (region.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
- support_region_kernel<<<gridSize, blockSize, 0, stream>>>(colour, region, tau, v_max, h_max);
+ if (sym) support_region_kernel<uchar4, true><<<gridSize, blockSize, 0, stream>>>(colour, region, tau, v_max, h_max);
+ else support_region_kernel<uchar4, false><<<gridSize, blockSize, 0, stream>>>(colour, region, tau, v_max, h_max);
cudaSafeCall( cudaGetLastError() );
@@ -118,12 +127,13 @@ void ftl::cuda::support_region(
float tau,
int v_max,
int h_max,
+ bool sym,
cudaStream_t stream) {
const dim3 gridSize((region.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (region.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
- support_region_kernel<<<gridSize, blockSize, 0, stream>>>(depth, region, tau, v_max, h_max);
+ support_region_kernel<float, true><<<gridSize, blockSize, 0, stream>>>(depth, region, tau, v_max, h_max);
cudaSafeCall( cudaGetLastError() );
diff --git a/components/operators/src/segmentation_cuda.hpp b/components/operators/src/segmentation_cuda.hpp
index 5ef55abb1..c2cb390d9 100644
--- a/components/operators/src/segmentation_cuda.hpp
+++ b/components/operators/src/segmentation_cuda.hpp
@@ -9,13 +9,13 @@ namespace cuda {
void support_region(
ftl::cuda::TextureObject<uchar4> &colour,
ftl::cuda::TextureObject<uchar4> ®ion,
- float tau, int v_max, int h_max,
+ float tau, int v_max, int h_max, bool sym,
cudaStream_t stream);
void support_region(
ftl::cuda::TextureObject<float> &depth,
ftl::cuda::TextureObject<uchar4> ®ion,
- float tau, int v_max, int h_max,
+ float tau, int v_max, int h_max, bool sym,
cudaStream_t stream);
void vis_support_region(
--
GitLab