Changes to normals and mask filtering

applications/reconstruct/src/ilw/correspondence.cu

@@ -176,7 +176,8 @@ __global__ void mask_filter_kernel(
 	const int y = blockIdx.y*blockDim.y + threadIdx.y;
 
 	if (x >= 0 && x < depth.width() && y >=0 && y < depth.height()) {
-		if (mask.tex2D(x,y) > 0) {
+		Mask m(mask.tex2D(x,y));
+		if (m.isFilled()) {
 			depth(x,y) = MINF;
 		}
 	}

applications/reconstruct/src/ilw/fill.cu

@@ -65,7 +65,7 @@ void ftl::cuda::preprocess_depth(
 	const dim3 gridSize((depth_in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
 	const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
 
-	preprocess_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_in, depth_out, colour, mask, camera, params);
+	preprocess_kernel<10><<<gridSize, blockSize, 0, stream>>>(depth_in, depth_out, colour, mask, camera, params);
 
 	cudaSafeCall( cudaGetLastError() );
 }

applications/reconstruct/src/ilw/ilw.cpp

@@ -170,7 +170,7 @@ bool ILW::process(ftl::rgbd::FrameSet &fs) {
 			
         auto &t = f.createTexture<float4>(Channel::Points, Format<float4>(f.get<GpuMat>(Channel::Colour).size()));
         auto pose = MatrixConversion::toCUDA(s->getPose().cast<float>()); //.inverse());
-        ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, 0, stream_);
+        ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, discon_mask_, stream_);
     }
 
 	cudaSafeCall(cudaStreamSynchronize(stream_));
@@ -320,11 +320,6 @@ bool ILW::_phase1(ftl::rgbd::FrameSet &fs, int win, cudaStream_t stream) {
                 LOG(ERROR) << "Exception in correspondence: " << e.what();
             }
 
-			if (f1.hasChannel(Channel::Mask)) {
-				ftl::cuda::mask_filter(f1.getTexture<float>(Channel::Depth2),
-					f1.getTexture<int>(Channel::Mask), stream_);
-			}
-
             //LOG(INFO) << "Correspondences done... " << i;
         }
     }
@@ -357,6 +352,11 @@ bool ILW::_phase2(ftl::rgbd::FrameSet &fs, float rate, cudaStream_t stream) {
              motion_window_,
              stream
         );
+
+        if (f.hasChannel(Channel::Mask)) {
+			ftl::cuda::mask_filter(f.getTexture<float>(Channel::Depth),
+				f.getTexture<int>(Channel::Mask), stream_);
+		}
     }
 
     return true;

components/renderers/cpp/src/normals.cu

@@ -77,6 +77,8 @@ __global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms,
     float3 nsum = make_float3(0.0f);
     float contrib = 0.0f;
 
+    output(x,y) = make_float4(0.0f,0.0f,0.0f,0.0f);
+
     if (p0.x == MINF) return;
 
     for (int v=-RADIUS; v<=RADIUS; ++v) {
@@ -120,6 +122,8 @@ __global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms,
     float3 nsum = make_float3(0.0f);
     float contrib = 0.0f;
 
+    output(x,y) = make_float4(0.0f,0.0f,0.0f,0.0f);
+
     if (p0.z < camera.minDepth || p0.z > camera.maxDepth) return;
 
     for (int v=-RADIUS; v<=RADIUS; ++v) {
@@ -159,6 +163,8 @@ __global__ void smooth_normals_kernel<0>(ftl::cuda::TextureObject<float4> norms,
 
     if(x >= depth.width() || y >= depth.height()) return;
 
+    output(x,y) = make_float4(0.0f,0.0f,0.0f,0.0f);
+
     const float3 p0 = camera.screenToCam(x,y, (float)depth.tex2D((int)x,(int)y) / 1000.0f);
 
     if (p0.z < camera.minDepth || p0.z > camera.maxDepth) return;