From daeaafab9d18cce98f291920539872af921062f2 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Mon, 22 Jul 2019 17:39:02 +0300
Subject: [PATCH] Revert "Switch to storing point cloud from MLS step"
This reverts commit 55795d670990759b89947f18f64266f2154544f7.
---
.../reconstruct/include/ftl/depth_camera.hpp | 2 --
applications/reconstruct/src/depth_camera.cpp | 11 +++-------
applications/reconstruct/src/depth_camera.cu | 22 +++++++++----------
.../reconstruct/src/depth_camera_cuda.hpp | 4 ++--
applications/reconstruct/src/dibr.cu | 2 +-
applications/reconstruct/src/integrators.cu | 4 ++--
.../reconstruct/src/scene_rep_hash_sdf.cu | 2 +-
applications/reconstruct/src/voxel_scene.cpp | 6 ++---
8 files changed, 23 insertions(+), 30 deletions(-)
diff --git a/applications/reconstruct/include/ftl/depth_camera.hpp b/applications/reconstruct/include/ftl/depth_camera.hpp
index 379e73110..a62f6f9d1 100644
--- a/applications/reconstruct/include/ftl/depth_camera.hpp
+++ b/applications/reconstruct/include/ftl/depth_camera.hpp
@@ -24,7 +24,6 @@ namespace voxhash {
struct __align__(16) DepthCameraCUDA {
cudaTextureObject_t depth;
cudaTextureObject_t depth2;
- cudaTextureObject_t point;
cudaTextureObject_t colour;
cudaTextureObject_t normal;
DepthCameraParams params;
@@ -55,7 +54,6 @@ struct DepthCamera {
cv::cuda::GpuMat *normal_mat_;
ftl::cuda::TextureObject<float> *depth_tex_;
ftl::cuda::TextureObject<float> *depth2_tex_;
- ftl::cuda::TextureObject<float4> *point_tex_;
ftl::cuda::TextureObject<uchar4> *colour_tex_;
ftl::cuda::TextureObject<float4> *normal_tex_;
diff --git a/applications/reconstruct/src/depth_camera.cpp b/applications/reconstruct/src/depth_camera.cpp
index f75fec99a..b2dd2671a 100644
--- a/applications/reconstruct/src/depth_camera.cpp
+++ b/applications/reconstruct/src/depth_camera.cpp
@@ -15,26 +15,22 @@ DepthCamera::DepthCamera() {
depth2_tex_ = nullptr;
colour_tex_ = nullptr;
normal_tex_ = nullptr;
- point_tex_ = nullptr;
}
void DepthCamera::alloc(const DepthCameraParams& params, bool withNormals) { //! todo resizing???
depth_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_32FC1);
depth2_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);
- point_mat_ = new cv::cuda::GpuMat(params.m_imageHeight, params.m_imageWidth, CV_32FC4);
depth_tex_ = new ftl::cuda::TextureObject<float>((cv::cuda::PtrStepSz<float>)*depth_mat_);
depth2_tex_ = new ftl::cuda::TextureObject<float>((cv::cuda::PtrStepSz<float>)*depth2_mat_);
- point_tex_ = new ftl::cuda::TextureObject<float4>((cv::cuda::PtrStepSz<float4>)*point_mat_);
colour_tex_ = new ftl::cuda::TextureObject<uchar4>((cv::cuda::PtrStepSz<uchar4>)*colour_mat_);
data.depth = depth_tex_->cudaTexture();
data.depth2 = depth2_tex_->cudaTexture();
data.colour = colour_tex_->cudaTexture();
- data.point = point_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();
@@ -52,7 +48,6 @@ void DepthCamera::free() {
delete depth_tex_;
delete colour_tex_;
if (normal_tex_) delete normal_tex_;
- if (point_tex_) delete normal_tex_;
}
void DepthCamera::updateData(const cv::Mat &depth, const cv::Mat &rgb, cv::cuda::Stream &stream) {
@@ -64,6 +59,6 @@ void DepthCamera::updateData(const cv::Mat &depth, const cv::Mat &rgb, cv::cuda:
}
void DepthCamera::_computeNormals(cudaStream_t stream) {
- //ftl::cuda::point_cloud((float3*)point_mat_->data, data, stream);
- //ftl::cuda::compute_normals((float3*)point_mat_->data, normal_tex_, stream);
+ ftl::cuda::point_cloud((float3*)point_mat_->data, data, stream);
+ ftl::cuda::compute_normals((float3*)point_mat_->data, normal_tex_, stream);
}
diff --git a/applications/reconstruct/src/depth_camera.cu b/applications/reconstruct/src/depth_camera.cu
index db1597fef..47ecf8822 100644
--- a/applications/reconstruct/src/depth_camera.cu
+++ b/applications/reconstruct/src/depth_camera.cu
@@ -12,17 +12,17 @@ using ftl::voxhash::HashParams;
extern __constant__ ftl::voxhash::DepthCameraCUDA c_cameras[MAX_CAMERAS];
-__global__ void clear_depth_kernel(ftl::cuda::TextureObject<float4> depth) {
+__global__ void clear_depth_kernel(ftl::cuda::TextureObject<float> depth) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x < depth.width() && y < depth.height()) {
- depth(x,y) = make_float4(MINF, MINF, MINF, MINF); //PINF;
+ depth(x,y) = 1000.0f; //PINF;
//colour(x,y) = make_uchar4(76,76,82,0);
}
}
-void ftl::cuda::clear_depth(const ftl::cuda::TextureObject<float4> &depth, cudaStream_t stream) {
+void ftl::cuda::clear_depth(const ftl::cuda::TextureObject<float> &depth, cudaStream_t stream) {
const dim3 clear_gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 clear_blockSize(T_PER_BLOCK, T_PER_BLOCK);
clear_depth_kernel<<<clear_gridSize, clear_blockSize, 0, stream>>>(depth);
@@ -66,7 +66,7 @@ __device__ float colordiffFloat2(const uchar4 &pa, const uchar4 &pb) {
#define WINDOW_RADIUS 5
-__global__ void mls_smooth_kernel(ftl::cuda::TextureObject<float4> output, HashData hashData, HashParams hashParams, int numcams, int cam) {
+__global__ void mls_smooth_kernel(ftl::cuda::TextureObject<float> output, HashData hashData, HashParams hashParams, int numcams, int cam) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
@@ -120,18 +120,18 @@ __global__ void mls_smooth_kernel(ftl::cuda::TextureObject<float4> output, HashD
//float sdf = wnorm.x * (mPos.x - wpos.x) + wnorm.y * (mPos.y - wpos.y) + wnorm.z * (mPos.z - wpos.z);
//wpos = wpos - (wnorm * (mPos - wpos));
- //wpos = mainCamera.poseInverse * wpos;
- //const uint2 screenPos = make_uint2(mainCamera.params.cameraToKinectScreenInt(wpos));
- //if (screenPos.x < output.width() && screenPos.y < output.height()) {
- //output(screenPos.x,screenPos.y) = (weights >= hashParams.m_confidenceThresh) ? wpos.z : 1000.0f;
- output(x,y) = (weights >= hashParams.m_confidenceThresh) ? make_float4(wpos, 0.0f) : make_float4(0.0f, 0.0f, 1000.0f, 0.0f);
- //}
+ wpos = mainCamera.poseInverse * wpos;
+ const uint2 screenPos = make_uint2(mainCamera.params.cameraToKinectScreenInt(wpos));
+ if (screenPos.x < output.width() && screenPos.y < output.height()) {
+ output(screenPos.x,screenPos.y) = (weights >= hashParams.m_confidenceThresh) ? wpos.z : 1000.0f;
+ //output(x,y) = (weights >= hashParams.m_confidenceThresh) ? wpos.z : 1000.0f;
+ }
}
}
}
-void ftl::cuda::mls_smooth(TextureObject<float4> &output, const HashData &hashData, const HashParams &hashParams, int numcams, int cam, cudaStream_t stream) {
+void ftl::cuda::mls_smooth(TextureObject<float> &output, const HashData &hashData, const HashParams &hashParams, int numcams, int cam, cudaStream_t stream) {
const dim3 gridSize((output.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (output.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
diff --git a/applications/reconstruct/src/depth_camera_cuda.hpp b/applications/reconstruct/src/depth_camera_cuda.hpp
index 562f4e959..90090f643 100644
--- a/applications/reconstruct/src/depth_camera_cuda.hpp
+++ b/applications/reconstruct/src/depth_camera_cuda.hpp
@@ -7,9 +7,9 @@
namespace ftl {
namespace cuda {
-void clear_depth(const TextureObject<float4> &depth, cudaStream_t stream);
+void clear_depth(const TextureObject<float> &depth, cudaStream_t stream);
-void mls_smooth(TextureObject<float4> &output, const ftl::voxhash::HashData &hashData, const ftl::voxhash::HashParams &hashParams, int numcams, int cam, cudaStream_t stream);
+void mls_smooth(TextureObject<float> &output, const ftl::voxhash::HashData &hashData, const ftl::voxhash::HashParams &hashParams, int numcams, int cam, cudaStream_t stream);
void point_cloud(float3* output, const ftl::voxhash::DepthCameraCUDA &depthCameraData, cudaStream_t stream);
diff --git a/applications/reconstruct/src/dibr.cu b/applications/reconstruct/src/dibr.cu
index 85965d6ee..d647e8f31 100644
--- a/applications/reconstruct/src/dibr.cu
+++ b/applications/reconstruct/src/dibr.cu
@@ -94,7 +94,7 @@ __global__ void dibr_kernel_rev(
const unsigned int cy = screenPos.y;
if (cx < camera.params.m_imageWidth && cy < camera.params.m_imageHeight) {
- float d = tex2D<float>(camera.depth2, (int)cx, (int)cy);
+ float d = tex2D<float>(camera.depth, (int)cx, (int)cy);
float camdiff = fabs(camPos.z-d);
if (camdiff < 0.1f) {
colour_out(x,y) = tex2D<uchar4>(camera.colour,cx,cy);
diff --git a/applications/reconstruct/src/integrators.cu b/applications/reconstruct/src/integrators.cu
index ae60e4a6e..d37591036 100644
--- a/applications/reconstruct/src/integrators.cu
+++ b/applications/reconstruct/src/integrators.cu
@@ -246,10 +246,10 @@ __global__ void integrateMLSKernel(HashData hashData, HashParams hashParams, int
for (int v=-WINDOW_RADIUS; v<=WINDOW_RADIUS; ++v) {
for (int u=-WINDOW_RADIUS; u<=WINDOW_RADIUS; ++u) {
if (screenPos.x+u < width && screenPos.y+v < height) { //on screen
- float3 camPos = make_float3(tex2D<float4>(camera.point, screenPos.x+u, screenPos.y+v));
+ float depth = tex2D<float>(camera.depth, screenPos.x+u, screenPos.y+v);
//float4 normal = tex2D<float4>(camera.normal, screenPos.x+u, screenPos.y+v);
- const float3 worldPos = camPos; //camera.pose * camPos; //camera.params.kinectDepthToSkeleton(screenPos.x+u, screenPos.y+v, depth);
+ const float3 worldPos = camera.pose * camera.params.kinectDepthToSkeleton(screenPos.x+u, screenPos.y+v, depth);
const float weight = spatialWeighting(length(pfb - worldPos));
wpos += weight*worldPos;
diff --git a/applications/reconstruct/src/scene_rep_hash_sdf.cu b/applications/reconstruct/src/scene_rep_hash_sdf.cu
index 001149494..927e22eed 100644
--- a/applications/reconstruct/src/scene_rep_hash_sdf.cu
+++ b/applications/reconstruct/src/scene_rep_hash_sdf.cu
@@ -182,7 +182,7 @@ __global__ void allocKernel(HashData hashData, HashParams hashParams, int camnum
if (x < cameraParams.m_imageWidth && y < cameraParams.m_imageHeight)
{
- float d = tex2D<float>(camera.depth2, x, y);
+ float d = tex2D<float>(camera.depth, x, y);
//if (d == MINF || d < cameraParams.m_sensorDepthWorldMin || d > cameraParams.m_sensorDepthWorldMax) return;
if (d == MINF || d == 0.0f) return;
diff --git a/applications/reconstruct/src/voxel_scene.cpp b/applications/reconstruct/src/voxel_scene.cpp
index bbd5e11be..424783e84 100644
--- a/applications/reconstruct/src/voxel_scene.cpp
+++ b/applications/reconstruct/src/voxel_scene.cpp
@@ -141,7 +141,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);
}
}
@@ -357,8 +357,8 @@ void SceneRep::_compactifyAllocated() {
void SceneRep::_integrateDepthMaps() {
for (size_t i=0; i<cameras_.size(); ++i) {
- ftl::cuda::clear_depth(*(cameras_[i].gpu.point_tex_), integ_stream_);
- ftl::cuda::mls_smooth(*(cameras_[i].gpu.point_tex_), m_hashData, m_hashParams, cameras_.size(), i, integ_stream_);
+ ftl::cuda::clear_depth(*(cameras_[i].gpu.depth_tex_), integ_stream_);
+ ftl::cuda::mls_smooth(*(cameras_[i].gpu.depth_tex_), m_hashData, m_hashParams, cameras_.size(), i, integ_stream_);
}
ftl::cuda::integrateDepthMaps(m_hashData, m_hashParams, cameras_.size(), integ_stream_);
}
--
GitLab