fix a race condition in add_element

include/rxmesh/cavity_manager.cuh

@@ -332,7 +332,7 @@ struct CavityManager
  * these spots in hashtable calibration
  */
  __device__ __inline__ uint16_t add_element(Bitmask& active_bitmask,
- uint16_t* num_elements,
+ uint32_t* num_elements,
  const uint16_t capacity,
  const Bitmask& in_cavity,
  const Bitmask& owned,
@@ -641,7 +641,9 @@ struct CavityManager
 
  // store the number of elements. we use pointers since the number of mesh
  // elements could change
- uint16_t *m_s_num_vertices, *m_s_num_edges, *m_s_num_faces;
+ // while they should be represented using uint16_t, we use uint32_t since we
+ // need to use atomicMax which is only supported 32 and 64 bit
+ uint32_t *m_s_num_vertices, *m_s_num_edges, *m_s_num_faces;
 
  // 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)

include/rxmesh/cavity_manager_impl.cuh

@@ -10,7 +10,7 @@ __device__ __inline__ CavityManager<blockThreads, cop>::CavityManager(
  __shared__ uint32_t s_patch_id;
 
 
- __shared__ uint16_t counts[3];
+ __shared__ uint32_t counts[3];
  m_s_num_vertices = counts + 0;
  m_s_num_edges = counts + 1;
  m_s_num_faces = counts + 2;
@@ -1050,35 +1050,35 @@ __device__ __inline__ FaceHandle CavityManager<blockThreads, cop>::add_face(
 template <uint32_t blockThreads, CavityOp cop>
 __device__ __inline__ uint16_t CavityManager<blockThreads, cop>::add_element(
  Bitmask& active_bitmask,
- uint16_t* num_elements,
+ uint32_t* num_elements,
  const uint16_t capacity,
  const Bitmask& in_cavity,
  const Bitmask& owned,
  bool avoid_not_owned_in_cavity)
 {
- // for (uint16_t i = 0; i < capacity; ++i) {
- const uint16_t count = num_elements[0];
- for (uint16_t i = 0; i < count; ++i) {
+ assert(capacity == in_cavity.size());
+
+ uint16_t found = INVALID16;
+ for (uint16_t i = 0; i < capacity; ++i) {
  if (!avoid_not_owned_in_cavity && in_cavity(i) && !owned(i)) {
  continue;
  }
  if (avoid_not_owned_in_cavity && in_cavity(i) /*&& !owned(i)*/) {
  continue;
  }
  if (active_bitmask.try_set(i)) {
- return i;
+ found = i;
+ break;
  }
  }
 
- // TODO there is a race condition here w.r.t active_bitmask
- uint16_t ret = atomicAdd(num_elements, uint16_t(1));
- if (ret >= capacity) {
- return INVALID16;
- } else {
- assert(ret < active_bitmask.size());
- active_bitmask.set(ret, true);
- return ret;
+
+ if (found != INVALID16) {
+ ::atomicMax(num_elements, found + 1);
+ assert(found < active_bitmask.size());
  }
+
+ return found;
 }