diff --git a/applications/reconstruct/src/ilw/correspondence.cu b/applications/reconstruct/src/ilw/correspondence.cu index b13702f203af7809d566a2c9f58eb7f056ba29b3..8fb9d94b75df5a5d2687fa81da01d999f2bcfeeb 100644 --- a/applications/reconstruct/src/ilw/correspondence.cu +++ b/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; } } diff --git a/applications/reconstruct/src/ilw/fill.cu b/applications/reconstruct/src/ilw/fill.cu index f5545f18e7ca0a2814563a730589fea09ad70ee8..4109c1034d045ecdede9e28582cd19df670f5a63 100644 --- a/applications/reconstruct/src/ilw/fill.cu +++ b/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() ); } diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp index 49564f695235a490eb09655d3bae983671f5b4ac..93fd75a188b5e48e48e387c5ddf96ea43bcb1eec 100644 --- a/applications/reconstruct/src/ilw/ilw.cpp +++ b/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; diff --git a/components/renderers/cpp/src/normals.cu b/components/renderers/cpp/src/normals.cu index 7b385a8c5c8713ed8bc0469ead391061f874cce0..14357d02ed00f212577e463051d315e6e5a90078 100644 --- a/components/renderers/cpp/src/normals.cu +++ b/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;