Merge pull request #36 from porumbes/dyn

Shortest edge collapse app using a preliminary cuCollections priority queue

Author: Ahdhn

.gitignore

@@ -12,6 +12,7 @@ input/*
 !input/sphere1.obj
 !input/bunnyhead.obj
 build/
+build_debug/
 include/rxmesh/util/git_sha1.cpp
 .vscode/
-scripts/*.log
+scripts/*.log

CMakeLists.txt

@@ -79,6 +79,22 @@ FetchContent_Declare(cereal
 )
 FetchContent_Populate(cereal)
 
+# Package Management
+# TODO: Consider using CPM for the various libraries above 
+include(cmake/CPM.cmake) 
+
+# Add cuCollection with priority queue. This should eventually come from
+# NVIDIA.
+CPMAddPackage(
+    NAME cuco
+    GITHUB_REPOSITORY andrewbriand/cuCollections
+    GIT_TAG d58dd9fedde721a264c8ae960f7393a3a3b08c58
+    OPTIONS
+       "BUILD_TESTS OFF"
+       "BUILD_BENCHMARKS OFF"
+       "BUILD_EXAMPLES OFF"
+)
+
 # Auto-detect GPU architecture
 include("cmake/AutoDetectCudaArch.cmake")
 
@@ -133,7 +149,7 @@ set(cxx_flags
 
 set(MSVC_XCOMPILER_FLAGS "/openmp:experimental /MP /std:c++17 /Zi")
 set(cuda_flags
-    -Xcompiler=$<$<CXX_COMPILER_ID:GNU>:-Wall -fopenmp -O3 -Wno-unused-function>
+    -Xcompiler=$<$<CXX_COMPILER_ID:GNU>:-rdynamic -Wall -fopenmp -O3 -Wno-unused-function>
 	-Xcompiler=$<$<CXX_COMPILER_ID:MSVC>:${MSVC_XCOMPILER_FLAGS}>
     #Disables warning
     #177-D "function XXX was declared but never referenced"

apps/CMakeLists.txt

@@ -4,12 +4,13 @@ add_subdirectory(MCF)
 add_subdirectory(Geodesic)
 add_subdirectory(Delaunay)
 add_subdirectory(GaussianCurvature)
-add_subdirectory(XPBD )
+add_subdirectory(XPBD)
 #add_subdirectory(Simplification)
 add_subdirectory(ShortestEdgeCollapse)
 add_subdirectory(Remesh)
+add_subdirectory(SECPriority)
+add_subdirectory(SurfaceTracking)
 add_subdirectory(SurfaceTracking)
 add_subdirectory(SCP)
 add_subdirectory(ARAP)
-add_subdirectory(Heat)
-
+add_subdirectory(Heat)

apps/SECPriority/CMakeLists.txt

@@ -0,0 +1,41 @@
+add_executable(SECPriority)
+
+set(SOURCE_LIST
+    #main.cu
+    secp.cu
+	secp_rxmesh.cuh
+	secp_kernels.cuh
+)
+
+set(COMMON_LIST    
+    ../common/openmesh_trimesh.h
+	../common/openmesh_report.h
+)
+
+target_sources(SECPriority 
+    PRIVATE
+    ${SOURCE_LIST} ${COMMON_LIST}
+)
+
+if (WIN32)
+    target_compile_definitions(SECPriority
+      PRIVATE _USE_MATH_DEFINES 
+      PRIVATE NOMINMAX
+      PRIVATE _CRT_SECURE_NO_WARNINGS)
+endif()
+
+set_target_properties(SECPriority PROPERTIES FOLDER "apps")
+
+set_property(TARGET SECPriority PROPERTY CUDA_SEPARABLE_COMPILATION ON)
+
+source_group(TREE ${CMAKE_CURRENT_LIST_DIR} PREFIX "SECPriority" FILES ${SOURCE_LIST})
+
+target_link_libraries(SECPriority     
+    PRIVATE RXMesh
+    PRIVATE gtest_main
+	PRIVATE OpenMeshCore
+    PRIVATE OpenMeshTools
+    PRIVATE cuco
+)
+
+#gtest_discover_tests( SECPriority )

apps/SECPriority/main.cu

@@ -0,0 +1,96 @@
+#include <cuco/priority_queue.cuh>
+#include <cuco/detail/pair.cuh>
+
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+
+#include <cooperative_groups.h>
+#include <cuda_runtime.h>
+
+#include <map>
+#include <vector>
+
+#include <iostream>
+#include <random>
+
+using namespace cuco;
+namespace cg = cooperative_groups;
+
+// grab some bits from priority queue tests and benchmarks
+
+// -- simulate reading the mesh, computing edge length
+// -- cuco:pair<float, uint32_t>
+//
+// setup pair_less template
+//
+// setup device function to pop items from queue
+//
+
+template <typename T>
+struct pair_less 
+{
+    __host__ __device__ bool operator()(const T& a, const T& b) const
+    {
+        return a.first < b.first;
+    }
+};
+
+template <typename PairType, typename OutputIt>
+void generate_kv_pairs_uniform(OutputIt output_begin, OutputIt output_end)
+{
+  std::random_device rd;
+  std::mt19937 gen{rd()};
+
+  const auto num_keys = std::distance(output_begin, output_end);
+  for(auto i = 0; i < num_keys; i++)
+  {
+    output_begin[i] = {static_cast<typename PairType::first_type>(gen()),
+                       static_cast<typename PairType::second_type>(i)};
+  }
+}
+
+void sp_pair()
+{
+  // Setup the cuco::priority_queue
+  const size_t insertion_size = 200;
+  const size_t deletion_size  = 100;
+  using PairType              = cuco::pair<float, uint32_t>;
+  using Compare               = pair_less<PairType>;
+
+  cuco::priority_queue<PairType, Compare> pq(insertion_size);
+
+  // Generate data for the queue
+  std::vector<PairType> h_pairs(insertion_size);
+  generate_kv_pairs_uniform<PairType>(h_pairs.begin(), h_pairs.end());
+
+  for(auto i = 0; i < h_pairs.size(); i++)
+  {
+    std::cout << "Priority: " << h_pairs[i].first 
+              << "\tID: " << h_pairs[i].second << "\n";
+  }
+
+  // Fill the priority queue
+  thrust::device_vector<PairType> d_pairs(h_pairs);
+  pq.push(d_pairs.begin(), d_pairs.end());
+  cudaDeviceSynchronize();
+
+  // Pop the priority queue
+  thrust::device_vector<PairType> d_popped(deletion_size);
+  pq.pop(d_popped.begin(), d_popped.end());
+  cudaDeviceSynchronize();
+
+  std::cout << "-----After Pop-----\n";
+  thrust::host_vector<PairType> h_popped(d_popped);
+  for(auto i = 0; i < h_popped.size(); i++)
+  {
+    std::cout << "Priority: " << h_popped[i].first 
+              << "\tID: " << h_popped[i].second << "\n";
+  }
+}
+
+int main(int argc, char* argv[])
+{
+  sp_pair();
+
+  return 0;
+}

apps/SECPriority/secp.cu

@@ -0,0 +1,102 @@
+#include "gtest/gtest.h"
+#include "rxmesh/util/log.h"
+#include "rxmesh/util/macros.h"
+#include "rxmesh/util/util.h"
+
+#include <filesystem>
+
+struct arg
+{
+    std::string obj_file_name = STRINGIFY(INPUT_DIR) "dragon.obj";
+    std::string output_folder = STRINGIFY(OUTPUT_DIR);
+    float       target        = 0.1;
+    float       edgefrac = 0.1;
+    uint32_t    device_id     = 0;
+    char**      argv;
+    int         argc;
+} Arg;
+
+#include "secp_rxmesh.cuh"
+
+TEST(Apps, SECPriority)
+{
+    using namespace rxmesh;
+
+    // Select device
+    cuda_query(Arg.device_id);
+
+    // RXMeshDynamic rx(Arg.obj_file_name);
+
+    const std::string p_file = STRINGIFY(OUTPUT_DIR) +
+                               extract_file_name(Arg.obj_file_name) +
+                               "_patches";
+    RXMeshDynamic rx(Arg.obj_file_name, p_file);
+    if (!std::filesystem::exists(p_file)) {
+        rx.save(p_file);
+    }
+
+    ASSERT_TRUE(rx.is_edge_manifold());
+
+    ASSERT_TRUE(rx.is_closed());
+
+    uint32_t final_num_vertices = Arg.target * rx.get_num_vertices();
+
+    secp_rxmesh(rx, final_num_vertices, Arg.edgefrac);
+}
+
+
+int main(int argc, char** argv)
+{
+    using namespace rxmesh;
+    Log::init();
+
+    ::testing::InitGoogleTest(&argc, argv);
+    Arg.argv = argv;
+    Arg.argc = argc;
+
+
+    if (argc > 1) {
+        if (cmd_option_exists(argv, argc + argv, "-h")) {
+            // clang-format off
+            RXMESH_INFO("\nUsage: SECPriority.exe < -option X>\n"
+                        " -h:          Display this massage and exit\n"
+                        " -input:      Input file. Input file should be under the input/ subdirectory\n"
+                        "              Default is {} \n"
+                        "              Hint: Only accept OBJ files\n"
+                        " -target:     The fraction of output #vertices from the input\n"
+                        " -edgefrac:   The fraction of edges to collapse in a round\n"
+                        " -o:          JSON file output folder. Default is {} \n"
+                        " -device_id:  GPU device ID. Default is {}",
+            Arg.obj_file_name, Arg.output_folder, Arg.device_id);
+            // clang-format on
+            exit(EXIT_SUCCESS);
+        }
+
+        if (cmd_option_exists(argv, argc + argv, "-input")) {
+            Arg.obj_file_name =
+                std::string(get_cmd_option(argv, argv + argc, "-input"));
+        }
+        if (cmd_option_exists(argv, argc + argv, "-o")) {
+            Arg.output_folder =
+                std::string(get_cmd_option(argv, argv + argc, "-o"));
+        }
+        if (cmd_option_exists(argv, argc + argv, "-device_id")) {
+            Arg.device_id =
+                atoi(get_cmd_option(argv, argv + argc, "-device_id"));
+        }
+        if (cmd_option_exists(argv, argc + argv, "-target")) {
+            Arg.target = atof(get_cmd_option(argv, argv + argc, "-target"));
+        }
+        if (cmd_option_exists(argv, argc + argv, "-edgefrac")) {
+            Arg.edgefrac = atof(get_cmd_option(argv, argv + argc, "-edgefrac"));
+        }
+    }
+
+    RXMESH_TRACE("input= {}", Arg.obj_file_name);
+    RXMESH_TRACE("output_folder= {}", Arg.output_folder);
+    RXMESH_TRACE("device_id= {}", Arg.device_id);
+    RXMESH_TRACE("target= {}", Arg.target);
+    RXMESH_TRACE("edgefrac= {}", Arg.edgefrac);
+
+    return RUN_ALL_TESTS();
+}