fix minor bug with the scheduler

owensgroup · Aug 16, 2023 · 99c6b45 · 99c6b45
1 parent deb9d44
commit 99c6b45
Show file tree

Hide file tree

Showing 3 changed files with 32 additions and 6 deletions.
diff --git 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
@@ -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
@@ -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
  */