Fix share memory allocation

owensgroup · Jan 7, 2025 · a82a46e · a82a46e
1 parent 3c19fb4
commit a82a46e
Show file tree

Hide file tree

Showing 11 changed files with 119 additions and 72 deletions.
diff --git a/apps/Remesh/collapse.cuh b/apps/Remesh/collapse.cuh
@@ -31,7 +31,7 @@ __global__ static void __launch_bounds__(blockThreads)
 
     // a bitmask that indicates which edge we want to flip
     // we also use it to mark updated edges (for edge_status)
-    Bitmask edge_mask(cavity.patch_info().edges_capacity[0], shrd_alloc);
+    Bitmask edge_mask(cavity.patch_info().edges_capacity, shrd_alloc);
     edge_mask.reset(block);
 
     uint32_t shmem_before = shrd_alloc.get_allocated_size_bytes();

diff --git a/apps/Remesh/flip.cuh b/apps/Remesh/flip.cuh
@@ -54,7 +54,7 @@ __global__ static void __launch_bounds__(blockThreads)
 
     // a bitmask that indicates which edge we want to flip
     // we also used it to mark the new edges
-    Bitmask edge_mask(cavity.patch_info().edges_capacity[0], shrd_alloc);
+    Bitmask edge_mask(cavity.patch_info().edges_capacity, shrd_alloc);
     edge_mask.reset(block);
 
     uint32_t shmem_before = shrd_alloc.get_allocated_size_bytes();

diff --git a/apps/Remesh/remesh_rxmesh.cuh b/apps/Remesh/remesh_rxmesh.cuh
@@ -389,10 +389,9 @@ inline void remesh_rxmesh(rxmesh::RXMeshDynamic& rx)
                 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("Flip Total Time {} (ms)", timers.elapsed_millis("FlipTotal"));
     RXMESH_INFO("Smooth Total Time {} (ms)",
-                timers.elapsed_millis("SmoothTotal"));   
+                timers.elapsed_millis("SmoothTotal"));
 
     report.add_member("split_time_ms", timers.elapsed_millis("SplitTotal"));
     report.add_member("collapse_time_ms",

diff --git a/apps/Remesh/smoothing.cuh b/apps/Remesh/smoothing.cuh
@@ -23,7 +23,7 @@ __global__ static void __launch_bounds__(blockThreads)
             return;
         }
 
-        if (v_boundary(v_id)) {
+        if (v_boundary(v_id) || iter.size() == 0) {
             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);

diff --git a/apps/SECHistogram/sec_kernels.cuh b/apps/SECHistogram/sec_kernels.cuh
@@ -24,7 +24,7 @@ __global__ static void sec(rxmesh::Context            context,
     // we first use this mask to set the edge we want to collapse (and then
     // filter them). Then after cavity.prologue, we reuse this bitmask to mark
     // the newly added edges
-    Bitmask edge_mask(cavity.patch_info().edges_capacity[0], shrd_alloc);
+    Bitmask edge_mask(cavity.patch_info().edges_capacity, shrd_alloc);
     edge_mask.reset(block);
 
     // we use this bitmask to mark the other end of to-be-collapse edge during

diff --git a/apps/SECPriority/secp_kernels.cuh b/apps/SECPriority/secp_kernels.cuh
@@ -24,7 +24,7 @@ __global__ static void secp(rxmesh::Context             context,
     // we first use this mask to set the edge we want to collapse (and then
     // filter them). Then after cavity.prologue, we reuse this bitmask to mark
     // the newly added edges
-    Bitmask edge_mask(cavity.patch_info().edges_capacity[0], shrd_alloc);
+    Bitmask edge_mask(cavity.patch_info().edges_capacity, shrd_alloc);
     edge_mask.reset(block);
 
     // we use this bitmask to mark the other end of to-be-collapse edge during

diff --git a/include/rxmesh/cavity_manager2.cuh b/include/rxmesh/cavity_manager2.cuh
@@ -1014,6 +1014,7 @@ struct CavityManager2
 
     // store the boundary edges of all cavities in compact format (similar to
     // CSR for sparse matrices using m_s_cavity_size_prefix but no value ptr)
+    // overlap with m_s_boudary_edges_cavity_id
     uint16_t* m_s_cavity_boundary_edges;
 
     // patch stash stored in shared memory