From d13531bf9a7da236b3cb0bd2c581296d4d200da5 Mon Sep 17 00:00:00 2001
From: Ahmed Mahmoud <ahmed.mahmoud@autodesk.com>
Date: Fri, 3 Nov 2023 18:04:30 -0400
Subject: [PATCH] store_hashtable while doing cavity fill-in

---
 include/rxmesh/cavity_manager.cuh      |  6 ++++
 include/rxmesh/cavity_manager_impl.cuh | 39 +++++++++++++++++---------
 2 files changed, 31 insertions(+), 14 deletions(-)

diff --git a/include/rxmesh/cavity_manager.cuh b/include/rxmesh/cavity_manager.cuh
index ed9d0ee2..32d7dc84 100644
--- 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
index 41f6956b..394b4834 100644
--- 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]) {