store_hashtable while doing cavity fill-in

owensgroup · Nov 3, 2023 · d13531b · d13531b
1 parent 8222135
commit d13531b
Show file tree

Hide file tree

Showing 2 changed files with 31 additions and 14 deletions.
diff --git a/include/rxmesh/cavity_manager.cuh b/include/rxmesh/cavity_manager.cuh
@@ -251,6 +251,12 @@ struct CavityManager
  __device__ __inline__ void load_hashtable(
  cooperative_groups::thread_block& block);
 
+ /**
+ * @brief store all hashtable from shared memory to global memory
+ */
+ __device__ __inline__ void store_hashtable(
+ cooperative_groups::thread_block& block);
+
  /**
  * @brief propage the cavity ID from the seeds to indicent/adjacent elements
  * based on CavityOp

diff --git a/include/rxmesh/cavity_manager_impl.cuh b/include/rxmesh/cavity_manager_impl.cuh
@@ -541,6 +541,11 @@ __device__ __inline__ bool CavityManager<blockThreads, cop>::prologue(
  m_s_fill_in_f.reset(block);
  block.sync();
 
+ // store hashtable now so we could re-use the shared memory
+ // overlap this memcpy with the user operations that mostly happens in
+ // shared memory
+ store_hashtable(block);
+
  return true;
 }
 
@@ -1441,6 +1446,26 @@ __device__ __inline__ void CavityManager<blockThreads, cop>::load_hashtable(
  m_s_table_f, true, m_s_table_stash_f);
 }
 
+template <uint32_t blockThreads, CavityOp cop>
+__device__ __inline__ void CavityManager<blockThreads, cop>::store_hashtable(
+ cooperative_groups::thread_block& block)
+{
+
+ // store hashtable
+ m_patch_info.lp_v.write_to_global_memory<blockThreads>(m_s_table_v,
+ m_s_table_stash_v);
+ m_patch_info.lp_e.write_to_global_memory<blockThreads>(m_s_table_e,
+ m_s_table_stash_e);
+ m_patch_info.lp_f.write_to_global_memory<blockThreads>(m_s_table_f,
+ m_s_table_stash_f);
+
+ // patch stash
+ for (uint32_t i = threadIdx.x; i < PatchStash::stash_size;
+ i += blockThreads) {
+ m_patch_info.patch_stash.m_stash[i] = m_s_patch_stash.m_stash[i];
+ }
+}
+
 
 template <uint32_t blockThreads, CavityOp cop>
 __device__ __inline__ void CavityManager<blockThreads, cop>::push()
@@ -3181,20 +3206,6 @@ __device__ __inline__ void CavityManager<blockThreads, cop>::epilogue(
  detail::store<blockThreads>(m_s_active_mask_f.m_bitmask,
  DIVIDE_UP(m_s_num_faces[0], 32),
  m_patch_info.active_mask_f);
-
- // store hashtable
- m_patch_info.lp_v.write_to_global_memory<blockThreads>(
- m_s_table_v, m_s_table_stash_v);
- m_patch_info.lp_e.write_to_global_memory<blockThreads>(
- m_s_table_e, m_s_table_stash_e);
- m_patch_info.lp_f.write_to_global_memory<blockThreads>(
- m_s_table_f, m_s_table_stash_f);
-
- // patch stash
- for (uint32_t i = threadIdx.x; i < PatchStash::stash_size;
- i += blockThreads) {
- m_patch_info.patch_stash.m_stash[i] = m_s_patch_stash.m_stash[i];
- }
  }
 
  if (m_s_should_slice[0]) {