Commit c11dfb6b authored by Tat Dat Duong's avatar Tat Dat Duong
Browse files

chore: replace fences in latch

parent 5ece8de7

src/BNodeLatch/BNodeLatchCuda.cu

+3 −6
Original line number Diff line number Diff line
@@ -31,24 +31,21 @@ struct Latch<BNode<KeyType, ValueType, Order, KeyInf>, Devices::Cuda> { 
  __device__ bool attempt(BTreeNode *volatile node) {

    if (atomicCAS(&node->mWriteLock, false, true) == false) {

      DBG_LATCH_ATTEMPT(node);

      __threadfence();

      return true;

    }

    __threadfence();

    return false;

  }



  __device__ void release(BTreeNode *volatile node) {

    __threadfence();

    DBG_LATCH_RELEASE(node);

    atomicAnd(&node->mWriteLock, false);

  }



#else  // HAVE_CUDA

  void acquire(volatile BTreeNode *node) { throw TNL::Exceptions::CudaSupportMissing(); }

  void acquire(BTreeNode *volatile node) { throw TNL::Exceptions::CudaSupportMissing(); }



  bool attempt(volatile BTreeNode *node) { throw TNL::Exceptions::CudaSupportMissing(); }

  bool attempt(BTreeNode *volatile node) { throw TNL::Exceptions::CudaSupportMissing(); }



  void release(volatile BTreeNode *node) { throw TNL::Exceptions::CudaSupportMissing(); }

  void release(BTreeNode *volatile node) { throw TNL::Exceptions::CudaSupportMissing(); }

#endif // HAVE_CUDA

};
 
 No newline at end of file

src/BTree/BTreeCuda.cu

+19 −1
Original line number Diff line number Diff line
@@ -247,6 +247,8 @@ public: 
    BTreeNode *prev = nullptr, *curr = nullptr;



    do {

      __threadfence();



      BTreeNode *tmpPrev = prev;

      if (curr == nullptr) {

        curr = root;
@@ -255,7 +257,6 @@ public: 
        curr = scan(curr, key, warp);

      }



      __threadfence();

      if (curr->mWriteLock) {

        return false;

      }
@@ -270,6 +271,8 @@ public: 
          return false;

        }



        __threadfence();



        size_t medianIdx = curr->size() / 2;

        size_t splitIdx = medianIdx + (!curr->leaf());

        KeyType medianKey = curr->mKeys[medianIdx];
@@ -277,6 +280,8 @@ public: 
        if (curr == root) {

          BTreeNode *split = splitNode(curr, alloc, splitIdx, medianIdx, warp);

          increaseTreeHeight(curr, medianKey, split, alloc, warp);



          __threadfence();

          releaseLatchWarp(curr, latch, warp);

        } else {

          // latch the parent too
@@ -285,6 +290,8 @@ public: 
            return false;

          }



          __threadfence();



          // verify, if the parent is actually our parent

          if (verifyParent(curr, prev, warp) == false) {

            releaseLatchWarp(curr, latch, warp);
@@ -293,9 +300,13 @@ public: 
          }



          BTreeNode *split = splitNode(curr, alloc, splitIdx, medianIdx, warp);



          __threadfence();

          releaseLatchWarp(curr, latch, warp);



          insertIntoFreeNode(prev, medianKey, 0, split, warp);



          __threadfence();

          releaseLatchWarp(prev, latch, warp);

        }
@@ -313,6 +324,8 @@ public: 
      return false;

    }



    __threadfence();



    // here we don't check if the node has subsequently become full

    if (curr->size() + !curr->leaf() >= Order) {

      releaseLatchWarp(curr, latch, warp);
@@ -321,6 +334,7 @@ public: 


    insertIntoFreeNode(curr, key, value, nullptr, warp);



    __threadfence();

    releaseLatchWarp(curr, latch, warp);

    return true;

  }
@@ -350,6 +364,8 @@ public: 
      return false;

    }



    __threadfence();



    size_t size = cursor->size();

    uint32_t ballot = warp.ballot(cursor->mKeys[rank] >= key);

    size_t keyIdx = __ffs(ballot) - 1;
@@ -366,6 +382,8 @@ public: 
      DBG_DELETE(cursor, keyIdx);

    }



    __threadfence();



    releaseLatchWarp(cursor, latch, warp);

    return true;

  }