From 7aa42052c1644dea7f486a8a151cc91d3be5c2a1 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Sun, 21 Jul 2019 09:46:08 +0300
Subject: [PATCH] Tidy params, add normal alloc
---
.../reconstruct/include/ftl/depth_camera.hpp | 21 ++++++-
.../include/ftl/voxel_hash_params.hpp | 19 +------
applications/reconstruct/src/integrators.cu | 57 ++++++-------------
.../reconstruct/src/scene_rep_hash_sdf.cu | 35 ------------
applications/reconstruct/src/voxel_scene.cpp | 36 ++++--------
5 files changed, 52 insertions(+), 116 deletions(-)
diff --git a/applications/reconstruct/include/ftl/depth_camera.hpp b/applications/reconstruct/include/ftl/depth_camera.hpp
index f40581b65..7553f8aa8 100644
--- a/applications/reconstruct/include/ftl/depth_camera.hpp
+++ b/applications/reconstruct/include/ftl/depth_camera.hpp
@@ -24,6 +24,7 @@ namespace voxhash {
struct __align__(16) DepthCameraCUDA {
cudaTextureObject_t depth;
cudaTextureObject_t colour;
+ cudaTextureObject_t normal;
DepthCameraParams params;
float4x4 pose;
float4x4 poseInverse;
@@ -44,12 +45,15 @@ struct DepthCamera {
depth_mat_ = nullptr;
colour_mat_ = nullptr;
+ point_mat_ = nullptr;
+ normal_mat_ = nullptr;
depth_tex_ = nullptr;
colour_tex_ = nullptr;
+ normal_tex_ = nullptr;
}
__host__
- void alloc(const DepthCameraParams& params) { //! todo resizing???
+ void alloc(const DepthCameraParams& params, bool withNormals=false) { //! todo resizing???
depth_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_32FC1);
colour_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_8UC4);
depth_tex_ = new ftl::cuda::TextureObject<float>((cv::cuda::PtrStepSz<float>)*depth_mat_);
@@ -57,6 +61,15 @@ struct DepthCamera {
data.depth = depth_tex_->cudaTexture();
data.colour = colour_tex_->cudaTexture();
data.params = params;
+
+ if (withNormals) {
+ point_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_32FC3);
+ normal_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_32FC4);
+ normal_tex_ = new ftl::cuda::TextureObject<float4>((cv::cuda::PtrStepSz<float4>)*normal_mat_);
+ data.normal = normal_tex_->cudaTexture();
+ } else {
+ data.normal = 0;
+ }
}
//__host__
@@ -74,16 +87,22 @@ struct DepthCamera {
void free() {
if (depth_mat_) delete depth_mat_;
if (colour_mat_) delete colour_mat_;
+ if (point_mat_) delete point_mat_;
+ if (normal_mat_) delete normal_mat_;
delete depth_tex_;
delete colour_tex_;
+ if (normal_tex_) delete normal_tex_;
}
// TODO(Nick) Should not need to pass all these pointers to device
cv::cuda::GpuMat *depth_mat_;
cv::cuda::GpuMat *colour_mat_;
+ cv::cuda::GpuMat *point_mat_;
+ cv::cuda::GpuMat *normal_mat_;
ftl::cuda::TextureObject<float> *depth_tex_;
ftl::cuda::TextureObject<uchar4> *colour_tex_;
+ ftl::cuda::TextureObject<float4> *normal_tex_;
//cudaTextureObject_t depth_obj_;
//cudaTextureObject_t colour_obj_;
diff --git a/applications/reconstruct/include/ftl/voxel_hash_params.hpp b/applications/reconstruct/include/ftl/voxel_hash_params.hpp
index cb94404c3..aa19368ea 100644
--- a/applications/reconstruct/include/ftl/voxel_hash_params.hpp
+++ b/applications/reconstruct/include/ftl/voxel_hash_params.hpp
@@ -17,30 +17,17 @@ struct __align__(16) HashParams {
HashParams() {
}
- float4x4 m_rigidTransform;
- float4x4 m_rigidTransformInverse;
-
unsigned int m_hashNumBuckets;
- unsigned int m_deprecated1;
- unsigned int m_deprecated2; //m_hashMaxCollisionLinkedListSize;
- unsigned int m_numSDFBlocks;
-
- int m_SDFBlockSize;
float m_virtualVoxelSize;
- unsigned int m_numOccupiedBlocks; //occupied blocks in the viewing frustum
-
float m_maxIntegrationDistance;
float m_truncScale;
float m_truncation;
unsigned int m_integrationWeightSample;
unsigned int m_integrationWeightMax;
- float3 m_streamingVoxelExtents;
- int3 m_streamingGridDimensions;
- int3 m_streamingMinGridPos;
- unsigned int m_streamingInitialChunkListSize;
- uint2 m_dummy;
-
+ float3 m_minBounds;
+ float3 m_maxBounds;
+ float m_spatialSmoothing;
};
} // namespace voxhash
diff --git a/applications/reconstruct/src/integrators.cu b/applications/reconstruct/src/integrators.cu
index 1259210b4..ce51c185d 100644
--- a/applications/reconstruct/src/integrators.cu
+++ b/applications/reconstruct/src/integrators.cu
@@ -7,6 +7,8 @@
#include <ftl/cuda_common.hpp>
#define T_PER_BLOCK 8
+#define NUM_CUDA_BLOCKS 10000
+#define WARP_SIZE 32
using ftl::voxhash::HashData;
using ftl::voxhash::HashParams;
@@ -15,45 +17,13 @@ using ftl::voxhash::HashEntry;
using ftl::voxhash::HashEntryHead;
using ftl::voxhash::FREE_ENTRY;
+extern __constant__ ftl::voxhash::DepthCameraCUDA c_cameras[MAX_CAMERAS];
+extern __constant__ HashParams c_hashParams;
+
__device__ float4 make_float4(uchar4 c) {
return make_float4(static_cast<float>(c.x), static_cast<float>(c.y), static_cast<float>(c.z), static_cast<float>(c.w));
}
-inline __device__ uchar4 bilinearFilterColor(const DepthCameraParams& cameraParams, const float2& screenPos, cudaTextureObject_t colorTextureRef) {
- //const DepthCameraParams& cameraParams = c_depthCameraParams;
- const int imageWidth = cameraParams.m_imageWidth;
- const int imageHeight = cameraParams.m_imageHeight;
- const int2 p00 = make_int2(screenPos.x+0.5f, screenPos.y+0.5f);
- const int2 dir = sign(make_float2(screenPos.x - p00.x, screenPos.y - p00.y));
-
- const int2 p01 = p00 + make_int2(0.0f, dir.y);
- const int2 p10 = p00 + make_int2(dir.x, 0.0f);
- const int2 p11 = p00 + make_int2(dir.x, dir.y);
-
- const float alpha = (screenPos.x - p00.x)*dir.x;
- const float beta = (screenPos.y - p00.y)*dir.y;
-
- float4 s0 = make_float4(0.0f, 0.0f, 0.0f, 0.0f); float w0 = 0.0f;
- if(p00.x >= 0 && p00.x < imageWidth && p00.y >= 0 && p00.y < imageHeight) { uchar4 v00 = tex2D<uchar4>(colorTextureRef, p00.x, p00.y); if(v00.x != 0) { s0 += (1.0f-alpha)*make_float4(v00); w0 += (1.0f-alpha); } }
- if(p10.x >= 0 && p10.x < imageWidth && p10.y >= 0 && p10.y < imageHeight) { uchar4 v10 = tex2D<uchar4>(colorTextureRef, p10.x, p10.y); if(v10.x != 0) { s0 += alpha *make_float4(v10); w0 += alpha ; } }
-
- float4 s1 = make_float4(0.0f, 0.0f, 0.0f, 0.0f); float w1 = 0.0f;
- if(p01.x >= 0 && p01.x < imageWidth && p01.y >= 0 && p01.y < imageHeight) { uchar4 v01 = tex2D<uchar4>(colorTextureRef, p01.x, p01.y); if(v01.x != 0) { s1 += (1.0f-alpha)*make_float4(v01); w1 += (1.0f-alpha);} }
- if(p11.x >= 0 && p11.x < imageWidth && p11.y >= 0 && p11.y < imageHeight) { uchar4 v11 = tex2D<uchar4>(colorTextureRef, p11.x, p11.y); if(v11.x != 0) { s1 += alpha *make_float4(v11); w1 += alpha ;} }
-
- const float4 p0 = s0/w0;
- const float4 p1 = s1/w1;
-
- float4 ss = make_float4(0.0f, 0.0f, 0.0f, 0.0f); float ww = 0.0f;
- if(w0 > 0.0f) { ss += (1.0f-beta)*p0; ww += (1.0f-beta); }
- if(w1 > 0.0f) { ss += beta *p1; ww += beta ; }
-
- if(ww > 0.0f) {
- ss /= ww;
- return make_uchar4(ss.x,ss.y,ss.z,ss.w);
- } else return make_uchar4(0, 0, 0, 0);
-}
-
__device__ float colourDistance(const uchar4 &c1, const uchar3 &c2) {
float x = c1.x-c2.x;
float y = c1.y-c2.y;
@@ -69,7 +39,17 @@ __device__ bool colordiff(const uchar4 &pa, const uchar3 &pb, float epsilon) {
return sqrt(x_2 - p_2) < epsilon;
}
-#define NUM_CUDA_BLOCKS 10000
+/*
+ * Guennebaud, G.; Gross, M. Algebraic point set surfaces. ACMTransactions on Graphics Vol. 26, No. 3, Article No. 23, 2007.
+ * Used in: FusionMLS: Highly dynamic 3D reconstruction with consumer-grade RGB-D cameras
+ */
+__device__ float spatialWeighting(float r) {
+ const float h = c_hashParams.m_spatialSmoothing;
+ float rh = r / h;
+ rh = 1.0f - rh*rh;
+ rh = rh*rh*rh*rh;
+ return (rh < h) ? rh : 0.0f;
+}
@@ -141,10 +121,6 @@ __device__ bool colordiff(const uchar4 &pa, const uchar3 &pb, float epsilon) {
} // Stride loop
}*/
-extern __constant__ ftl::voxhash::DepthCameraCUDA c_cameras[MAX_CAMERAS];
-
-#define WARP_SIZE 32
-
__global__ void integrateDepthMapsKernel(HashData hashData, HashParams hashParams, int numcams) {
__shared__ uint all_warp_ballot;
__shared__ uint voxels[16];
@@ -254,6 +230,7 @@ __global__ void integrateDepthMapsKernel(HashData hashData, HashParams hashParam
__syncthreads();
// Work out if block is occupied or not and save voxel masks
+ // TODO:(Nick) Is it faster to do this in a separate garbage kernel?
if (i < 16) {
const uint v = voxels[i];
hashData.d_hashCompactified[bi]->voxels[i] = v;
diff --git a/applications/reconstruct/src/scene_rep_hash_sdf.cu b/applications/reconstruct/src/scene_rep_hash_sdf.cu
index 5c5ebf6fd..927e22eed 100644
--- a/applications/reconstruct/src/scene_rep_hash_sdf.cu
+++ b/applications/reconstruct/src/scene_rep_hash_sdf.cu
@@ -168,41 +168,6 @@ extern "C" void resetHashBucketMutexCUDA(HashData& hashData, const HashParams& h
}
-__device__
-unsigned int linearizeChunkPos(const int3& chunkPos)
-{
- int3 p = chunkPos-c_hashParams.m_streamingMinGridPos;
- return p.z * c_hashParams.m_streamingGridDimensions.x * c_hashParams.m_streamingGridDimensions.y +
- p.y * c_hashParams.m_streamingGridDimensions.x +
- p.x;
-}
-
-__device__
-int3 worldToChunks(const float3& posWorld)
-{
- float3 p;
- p.x = posWorld.x/c_hashParams.m_streamingVoxelExtents.x;
- p.y = posWorld.y/c_hashParams.m_streamingVoxelExtents.y;
- p.z = posWorld.z/c_hashParams.m_streamingVoxelExtents.z;
-
- float3 s;
- s.x = (float)sign(p.x);
- s.y = (float)sign(p.y);
- s.z = (float)sign(p.z);
-
- return make_int3(p+s*0.5f);
-}
-
-__device__
-bool isSDFBlockStreamedOut(const int3& sdfBlock, const HashData& hashData, const unsigned int* d_bitMask) //TODO MATTHIAS (-> move to HashData)
-{
- float3 posWorld = hashData.virtualVoxelPosToWorld(hashData.SDFBlockToVirtualVoxelPos(sdfBlock)); // sdfBlock is assigned to chunk by the bottom right sample pos
-
- uint index = linearizeChunkPos(worldToChunks(posWorld));
- uint nBitsInT = 32;
- return ((d_bitMask[index/nBitsInT] & (0x1 << (index%nBitsInT))) != 0x0);
-}
-
// Note: bitMask used for Streaming out code... could be set to nullptr if not streaming out
// Note: Allocations might need to be around fat rays since multiple voxels could correspond
// to same depth map pixel at larger distances.
diff --git a/applications/reconstruct/src/voxel_scene.cpp b/applications/reconstruct/src/voxel_scene.cpp
index 51f9bc75b..861e480e2 100644
--- a/applications/reconstruct/src/voxel_scene.cpp
+++ b/applications/reconstruct/src/voxel_scene.cpp
@@ -138,7 +138,7 @@ int SceneRep::upload() {
cam.params.m_imageHeight = in->parameters().height;
cam.params.m_sensorDepthWorldMax = in->parameters().maxDepth;
cam.params.m_sensorDepthWorldMin = in->parameters().minDepth;
- cam.gpu.alloc(cam.params);
+ cam.gpu.alloc(cam.params, true);
}
}
@@ -179,6 +179,8 @@ int SceneRep::upload() {
//allocate all hash blocks which are corresponding to depth map entries
_alloc(i, cv::cuda::StreamAccessor::getStream(cam.stream));
+
+ // Calculate normals
}
// Must have finished all allocations and rendering before next integration
@@ -237,21 +239,11 @@ void SceneRep::garbage() {
m_numIntegratedFrames++;
}*/
-void SceneRep::setLastRigidTransform(const Eigen::Matrix4f& lastRigidTransform) {
- m_hashParams.m_rigidTransform = MatrixConversion::toCUDA(lastRigidTransform);
- m_hashParams.m_rigidTransformInverse = MatrixConversion::toCUDA(lastRigidTransform.inverse()); //m_hashParams.m_rigidTransform.getInverse();
-}
-
/*void SceneRep::setLastRigidTransformAndCompactify(const Eigen::Matrix4f& lastRigidTransform, const DepthCameraData& depthCameraData) {
setLastRigidTransform(lastRigidTransform);
_compactifyHashEntries();
}*/
-
-const Eigen::Matrix4f SceneRep::getLastRigidTransform() const {
- return MatrixConversion::toEigen(m_hashParams.m_rigidTransform);
-}
-
/* Nick: To reduce weights between frames */
void SceneRep::nextFrame() {
if (do_reset_) {
@@ -269,10 +261,6 @@ void SceneRep::nextFrame() {
//! resets the hash to the initial state (i.e., clears all data)
void SceneRep::reset() {
m_numIntegratedFrames = 0;
-
- //m_hashParams.m_rigidTransform.setIdentity();
- //m_hashParams.m_rigidTransformInverse.setIdentity();
- m_hashParams.m_numOccupiedBlocks = 0;
m_hashData.updateParams(m_hashParams);
resetCUDA(m_hashData, m_hashParams);
}
@@ -288,22 +276,22 @@ HashParams SceneRep::_parametersFromConfig() {
HashParams params;
// First camera view is set to identity pose to be at the centre of
// the virtual coordinate space.
- params.m_rigidTransform.setIdentity();
- params.m_rigidTransformInverse.setIdentity();
params.m_hashNumBuckets = value("hashNumBuckets", 100000);
- params.m_SDFBlockSize = SDF_BLOCK_SIZE;
- params.m_numSDFBlocks = value("hashNumSDFBlocks",500000);
params.m_virtualVoxelSize = value("SDFVoxelSize", 0.006f);
params.m_maxIntegrationDistance = value("SDFMaxIntegrationDistance", 10.0f);
params.m_truncation = value("SDFTruncation", 0.1f);
params.m_truncScale = value("SDFTruncationScale", 0.01f);
params.m_integrationWeightSample = value("SDFIntegrationWeightSample", 10);
params.m_integrationWeightMax = value("SDFIntegrationWeightMax", 255);
- // Note (Nick): We are not streaming voxels in/out of GPU
- //params.m_streamingVoxelExtents = MatrixConversion::toCUDA(gas.s_streamingVoxelExtents);
- //params.m_streamingGridDimensions = MatrixConversion::toCUDA(gas.s_streamingGridDimensions);
- //params.m_streamingMinGridPos = MatrixConversion::toCUDA(gas.s_streamingMinGridPos);
- //params.m_streamingInitialChunkListSize = gas.s_streamingInitialChunkListSize;
+ params.m_spatialSmoothing = value("spatialSmoothing", 0.04f); // 4cm
+ params.m_maxBounds = make_float3(
+ value("bbox_x_max", 2.0f),
+ value("bbox_y_max", 2.0f),
+ value("bbox_z_max", 2.0f));
+ params.m_minBounds = make_float3(
+ value("bbox_x_min", -2.0f),
+ value("bbox_y_min", -2.0f),
+ value("bbox_z_min", -2.0f));
return params;
}
--
GitLab