diff --git a/applications/reconstruct/include/ftl/ray_cast_sdf.hpp b/applications/reconstruct/include/ftl/ray_cast_sdf.hpp index d796b1cfc209d3f69c378cacb7fd4a98b638c0ae..7fb9b940fed69dc15c5c42247f71a420d9f2699b 100644 --- a/applications/reconstruct/include/ftl/ray_cast_sdf.hpp +++ b/applications/reconstruct/include/ftl/ray_cast_sdf.hpp @@ -42,7 +42,7 @@ public: return params; } - void render(const ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams, const DepthCameraParams& cameraParams, const Eigen::Matrix4f& lastRigidTransform); + void render(ftl::voxhash::HashData& hashData, ftl::voxhash::HashParams& hashParams, const DepthCameraParams& cameraParams, const Eigen::Matrix4f& lastRigidTransform); const RayCastData& getRayCastData(void) { return m_data; @@ -60,6 +60,8 @@ private: void create(const RayCastParams& params); void destroy(void); + void compactifyHashEntries(ftl::voxhash::HashData& hashData, ftl::voxhash::HashParams& hashParams); + void rayIntervalSplatting(const ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams, const Eigen::Matrix4f& lastRigidTransform); // rasterize RayCastParams m_params; diff --git a/applications/reconstruct/include/ftl/scene_rep_hash_sdf.hpp b/applications/reconstruct/include/ftl/scene_rep_hash_sdf.hpp index 09fd3d1a7a01102ff01a855fc2e58105978cfb8c..55303cef06bfc58a39e8ca780e3dfad03536790e 100644 --- a/applications/reconstruct/include/ftl/scene_rep_hash_sdf.hpp +++ b/applications/reconstruct/include/ftl/scene_rep_hash_sdf.hpp @@ -98,7 +98,7 @@ class SceneRep : public ftl::Configurable { alloc(depthCameraData, depthCameraParams, d_bitMask); //generate a linear hash array with only occupied entries - compactifyHashEntries(depthCameraData); + compactifyHashEntries(); //volumetrically integrate the depth data into the depth SDFBlocks integrateDepthMap(depthCameraData, depthCameraParams); @@ -115,7 +115,7 @@ class SceneRep : public ftl::Configurable { void setLastRigidTransformAndCompactify(const Eigen::Matrix4f& lastRigidTransform, const DepthCameraData& depthCameraData) { setLastRigidTransform(lastRigidTransform); - compactifyHashEntries(depthCameraData); + compactifyHashEntries(); } @@ -145,7 +145,7 @@ class SceneRep : public ftl::Configurable { return m_hashData; } - const HashParams& getHashParams() const { + HashParams& getHashParams() { return m_hashParams; } @@ -312,7 +312,7 @@ private: } - void compactifyHashEntries(const DepthCameraData& depthCameraData) { + void compactifyHashEntries() { //const DepthCameraData& depthCameraData) { //Start Timing //if(GlobalAppState::get().s_timingsDetailledEnabled) { cutilSafeCall(cudaDeviceSynchronize()); m_timer.start(); } diff --git a/applications/reconstruct/src/ray_cast_sdf.cpp b/applications/reconstruct/src/ray_cast_sdf.cpp index c3723429b914629da77b29c53970dd44f39b01ab..0f792b4c3608b1c113631bdb9b791595bd50e007 100644 --- a/applications/reconstruct/src/ray_cast_sdf.cpp +++ b/applications/reconstruct/src/ray_cast_sdf.cpp @@ -36,7 +36,17 @@ void CUDARayCastSDF::destroy(void) //m_rayIntervalSplatting.OnD3D11DestroyDevice(); } -void CUDARayCastSDF::render(const ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams, const DepthCameraParams& cameraParams, const Eigen::Matrix4f& lastRigidTransform) +extern "C" unsigned int compactifyHashAllInOneCUDA(ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams); + + +void CUDARayCastSDF::compactifyHashEntries(ftl::voxhash::HashData& hashData, ftl::voxhash::HashParams& hashParams) { //const DepthCameraData& depthCameraData) { + + hashParams.m_numOccupiedBlocks = compactifyHashAllInOneCUDA(hashData, hashParams); //this version uses atomics over prefix sums, which has a much better performance + std::cout << "Ray blocks = " << hashParams.m_numOccupiedBlocks << std::endl; + hashData.updateParams(hashParams); //make sure numOccupiedBlocks is updated on the GPU +} + +void CUDARayCastSDF::render(ftl::voxhash::HashData& hashData, ftl::voxhash::HashParams& hashParams, const DepthCameraParams& cameraParams, const Eigen::Matrix4f& lastRigidTransform) { updateConstantDepthCameraParams(cameraParams); //rayIntervalSplatting(hashData, hashParams, lastRigidTransform); @@ -48,6 +58,8 @@ void CUDARayCastSDF::render(const ftl::voxhash::HashData& hashData, const ftl::v m_params.m_viewMatrixInverse = MatrixConversion::toCUDA(lastRigidTransform); m_data.updateParams(m_params); + compactifyHashEntries(hashData, hashParams); + if (hash_render_) nickRenderCUDA(hashData, hashParams, m_data, m_params); else renderCS(hashData, m_data, m_params); diff --git a/applications/reconstruct/src/ray_cast_sdf.cu b/applications/reconstruct/src/ray_cast_sdf.cu index b28bea8acfd2ba415c2e08182a09a1ec58e9bb28..b0c6f31124760b71c97a0c39c4c38b29d3c9ac40 100644 --- a/applications/reconstruct/src/ray_cast_sdf.cu +++ b/applications/reconstruct/src/ray_cast_sdf.cu @@ -298,6 +298,8 @@ __global__ void nickRenderKernel(ftl::voxhash::HashData hashData, RayCastData ra const float2 screenPosf = DepthCameraData::cameraToKinectScreenFloat(camPos); const uint2 screenPos = make_uint2(make_int2(screenPosf)); // + make_float2(0.5f, 0.5f) + if (camPos.z < 0.0f) return; + /*if (screenPos.x < params.m_width && screenPos.y < params.m_height && rayCastData.d_depth[(screenPos.y)*params.m_width+screenPos.x] > camPos.z) { rayCastData.d_depth[(screenPos.y)*params.m_width+screenPos.x] = camPos.z;