diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 0dd0a51912e10157a44ea26d793d8364dd60a441..42d339fd03e43d12b88e5ba4acfdadf2b7406947 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/applications/reconstruct/src/main.cpp
@@ -331,9 +331,10 @@ static void run(ftl::Configurable *root) {
//pipeline1->append<ftl::operators::ScanFieldFill>("filling"); // Generate a smoothing channel
pipeline1->append<ftl::operators::CrossSupport>("cross");
pipeline1->append<ftl::operators::DiscontinuityMask>("discontinuity");
+ pipeline1->append<ftl::operators::CrossSupport>("cross2")->set("discon_support", true);
pipeline1->append<ftl::operators::CullDiscontinuity>("remove_discontinuity");
//pipeline1->append<ftl::operators::AggreMLS>("mls"); // Perform MLS (using smoothing channel)
- pipeline1->append<ftl::operators::VisCrossSupport>("viscross")->set("enabled", false);
+ pipeline1->append<ftl::operators::VisCrossSupport>("viscross"); //->set("enabled", false);
pipeline1->append<ftl::operators::MultiViewMLS>("mvmls");
// Alignment
diff --git a/components/operators/src/segmentation.cpp b/components/operators/src/segmentation.cpp
index fb015b6b16e1aa7e1dcac9735103c8c6af41567e..9bf7605e1ea9ef6b5dd73a4c3078ba70bb61abb9 100644
--- a/components/operators/src/segmentation.cpp
+++ b/components/operators/src/segmentation.cpp
@@ -14,18 +14,17 @@ CrossSupport::~CrossSupport() {
}
bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
- bool use_depth = config()->value("depth_region", false);
+ bool use_mask = config()->value("discon_support", false);
- if (use_depth) {
+ if (use_mask) {
ftl::cuda::support_region(
- in.createTexture<float>(Channel::Depth),
+ in.createTexture<int>(Channel::Mask),
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),
- config()->value("symmetric", true), stream
+ config()->value("symmetric", false), stream
);
- } //else {
+ } else {
ftl::cuda::support_region(
in.createTexture<uchar4>(Channel::Colour),
out.createTexture<uchar4>(Channel::Support1, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())),
@@ -34,7 +33,7 @@ bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd:
config()->value("h_max", 5),
config()->value("symmetric", true), stream
);
- //}
+ }
return true;
}
diff --git a/components/operators/src/segmentation.cu b/components/operators/src/segmentation.cu
index c16e647931f7d4c4d92551f44f6242fc415bd91d..aaa81e2ad98b5171572b6768754b93f32abcd1eb 100644
--- a/components/operators/src/segmentation.cu
+++ b/components/operators/src/segmentation.cu
@@ -1,8 +1,10 @@
#include "segmentation_cuda.hpp"
+#include "mask_cuda.hpp"
#define T_PER_BLOCK 8
using ftl::cuda::TextureObject;
+using ftl::cuda::Mask;
template <typename T>
__device__ inline float cross(T p1, T p2);
@@ -89,6 +91,64 @@ __device__ uchar4 calculate_support_region(const TextureObject<T> &img, int x, i
return result;
}
+__device__ uchar4 calculate_support_region(const TextureObject<int> &img, int x, int y, int v_max, int h_max) {
+ int x_min = max(0, x - h_max);
+ int x_max = min(img.width()-1, x + h_max);
+ int y_min = max(0, y - v_max);
+ int y_max = min(img.height()-1, y + v_max);
+
+ uchar4 result = make_uchar4(0, 0, 0, 0);
+
+ Mask m1(img.tex2D(x,y));
+
+ int u;
+ for (u=x-1; u >= x_min; --u) {
+ Mask m2(img.tex2D(u,y));
+ if (m2.isDiscontinuity()) {
+ result.x = x - u - 1;
+ break;
+ }
+ }
+ if (u < x_min) result.x = x - x_min;
+
+ for (u=x+1; u <= x_max; ++u) {
+ Mask m2(img.tex2D(u,y));
+ if (m2.isDiscontinuity()) {
+ result.y = u - x - 1;
+ break;
+ }
+ }
+ if (u > x_max) result.y = x_max - x;
+
+ int v;
+ for (v=y-1; v >= y_min; --v) {
+ Mask m2(img.tex2D(x,v));
+ if (m2.isDiscontinuity()) {
+ result.z = y - v - 1;
+ break;
+ }
+ }
+ if (v < y_min) result.z = y - y_min;
+
+ for (v=y+1; v <= y_max; ++v) {
+ Mask m2(img.tex2D(x,v));
+ if (m2.isDiscontinuity()) {
+ result.w = v - y - 1;
+ break;
+ }
+ }
+ if (v > y_max) result.w = y_max - y;
+
+ // Make symetric left/right and up/down
+ if (false) {
+ 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, 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;
@@ -99,6 +159,15 @@ __global__ void support_region_kernel(TextureObject<T> img, TextureObject<uchar4
region(x,y) = calculate_support_region<T,SYM>(img, x, y, tau, v_max, h_max);
}
+__global__ void support_region_kernel(TextureObject<int> img, TextureObject<uchar4> region, 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, v_max, h_max);
+}
+
void ftl::cuda::support_region(
ftl::cuda::TextureObject<uchar4> &colour,
ftl::cuda::TextureObject<uchar4> ®ion,
@@ -142,6 +211,26 @@ void ftl::cuda::support_region(
#endif
}
+void ftl::cuda::support_region(
+ ftl::cuda::TextureObject<int> &mask,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ 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>>>(mask, region, v_max, h_max);
+ cudaSafeCall( cudaGetLastError() );
+
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
+
__global__ void vis_support_region_kernel(TextureObject<uchar4> colour, TextureObject<uchar4> region, uchar4 bcolour, uchar4 acolour,
int ox, int oy, int dx, int dy) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
diff --git a/components/operators/src/segmentation_cuda.hpp b/components/operators/src/segmentation_cuda.hpp
index c2cb390d9c0ee62a33127eec1720cf0d6fee8cae..1383489337dc968a33ec630facb597920518e2bf 100644
--- a/components/operators/src/segmentation_cuda.hpp
+++ b/components/operators/src/segmentation_cuda.hpp
@@ -18,6 +18,12 @@ void support_region(
float tau, int v_max, int h_max, bool sym,
cudaStream_t stream);
+void support_region(
+ ftl::cuda::TextureObject<int> &mask,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ int v_max, int h_max, bool sym,
+ cudaStream_t stream);
+
void vis_support_region(
ftl::cuda::TextureObject<uchar4> &colour,
ftl::cuda::TextureObject<uchar4> ®ion,