Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Shortest edge collapse app using a preliminary cuCollections priority queue #36

Merged
merged 24 commits into from
Aug 9, 2024
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
24 commits
Select commit Hold shift + click to select a range
cdf6537
Add cmake package manager, CPM
porumbes Apr 17, 2024
21f5b34
Add cuColletions with priority queue
porumbes Apr 17, 2024
f02f351
Add initial shortest edge collapse with priority queue.
porumbes Apr 17, 2024
b625753
Initial rxmesh priority queue-based edge collapse
porumbes Apr 18, 2024
9672c5c
Merge branch 'dyn' of https://github.com/owensgroup/RXMesh into dyn
porumbes Apr 25, 2024
1c0a4ba
ignore build_debug directory
porumbes May 7, 2024
f825a1d
Add initial cuCollections priority queue
porumbes May 7, 2024
4a91e8a
Merge branch 'dyn' of https://github.com/owensgroup/RXMesh into dyn
porumbes May 7, 2024
c889d45
Encode patch and local into 32 bits, clean up some dead code
porumbes May 7, 2024
65eb418
Add kernel to pop and mark edges to be collapsed
porumbes May 8, 2024
3905344
Remove dead code
porumbes May 8, 2024
352c8e7
Use priority queue to actually simplify
porumbes May 9, 2024
4c25923
Merge branch 'dyn' of https://github.com/owensgroup/RXMesh into dyn
porumbes May 9, 2024
35fea44
Update to match histogram and add reporting
porumbes May 9, 2024
7a0bf73
Use DIVIDE_UP for number of kernel blocks
porumbes May 14, 2024
f22008b
Merge branch 'dyn' of https://github.com/owensgroup/RXMesh into dyn
porumbes May 14, 2024
18925fc
Fix edge attr reset bug that swapped value with device
porumbes May 14, 2024
3cae5d6
Add debug function to view edges to be collaped
porumbes May 15, 2024
23145d8
Add a bash script to test a range of target vertex counts
porumbes May 15, 2024
452bb39
Render edges to collapse with patch boundaries
porumbes May 16, 2024
b6e4cb1
Add edgefrac(tion) to collapse per round
porumbes May 16, 2024
024a11b
Add bash script to sweep edgefrac(tion)
porumbes May 16, 2024
f82c9b9
Merge branch 'dyn' of https://github.com/owensgroup/RXMesh into dyn
porumbes May 19, 2024
4c49be4
Merge branch 'main' into dyn
Ahdhn Aug 9, 2024
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion .gitignore
Original file line number Diff line number Diff line change
Expand Up @@ -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
18 changes: 17 additions & 1 deletion CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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")

Expand Down Expand Up @@ -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"
Expand Down
7 changes: 4 additions & 3 deletions apps/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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)
41 changes: 41 additions & 0 deletions apps/SECPriority/CMakeLists.txt
Original file line number Diff line number Diff line change
@@ -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 )
96 changes: 96 additions & 0 deletions apps/SECPriority/main.cu
Original file line number Diff line number Diff line change
@@ -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;
}
102 changes: 102 additions & 0 deletions apps/SECPriority/secp.cu
Original file line number Diff line number Diff line change
@@ -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();
}
Loading
Loading