From 99c6b45a5c7aa75f53989a705903ae25be40d298 Mon Sep 17 00:00:00 2001
From: ahmed <ahmahmoud@ucdavis.edu>
Date: Tue, 15 Aug 2023 20:25:18 -0400
Subject: [PATCH] fix minor bug with the scheduler

---
 include/rxmesh/cavity_manager.cuh      |  8 +++++---
 include/rxmesh/cavity_manager_impl.cuh | 16 +++++++++++++---
 include/rxmesh/patch_scheduler.cuh     | 14 ++++++++++++++
 3 files changed, 32 insertions(+), 6 deletions(-)

diff --git a/include/rxmesh/cavity_manager.cuh b/include/rxmesh/cavity_manager.cuh
index 94e60828..2c15a11e 100644
--- a/include/rxmesh/cavity_manager.cuh
+++ b/include/rxmesh/cavity_manager.cuh
@@ -354,6 +354,8 @@ struct CavityManager
      */
     __device__ __inline__ void push();
 
+    __device__ __inline__ void push(const uint32_t pid);
+
     /**
      * @brief release the lock of this patch
      */
@@ -679,9 +681,9 @@ struct CavityManager
     bool*      m_s_should_slice;
     ShmemMutex m_s_patch_stash_mutex;
 
-    //LPPair*  m_s_table_q;
-    //LPPair*  m_s_table_stash_q;
-    //uint32_t m_s_table_q_size;    
+    // LPPair*  m_s_table_q;
+    // LPPair*  m_s_table_stash_q;
+    // uint32_t m_s_table_q_size;
 };
 
 }  // namespace rxmesh
diff --git a/include/rxmesh/cavity_manager_impl.cuh b/include/rxmesh/cavity_manager_impl.cuh
index 3f585210..c9ce1fb9 100644
--- a/include/rxmesh/cavity_manager_impl.cuh
+++ b/include/rxmesh/cavity_manager_impl.cuh
@@ -54,7 +54,7 @@ __device__ __inline__ CavityManager<blockThreads, cop>::CavityManager(
 
             if (!locked) {
                 // if we can not, we add it again to the queue
-                push();
+                push(s_patch_id);
 
                 // and signal other threads to also exit
                 s_patch_id = INVALID32;
@@ -235,8 +235,8 @@ CavityManager<blockThreads, cop>::alloc_shared_memory(
     // m_s_table_q = shrd_alloc.alloc<LPPair>(m_s_table_q_size);
 
     //__shared__ LPPair st_q[LPHashTable::stash_size];
-    //m_s_table_stash_q = st_q;
-    //fill_n<blockThreads>(
+    // m_s_table_stash_q = st_q;
+    // fill_n<blockThreads>(
     //    m_s_table_stash_q, uint16_t(LPHashTable::stash_size), LPPair());
 
     // lp stash
@@ -1215,6 +1215,16 @@ __device__ __inline__ void CavityManager<blockThreads, cop>::push()
     }
 }
 
+template <uint32_t blockThreads, CavityOp cop>
+__device__ __inline__ void CavityManager<blockThreads, cop>::push(
+    const uint32_t pid)
+{
+    if (threadIdx.x == 0) {
+        bool ret = m_context.m_patch_scheduler.push(pid);
+        assert(ret);
+    }
+}
+
 
 template <uint32_t blockThreads, CavityOp cop>
 __device__ __forceinline__ bool CavityManager<blockThreads, cop>::lock(
diff --git a/include/rxmesh/patch_scheduler.cuh b/include/rxmesh/patch_scheduler.cuh
index 8d1f35e9..302b1d5b 100644
--- a/include/rxmesh/patch_scheduler.cuh
+++ b/include/rxmesh/patch_scheduler.cuh
@@ -31,6 +31,7 @@ struct PatchScheduler
 #ifdef PROCESS_SINGLE_PATCH
         return true;
 #else
+        assert(pid != INVALID32);
         if (::atomicAdd(count, 1) < static_cast<int>(capacity)) {
             int pos = ::atomicAdd(back, 1) % capacity;
 
@@ -123,6 +124,19 @@ struct PatchScheduler
         CUDA_ERROR(cudaMalloc((void**)&list, sizeof(uint32_t) * capacity));
     }
 
+    __host__ void print_list() const
+    {
+        std::vector<uint32_t> h_list(capacity);
+        CUDA_ERROR(cudaMemcpy(h_list.data(),
+                              list,
+                              h_list.size() * sizeof(uint32_t),
+                              cudaMemcpyDeviceToHost));
+        for (uint32_t i = 0; i < h_list.size(); ++i) {
+            printf("\n list[%u]= %u", i, h_list[i]);
+        }
+    }
+
+
     /**
      * @brief free all the memories
      */