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

chore: add logging

parent 9afe73c1

src/btree.hpp

+2 −0
Original line number Diff line number Diff line
@@ -86,6 +86,7 @@ __cuda_callable__ void Insert(BNode *root, KeyType key, ValueType value, 
  BNode *cursor = FindLeaf<Device>(root, key, &latch);



  while (cursor != nullptr) {

    printf("Continuing with inserting %lu\n", toInsertKey);

    size_t keyIdx = upperBound(cursor->mKeys, cursor->size(), toInsertKey);



    // check if inserting won't cause rebalance, if so, early exit
@@ -169,6 +170,7 @@ __cuda_callable__ void Insert(BNode *root, KeyType key, ValueType value, 
    }



    // if we're splitting a root node, we need to create a new root directly

    // TODO: in this scenario, we need to restart the search process

    if (cursor == root) {

      BNode *newLeft = allocator.allocate();

      newLeft->init(cursor);

src/latch.hpp

+6 −1
Original line number Diff line number Diff line
@@ -34,14 +34,19 @@ template <> struct Latch<Devices::Cuda> { 
  __device__ void acquire(BNode *node) {

    bool locked = true;



    printf("Acquiring latch for %lu\n", (unsigned long)node);

    while (locked) {

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

        printf("Latch acquired for %lu\n", (unsigned long)node);

        break;

      }

    }

  }



  __device__ void release(BNode *node) { atomicAnd(&node->mWriteLock, false); }

  __device__ void release(BNode *node) {

    printf("Releasing latch for %lu\n", (unsigned long)node);

    atomicAnd(&node->mWriteLock, false);

  }



#else  // HAVE_CUDA

  void acquire(BNode *node) {

test/device.cu

+3 −5
Original line number Diff line number Diff line
@@ -35,8 +35,8 @@ __device__ void bulkInsertWarp(bool &toInsert, KeyType key, ValueType value, 
  while ((queue = warp.ballot(toInsert))) {

    toInsert = false;



    BTree::Insert<Devices::Cuda>(root, key, value, allocator, latch);

    printf("Inserting %lu %lu\n", key, value);

    BTree::Insert<Devices::Cuda>(root, key, value, allocator, latch);

  }

};
@@ -57,7 +57,7 @@ bulkInsert(Containers::ArrayView<std::add_const_t<KeyType>, Devices::Cuda> keys, 
    value = values.getElement(threadId);

  }



  bulkInsertWarp(toInsert, key, value, allocator, latch);

  BTree::Insert<Devices::Cuda>(root, key, value, allocator, latch);

}



__global__ void
@@ -75,9 +75,7 @@ bulkQuery(Containers::ArrayView<std::add_const_t<KeyType>, Devices::Cuda> keys, 
  }

}



__global__ void devicePrintTree() {

  printTree(root);

}

__global__ void devicePrintTree() { printTree(root); }



int main(void) {