Add radius option and correct coord bug

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() );

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(