[add]: SurfaceMeshOperation

CMakeLists.txt

@@ -72,37 +72,38 @@ if(NOT MSVC)
 endif()
 
 include_directories(external/probabilistic-quadrics)
-
+add_subdirectory(external/polyscope)
 
 FILE(GLOB CU_UTILS_SRC_IMPL
- "*.cu"
- "**/*.cu"
- "**/**/*.cu"
+ cuMesh/*.cu
+ cuMesh/**/*.cu
  )
-FILE(GLOB CU_UTILS_SRC ./*.cpp
- ./**/*.cpp
- ./**/**/*.cpp)
-
-message(STATUS >>>>> ${CU_UTILS_SRC_IMPL})
+FILE(GLOB CU_UTILS_SRC
+ cuMesh/*.cpp
+ cuMesh/**/*.cpp)
 
 set(CUDA_LINK_LIBRARIES_KEYWORD PRIVATE)
 
-include_directories(${CMAKE_CURRENT_SOURCE_DIR})
+include_directories(${PROJECT_SOURCE_DIR})
 
-cuda_add_library(cuMesh_Impl ${CU_UTILS_SRC_IMPL} cuMesh/cuda/UpdateTopology.cuh cuMesh/cuda/HalfEdgeNavigators.cuh)
+cuda_add_library(cuMesh_Impl ${CU_UTILS_SRC_IMPL} )
 TARGET_LINK_LIBRARIES(cuMesh_Impl
  PRIVATE ${CUDA_LIBRARIES}
  PRIVATE ${CUDA_CUDART_LIBRARY}
  PRIVATE ${CUDA_cublas_LIBRARY}
  )
 
 
-add_library(cuMesh SHARED ${CU_UTILS_SRC} tests/TestCudaConvertHedge.cpp)
+message(STATUS ${CU_UTILS_SRC})
+add_library(cuMesh SHARED ${CU_UTILS_SRC} )
 target_link_libraries(cuMesh
  PRIVATE cuMesh_Impl
  PRIVATE ${CUDA_LIBRARIES}
  PRIVATE ${CUDA_CUDART_LIBRARY}
  PRIVATE ${CUDA_cublas_LIBRARY}
+ PUBLIC polyscope
  )
 
+
+
 add_subdirectory(tests)

cuMesh/HalfEdgeElementType.cuh

@@ -0,0 +1,84 @@
+//
+// Created by pupa on 2021/11/11.
+//
+#pragma once
+
+#include <thrust/device_vector.h>
+
+namespace cuMesh {
+
+using Index = uint32_t;
+#define NonIndex 0xffffffff
+
+ struct VFMeshData{
+ enum class Tag{
+ NONE = 0x0000,
+ USER0 = 0x0200, // First user bit
+ NOTREAD = 0x0002, // This bit indicate that the vertex of the mesh is not readable
+ NOTWRITE = 0x0004, // This bit indicate that the vertex is not modifiable
+ MODIFIED = 0x0008, // This bit indicate that the vertex is modified
+ VISITED = 0x0010, // This bit can be used to mark the visited vertex
+ SELECTED = 0x0020, // This bit can be used to select
+ BORDER = 0x0100, // Border Flag
+ DELETED = 0x8000, // This first bit indicate that the vertex is deleted from the mesh
+ };
+
+ thrust::device_vector<float3> _vertices;
+ thrust::device_vector<uint3> _facets;
+ thrust::device_vector<Tag> _v_tags;
+ thrust::device_vector<Tag> _f_tags;
+ thrust::device_vector<Index> _hedges_twin;
+
+ void Update(const thrust::host_vector<float3>& V) {
+ _vertices = V;
+ _v_tags = thrust::host_vector<Tag>(_vertices.size(), Tag::NONE);
+ }
+
+ void Update(const thrust::host_vector<uint3>& F) {
+ _facets = F ;
+ _f_tags = thrust::host_vector<Tag>(_facets.size(), Tag::NONE);
+ _hedges_twin = thrust::host_vector<Index>(_facets.size() * 3, NonIndex);
+ }
+
+ void Update(decltype(_vertices)::iterator begin, decltype(_vertices)::iterator end) {
+ _vertices = decltype(_vertices) (begin, end);
+ _v_tags = thrust::host_vector<Tag>(_vertices.size(), Tag::NONE);
+ }
+
+ void Update(decltype(_facets)::iterator begin, decltype(_facets)::iterator end) {
+ _facets = decltype(_facets) (begin, end);
+ _f_tags = thrust::host_vector<Tag>(_facets.size(), Tag::NONE);
+ _hedges_twin = thrust::host_vector<Index>(_facets.size() * 3, NonIndex);
+ }
+
+ std::vector<std::array<double, 3>> GetVertexPositions() {
+ std::vector<std::array<double, 3>> positions(_vertices.size());
+ thrust::host_vector<float3> vertices = _vertices;
+ for(int i = 0; i < vertices.size(); i++)
+ positions[i] = {vertices[i].x, vertices[i].y, vertices[i].z};
+ return positions;
+ }
+
+ std::vector<std::vector<size_t>> GetTriangleIndices() {
+ std::vector<std::vector<size_t>> triangles(_facets.size(), std::vector<size_t>(3, 0));
+ thrust::host_vector<uint3> facets = _facets;
+ for(int i = 0; i < facets.size(); i++){
+ triangles[i][0] = facets[i].x;
+ triangles[i][1] = facets[i].y;
+ triangles[i][2] = facets[i].z;
+ }
+ return triangles;
+ }
+
+
+ struct PointerHub {
+ float3 *_vert{nullptr};
+ uint3 *_face{nullptr};
+ Tag *_v_tag{nullptr};
+ Tag *_f_tag{nullptr};
+ Index *_hedge_twin{nullptr};
+ uint32_t n_v{0}, n_f{0}, n_h{0};
+ };
+ };
+}
+

cuMesh/HalfEdgeNavigators.cuh

@@ -0,0 +1,61 @@
+//
+// Created by pupa on 2021/11/11.
+//
+#pragma once
+
+#include <cuMesh/HalfEdgeNavigators.cuh>
+#include <helper_math.h>
+
+#define blk_size 32
+inline uint32_t GetGridDim(const uint32_t& n_thread) { return (n_thread - 1 + blk_size) / blk_size; }
+
+namespace cuMesh {
+ struct MeshNavigators: VFMeshData::PointerHub {
+ __host__ MeshNavigators(VFMeshData& meshData) {
+ n_v = meshData._vertices.size();
+ n_f = meshData._facets.size();
+ n_h = meshData._facets.size() * 3;
+ _vert = meshData._vertices.data().get();
+ _face = meshData._facets.data().get();
+ _hedge_twin = meshData._hedges_twin.data().get();
+
+ _v_tag = meshData._v_tags.data().get();
+ _f_tag = meshData._f_tags.data().get();
+ }
+ // Hedge -> Face
+ __host__ __device__ __forceinline__ Index static hedge2face(Index hid) { return hid == NonIndex? hid : hid/3 ; }
+
+ // Hedge -> Hedge
+ __host__ __device__ __forceinline__ Index static next_halfedge(Index hid) { return hid / 3 * 3 + (hid + 1) % 3 ;}
+ __host__ __device__ __forceinline__ Index static prev_halfedge(Index hid) { return hid / 3 * 3 + (hid + 2) % 3 ;}
+ __host__ __device__ __forceinline__ Index twin_halfedge(Index hid) { return hid == NonIndex ? hid : _hedge_twin[hid];}
+
+ // Hedge -> Vertex
+ __host__ __device__ __forceinline__ Index tip_vertex(Index hid) {
+ return ((Index*)(_face + hedge2face(hid)))[(hid + 1) % 3];
+ }
+ __host__ __device__ __forceinline__ Index tail_vertex(Index hid) {
+ return ((Index*)(_face + hedge2face(hid)))[hid % 3];
+ }
+
+
+ // Face -> Face
+ __host__ __device__ __forceinline__ Index adjacent_face(Index fid, size_t i_adj) {
+ return hedge2face(twin_halfedge(fid * 3 + i_adj ));
+ }
+
+ // Face -> Vertex
+ __host__ __device__ __forceinline__ Index vertex(Index fid, size_t i_adj) {
+ return tail_vertex(fid*3+i_adj);
+ }
+
+
+ // Geometry Query
+ __host__ __device__ __forceinline__ float3 triangle_centroid(Index fid) {
+ return (_vert[_face[fid].x]+_vert[_face[fid].y]+_vert[_face[fid].z])/3.0f;
+ }
+
+ __host__ __device__ __forceinline__ float3& position(Index vid) { return _vert[vid]; }
+ };
+}
+

cuMesh/algorithm/SurfaceMeshOperation.cuh

@@ -0,0 +1,37 @@
+//
+// Created by pupa on 2021/11/13.
+//
+#pragma once
+#include <cuMesh/HalfEdgeNavigators.cuh>
+
+
+struct Quadric;
+
+namespace cuMesh {
+
+
+namespace Operation{
+
+
+
+struct ComputeEdgeLength: MeshNavigators {
+ explicit __host__ ComputeEdgeLength(VFMeshData& meshData): MeshNavigators(meshData){}
+ __device__ float operator()(Index hid) ;
+};
+
+struct ComputeFaceNormal: MeshNavigators {
+ explicit __host__ ComputeFaceNormal(VFMeshData& meshData): MeshNavigators(meshData){}
+ __device__ float3 operator()(Index fid) ;
+};
+
+struct ComputeFaceQuadric: MeshNavigators {
+ explicit __host__ ComputeFaceQuadric(VFMeshData& meshData): MeshNavigators(meshData){}
+ __device__ Quadric operator()(Index& fid,float3& normal) ;
+};
+
+
+
+
+
+}
+}

cuMesh/algorithm/SurfaceSmoothing.cuh

@@ -0,0 +1,20 @@
+//
+// Created by pupa on 2021/11/13.
+//
+#pragma once
+#include <cuMesh/HalfEdgeElementType.cuh>
+#include <probabilistic-quadrics.cuh>
+
+namespace cuMesh {
+
+class MeshNavigators;
+
+namespace Smoothing {
+
+void PQGFSmoothing(VFMeshData& mesh_data);
+
+}
+
+}
+
+

cuMesh/algorithm/UpdateTopology.cuh

@@ -0,0 +1,35 @@
+//
+// Created by pupa on 2021/11/11.
+//
+#pragma once
+
+#include <cuMesh/HalfEdgeElementType.cuh>
+
+#include <thrust/sort.h>
+
+namespace cuMesh {
+
+class MeshNavigators;
+
+namespace Topology {
+
+void UpdateFaceFace(VFMeshData& mesh_data, bool is_manifold=false);
+
+}
+
+}
+
+
+namespace Kernel {
+ using cuMesh::MeshNavigators;
+
+ __global__ void CountVertexDegree(const uint3* face, uint32_t n_f, uint32_t * v_degree);
+
+ __global__ void FillOutGoingHedge(const uint3* face, uint32_t n_f, uint32_t * v_offset, uint32_t* v_hedges);
+
+ __global__ void FindHedgeTwin(MeshNavigators navigators, const uint32_t *v_offset, const uint32_t* v_hedges);
+
+ __global__ void TagNonManifoldTriangle(MeshNavigators navigators);
+
+ __global__ void CountBoundaryEdges(MeshNavigators navigators, uint32_t* n_boundary_hedges);
+}