From 9d22d5a7f00b84180b58c0885a1d5ab71f92eada Mon Sep 17 00:00:00 2001 From: Nicolas Pope <nwpope@utu.fi> Date: Sat, 5 Oct 2019 17:20:04 +0300 Subject: [PATCH] Pass params to move points --- applications/reconstruct/src/ilw/ilw.cpp | 3 ++- applications/reconstruct/src/ilw/ilw.cu | 18 ++++++++++-------- applications/reconstruct/src/ilw/ilw_cuda.hpp | 1 + 3 files changed, 13 insertions(+), 9 deletions(-) diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp index 3feb23dff..b9931d92b 100644 --- a/applications/reconstruct/src/ilw/ilw.cpp +++ b/applications/reconstruct/src/ilw/ilw.cpp @@ -96,7 +96,7 @@ bool ILW::process(ftl::rgbd::FrameSet &fs, cudaStream_t stream) { _phase0(fs, stream); - params_.range = 0.05f; + params_.range = 0.1f; for (int i=0; i<iterations_; ++i) { int win; @@ -243,6 +243,7 @@ bool ILW::_phase2(ftl::rgbd::FrameSet &fs, float rate, cudaStream_t stream) { f.getTexture<float4>(Channel::EnergyVector), fs.sources[i]->parameters(), pose, + params_, rate, motion_window_, stream diff --git a/applications/reconstruct/src/ilw/ilw.cu b/applications/reconstruct/src/ilw/ilw.cu index 2c8d3199b..c0b14dd8a 100644 --- a/applications/reconstruct/src/ilw/ilw.cu +++ b/applications/reconstruct/src/ilw/ilw.cu @@ -194,7 +194,8 @@ __global__ void move_points_kernel( ftl::cuda::TextureObject<float> d, ftl::cuda::TextureObject<float4> ev, ftl::rgbd::Camera camera, - float4x4 pose, + float4x4 pose, + ftl::cuda::ILWParams params, float rate) { const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; @@ -218,7 +219,7 @@ __global__ void move_points_kernel( //if (pn.x == MINF) continue; if (vecn.x == 0.0f) continue; - const float s = ftl::cuda::weighting(fabs(vec0.x - vecn.x), 0.04f); + const float s = ftl::cuda::weighting(fabs(vec0.x - vecn.x), params.range); contrib += vecn.w * s; delta += vecn.w * s * vecn.z; } @@ -238,7 +239,8 @@ void ftl::cuda::move_points( ftl::cuda::TextureObject<float> &d, ftl::cuda::TextureObject<float4> &v, const ftl::rgbd::Camera &camera, - const float4x4 &pose, + const float4x4 &pose, + const ftl::cuda::ILWParams ¶ms, float rate, int radius, cudaStream_t stream) { @@ -247,11 +249,11 @@ void ftl::cuda::move_points( const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); switch (radius) { - case 9 : move_points_kernel<9><<<gridSize, blockSize, 0, stream>>>(d,v,camera, pose,rate); break; - case 5 : move_points_kernel<5><<<gridSize, blockSize, 0, stream>>>(d,v,camera, pose,rate); break; - case 3 : move_points_kernel<3><<<gridSize, blockSize, 0, stream>>>(d,v,camera, pose,rate); break; - case 1 : move_points_kernel<1><<<gridSize, blockSize, 0, stream>>>(d,v,camera, pose,rate); break; - case 0 : move_points_kernel<0><<<gridSize, blockSize, 0, stream>>>(d,v,camera, pose,rate); break; + case 9 : move_points_kernel<9><<<gridSize, blockSize, 0, stream>>>(d,v,camera, pose, params, rate); break; + case 5 : move_points_kernel<5><<<gridSize, blockSize, 0, stream>>>(d,v,camera, pose, params, rate); break; + case 3 : move_points_kernel<3><<<gridSize, blockSize, 0, stream>>>(d,v,camera, pose, params, rate); break; + case 1 : move_points_kernel<1><<<gridSize, blockSize, 0, stream>>>(d,v,camera, pose, params, rate); break; + case 0 : move_points_kernel<0><<<gridSize, blockSize, 0, stream>>>(d,v,camera, pose, params, rate); break; } cudaSafeCall( cudaGetLastError() ); diff --git a/applications/reconstruct/src/ilw/ilw_cuda.hpp b/applications/reconstruct/src/ilw/ilw_cuda.hpp index 503d1bf4c..e4a099031 100644 --- a/applications/reconstruct/src/ilw/ilw_cuda.hpp +++ b/applications/reconstruct/src/ilw/ilw_cuda.hpp @@ -43,6 +43,7 @@ void move_points( ftl::cuda::TextureObject<float4> &v, const ftl::rgbd::Camera &camera, const float4x4 &pose, + const ILWParams ¶ms, float rate, int radius, cudaStream_t stream -- GitLab