debugging

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<float, blockThreads>);
  delaunay_edge_flip<float, blockThreads>
- <<<launch_box.blocks,
- launch_box.num_threads,
- launch_box.smem_bytes_dyn>>>(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;

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