Skip to content

Commit

Permalink
dense: refactor PatchMatchCUDA::Camera into a stand alone structure
Browse files Browse the repository at this point in the history
  • Loading branch information
cdcseacave committed Aug 22, 2024
1 parent a759b87 commit 8e42439
Show file tree
Hide file tree
Showing 16 changed files with 340 additions and 215 deletions.
2 changes: 1 addition & 1 deletion apps/DensifyPointCloud/DensifyPointCloud.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -109,7 +109,7 @@ bool Application::Initialize(size_t argc, LPCTSTR* argv)
), "verbosity level")
#endif
#ifdef _USE_CUDA
("cuda-device", boost::program_options::value(&CUDA::desiredDeviceID)->default_value(-1), "CUDA device number to be used for depth-map estimation (-2 - CPU processing, -1 - best GPU, >=0 - device index)")
("cuda-device", boost::program_options::value(&SEACAVE::CUDA::desiredDeviceID)->default_value(-1), "CUDA device number to be used for depth-map estimation (-2 - CPU processing, -1 - best GPU, >=0 - device index)")
#endif
;

Expand Down
2 changes: 1 addition & 1 deletion apps/ReconstructMesh/ReconstructMesh.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -118,7 +118,7 @@ bool Application::Initialize(size_t argc, LPCTSTR* argv)
), "verbosity level")
#endif
#ifdef _USE_CUDA
("cuda-device", boost::program_options::value(&CUDA::desiredDeviceID)->default_value(-1), "CUDA device number to be used to reconstruct the mesh (-2 - CPU processing, -1 - best GPU, >=0 - device index)")
("cuda-device", boost::program_options::value(&SEACAVE::CUDA::desiredDeviceID)->default_value(-1), "CUDA device number to be used to reconstruct the mesh (-2 - CPU processing, -1 - best GPU, >=0 - device index)")
#endif
;

Expand Down
4 changes: 2 additions & 2 deletions apps/RefineMesh/RefineMesh.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -108,7 +108,7 @@ bool Application::Initialize(size_t argc, LPCTSTR* argv)
), "verbosity level")
#endif
#ifdef _USE_CUDA
("cuda-device", boost::program_options::value(&CUDA::desiredDeviceID)->default_value(-2), "CUDA device number to be used for mesh refinement (-2 - CPU processing, -1 - best GPU, >=0 - device index)")
("cuda-device", boost::program_options::value(&SEACAVE::CUDA::desiredDeviceID)->default_value(-2), "CUDA device number to be used for mesh refinement (-2 - CPU processing, -1 - best GPU, >=0 - device index)")
#endif
;

Expand Down Expand Up @@ -229,7 +229,7 @@ int main(int argc, LPCTSTR* argv)
}
TD_TIMER_START();
#ifdef _USE_CUDA
if (CUDA::desiredDeviceID < -1 ||
if (SEACAVE::CUDA::desiredDeviceID < -1 ||
!scene.RefineMeshCUDA(OPT::nResolutionLevel, OPT::nMinResolution, OPT::nMaxViews,
OPT::fDecimateMesh, OPT::nCloseHoles, OPT::nEnsureEdgeSize,
OPT::nMaxFaceArea,
Expand Down
2 changes: 1 addition & 1 deletion apps/TextureMesh/TextureMesh.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -110,7 +110,7 @@ bool Application::Initialize(size_t argc, LPCTSTR* argv)
), "verbosity level")
#endif
#ifdef _USE_CUDA
("cuda-device", boost::program_options::value(&CUDA::desiredDeviceID)->default_value(-1), "CUDA device number to be used to texture the mesh (-2 - CPU processing, -1 - best GPU, >=0 - device index)")
("cuda-device", boost::program_options::value(&SEACAVE::CUDA::desiredDeviceID)->default_value(-1), "CUDA device number to be used to texture the mesh (-2 - CPU processing, -1 - best GPU, >=0 - device index)")
#endif
;

Expand Down
18 changes: 7 additions & 11 deletions libs/Common/Config.h
Original file line number Diff line number Diff line change
Expand Up @@ -206,9 +206,6 @@
# define FORCEINLINE inline
#endif

#ifndef _SUPPORT_CPP11
# define constexpr inline
#endif
#ifdef _SUPPORT_CPP17
# undef MAYBEUNUSED
# define MAYBEUNUSED [[maybe_unused]]
Expand All @@ -224,37 +221,36 @@

#ifdef _MSC_VER
#define _DEBUGINFO
#define _CRTDBG_MAP_ALLOC //enable this to show also the filename (DEBUG_NEW should also be defined in each file)
#define _CRTDBG_MAP_ALLOC //enable this to show also the filename (DEBUG_NEW should also be defined in each file)
#include <cstdlib>
#include <crtdbg.h>
#ifdef _INC_CRTDBG
#define ASSERT(exp) {if (!(exp) && 1 == _CrtDbgReport(_CRT_ASSERT, __FILE__, __LINE__, NULL, #exp)) _CrtDbgBreak();}
#define ASSERT(exp, ...) {if (!(exp) && 1 == _CrtDbgReport(_CRT_ASSERT, __FILE__, __LINE__, NULL, #exp)) _CrtDbgBreak();}
#else
#define ASSERT(exp) {if (!(exp)) __debugbreak();}
#define ASSERT(exp, ...) {if (!(exp)) __debugbreak();}
#endif // _INC_CRTDBG
#define TRACE(...) {TCHAR buffer[2048]; _sntprintf(buffer, 2048, __VA_ARGS__); OutputDebugString(buffer);}
#else // _MSC_VER
#include <assert.h>
#define ASSERT(exp) assert(exp)
#define ASSERT(exp, ...) assert(exp)
#define TRACE(...)
#endif // _MSC_VER

#else

#ifdef _RELEASE
#define ASSERT(exp)
#define ASSERT(exp, ...)
#else
#ifdef _MSC_VER
#define ASSERT(exp) {if (!(exp)) __debugbreak();}
#define ASSERT(exp, ...) {if (!(exp)) __debugbreak();}
#else // _MSC_VER
#define ASSERT(exp) {if (!(exp)) __builtin_trap();}
#define ASSERT(exp, ...) {if (!(exp)) __builtin_trap();}
#endif // _MSC_VER
#endif
#define TRACE(...)

#endif // _DEBUG

#define ASSERTM(exp, msg) ASSERT(exp)

namespace SEACAVE_ASSERT
{
Expand Down
6 changes: 0 additions & 6 deletions libs/Common/FastDelegate.h
Original file line number Diff line number Diff line change
Expand Up @@ -26,12 +26,6 @@
#include <type_traits>
#include <utility>

// VC work around for constexpr and noexcept: VC2013 and below do not support these 2 keywords
#if defined(_MSC_VER) && (_MSC_VER <= 1800)
#define constexpr const
#define noexcept throw()
#endif

namespace fastdelegate
{

Expand Down
6 changes: 3 additions & 3 deletions libs/Common/UtilCUDA.h
Original file line number Diff line number Diff line change
Expand Up @@ -67,8 +67,8 @@ inline CUresult __reportCudaError(CUresult result, LPCSTR errorMessage) {
ASSERT("CudaError" == NULL);
return result;
}
#define reportCudaError(val) CUDA::__reportCudaError(val, #val)
#define checkCudaError(val) { const CUresult ret(CUDA::__reportCudaError(val, #val)); if (ret != CUDA_SUCCESS) return ret; }
#define reportCudaError(val) SEACAVE::CUDA::__reportCudaError(val, #val)
#define checkCudaError(val) { const CUresult ret(SEACAVE::CUDA::__reportCudaError(val, #val)); if (ret != CUDA_SUCCESS) return ret; }

// outputs the proper CUDA error code and abort in the event that a CUDA host call returns an error
inline void __ensureCudaResult(CUresult result, LPCSTR errorMessage) {
Expand All @@ -77,7 +77,7 @@ inline void __ensureCudaResult(CUresult result, LPCSTR errorMessage) {
ASSERT("CudaAbort" == NULL);
exit(EXIT_FAILURE);
}
#define ensureCudaResult(val) CUDA::__ensureCudaResult(val, #val)
#define ensureCudaResult(val) SEACAVE::CUDA::__ensureCudaResult(val, #val)
/*----------------------------------------------------------------*/

// rounds up addr to the align boundary
Expand Down
6 changes: 6 additions & 0 deletions libs/Common/UtilCUDADevice.h
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,12 @@

// D E F I N E S ///////////////////////////////////////////////////

#if __CUDA_ARCH__ > 0
#define __CDC__CUDA__ARCH__ 1
#else
#undef __CDC__CUDA__ARCH__
#endif

#ifndef VERBOSE
#define DEFINE_VERBOSE 1
#define VERBOSE(...) fprintf(stderr, __VA_ARGS__)
Expand Down
171 changes: 171 additions & 0 deletions libs/MVS/CameraCUDA.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,171 @@
/*
* CameraCUDA.h
*
* Copyright (c) 2014-2024 SEACAVE
*
* Author(s):
*
* cDc <[email protected]>
*
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Affero General Public License as published by
* the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* This program is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU Affero General Public License for more details.
*
* You should have received a copy of the GNU Affero General Public License
* along with this program. If not, see <http://www.gnu.org/licenses/>.
*
*
* Additional Terms:
*
* You are required to preserve legal notices and author attributions in
* that material or in the Appropriate Legal Notices displayed by works
* containing it.
*/

#ifndef _MVS_CAMERACUDA_H_
#define _MVS_CAMERACUDA_H_


// I N C L U D E S /////////////////////////////////////////////////

#define _USE_MATH_DEFINES
#include <cmath>
#include <cstdint>
#include <float.h>
#include <string>
#include <vector>

// Eigen
#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int
#include <Eigen/Dense>

// CUDA toolkit
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>
#include <cuda_texture_types.h>
#include <curand_kernel.h>
#include <vector_types.h>

#include "../Common/UtilCUDADevice.h"


// D E F I N E S ///////////////////////////////////////////////////


// S T R U C T S ///////////////////////////////////////////////////

namespace MVS {

namespace CUDA {

typedef Eigen::Matrix<int,2,1> Point2i;
typedef Eigen::Matrix<float,2,1> Point2;
typedef Eigen::Matrix<float,3,1> Point3;
typedef Eigen::Matrix<float,4,1> Point4;
typedef Eigen::Matrix<float,3,3> Matrix3;

// Linear camera model
struct LinearCameraModel {
Point2 f; // focal length
Point2 p; // principal point

__host__ __device__ LinearCameraModel() {}
__host__ __device__ LinearCameraModel(float fx, float fy, float cx, float cy) :
f(fx, fy), p(cx, cy) {}
__host__ __device__ LinearCameraModel(const Matrix3& K) :
f(K(0,0), K(1,1)), p(K(0,2), K(1,2)) { ASSERT(K(0,1) == 0); }

__host__ __device__ inline Matrix3 K() const {
Matrix3 M; M <<
f.x(), 0, p.x(),
0, f.y(), p.y(),
0, 0, 1;
return M;
}

// transform a point in image space to camera space
__host__ __device__ inline Point2 NormalizePoint(const Point2& x) const {
return Point2(
(x.x() - p.x()) / f.x(),
(x.y() - p.y()) / f.y());
}

// project a point in camera space to image space
__host__ __device__ inline Point2 TransformPointC2I(const Point3& X) const {
return Point2(
f.x() * X.x() / X.z() + p.x(),
f.y() * X.y() / X.z() + p.y());
}

// back-project a point in image space to camera space
__host__ __device__ inline Point3 TransformPointI2C(const Point2& x, const float depth = 1.f) const {
return Point3(
depth * (x.x() - p.x()) / f.x(),
depth * (x.y() - p.y()) / f.y(),
depth);
}

// compute camera ray direction for the given pixel
__host__ __device__ inline Point3 ViewDirection(const Point2i& x) const {
return TransformPointI2C(x.cast<float>()).normalized();
}
};
/*----------------------------------------------------------------*/

// Camera pose
struct Pose {
Matrix3 R; // rotation matrix
Point3 C; // camera center

__host__ __device__ Pose() {}
__host__ __device__ Pose(const Matrix3& R, const Point3& C) :
R(R), C(C) {}

// transform a 3D point from world space to camera space
__host__ __device__ inline Point3 TransformPointW2C(const Point3& X) const {
return R * (X - C);
}

// transform a 3D point in camera space to world space
__host__ __device__ inline Point3 TransformPointC2W(const Point3& X) const {
return R.transpose() * X + C;
}
};
/*----------------------------------------------------------------*/

// Camera view
struct Camera {
LinearCameraModel model;
Pose pose;
Point2i size;

__host__ __device__ Camera() {}
__host__ __device__ Camera(const LinearCameraModel& model, const Pose& pose, int width=0, int height=0) :
model(model), pose(pose), size(width, height) {}
__host__ __device__ Camera(const Matrix3& K, const Matrix3& R, const Point3& C, int width=0, int height=0) :
model(K), pose(R, C), size(width, height) {}

// project a 3D point in world space to image space
__host__ __device__ inline Point2 TransformPointW2I(const Point3& X) const {
return model.TransformPointC2I(pose.TransformPointW2C(X));
}

// back-project a point in image space to 3D point in world space
__host__ __device__ inline Point3 TransformPointI2W(const Point2& x, const float depth = 1.f) const {
return pose.TransformPointC2W(model.TransformPointI2C(x, depth));
}
};
/*----------------------------------------------------------------*/

} // namespace CUDA

} // namespace MVS

#endif // _MVS_CAMERACUDA_H_
4 changes: 2 additions & 2 deletions libs/MVS/Mesh.h
Original file line number Diff line number Diff line change
Expand Up @@ -137,13 +137,13 @@ class MVS_API Mesh
Image8U3Arr texturesDiffuse; // textures containing the diffuse color (optional)

#ifdef _USE_CUDA
static CUDA::KernelRT kernelComputeFaceNormal;
static SEACAVE::CUDA::KernelRT kernelComputeFaceNormal;
#endif

public:
#ifdef _USE_CUDA
inline Mesh() {
InitKernels(CUDA::desiredDeviceID);
InitKernels(SEACAVE::CUDA::desiredDeviceID);
}
#endif

Expand Down
Loading

0 comments on commit 8e42439

Please sign in to comment.