From 8c84c12a7495b2d77eeecb6a7b7e1ee0c9a10bd4 Mon Sep 17 00:00:00 2001 From: Nicolas Pope <nicolas.pope@utu.fi> Date: Sat, 29 Feb 2020 20:11:11 +0200 Subject: [PATCH] Resolves #337 MLS block errors --- components/operators/src/fusion/mvmls.cpp | 1 + components/operators/src/mls.cu | 8 +++++--- components/operators/src/smoothing.cpp | 1 + components/operators/src/smoothing_cuda.hpp | 1 + 4 files changed, 8 insertions(+), 3 deletions(-) diff --git a/components/operators/src/fusion/mvmls.cpp b/components/operators/src/fusion/mvmls.cpp index 55b992957..41e0a2eef 100644 --- a/components/operators/src/fusion/mvmls.cpp +++ b/components/operators/src/fusion/mvmls.cpp @@ -526,6 +526,7 @@ bool MultiViewMLS::apply(ftl::rgbd::FrameSet &in, ftl::rgbd::FrameSet &out, cuda f.getTexture<half4>(Channel::Normals), *centroid_vert_[i], f.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(size)), + f.getTexture<float>(Channel::Depth), f.getLeftCamera(), stream ); diff --git a/components/operators/src/mls.cu b/components/operators/src/mls.cu index afe15d559..ee3ce4125 100644 --- a/components/operators/src/mls.cu +++ b/components/operators/src/mls.cu @@ -559,6 +559,7 @@ __global__ void mls_adjust_depth_kernel( TextureObject<half4> normals_in, TextureObject<float4> centroid_in, // Virtual depth map TextureObject<float> depth_out, + TextureObject<float> depth_in, ftl::rgbd::Camera camera) { const int x = blockIdx.x*blockDim.x + threadIdx.x; @@ -568,8 +569,8 @@ __global__ void mls_adjust_depth_kernel( const float3 aX = make_float3(centroid_in(x,y)); const float3 nX = make_float3(normals_in(x,y)); - //float d0 = depth_in.tex2D(x, y); - depth_out(x,y) = aX.z; + float d0 = depth_in.tex2D(x, y); + depth_out(x,y) = d0; if (aX.z > camera.minDepth && aX.z < camera.maxDepth) { float3 X = camera.screenToCam((int)(x),(int)(y),aX.z); @@ -663,6 +664,7 @@ void ftl::cuda::mls_adjust_depth( ftl::cuda::TextureObject<half4> &normals_in, ftl::cuda::TextureObject<float4> ¢roid_in, ftl::cuda::TextureObject<float> &depth_out, + ftl::cuda::TextureObject<float> &depth_in, const ftl::rgbd::Camera &camera, cudaStream_t stream) { @@ -672,7 +674,7 @@ void ftl::cuda::mls_adjust_depth( const dim3 gridSize((depth_out.width() + THREADS_X - 1)/THREADS_X, (depth_out.height() + THREADS_Y - 1)/THREADS_Y); const dim3 blockSize(THREADS_X, THREADS_Y); - mls_adjust_depth_kernel<<<gridSize, blockSize, 0, stream>>>(normals_in, centroid_in, depth_out, camera); + mls_adjust_depth_kernel<<<gridSize, blockSize, 0, stream>>>(normals_in, centroid_in, depth_out, depth_in, camera); cudaSafeCall( cudaGetLastError() ); diff --git a/components/operators/src/smoothing.cpp b/components/operators/src/smoothing.cpp index 49ca3f4b9..739162622 100644 --- a/components/operators/src/smoothing.cpp +++ b/components/operators/src/smoothing.cpp @@ -307,6 +307,7 @@ bool AggreMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t s in.createTexture<half4>(Channel::Normals), centroid_vert_, temp_.createTexture<float>(Channel::Depth, ftl::rgbd::Format<float>(size)), + in.getTexture<float>(Channel::Depth), in.getLeftCamera(), stream ); diff --git a/components/operators/src/smoothing_cuda.hpp b/components/operators/src/smoothing_cuda.hpp index 06c6d713f..5454d91e2 100644 --- a/components/operators/src/smoothing_cuda.hpp +++ b/components/operators/src/smoothing_cuda.hpp @@ -46,6 +46,7 @@ void mls_adjust_depth( ftl::cuda::TextureObject<half4> &normals_in, ftl::cuda::TextureObject<float4> ¢roid_in, ftl::cuda::TextureObject<float> &depth_out, + ftl::cuda::TextureObject<float> &depth_in, const ftl::rgbd::Camera &camera, cudaStream_t stream); -- GitLab