Iterate smoothing kernel

applications/reconstruct/src/filters/smoothing.cu

@@ -52,30 +52,32 @@ void ftl::cuda::depth_smooth(
 		ftl::cuda::TextureObject<uchar4> &colour_in,
 		ftl::cuda::TextureObject<float> &depth_out,
 		const ftl::rgbd::Camera &camera,
-		int radius, float factor, float thresh, cudaStream_t stream) {
+		int radius, float factor, float thresh, int iters, cudaStream_t stream) {
 
 	const dim3 gridSize((depth_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
 	const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
 
-	switch (radius) {
-    case 5 :	depth_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
-	case 4 :	depth_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
-	case 3 :	depth_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
-	case 2 :	depth_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
-	case 1 :	depth_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
-	default:	break;
-	}
-	cudaSafeCall( cudaGetLastError() );
+	for (int n=0; n<iters; ++n) {
+		switch (radius) {
+		case 5 :	depth_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
+		case 4 :	depth_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
+		case 3 :	depth_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
+		case 2 :	depth_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
+		case 1 :	depth_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
+		default:	break;
+		}
+		cudaSafeCall( cudaGetLastError() );
 
-	switch (radius) {
-	case 5 :	depth_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
-	case 4 :	depth_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
-	case 3 :	depth_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
-	case 2 :	depth_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
-	case 1 :	depth_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
-	default:	break;
+		switch (radius) {
+		case 5 :	depth_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
+		case 4 :	depth_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
+		case 3 :	depth_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
+		case 2 :	depth_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
+		case 1 :	depth_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
+		default:	break;
+		}
+		cudaSafeCall( cudaGetLastError() );
 	}
-	cudaSafeCall( cudaGetLastError() );
 
 #ifdef _DEBUG
 	cudaSafeCall(cudaDeviceSynchronize());

applications/reconstruct/src/filters/smoothing.hpp

@@ -12,7 +12,7 @@ void depth_smooth(
 	ftl::cuda::TextureObject<uchar4> &colour_in,
 	ftl::cuda::TextureObject<float> &depth_out,
 	const ftl::rgbd::Camera &camera,
-	int radius, float factor, float thresh,
+	int radius, float factor, float thresh, int iters,
 	cudaStream_t stream);
 
 }

applications/reconstruct/src/main.cpp

@@ -245,9 +245,11 @@ static void run(ftl::Configurable *root) {
 
 	bool busy = false;
 
+	auto *filter = ftl::config::create<ftl::Configurable>(root, "filters");
+
 	group->setLatency(4);
 	group->setName("ReconGroup");
-	group->sync([splat,virt,&busy,&slave,&scene_A,&scene_B,&align,controls](ftl::rgbd::FrameSet &fs) -> bool {
+	group->sync([splat,virt,&busy,&slave,&scene_A,&scene_B,&align,controls,filter](ftl::rgbd::FrameSet &fs) -> bool {
 		//cudaSetDevice(scene->getCUDADevice());
 
 		//if (slave.isPaused()) return true;
@@ -262,7 +264,7 @@ static void run(ftl::Configurable *root) {
 		// Swap the entire frameset to allow rapid return
 		fs.swapTo(scene_A);
 
-		ftl::pool.push([&scene_B,&scene_A,&busy,&slave,&align](int id) {
+		ftl::pool.push([&scene_B,&scene_A,&busy,&slave,&align, filter](int id) {
 			//cudaSetDevice(scene->getCUDADevice());
 			// TODO: Release frameset here...
 			//cudaSafeCall(cudaStreamSynchronize(scene->getIntegrationStream()));
@@ -270,9 +272,10 @@ static void run(ftl::Configurable *root) {
 			UNIQUE_LOCK(scene_A.mtx, lk);
 
 			cv::cuda::GpuMat tmp;
-			float factor = align->value("smooth_factor", 0.04f);
-			float colour_limit = align->value("colour_limit", 50.0f);
-			bool do_smooth = align->value("pre_smooth", false);
+			float factor = filter->value("smooth_factor", 0.04f);
+			float colour_limit = filter->value("colour_limit", 50.0f);
+			bool do_smooth = filter->value("pre_smooth", false);
+			int iters = filter->value("iterations", 1);
 
 			if (do_smooth) {
 				// Presmooth...
@@ -295,7 +298,7 @@ static void run(ftl::Configurable *root) {
 						f.createTexture<uchar4>(Channel::Colour),
 						f.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(f.get<cv::cuda::GpuMat>(Channel::Depth).size())),
 						s->parameters(),
-						1, factor, colour_limit, 0
+						1, factor, colour_limit, iters, 0
 					);
 				}
 			}