diff --git a/components/filters/src/smoothing.cpp b/components/filters/src/smoothing.cpp
index 517179976fc5e5a3db73e769f5ad8e0cdb796241..e2805174544875aaec56cc8c092a31e5173dfaa3 100644
--- a/components/filters/src/smoothing.cpp
+++ b/components/filters/src/smoothing.cpp
@@ -71,24 +71,28 @@ void MLSSmoother::smooth(ftl::rgbd::Frame &f, ftl::rgbd::Source *s) {
 	bool do_smooth = value("mls_smooth", false);
 	if (!do_smooth) return;
 
-	if (!f.hasChannel(Channel::Normals)) {
-		ftl::cuda::normals(
-			f.createTexture<float4>(Channel::Normals, ftl::rgbd::Format<float4>(f.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+	float thresh = value("mls_threshold", 0.04f);
+	int iters = value("mls_iterations", 1);
+
+	for (int i=0; i<iters; ++i) {
+		if (i > 0 || !f.hasChannel(Channel::Normals)) {
+			ftl::cuda::normals(
+				f.createTexture<float4>(Channel::Normals, ftl::rgbd::Format<float4>(f.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+				f.createTexture<float>(Channel::Depth),
+				s->parameters(), 0
+			);
+		}
+
+		ftl::cuda::mls_smooth(
+			f.createTexture<float4>(Channel::Normals),
 			f.createTexture<float>(Channel::Depth),
-			s->parameters(), 0
+			f.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(f.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+			thresh,
+			s->parameters(),
+			0
 		);
-	}
 
-	float thresh = value("mls_threshold", 0.04f);
-	ftl::cuda::mls_smooth(
-		f.createTexture<float4>(Channel::Normals),
-		f.createTexture<float>(Channel::Depth),
-		f.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(f.get<cv::cuda::GpuMat>(Channel::Depth).size())),
-		thresh,
-		s->parameters(),
-		0
-	);
-
-	f.swapChannels(Channel::Depth, Channel::Depth2);
+		f.swapChannels(Channel::Depth, Channel::Depth2);
+	}
 }
 
diff --git a/components/filters/src/smoothing.cu b/components/filters/src/smoothing.cu
index 0ba749d25c57e23496df0a59b5428dea0a2c0d27..d7d2a8a3fcccc903c70f2cb39adfa169ee20acd7 100644
--- a/components/filters/src/smoothing.cu
+++ b/components/filters/src/smoothing.cu
@@ -79,7 +79,7 @@ void ftl::cuda::mls_smooth(
 	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);
 
-	mls_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(normals, depth_in, depth_out, smoothing, camera);
+	mls_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(normals, depth_in, depth_out, smoothing, camera);
 	cudaSafeCall( cudaGetLastError() );