From e86f750a5e574a70f452ce89b108038e479c670e Mon Sep 17 00:00:00 2001 From: ahmed Date: Mon, 22 May 2023 22:37:27 -0400 Subject: [PATCH] debugging --- apps/Delaunay/delaunay_rxmesh.cuh | 25 +++-- include/rxmesh/cavity_manager.cuh | 15 ++- include/rxmesh/cavity_manager_impl.cuh | 146 +++++++++++++++---------- include/rxmesh/rxmesh.h | 2 +- include/rxmesh/rxmesh_dynamic.cu | 17 +-- include/rxmesh/rxmesh_dynamic.h | 2 +- tests/RXMesh_test/test_dynamic.cuh | 71 ++++++------ 7 files changed, 163 insertions(+), 115 deletions(-) diff --git a/apps/Delaunay/delaunay_rxmesh.cuh b/apps/Delaunay/delaunay_rxmesh.cuh index 04825c58..ec6972ee 100644 --- a/apps/Delaunay/delaunay_rxmesh.cuh +++ b/apps/Delaunay/delaunay_rxmesh.cuh @@ -16,7 +16,9 @@ __global__ static void delaunay_edge_flip(rxmesh::Context context, if (pid == INVALID32) { return; } - + if (threadIdx.x == 0) { + printf("\n working on p=%u", pid); + } // edge diamond query to and calc delaunay condition auto is_delaunay = [&](const EdgeHandle& eh, const VertexIterator& iter) { // iter[0] and iter[2] are the edge two vertices @@ -115,7 +117,7 @@ __global__ static void delaunay_edge_flip(rxmesh::Context context, }); if (threadIdx.x == 0) { - printf("\n p= %u", cavity.patch_id()); + printf("\n success p= %u\n", cavity.patch_id()); } } cavity.epilogue(block); @@ -130,6 +132,7 @@ inline bool delaunay_rxmesh(rxmesh::RXMeshDynamic& rx) const uint32_t num_edges = rx.get_num_edges(); const uint32_t num_faces = rx.get_num_faces(); + #if USE_POLYSCOPE rx.polyscope_render_vertex_patch(); rx.polyscope_render_edge_patch(); @@ -140,6 +143,7 @@ inline bool delaunay_rxmesh(rxmesh::RXMeshDynamic& rx) EXPECT_TRUE(rx.validate()); + GPUTimer timer; timer.start(); int iter = 0; @@ -150,14 +154,17 @@ inline bool delaunay_rxmesh(rxmesh::RXMeshDynamic& rx) launch_box, (void*)delaunay_edge_flip); delaunay_edge_flip - <<>>(rx.get_context(), *coords); + <<<1, launch_box.num_threads, launch_box.smem_bytes_dyn>>>( + rx.get_context(), *coords); CUDA_ERROR(cudaDeviceSynchronize()); + RXMESH_INFO("slicing"); rx.slice_patches(*coords); + CUDA_ERROR(cudaDeviceSynchronize()); + RXMESH_INFO("cleanup"); rx.cleanup(); + CUDA_ERROR(cudaDeviceSynchronize()); timer.stop(); CUDA_ERROR(cudaDeviceSynchronize()); @@ -168,7 +175,7 @@ inline bool delaunay_rxmesh(rxmesh::RXMeshDynamic& rx) rx.update_host(); coords->move(DEVICE, HOST); - + EXPECT_EQ(num_vertices, rx.get_num_vertices()); EXPECT_EQ(num_edges, rx.get_num_edges()); EXPECT_EQ(num_faces, rx.get_num_faces()); @@ -181,14 +188,18 @@ inline bool delaunay_rxmesh(rxmesh::RXMeshDynamic& rx) rx.update_polyscope(); auto ps_mesh = rx.get_polyscope_mesh(); + ps_mesh->updateVertexPositions(*coords); ps_mesh->setEnabled(false); + for (uint32_t p = 0; p < rx.get_num_patches(); ++p) { + rx.render_patch(p)->setEnabled(false); + } rx.polyscope_render_vertex_patch(); rx.polyscope_render_edge_patch(); rx.polyscope_render_face_patch(); - // polyscope::show(); + polyscope::show(); #endif } return true; diff --git a/include/rxmesh/cavity_manager.cuh b/include/rxmesh/cavity_manager.cuh index 754041e9..b4821c34 100644 --- a/include/rxmesh/cavity_manager.cuh +++ b/include/rxmesh/cavity_manager.cuh @@ -22,7 +22,7 @@ struct CavityManager __device__ __inline__ CavityManager() : m_write_to_gmem(true), m_s_num_cavities(nullptr), - m_s_cavity_size_prefix(nullptr), + m_s_cavity_size_prefix(nullptr), m_s_readd_to_queue(nullptr), m_s_ev(nullptr), m_s_fe(nullptr), @@ -195,7 +195,7 @@ struct CavityManager */ __device__ __inline__ void alloc_shared_memory( cooperative_groups::thread_block& block, - ShmemAllocator& shrd_alloc); + ShmemAllocator& shrd_alloc); /** * @brief load hashtable into shared memory @@ -268,7 +268,8 @@ struct CavityManager Bitmask& active_bitmask, Bitmask& in_cavity, const uint16_t* element_cavity_id, - const uint16_t num_elements); + const uint16_t num_elements, + bool debug = false); /** * @brief construct the cavities boundary loop for all cavities created in @@ -287,9 +288,12 @@ struct CavityManager * @brief find the index of the next element to add. We do this by * atomically attempting to set the active_bitmask until we find an element * where we successfully flipped its status from inactive to active. + * If we fail, we just atomically increment num_elements and set the + * corresponding bit in active_bitmask */ __device__ __inline__ uint16_t add_element(Bitmask active_bitmask, - const uint16_t num_elements); + uint16_t* num_elements, + const uint16_t capacity); /** * @brief enqueue patch in the patch scheduler so that it can be scheduled @@ -449,7 +453,8 @@ struct CavityManager const Bitmask& s_ownership_change, const LPPair* s_table, const LPPair* s_stash, - Bitmask& s_owned_bitmask); + Bitmask& s_owned_bitmask, + bool debug = false); /** * @brief update an attribute such that it can be used after the topology diff --git a/include/rxmesh/cavity_manager_impl.cuh b/include/rxmesh/cavity_manager_impl.cuh index 79c9772b..26386f14 100644 --- a/include/rxmesh/cavity_manager_impl.cuh +++ b/include/rxmesh/cavity_manager_impl.cuh @@ -31,7 +31,8 @@ __device__ __inline__ CavityManager::CavityManager( m_s_num_cavities[0] = 0; // get a patch - s_patch_id = m_context.m_patch_scheduler.pop(); + // s_patch_id = m_context.m_patch_scheduler.pop(); + s_patch_id = 0; if (s_patch_id != INVALID32) { if (m_context.m_patches_info[s_patch_id].patch_id == INVALID32) { @@ -556,8 +557,11 @@ CavityManager::clear_bitmask_if_in_cavity() m_s_num_vertices[0]); clear_bitmask_if_in_cavity( m_s_active_mask_e, m_s_in_cavity_e, m_s_cavity_id_e, m_s_num_edges[0]); - clear_bitmask_if_in_cavity( - m_s_active_mask_f, m_s_in_cavity_f, m_s_cavity_id_f, m_s_num_faces[0]); + clear_bitmask_if_in_cavity(m_s_active_mask_f, + m_s_in_cavity_f, + m_s_cavity_id_f, + m_s_num_faces[0], + true); } @@ -567,11 +571,16 @@ CavityManager::clear_bitmask_if_in_cavity( Bitmask& active_bitmask, Bitmask& in_cavity, const uint16_t* element_cavity_id, - const uint16_t num_elements) + const uint16_t num_elements, + bool debug) { for (uint16_t b = threadIdx.x; b < num_elements; b += blockThreads) { if (element_cavity_id[b] != INVALID16) { active_bitmask.reset(b, true); + if (debug) { + printf("\n resetting %u", b); + } + // we don't reset owned bitmask since we use it in find_copy in_cavity.set(b, true); assert(!active_bitmask(b)); } @@ -806,16 +815,10 @@ template __device__ __inline__ VertexHandle CavityManager::add_vertex() { - // First try to reuse a vertex in the cavity or a deleted vertex - uint16_t v_id = add_element(m_s_active_mask_v, m_s_num_vertices[0]); - - if (v_id == INVALID16) { - // if this fails, then add a new vertex to the mesh - v_id = atomicAdd(m_s_num_vertices, 1); - assert(v_id < m_patch_info.vertices_capacity[0]); - } - assert(!m_s_active_mask_v(v_id)); + uint16_t v_id = add_element( + m_s_active_mask_v, m_s_num_vertices, m_patch_info.vertices_capacity[0]); + assert(m_s_active_mask_v(v_id)); m_s_owned_mask_v.set(v_id, true); return {m_patch_info.patch_id, v_id}; } @@ -829,14 +832,10 @@ __device__ __inline__ DEdgeHandle CavityManager::add_edge( assert(src.patch_id() == m_patch_info.patch_id); assert(dest.patch_id() == m_patch_info.patch_id); - // First try to reuse an edge in the cavity or a deleted edge - uint16_t e_id = add_element(m_s_active_mask_e, m_s_num_edges[0]); - if (e_id == INVALID16) { - // if this fails, then add a new edge to the mesh - e_id = atomicAdd(m_s_num_edges, 1); - assert(e_id < m_patch_info.edges_capacity[0]); - } + uint16_t e_id = add_element( + m_s_active_mask_e, m_s_num_edges, m_patch_info.edges_capacity[0]); + assert(m_s_active_mask_e(e_id)); assert(m_s_active_mask_v(src.local_id())); assert(m_s_active_mask_v(dest.local_id())); @@ -857,14 +856,12 @@ __device__ __inline__ FaceHandle CavityManager::add_face( assert(e1.patch_id() == m_patch_info.patch_id); assert(e2.patch_id() == m_patch_info.patch_id); - // First try to reuse a face in the cavity or a deleted face - uint16_t f_id = add_element(m_s_active_mask_f, m_s_num_faces[0]); + uint16_t f_id = add_element( + m_s_active_mask_f, m_s_num_faces, m_patch_info.faces_capacity[0]); - if (f_id == INVALID16) { - // if this fails, then add a new face to the mesh - f_id = atomicAdd(m_s_num_faces, 1); - assert(f_id < m_patch_info.faces_capacity[0]); - } + printf("\n add_face %u", f_id); + + assert(m_s_active_mask_f(f_id)); m_s_fe[3 * f_id + 0] = e0.local_id(); m_s_fe[3 * f_id + 1] = e1.local_id(); @@ -884,16 +881,22 @@ __device__ __inline__ FaceHandle CavityManager::add_face( template __device__ __inline__ uint16_t CavityManager::add_element( Bitmask active_bitmask, - const uint16_t num_elements) + uint16_t* num_elements, + const uint16_t capacity) { - - for (uint16_t i = 0; i < num_elements; ++i) { + const uint16_t count = num_elements[0]; + for (uint16_t i = 0; i < count; ++i) { if (active_bitmask.try_set(i)) { return i; } } - - return INVALID16; + uint16_t ret = atomicAdd(num_elements, uint16_t(1)); + if (ret >= capacity) { + return INVALID16; + } else { + active_bitmask.set(ret, true); + return ret; + } } @@ -919,6 +922,7 @@ template __device__ __inline__ void CavityManager::push() { if (threadIdx.x == 0) { + printf("\n readding %u\n", patch_id()); m_context.m_patch_scheduler.push(m_patch_info.patch_id); } } @@ -1513,16 +1517,15 @@ __device__ __inline__ LPPair CavityManager::migrate_vertex( assert(m_context.m_patches_info[o].is_owned(LocalVertexT(vq))); if (vp == INVALID16) { - - vp = atomicAdd(m_s_num_vertices, 1u); - if (vp >= m_patch_info.vertices_capacity[0]) { + vp = add_element(m_s_active_mask_v, + m_s_num_vertices, + m_patch_info.vertices_capacity[0]); + if (vp == INVALID16) { m_s_should_slice[0] = true; return ret; } - // assert(vp < m_patch_info.vertices_capacity[0]); - // activate the vertex in the bit mask - m_s_active_mask_v.set(vp, true); + // active bitmask is set in add_element // since it is owned by some other patch m_s_owned_mask_v.reset(vp, true); @@ -1585,9 +1588,10 @@ __device__ __inline__ LPPair CavityManager::migrate_edge( assert(m_context.m_patches_info[o].is_owned(LocalEdgeT(eq))); if (ep == INVALID16) { - ep = atomicAdd(m_s_num_edges, 1u); - - if (ep >= m_patch_info.edges_capacity[0]) { + ep = add_element(m_s_active_mask_e, + m_s_num_edges, + m_patch_info.edges_capacity[0]); + if (ep == INVALID16) { m_s_should_slice[0] = true; return ret; } @@ -1624,8 +1628,7 @@ __device__ __inline__ LPPair CavityManager::migrate_edge( m_s_ev[2 * ep + 0] = v0p; m_s_ev[2 * ep + 1] = v1p; - // activate the edge in the bitmask - m_s_active_mask_e.set(ep, true); + // active bitmask is set in add_element // since it is owned by some other patch m_s_owned_mask_e.reset(ep, true); @@ -1690,9 +1693,11 @@ __device__ __inline__ LPPair CavityManager::migrate_face( assert(m_context.m_patches_info[o].is_owned(LocalFaceT(fq))); if (fp == INVALID16) { - fp = atomicAdd(m_s_num_faces, 1u); - - if (fp >= m_patch_info.faces_capacity[0]) { + fp = add_element(m_s_active_mask_f, + m_s_num_faces, + m_patch_info.faces_capacity[0]); + printf("\n migrate_Face = %u", fp); + if (fp == INVALID16) { m_s_should_slice[0] = true; return ret; } @@ -1730,8 +1735,7 @@ __device__ __inline__ LPPair CavityManager::migrate_face( m_s_fe[3 * fp + 1] = (e1p << 1) | d1; m_s_fe[3 * fp + 2] = (e2p << 1) | d2; - // activate the face in the bitmask - m_s_active_mask_f.set(fp, true); + // active bitmask is set in add_element // since it is owned by some other patch m_s_owned_mask_f.reset(fp, true); @@ -1849,9 +1853,16 @@ __device__ __inline__ uint16_t CavityManager::find_copy( const HandleT handle = m_patch_info.find( i, s_table, s_stash, m_s_patch_stash); - assert(handle.is_valid()); + + // These assertion does not work any more since we change the active + // and owned mask when we add new elements. So, a thread A might set + // the bit for the active mask and reset the owned for element X + // before adding it to the hashtable leading to another thread B + // looking for it without finding it + + /*assert(handle.is_valid()); assert(handle.patch_id() != INVALID32); - assert(handle.local_id() != INVALID16); + assert(handle.local_id() != INVALID16);*/ if (handle.patch_id() == src_patch && handle.local_id() == lid) { return i; @@ -1868,24 +1879,25 @@ __device__ __inline__ void CavityManager::change_ownership( { change_ownership(block, m_s_num_vertices[0], - m_s_ownership_change_mask_v, + m_s_ownership_change_mask_v, m_s_table_v, m_s_table_stash_v, m_s_owned_mask_v); change_ownership(block, m_s_num_edges[0], - m_s_ownership_change_mask_e, + m_s_ownership_change_mask_e, m_s_table_e, m_s_table_stash_e, m_s_owned_mask_e); change_ownership(block, m_s_num_faces[0], - m_s_ownership_change_mask_f, + m_s_ownership_change_mask_f, m_s_table_f, m_s_table_stash_f, - m_s_owned_mask_f); + m_s_owned_mask_f, + true); } @@ -1894,16 +1906,16 @@ template __device__ __inline__ void CavityManager::change_ownership( cooperative_groups::thread_block& block, const uint16_t num_elements, - const Bitmask& s_ownership_change, + const Bitmask& s_ownership_change, const LPPair* s_table, const LPPair* s_stash, - Bitmask& s_owned_bitmask) + Bitmask& s_owned_bitmask, + bool debug) { for (uint16_t vp = threadIdx.x; vp < num_elements; vp += blockThreads) { if (s_ownership_change(vp)) { - - assert(!m_patch_info.is_owned(HandleT::LocalT(vp))); + assert(!s_owned_bitmask(vp)); const HandleT h = m_patch_info.find( vp, s_table, s_stash, m_s_patch_stash); @@ -1916,7 +1928,9 @@ __device__ __inline__ void CavityManager::change_ownership( // set the bitmask of this element in shared memory s_owned_bitmask.set(vp, true); - + if (debug) { + printf("\n owning %u", vp); + } // m_patch_info.get_lp().remove(vp); // ensure patch inclusion @@ -2135,7 +2149,19 @@ __device__ __inline__ void CavityManager::epilogue( } } else if (m_s_should_slice[0]) { if (threadIdx.x == 0) { - // printf("\n should slice =%u", patch_id()); + printf( + "\n should slice =%u, Vertices g(%u), s(%u), cap(%u), Edges " + "g(%u), s(%u), cap(%u), Face g(%u), s(%u), cap(%u), ", + patch_id(), + m_patch_info.num_vertices[0], + m_s_num_vertices[0], + m_patch_info.vertices_capacity[0], + m_patch_info.num_edges[0], + m_s_num_edges[0], + m_patch_info.edges_capacity[0], + m_patch_info.num_faces[0], + m_s_num_faces[0], + m_patch_info.faces_capacity[0]); m_context.m_patches_info[patch_id()].should_slice = true; } } diff --git a/include/rxmesh/rxmesh.h b/include/rxmesh/rxmesh.h index e876c0bd..e93f5ef7 100644 --- a/include/rxmesh/rxmesh.h +++ b/include/rxmesh/rxmesh.h @@ -291,7 +291,7 @@ class RXMesh const bool quite = false, const float capacity_factor = 1.2, const float patch_alloc_factor = 2.0, - const float lp_hashtable_load_factor = 0.7); + const float lp_hashtable_load_factor = 0.5); /** * @brief build different supporting data structure used to build RXMesh diff --git a/include/rxmesh/rxmesh_dynamic.cu b/include/rxmesh/rxmesh_dynamic.cu index 84cb58de..c31b9b5c 100644 --- a/include/rxmesh/rxmesh_dynamic.cu +++ b/include/rxmesh/rxmesh_dynamic.cu @@ -698,17 +698,20 @@ __inline__ __device__ void slice(Context& context, // if the element is active, owned by this patch and the new patch, then // we remove the ownership from this patch for (uint16_t v = threadIdx.x; v < num_vertices; v += blockThreads) { - if (s_active_v(v) && s_owned_v(v) && s_new_p_owned_v(v)) { + if ((s_active_v(v) && s_owned_v(v) && s_new_p_owned_v(v)) || + !s_active_v(v)) { s_owned_v.reset(v, true); } } for (uint16_t e = threadIdx.x; e < num_edges; e += blockThreads) { - if (s_active_e(e) && s_owned_e(e) && s_new_p_owned_e(e)) { + if ((s_active_e(e) && s_owned_e(e) && s_new_p_owned_e(e)) || + !s_active_e(e)) { s_owned_e.reset(e, true); } } for (uint16_t f = threadIdx.x; f < num_faces; f += blockThreads) { - if (s_active_f(f) && s_owned_f(f) && s_new_p_owned_f(f)) { + if ((s_active_f(f) && s_owned_f(f) && s_new_p_owned_f(f)) || + !s_active_f(f)) { s_owned_f.reset(f, true); } } @@ -2278,16 +2281,16 @@ void RXMeshDynamic::update_host() RXMESH_ERROR( "RXMeshDynamic::update_host error in updating host. m_num_edges " "{} does not match m_h_edge_prefix calculation {}", - this->m_num_faces, - m_h_face_prefix[m_num_patches]); + this->m_num_edges, + m_h_edge_prefix[m_num_patches]); } if (m_h_face_prefix[m_num_patches] != this->m_num_faces) { RXMESH_ERROR( "RXMeshDynamic::update_host error in updating host. m_num_faces " "{} does not match m_h_face_prefix calculation {}", - this->m_num_edges, - m_h_edge_prefix[m_num_patches]); + this->m_num_faces, + m_h_face_prefix[m_num_patches]); } const uint32_t patches_1_bytes = (m_num_patches + 1) * sizeof(uint32_t); diff --git a/include/rxmesh/rxmesh_dynamic.h b/include/rxmesh/rxmesh_dynamic.h index eea28355..ee205d88 100644 --- a/include/rxmesh/rxmesh_dynamic.h +++ b/include/rxmesh/rxmesh_dynamic.h @@ -146,7 +146,7 @@ __global__ static void slice_patches(Context context, __shared__ uint32_t s_new_patch_id; if (threadIdx.x == 0) { - // printf("\n*** slicing %u", pi.patch_id); + printf("\n*** slicing %u\n", pi.patch_id); s_new_patch_id = ::atomicAdd(context.m_num_patches, uint32_t(1)); assert(s_new_patch_id < context.m_max_num_patches); } diff --git a/tests/RXMesh_test/test_dynamic.cuh b/tests/RXMesh_test/test_dynamic.cuh index 355d1f30..2e67be99 100644 --- a/tests/RXMesh_test/test_dynamic.cuh +++ b/tests/RXMesh_test/test_dynamic.cuh @@ -126,8 +126,10 @@ TEST(RXMeshDynamic, RandomFlips) to_flip->reset(0, HOST); - const Config config = InteriorNotConflicting | InteriorConflicting | - OnRibbonNotConflicting | OnRibbonConflicting; + const Config config = /* InteriorNotConflicting | InteriorConflicting | + */ + OnRibbonNotConflicting /* | OnRibbonConflicting */ + ; rx.for_each_edge(HOST, [to_flip = *to_flip, config](const EdgeHandle eh) { @@ -172,7 +174,7 @@ TEST(RXMeshDynamic, RandomFlips) } - if (eh.patch_id() == 1) { + /* if (eh.patch_id() == 1) { if ((config & OnRibbonNotConflicting) == OnRibbonNotConflicting) { if (eh.local_id() == 383 || eh.local_id() == 324 || eh.local_id() == 355 || eh.local_id() == 340 || @@ -214,13 +216,13 @@ TEST(RXMeshDynamic, RandomFlips) to_flip(eh) = 1; } } - } + }*/ }); to_flip->move(HOST, DEVICE); - set_should_slice<<>>(rx.get_context()); + /* set_should_slice<<>>(rx.get_context()); rx.slice_patches(*coords, *to_flip); rx.cleanup(); CUDA_ERROR(cudaDeviceSynchronize()); @@ -230,7 +232,7 @@ TEST(RXMeshDynamic, RandomFlips) coords->move(DEVICE, HOST); to_flip->move(DEVICE, HOST); rx.update_polyscope(); - rx.get_polyscope_mesh()->updateVertexPositions(*coords); + rx.get_polyscope_mesh()->updateVertexPositions(*coords);*/ #if USE_POLYSCOPE @@ -243,7 +245,7 @@ TEST(RXMeshDynamic, RandomFlips) for (uint32_t p = 0; p < rx.get_num_patches(); ++p) { rx.render_patch(p)->setEnabled(false); } - // polyscope::show(); + polyscope::show(); #endif @@ -255,7 +257,7 @@ TEST(RXMeshDynamic, RandomFlips) LaunchBox launch_box; rx.prepare_launch_box( {}, launch_box, (void*)random_flips, false, true); - random_flips<<<<<1, launch_box.num_threads, launch_box.smem_bytes_dyn>>>( rx.get_context(), *coords, *to_flip); @@ -263,44 +265,45 @@ TEST(RXMeshDynamic, RandomFlips) rx.slice_patches(*coords, *to_flip); rx.cleanup(); - } - CUDA_ERROR(cudaDeviceSynchronize()); - rx.update_host(); + CUDA_ERROR(cudaDeviceSynchronize()); - coords->move(DEVICE, HOST); - to_flip->move(DEVICE, HOST); + rx.update_host(); + coords->move(DEVICE, HOST); + to_flip->move(DEVICE, HOST); - EXPECT_EQ(num_vertices, rx.get_num_vertices()); - EXPECT_EQ(num_edges, rx.get_num_edges()); - EXPECT_EQ(num_faces, rx.get_num_faces()); - EXPECT_TRUE(rx.validate()); + EXPECT_EQ(num_vertices, rx.get_num_vertices()); + EXPECT_EQ(num_edges, rx.get_num_edges()); + EXPECT_EQ(num_faces, rx.get_num_faces()); + + EXPECT_TRUE(rx.validate()); - /*rx.export_obj( - STRINGIFY(OUTPUT_DIR) "sphere3" + std::to_string(iter) + ".obj", - *coords); - rx.save(STRINGIFY(OUTPUT_DIR) "sphere3_patches" + - std::to_string(iter));*/ + /*rx.export_obj( + STRINGIFY(OUTPUT_DIR) "sphere3" + std::to_string(iter) + ".obj", + *coords); + rx.save(STRINGIFY(OUTPUT_DIR) "sphere3_patches" + + std::to_string(iter));*/ #if USE_POLYSCOPE - rx.update_polyscope(); - rx.polyscope_render_vertex_patch(); - rx.polyscope_render_edge_patch(); - rx.polyscope_render_face_patch(); + rx.update_polyscope(); + rx.polyscope_render_vertex_patch(); + rx.polyscope_render_edge_patch(); + rx.polyscope_render_face_patch(); - for (uint32_t p = 0; p < rx.get_num_patches(); ++p) { - rx.render_patch(p)->setEnabled(false); - } + for (uint32_t p = 0; p < rx.get_num_patches(); ++p) { + rx.render_patch(p)->setEnabled(false); + } - auto ps_mesh = rx.get_polyscope_mesh(); - ps_mesh->updateVertexPositions(*coords); - ps_mesh->addEdgeScalarQuantity("toFlip", *to_flip)->setMapRange({0, 2}); - ps_mesh->setEnabled(false); - // polyscope::show(); + auto ps_mesh = rx.get_polyscope_mesh(); + ps_mesh->updateVertexPositions(*coords); + ps_mesh->addEdgeScalarQuantity("toFlip", *to_flip)->setMapRange({0, 2}); + ps_mesh->setEnabled(false); + // polyscope::show(); #endif + } }