diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp index d8dd66b59323d5ddf5aa0f136f4c25f83213441a..d05453f9410ee52e8c5fed610656a42f662bed93 100644 --- a/applications/reconstruct/src/main.cpp +++ b/applications/reconstruct/src/main.cpp @@ -255,7 +255,7 @@ static void run(ftl::Configurable *root) { pipeline1->append<ftl::operators::HFSmoother>("hfnoise"); // Remove high-frequency noise pipeline1->append<ftl::operators::Normals>("normals"); // Estimate surface normals pipeline1->append<ftl::operators::SmoothChannel>("smoothing"); // Generate a smoothing channel - pipeline1->append<ftl::operators::SimpleMLS>("mls"); // Perform MLS (using smoothing channel) + pipeline1->append<ftl::operators::AdaptiveMLS>("mls"); // Perform MLS (using smoothing channel) // Alignment diff --git a/components/operators/src/mls.cu b/components/operators/src/mls.cu index 7e408b829fa39b15ad687ab8246886be3da3ee38..e030d1967695bd71760c0bc2af26706d21f0f2e3 100644 --- a/components/operators/src/mls.cu +++ b/components/operators/src/mls.cu @@ -206,3 +206,109 @@ void ftl::cuda::colour_mls_smooth( #endif } + +// ===== Adaptive MLS Smooth ===================================================== + +/* + * Smooth depth map using Moving Least Squares. This version uses colour + * similarity weights to adjust the spatial smoothing factor. It is naive in + * that similar colours may exist on both sides of an edge and colours at the + * level of single pixels can be subject to noise. Colour noise should first + * be removed from the image. + */ + template <int SEARCH_RADIUS> + __global__ void adaptive_mls_smooth_kernel( + TextureObject<float4> normals_in, + TextureObject<float4> normals_out, + TextureObject<float> depth_in, // Virtual depth map + TextureObject<float> depth_out, // Accumulated output + TextureObject<float> smoothing, + ftl::rgbd::Camera camera) { + + const int x = blockIdx.x*blockDim.x + threadIdx.x; + const int y = blockIdx.y*blockDim.y + threadIdx.y; + + if (x < 0 || y < 0 || x >= depth_in.width() || y >= depth_in.height()) return; + + float3 aX = make_float3(0.0f,0.0f,0.0f); + float3 nX = make_float3(0.0f,0.0f,0.0f); + float contrib = 0.0f; + + float d0 = depth_in.tex2D(x, y); + depth_out(x,y) = d0; + normals_out(x,y) = normals_in(x,y); + if (d0 < camera.minDepth || d0 > camera.maxDepth) return; + float3 X = camera.screenToCam((int)(x),(int)(y),d0); + + //uchar4 c0 = colour_in.tex2D(x, y); + float smooth = smoothing(x,y); + + // Neighbourhood + for (int v=-SEARCH_RADIUS; v<=SEARCH_RADIUS; ++v) { + for (int u=-SEARCH_RADIUS; u<=SEARCH_RADIUS; ++u) { + const float d = depth_in.tex2D(x+u, y+v); + if (d < camera.minDepth || d > camera.maxDepth) continue; + + // Point and normal of neighbour + const float3 Xi = camera.screenToCam((int)(x)+u,(int)(y)+v,d); + const float3 Ni = make_float3(normals_in.tex2D((int)(x)+u, (int)(y)+v)); + + if (Ni.x+Ni.y+Ni.z == 0.0f) continue; + + // Gauss approx weighting function using point distance + const float w = ftl::cuda::spatialWeighting(X,Xi,smooth*0.5f); + + aX += Xi*w; + nX += Ni*w; + contrib += w; + } + } + + if (contrib <= 0.0f) return; + + nX /= contrib; // Weighted average normal + aX /= contrib; // Weighted average point (centroid) + + // Signed-Distance Field function + float fX = nX.x * (X.x - aX.x) + nX.y * (X.y - aX.y) + nX.z * (X.z - aX.z); + + // Calculate new point using SDF function to adjust depth (and position) + X = X - nX * fX; + + //uint2 screen = camera.camToScreen<uint2>(X); + + //if (screen.x < depth_out.width() && screen.y < depth_out.height()) { + // depth_out(screen.x,screen.y) = X.z; + //} + depth_out(x,y) = X.z; + normals_out(x,y) = make_float4(nX / length(nX), 0.0f); +} + +void ftl::cuda::adaptive_mls_smooth( + ftl::cuda::TextureObject<float4> &normals_in, + ftl::cuda::TextureObject<float4> &normals_out, + ftl::cuda::TextureObject<float> &depth_in, + ftl::cuda::TextureObject<float> &depth_out, + ftl::cuda::TextureObject<float> &smoothing, + int radius, + const ftl::rgbd::Camera &camera, + 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: adaptive_mls_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break; + case 4: adaptive_mls_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break; + case 3: adaptive_mls_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break; + case 2: adaptive_mls_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break; + case 1: adaptive_mls_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break; + } + cudaSafeCall( cudaGetLastError() ); + + + #ifdef _DEBUG + cudaSafeCall(cudaDeviceSynchronize()); + #endif +} + diff --git a/components/operators/src/smoothing.cpp b/components/operators/src/smoothing.cpp index 77259b8de216fe022a657e87854446af950acf82..8e4f458bf9f996a6d38bef816ab7bcec75b69b11 100644 --- a/components/operators/src/smoothing.cpp +++ b/components/operators/src/smoothing.cpp @@ -6,6 +6,7 @@ using ftl::operators::HFSmoother; using ftl::operators::SimpleMLS; using ftl::operators::ColourMLS; +using ftl::operators::AdaptiveMLS; using ftl::operators::SmoothChannel; using ftl::codecs::Channel; using ftl::rgbd::Format; @@ -210,3 +211,43 @@ bool ColourMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So return true; } + +// ====== Adaptive MLS ========================================================= + +AdaptiveMLS::AdaptiveMLS(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) { + +} + +AdaptiveMLS::~AdaptiveMLS() { + +} + +bool AdaptiveMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) { + int iters = config()->value("mls_iterations", 1); + int radius = config()->value("mls_radius",2); + + if (!in.hasChannel(Channel::Normals)) { + LOG(ERROR) << "Required normals channel missing for MLS"; + return false; + } + + // FIXME: Assume in and out are the same frame. + for (int i=0; i<iters; ++i) { + ftl::cuda::adaptive_mls_smooth( + in.createTexture<float4>(Channel::Normals), + in.createTexture<float4>(Channel::Points, ftl::rgbd::Format<float4>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())), + in.createTexture<float>(Channel::Depth), + in.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())), + in.createTexture<float>(Channel::Smoothing), + radius, + s->parameters(), + 0 + ); + + in.swapChannels(Channel::Depth, Channel::Depth2); + in.swapChannels(Channel::Normals, Channel::Points); + } + + return true; +} + diff --git a/components/operators/src/smoothing_cuda.hpp b/components/operators/src/smoothing_cuda.hpp index 67de60801fb7294997e4ed55ff565313df99c313..53e8cf898c4e99875dc4a9881df58700ee31c127 100644 --- a/components/operators/src/smoothing_cuda.hpp +++ b/components/operators/src/smoothing_cuda.hpp @@ -29,6 +29,16 @@ void colour_mls_smooth( const ftl::rgbd::Camera &camera, cudaStream_t stream); +void adaptive_mls_smooth( + ftl::cuda::TextureObject<float4> &normals_in, + ftl::cuda::TextureObject<float4> &normals_out, + ftl::cuda::TextureObject<float> &depth_in, + ftl::cuda::TextureObject<float> &depth_out, + ftl::cuda::TextureObject<float> &smoothing, + int radius, + const ftl::rgbd::Camera &camera, + cudaStream_t stream); + void smooth_channel( ftl::cuda::TextureObject<uchar4> &colour_in, //ftl::cuda::TextureObject<float> &depth_in,