Skip to content

Commit

Permalink
remesh with open inputs
Browse files Browse the repository at this point in the history
  • Loading branch information
@Ahdhn
Ahdhn committed Jan 6, 2025
1 parent 99f6d0b commit 75b1c36
Show file tree
Hide file tree
Showing 6 changed files with 146 additions and 38 deletions.
43 changes: 38 additions & 5 deletions apps/Remesh/collapse.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -233,6 +233,7 @@ __global__ static void __launch_bounds__(blockThreads)
edge_collapse_1(rxmesh::Context context,
const rxmesh::VertexAttribute<T> coords,
rxmesh::EdgeAttribute<EdgeStatus> edge_status,
rxmesh::VertexAttribute<bool> v_boundary,
const T low_edge_len_sq,
const T high_edge_len_sq)
{
Expand Down Expand Up @@ -291,6 +292,12 @@ __global__ static void __launch_bounds__(blockThreads)
return;
}

// don't touch boundary vertices
if (v_boundary(v0) || v_boundary(v1) || v_boundary(v2) ||
v_boundary(v3)) {
return;
}

// degenerate cases
if (v0 == v1 || v0 == v2 || v0 == v3 || v1 == v2 || v1 == v3 ||
v2 == v3) {
Expand Down Expand Up @@ -321,7 +328,7 @@ __global__ static void __launch_bounds__(blockThreads)
}
};

// 1. mark edge that we want to collapse based on the edge lenght
// 1. mark edge that we want to collapse based on the edge length
Query<blockThreads> query(context, cavity.patch_id());
query.dispatch<Op::EVDiamond>(block, shrd_alloc, should_collapse);
block.sync();
Expand Down Expand Up @@ -375,7 +382,7 @@ __global__ static void __launch_bounds__(blockThreads)
shrd_alloc.dealloc(shrd_alloc.get_allocated_size_bytes() - shmem_before);

// create the cavity
if (cavity.prologue(block, shrd_alloc, coords, edge_status)) {
if (cavity.prologue(block, shrd_alloc, coords, edge_status, v_boundary)) {

is_updated.reset(block);
block.sync();
Expand Down Expand Up @@ -494,10 +501,11 @@ inline void collapse_short_edges(rxmesh::RXMeshDynamic& rx,
rxmesh::VertexAttribute<T>* coords,
rxmesh::EdgeAttribute<EdgeStatus>* edge_status,
rxmesh::EdgeAttribute<int8_t>* edge_link,
rxmesh::VertexAttribute<bool>* v_boundary,
const T low_edge_len_sq,
const T high_edge_len_sq,
rxmesh::Timers<rxmesh::GPUTimer> timers,
int* d_buffer)
rxmesh::Timers<rxmesh::GPUTimer>& timers,
int* d_buffer)
{
using namespace rxmesh;

Expand Down Expand Up @@ -547,6 +555,7 @@ inline void collapse_short_edges(rxmesh::RXMeshDynamic& rx,
rx.get_context(),
*coords,
*edge_status,
*v_boundary,
low_edge_len_sq,
high_edge_len_sq);

Expand All @@ -557,12 +566,36 @@ inline void collapse_short_edges(rxmesh::RXMeshDynamic& rx,
timers.stop("CollapseCleanup");

timers.start("CollapseSlice");
rx.slice_patches(*coords, *edge_status /*, *edge_link */);
rx.slice_patches(
*coords, *edge_status, *v_boundary /*, *edge_link */);
timers.stop("CollapseSlice");

timers.start("CollapseCleanup");
rx.cleanup();
timers.stop("CollapseCleanup");

bool show = false;
if (show) {

rx.update_host();
EXPECT_TRUE(rx.validate());

coords->move(DEVICE, HOST);
edge_status->move(DEVICE, HOST);
rx.update_polyscope();
auto ps_mesh = rx.get_polyscope_mesh();
ps_mesh->updateVertexPositions(*coords);
ps_mesh->setEnabled(false);

ps_mesh->addEdgeScalarQuantity("EdgeStatus", *edge_status);
ps_mesh->addVertexScalarQuantity("BoundaryV", *v_boundary);

rx.render_vertex_patch();
rx.render_edge_patch();
rx.render_face_patch()->setEnabled(false);

polyscope::show();
}
}

int remaining_work = is_done(rx, edge_status, d_buffer);
Expand Down
31 changes: 20 additions & 11 deletions apps/Remesh/flip.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -185,7 +185,8 @@ __global__ static void __launch_bounds__(blockThreads)
edge_flip_1(rxmesh::Context context,
const rxmesh::VertexAttribute<T> coords,
const rxmesh::VertexAttribute<uint8_t> v_valence,
rxmesh::EdgeAttribute<EdgeStatus> edge_status)
rxmesh::EdgeAttribute<EdgeStatus> edge_status,
rxmesh::VertexAttribute<bool> v_boundary)
{
using namespace rxmesh;

Expand Down Expand Up @@ -235,6 +236,12 @@ __global__ static void __launch_bounds__(blockThreads)
return;
}

// don't touch boundary vertices
if (v_boundary(iter[0]) || v_boundary(iter[1]) ||
v_boundary(iter[2]) || v_boundary(iter[3])) {
return;
}

// degenerate cases
if (iter[0] == iter[1] || iter[0] == iter[2] ||
iter[0] == iter[3] || iter[1] == iter[2] ||
Expand Down Expand Up @@ -332,7 +339,7 @@ __global__ static void __launch_bounds__(blockThreads)

shrd_alloc.dealloc(shrd_alloc.get_allocated_size_bytes() - shmem_before);

if (cavity.prologue(block, shrd_alloc, coords, edge_status)) {
if (cavity.prologue(block, shrd_alloc, coords, edge_status, v_boundary)) {

is_updated.reset(block);
block.sync();
Expand Down Expand Up @@ -378,7 +385,8 @@ inline void equalize_valences(rxmesh::RXMeshDynamic& rx,
rxmesh::VertexAttribute<uint8_t>* v_valence,
rxmesh::EdgeAttribute<EdgeStatus>* edge_status,
rxmesh::EdgeAttribute<int8_t>* edge_link,
rxmesh::Timers<rxmesh::GPUTimer> timers,
rxmesh::VertexAttribute<bool>* v_boundary,
rxmesh::Timers<rxmesh::GPUTimer>& timers,
int* d_buffer)
{

Expand Down Expand Up @@ -440,10 +448,14 @@ inline void equalize_valences(rxmesh::RXMeshDynamic& rx,

// link_condition(rx, edge_link);

edge_flip_1<T, blockThreads><<<lb_flip.blocks,
lb_flip.num_threads,
lb_flip.smem_bytes_dyn>>>(
rx.get_context(), *coords, *v_valence, *edge_status);
edge_flip_1<T, blockThreads>
<<<lb_flip.blocks,
lb_flip.num_threads,
lb_flip.smem_bytes_dyn>>>(rx.get_context(),
*coords,
*v_valence,
*edge_status,
*v_boundary);

timers.stop("Flip");

Expand All @@ -452,7 +464,7 @@ inline void equalize_valences(rxmesh::RXMeshDynamic& rx,
timers.stop("FlipCleanup");

timers.start("FlipSlice");
rx.slice_patches(*coords, *edge_status /*,edge_link*/);
rx.slice_patches(*coords, *edge_status, *v_boundary /*,edge_link*/);
timers.stop("FlipSlice");

timers.start("FlipCleanup");
Expand All @@ -466,9 +478,6 @@ inline void equalize_valences(rxmesh::RXMeshDynamic& rx,
break;
}
prv_remaining_work = remaining_work;
// RXMESH_INFO("num_flips {}, time {}",
// num_flips,
// app_time + slice_time + cleanup_time);
}
timers.stop("FlipTotal");

Expand Down
37 changes: 31 additions & 6 deletions apps/Remesh/remesh.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3,13 +3,17 @@
#include "rxmesh/util/macros.h"
#include "rxmesh/util/util.h"

#include "rxmesh/geometry_factory.h"
struct arg
{
std::string obj_file_name = STRINGIFY(INPUT_DIR) "sphere3.obj";
std::string output_folder = STRINGIFY(OUTPUT_DIR);
float relative_len = 1.0;
uint32_t num_iter = 3;
uint32_t device_id = 0;
std::string obj_file_name = STRINGIFY(INPUT_DIR) "cloth.obj";
std::string output_folder = STRINGIFY(OUTPUT_DIR);
uint32_t nx = 66;
uint32_t ny = 66;
float relative_len = 1.0;
int num_smooth_iters = 5;
uint32_t num_iter = 1;
uint32_t device_id = 0;
char** argv;
int argc;
} Arg;
Expand All @@ -23,6 +27,13 @@ TEST(Apps, Remesh)
// Select device
cuda_query(Arg.device_id);

// std::vector<std::vector<float>> verts;
// std::vector<std::vector<uint32_t>> fv;
// create_plane(verts, fv, Arg.nx, Arg.ny);
// RXMeshDynamic rx(fv);
// rx.add_vertex_coordinates(verts, "Coords");


RXMeshDynamic rx(Arg.obj_file_name);
// rx.save(STRINGIFY(OUTPUT_DIR) + extract_file_name(Arg.obj_file_name) +
// "_patches");
Expand All @@ -31,7 +42,12 @@ TEST(Apps, Remesh)
// STRINGIFY(OUTPUT_DIR) +
// extract_file_name(Arg.obj_file_name) + "_patches");
//
// ASSERT_TRUE(rx.is_closed());

ASSERT_TRUE(rx.is_edge_manifold());

// rx.export_obj("grid_" + std::to_string(Arg.nx) + "_" +
// std::to_string(Arg.ny) + ".obj",
// *rx.get_input_vertex_coordinates());

remesh_rxmesh(rx);
}
Expand Down Expand Up @@ -84,13 +100,22 @@ int main(int argc, char** argv)
Arg.relative_len =
std::stof(get_cmd_option(argv, argv + argc, "-relative_len"));
}

if (cmd_option_exists(argv, argc + argv, "-nx")) {
Arg.nx = atoi(get_cmd_option(argv, argv + argc, "-nx"));
}
if (cmd_option_exists(argv, argc + argv, "-ny")) {
Arg.ny = atoi(get_cmd_option(argv, argv + argc, "-ny"));
}
}

RXMESH_TRACE("input= {}", Arg.obj_file_name);
RXMESH_TRACE("output_folder= {}", Arg.output_folder);
RXMESH_TRACE("device_id= {}", Arg.device_id);
RXMESH_TRACE("num_iter= {}", Arg.num_iter);
RXMESH_TRACE("relative_len= {}", Arg.relative_len);
RXMESH_TRACE("nx= {}", Arg.nx);
RXMESH_TRACE("ny= {}", Arg.ny);


return RUN_ALL_TESTS();
Expand Down
21 changes: 20 additions & 1 deletion apps/Remesh/remesh_rxmesh.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -223,6 +223,7 @@ inline void remesh_rxmesh(rxmesh::RXMeshDynamic& rx)

auto v_valence = rx.add_vertex_attribute<uint8_t>("Valence", 1);

auto v_boundary = rx.add_vertex_attribute<bool>("BoundaryV", 1);

auto edge_len = rx.add_edge_attribute<float>("edgeLen", 1);
auto vertex_valence = rx.add_vertex_attribute<int>("vertexValence", 1);
Expand All @@ -237,6 +238,8 @@ inline void remesh_rxmesh(rxmesh::RXMeshDynamic& rx)

auto edge_link = rx.add_edge_attribute<int8_t>("edgeLink", 1);

rx.get_boundary_vertices(*v_boundary);

// edge_link->move(DEVICE, HOST);
// rx.get_polyscope_mesh()->addEdgeScalarQuantity("edgeLink", *edge_link);

Expand Down Expand Up @@ -300,6 +303,7 @@ inline void remesh_rxmesh(rxmesh::RXMeshDynamic& rx)
split_long_edges(rx,
coords.get(),
edge_status.get(),
v_boundary.get(),
high_edge_len_sq,
low_edge_len_sq,
timers,
Expand All @@ -310,6 +314,7 @@ inline void remesh_rxmesh(rxmesh::RXMeshDynamic& rx)
coords.get(),
edge_status.get(),
edge_link.get(),
v_boundary.get(),
low_edge_len_sq,
high_edge_len_sq,
timers,
Expand All @@ -322,11 +327,17 @@ inline void remesh_rxmesh(rxmesh::RXMeshDynamic& rx)
v_valence.get(),
edge_status.get(),
edge_link.get(),
v_boundary.get(),
timers,
d_buffer);

RXMESH_INFO(" Vertex Smoothing -- iter {}", iter);
tangential_relaxation(rx, coords.get(), new_coords.get(), timers);
tangential_relaxation(rx,
coords.get(),
new_coords.get(),
v_boundary.get(),
Arg.num_smooth_iters,
timers);
std::swap(new_coords, coords);
}

Expand Down Expand Up @@ -374,6 +385,14 @@ inline void remesh_rxmesh(rxmesh::RXMeshDynamic& rx)
edge_status->get_memory_mg() +
vertex_valence->get_memory_mg());

RXMESH_INFO("Split Total Time {} (ms)",
timers.elapsed_millis("SplitTotal"));
RXMESH_INFO("Collapse Total Time {} (ms)",
timers.elapsed_millis("CollapseTotal"));
RXMESH_INFO("Flip Total Time {} (ms)",
timers.elapsed_millis("FlipTotal"));
RXMESH_INFO("Smooth Total Time {} (ms)",
timers.elapsed_millis("SmoothTotal"));

report.add_member("split_time_ms", timers.elapsed_millis("SplitTotal"));
report.add_member("collapse_time_ms",
Expand Down
31 changes: 22 additions & 9 deletions apps/Remesh/smoothing.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,8 @@ template <typename T, uint32_t blockThreads>
__global__ static void __launch_bounds__(blockThreads)
vertex_smoothing(const rxmesh::Context context,
const rxmesh::VertexAttribute<T> coords,
rxmesh::VertexAttribute<T> new_coords)
rxmesh::VertexAttribute<T> new_coords,
rxmesh::VertexAttribute<bool> v_boundary)
{
// VV to compute vertex sum and normal
using namespace rxmesh;
Expand All @@ -22,6 +23,13 @@ __global__ static void __launch_bounds__(blockThreads)
return;
}

if (v_boundary(v_id)) {
new_coords(v_id, 0) = coords(v_id, 0);
new_coords(v_id, 1) = coords(v_id, 1);
new_coords(v_id, 2) = coords(v_id, 2);
return;
}

const vec3<T> v = coords.to_glm<3>(v_id);

// compute both vertex normal and the new position
Expand Down Expand Up @@ -94,10 +102,12 @@ __global__ static void __launch_bounds__(blockThreads)
}

template <typename T>
inline void tangential_relaxation(rxmesh::RXMeshDynamic& rx,
rxmesh::VertexAttribute<T>* coords,
rxmesh::VertexAttribute<T>* new_coords,
rxmesh::Timers<rxmesh::GPUTimer> timers)
inline void tangential_relaxation(rxmesh::RXMeshDynamic& rx,
rxmesh::VertexAttribute<T>* coords,
rxmesh::VertexAttribute<T>* new_coords,
rxmesh::VertexAttribute<bool>* v_boundary,
const int num_smooth_iters,
rxmesh::Timers<rxmesh::GPUTimer>& timers)
{
using namespace rxmesh;

Expand All @@ -111,10 +121,13 @@ inline void tangential_relaxation(rxmesh::RXMeshDynamic& rx,
true);

timers.start("SmoothTotal");
vertex_smoothing<T, blockThreads>
<<<launch_box.blocks,
launch_box.num_threads,
launch_box.smem_bytes_dyn>>>(rx.get_context(), *coords, *new_coords);
for (int i = 0; i < num_smooth_iters; ++i) {
vertex_smoothing<T, blockThreads><<<launch_box.blocks,
launch_box.num_threads,
launch_box.smem_bytes_dyn>>>(
rx.get_context(), *coords, *new_coords, *v_boundary);
std::swap(new_coords, coords);
}
timers.stop("SmoothTotal");

RXMESH_INFO("Relax time {} (ms)", timers.elapsed_millis("SmoothTotal"));
Expand Down
Loading

0 comments on commit 75b1c36

Please sign in to comment.