Commit 5f4d7b33 authored by Tat Dat Duong's avatar Tat Dat Duong
Browse files

chore: improve query performance of b+tree

parent 0bc68f79

implementation/src/BLinkTree/BLinkNode.hpp

+4 −8
Original line number Diff line number Diff line
@@ -9,10 +9,10 @@ template <typename KeyType, typename ValueType, size_t Order> struct BLinkNode { 
  BLinkNode *volatile mSibling;



  KeyType mHighKey;

  uint8_t mHighKeyFlag;

  uint16_t mHighKeyFlag;



  uint8_t mLeaf;

  uint8_t mSize;

  uint16_t mLeaf;

  uint16_t mSize;



  int mWriteLock = false;
@@ -20,11 +20,7 @@ template <typename KeyType, typename ValueType, size_t Order> struct BLinkNode { 
  ValueType mValues[Order];

  BLinkNode *volatile mChildren[Order];



  __cuda_callable__ inline bool leaf() { return this->mLeaf; }



  __cuda_callable__ inline size_t size() { return this->mSize; }



  __cuda_callable__ inline size_t childSize() {

  __cuda_callable__ inline uint16_t childSize() {

    if (this->mLeaf == false && this->mSize > 0) {

      return this->mSize + 1;

    }

implementation/src/BLinkTree/BLinkTreeCuda.cu

+15 −16
Original line number Diff line number Diff line
@@ -35,7 +35,7 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Cuda> { 
      return node->mSibling;

    }



    size_t size = node->size();

    size_t size = node->mSize;

    size_t childSize = node->childSize();

    size_t targetIndex = upperBoundWarp(node->mKeys, size, key, warp);
@@ -61,7 +61,7 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Cuda> { 
  template <typename Warp>

  __device__ static inline BNode *findLeaf(BNode *root, KeyType key, Warp &warp) {

    BNode *cursor = root;

    while (cursor != nullptr && cursor->leaf() == false) {

    while (cursor != nullptr && cursor->mLeaf == false) {

      cursor = scan(cursor, key, warp);

    }
@@ -85,9 +85,8 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Cuda> { 


    sibling = warp.shfl(sibling, 0);



    size_t size = cursor->size();

    size_t size = cursor->mSize;

    size_t childSize = cursor->childSize();

    bool isLeaf = sibling->leaf();



    if (rank >= siblingStart) {

      if (rank < childSize) {
@@ -119,9 +118,9 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Cuda> { 
  __device__ static inline void insertIntoFreeNode(BNode *curr, KeyType insertKey, ValueType value,

                                                   BNode *insertNode, Warp &warp) {

    auto rank = warp.thread_rank();

    size_t keyIdx = upperBoundWarp(curr->mKeys, curr->size(), insertKey, warp);

    size_t keyIdx = upperBoundWarp(curr->mKeys, curr->mSize, insertKey, warp);



    bool isLeaf = curr->leaf();

    bool isLeaf = curr->mLeaf;

    size_t childSize = curr->childSize();



    if (isLeaf) {
@@ -261,8 +260,8 @@ public: 


    BNode *next = root;

    while (next) {

      bool isLeaf = next->leaf();

      size_t size = next->size();

      bool isLeaf = next->mLeaf;

      size_t size = next->mSize;



      KeyType nextKey = next->mKeys[rank];

      uint32_t ballot = warp.ballot(nextKey >= key);
@@ -302,7 +301,7 @@ public: 
        curr = scan(curr, key, warp);

      }



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

      if (curr != nullptr && curr->mSize + !curr->mLeaf >= Order) {

        // if node is full, we haven't propagated, restart the operation

        if (prev && curr != root && prev->mSibling == curr) {

          return false;
@@ -312,8 +311,8 @@ public: 
          return false;

        }



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

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

        size_t medianIdx = curr->mSize / 2;

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

        KeyType medianKey = curr->mKeys[medianIdx];



        if (curr == root) {
@@ -342,7 +341,7 @@ public: 
          }



          // verify, if the parent has not subsequently went full (esp. when )

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

          if (prev->mSize + !prev->mLeaf >= Order) {

            releaseLatchWarp(curr, latch, warp);

            releaseLatchWarp(prev, latch, warp);

            return false;
@@ -362,7 +361,7 @@ public: 
          prev = tmpPrev;

        }

      }

    } while (curr != nullptr && curr->leaf() == false);

    } while (curr != nullptr && curr->mLeaf == false);



    if (curr == nullptr) {

      return false;
@@ -384,7 +383,7 @@ public: 
    }



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

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

    if (curr->mSize + !curr->mLeaf >= Order) {

      releaseLatchWarp(curr, latch, warp);

      return false;

    }
@@ -411,8 +410,8 @@ public: 
    __threadfence();



    // remove item from keys and remove the respective node

    size_t keyIdx = lowerBoundWarp(cursor->mKeys, cursor->size(), key, warp);

    if (keyIdx >= cursor->size() || cursor->mKeys[keyIdx] != key) {

    size_t keyIdx = lowerBoundWarp(cursor->mKeys, cursor->mSize, key, warp);

    if (keyIdx >= cursor->mSize || cursor->mKeys[keyIdx] != key) {

      releaseLatchWarp(cursor, latch, warp);

      return true;

    }

implementation/src/BLinkTree/BLinkTreeHost.hpp

+16 −16
Original line number Diff line number Diff line
@@ -32,7 +32,7 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Host> { 
      return node->mSibling;

    }



    size_t size = node->size();

    size_t size = node->mSize;

    size_t childSize = node->childSize();

    size_t targetIndex = upperBound(node->mKeys, size, key);
@@ -56,7 +56,7 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Host> { 


  __cuda_callable__ static inline BNode *findLeaf(BNode *root, KeyType key) {

    BNode *cursor = root;

    while (cursor != nullptr && cursor->leaf() == false) {

    while (cursor != nullptr && cursor->mLeaf == false) {

      cursor = scan(cursor, key);

    }
@@ -75,8 +75,8 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Host> { 
      _Operations::appendChild(sibling, node, i - siblingStart);

    }



    for (size_t i = siblingStart; i < cursor->size(); ++i) {

      if (sibling->leaf()) {

    for (size_t i = siblingStart; i < cursor->mSize; ++i) {

      if (sibling->mLeaf) {

        _Operations::appendKey(sibling, cursor->mKeys[i], cursor->mValues[i]);

      } else {

        _Operations::appendKey(sibling, cursor->mKeys[i]);
@@ -93,9 +93,9 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Host> { 


  __cuda_callable__ static inline void insertIntoFreeNode(BNode *curr, KeyType insertKey,

                                                          ValueType value, BNode *insertNode) {

    size_t keyIdx = upperBound(curr->mKeys, curr->size(), insertKey);

    size_t keyIdx = upperBound(curr->mKeys, curr->mSize, insertKey);



    if (curr->leaf()) {

    if (curr->mLeaf) {

      _Operations::insertChild(curr, keyIdx, insertNode, curr->childSize());

      _Operations::insertKey(curr, keyIdx, insertKey, value);

    } else {
@@ -115,8 +115,8 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Host> { 
      BNode *child = curr->mChildren[childKey];

      _Operations::appendChild(leftNode, child, childKey);



      if (childKey < curr->size()) {

        if (leftNode->leaf()) {

      if (childKey < curr->mSize) {

        if (leftNode->mLeaf) {

          _Operations::appendKey(leftNode, curr->mKeys[childKey], curr->mValues[childKey]);

        } else {

          _Operations::appendKey(leftNode, curr->mKeys[childKey]);
@@ -146,8 +146,8 @@ public: 
    if (leaf == nullptr)

      return false;



    size_t it = lowerBound(leaf->mKeys, leaf->size(), key);

    if (it < leaf->size() && leaf->mKeys[it] == key) {

    size_t it = lowerBound(leaf->mKeys, leaf->mSize, key);

    if (it < leaf->mSize && leaf->mKeys[it] == key) {

      result = leaf->mValues[it];

      return true;

    }
@@ -166,14 +166,14 @@ public: 
        curr = scan(curr, key);

      }



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

      if (curr != nullptr && curr->mSize + !curr->mLeaf >= Order) {

        // if node is full, we haven't propagated, restart the operation

        if (curr->mSibling == prev && curr != root) {

          return false;

        }



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

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

        size_t medianIdx = curr->mSize / 2;

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



        KeyType medianKey = curr->mKeys[medianIdx];

        BNode *split = nullptr;
@@ -188,7 +188,7 @@ public: 
          prev = tmpPrev;

        }

      }

    } while (curr != nullptr && curr->leaf() == false);

    } while (curr != nullptr && curr->mLeaf == false);



    if (curr == nullptr) {

      return false;
@@ -214,8 +214,8 @@ public: 
    }



    // remove item from keys and remove the respective node

    size_t keyIdx = lowerBound(cursor->mKeys, cursor->size(), key);

    if (keyIdx == cursor->size() || cursor->mKeys[keyIdx] != key) {

    size_t keyIdx = lowerBound(cursor->mKeys, cursor->mSize, key);

    if (keyIdx == cursor->mSize || cursor->mKeys[keyIdx] != key) {

      latch.release(cursor);

      return true;

    }

implementation/src/BLinkTree/BNodeOperations/BNodeOperationsCuda.cu

+8 −8
Original line number Diff line number Diff line
@@ -32,7 +32,7 @@ struct BLinkOperations<KeyType, ValueType, Order, TNL::Devices::Cuda> { 


  template <typename Warp>

  __device__ static void init(Warp &warp, Node *volatile node, Node *volatile cursor) {

    BLinkOperations::init(warp, node, cursor->leaf(), cursor->mSibling, cursor->mHighKeyFlag);

    BLinkOperations::init(warp, node, cursor->mLeaf, cursor->mSibling, cursor->mHighKeyFlag);

    if (warp.thread_rank() == 0) {

      node->mHighKey = cursor->mHighKey;

    }
@@ -41,7 +41,7 @@ struct BLinkOperations<KeyType, ValueType, Order, TNL::Devices::Cuda> { 
  template <typename Warp>

  __device__ static void insertKey(Node *volatile node, size_t index, KeyType key, ValueType value,

                                   Warp &warp) {

    TNL_ASSERT_TRUE(node->leaf(), "Inserting value is only supported in leaves");

    TNL_ASSERT_TRUE(node->mLeaf, "Inserting value is only supported in leaves");



    auto rank = warp.thread_rank();

    if (rank < Order) {
@@ -70,7 +70,7 @@ struct BLinkOperations<KeyType, ValueType, Order, TNL::Devices::Cuda> { 


  template <typename Warp>

  __device__ static void insertKey(Node *volatile node, size_t index, KeyType key, Warp &warp) {

    TNL_ASSERT_FALSE(node->leaf(), "Inserting with no value in an internal node");

    TNL_ASSERT_FALSE(node->mLeaf, "Inserting with no value in an internal node");



    auto rank = warp.thread_rank();

    if (rank < Order) {
@@ -94,23 +94,23 @@ struct BLinkOperations<KeyType, ValueType, Order, TNL::Devices::Cuda> { 


  template <typename Warp>

  __device__ static void appendKey(Node *volatile node, KeyType key, ValueType value, Warp &warp) {

    TNL_ASSERT_TRUE(node->leaf(), "Inserting value is only supported in leaves");

    TNL_ASSERT_TRUE(node->mLeaf, "Inserting value is only supported in leaves");



    auto rank = warp.thread_rank();

    if (rank == 0) {

      node->mKeys[node->size()] = key;

      node->mValues[node->size()] = value;

      node->mKeys[node->mSize] = key;

      node->mValues[node->mSize] = value;

      node->mSize += 1;

    }

  }



  template <typename Warp>

  __device__ static void appendKey(Node *volatile node, KeyType key, Warp &warp) {

    TNL_ASSERT_FALSE(node->leaf(), "Inserting with no value in an internal node");

    TNL_ASSERT_FALSE(node->mLeaf, "Inserting with no value in an internal node");

    auto rank = warp.thread_rank();



    if (rank == 0) {

      node->mKeys[node->size()] = key;

      node->mKeys[node->mSize] = key;

      node->mSize += 1;

    }

  }

implementation/src/BLinkTree/BNodeOperations/BNodeOperationsHost.hpp

+13 −13
Original line number Diff line number Diff line
@@ -31,35 +31,35 @@ struct BLinkOperations<KeyType, ValueType, Order, TNL::Devices::Host> { 
  }



  static void insertKey(Node *volatile node, size_t index, KeyType key, ValueType value) {

    TNL_ASSERT_TRUE(node->leaf(), "Inserting value is only supported in leaves");

    insert(node->mKeys, node->size(), index, key);

    insert(node->mValues, node->size(), index, value);

    TNL_ASSERT_TRUE(node->mLeaf, "Inserting value is only supported in leaves");

    insert(node->mKeys, node->mSize, index, key);

    insert(node->mValues, node->mSize, index, value);

    node->mSize += 1;

  }



  static void insertKey(Node *volatile node, size_t index, KeyType key) {

    TNL_ASSERT_FALSE(node->leaf(), "Inserting with no value in an internal node");

    insert(node->mKeys, node->size(), index, key);

    TNL_ASSERT_FALSE(node->mLeaf, "Inserting with no value in an internal node");

    insert(node->mKeys, node->mSize, index, key);

    node->mSize += 1;

  }



  static void appendKey(Node *volatile node, KeyType key, ValueType value) {

    TNL_ASSERT_TRUE(node->leaf(), "Inserting value is only supported in leaves");

    node->mKeys[node->size()] = key;

    node->mValues[node->size()] = value;

    TNL_ASSERT_TRUE(node->mLeaf, "Inserting value is only supported in leaves");

    node->mKeys[node->mSize] = key;

    node->mValues[node->mSize] = value;

    node->mSize += 1;

  }



  static void appendKey(Node *volatile node, KeyType key) {

    TNL_ASSERT_FALSE(node->leaf(), "Inserting with no value in an internal node");

    node->mKeys[node->size()] = key;

    TNL_ASSERT_FALSE(node->mLeaf, "Inserting with no value in an internal node");

    node->mKeys[node->mSize] = key;

    node->mSize += 1;

  }



  static void removeKey(Node *volatile node, size_t index) {

    if (node->size() > 0) {

      remove(node->mKeys, node->size(), index);

      remove(node->mValues, node->size(), index);

    if (node->mSize > 0) {

      remove(node->mKeys, node->mSize, index);

      remove(node->mValues, node->mSize, index);

      node->mSize -= 1;

    }

  }