chore: naming refactor (4cf5b6b7) · Commits · TNL / GPUBTree

implementation/src/BLinkTree/BNodeOperations/BNodeOperationsCuda.cu→implementation/src/BLinkTree/BLinkOperations/BLinkOperationsCuda.cu

+0 −0

File moved.

View file

implementation/src/BLinkTree/BNodeOperations/BNodeOperationsHost.hpp→implementation/src/BLinkTree/BLinkOperations/BLinkOperationsHost.hpp

+0 −0

File moved.

View file

implementation/src/BLinkTree/BNodeOperations/Default.hpp→implementation/src/BLinkTree/BLinkOperations/Default.hpp

+2 −2

Original line number	Diff line number	Diff line
		@@ -4,7 +4,7 @@
		template <typename KeyType, typename ValueType, size_t Order, typename Device>
		struct BLinkOperations {};

		#include "./BNodeOperationsHost.hpp"
		#include "./BLinkOperationsHost.hpp"
		#ifdef HAVE_CUDA
		#include "./BNodeOperationsCuda.cu"
		#include "./BLinkOperationsCuda.cu"
		#endif
		No newline at end of file

implementation/src/BLinkTree/BLinkTreeCuda.cu

+14 −14

Original line number	Diff line number	Diff line
		@@ -15,7 +15,7 @@
		#include "../Utils/Print.hpp"

		#include "../BNodeLatch/Default.hpp"
		#include "./BNodeOperations/Default.hpp"
		#include "./BLinkOperations/Default.hpp"
		#include "TNL/Assert.h"

		template <typename KeyType, typename ValueType, size_t Order, typename Device> class BLinkTree;
		@@ -26,7 +26,7 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Cuda> {
		using BNode = BLinkNode<KeyType, ValueType, Order>;
		using Latch = BNodeLatch<BNode, Device>;
		using Allocator = BumpAllocator<BNode, Device>;
		using _Operations = BLinkOperations<KeyType, ValueType, Order, Device>;
		using Operations = BLinkOperations<KeyType, ValueType, Order, Device>;

		template <typename Warp>
		__device__ static inline BNode scan(BNode node, KeyType key, Warp &warp) {
		@@ -56,7 +56,7 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Cuda> {
		BNode *sibling = (rank == 0) ? alloc.allocate() : nullptr;
		sibling = warp.shfl(sibling, 0);

		_Operations::init(warp, sibling, cursor);
		Operations::init(warp, sibling, cursor);
		warp.sync();

		uint16_t size = cursor->mSize;
		@@ -96,8 +96,8 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Cuda> {
		if (rank == 0) {
		DBG_INSERT_INTERNAL(curr, insertKey, insertNode, keyIdx, (keyIdx + 1));
		}
		_Operations::insertChild(curr, keyIdx + 1, insertNode, childSize, warp);
		_Operations::insertKey(curr, keyIdx, insertKey, warp);
		Operations::insertChild(curr, keyIdx + 1, insertNode, childSize, warp);
		Operations::insertKey(curr, keyIdx, insertKey, warp);

		warp.sync();
		}
		@@ -112,8 +112,8 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Cuda> {
		if (rank == 0) {
		DBG_INSERT_LEAF(curr, insertKey, keyIdx);
		}
		_Operations::insertChild(curr, keyIdx, insertNode, childSize, warp);
		_Operations::insertKey(curr, keyIdx, insertKey, value, warp);
		Operations::insertChild(curr, keyIdx, insertNode, childSize, warp);
		Operations::insertKey(curr, keyIdx, insertKey, value, warp);

		warp.sync();
		}
		@@ -130,7 +130,7 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Cuda> {
		}
		rootNode = warp.shfl(rootNode, 0);

		_Operations::init(warp, rootNode, false, nullptr, false);
		Operations::init(warp, rootNode, false, nullptr, false);
		warp.sync();

		if (rank == 0) {
		@@ -209,12 +209,12 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Cuda> {
		}

		public:
		using Node = BNode;
		using Operations = _Operations;
		using NodeType = BNode;
		using OperationsType = Operations;

		template <typename Warp> __device__ static inline BNode *init(Allocator &alloc, Warp &warp) {
		BNode *root = alloc.allocate();
		_Operations::init(warp, root, true, nullptr, false);
		Operations::init(warp, root, true, nullptr, false);

		DBG_INIT(root, Order);
		return root;
		@@ -379,8 +379,8 @@ public:
		return true;
		}

		_Operations::removeChild(cursor, keyIdx, cursor->childSize(), warp);
		_Operations::removeKey(cursor, keyIdx, warp);
		Operations::removeChild(cursor, keyIdx, cursor->childSize(), warp);
		Operations::removeKey(cursor, keyIdx, warp);
		warp.sync();

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

implementation/src/BLinkTree/BLinkTreeHost.hpp

+52 −53

Original line number	Diff line number	Diff line
		@@ -15,19 +15,19 @@

		#include "../BNodeLatch/Default.hpp"
		#include "./BLinkNode.hpp"
		#include "./BNodeOperations/Default.hpp"
		#include "./BLinkOperations/Default.hpp"

		template <typename KeyType, typename ValueType, size_t Order, typename Device> class BLinkTree;

		template <typename KeyType, typename ValueType, size_t Order>
		class BLinkTree<KeyType, ValueType, Order, Devices::Host> {
		using Device = Devices::Host;
		using BNode = BLinkNode<KeyType, ValueType, Order>;
		using Latch = BNodeLatch<BNode, Device>;
		using Allocator = BumpAllocator<BNode, Device>;
		using _Operations = BLinkOperations<KeyType, ValueType, Order, Device>;
		using Node = BLinkNode<KeyType, ValueType, Order>;
		using Latch = BNodeLatch<Node, Device>;
		using Allocator = BumpAllocator<Node, Device>;
		using Operations = BLinkOperations<KeyType, ValueType, Order, Device>;

		static inline BNode scan(BNode node, KeyType key) {
		static inline Node scan(Node node, KeyType key) {
		if (node->mHighKeyFlag && !(key < node->mHighKey)) {
		return node->mSibling;
		}
		@@ -40,12 +40,12 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Host> {
		return nullptr;
		}

		BNode *result = node->mChildren[targetIndex];
		Node *result = node->mChildren[targetIndex];
		return result;
		}

		static inline BNode moveSide(BNode cursor, KeyType key) {
		BNode tmp, result = cursor;
		static inline Node moveSide(Node cursor, KeyType key) {
		Node tmp, result = cursor;
		while (result != nullptr && result->mSibling != nullptr &&
		(tmp = scan(result, key)) == result->mSibling) {
		result = tmp;
		@@ -54,8 +54,8 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Host> {
		return result;
		}

		static inline BNode findLeaf(BNode root, KeyType key) {
		BNode *cursor = root;
		static inline Node findLeaf(Node root, KeyType key) {
		Node *cursor = root;
		while (cursor != nullptr && cursor->mLeaf == false) {
		cursor = scan(cursor, key);
		}
		@@ -64,22 +64,21 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Host> {
		return cursor;
		}

		static BNode splitNode(BNode cursor, Allocator &alloc, size_t siblingStart,
		size_t cursorCount) {
		static Node splitNode(Node cursor, Allocator &alloc, size_t siblingStart, size_t cursorCount) {

		BNode *sibling = alloc.allocate();
		_Operations::init(sibling, cursor);
		Node *sibling = alloc.allocate();
		Operations::init(sibling, cursor);

		for (size_t i = siblingStart; i < cursor->childSize(); ++i) {
		BNode *node = cursor->mChildren[i];
		_Operations::appendChild(sibling, node, i - siblingStart);
		Node *node = cursor->mChildren[i];
		Operations::appendChild(sibling, node, i - siblingStart);
		}

		for (size_t i = siblingStart; i < cursor->mSize; ++i) {
		if (sibling->mLeaf) {
		_Operations::appendKey(sibling, cursor->mKeys[i], cursor->mValues[i]);
		Operations::appendKey(sibling, cursor->mKeys[i], cursor->mValues[i]);
		} else {
		_Operations::appendKey(sibling, cursor->mKeys[i]);
		Operations::appendKey(sibling, cursor->mKeys[i]);
		}
		}

		@@ -91,10 +90,10 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Host> {
		return sibling;
		}

		static inline void insertIntoFreeInternal(BNode curr, KeyType insertKey, BNode insertNode) {
		static inline void insertIntoFreeInternal(Node curr, KeyType insertKey, Node insertNode) {
		size_t keyIdx = upperBound(curr->mKeys, curr->mSize, insertKey);
		_Operations::insertChild(curr, keyIdx + 1, insertNode, curr->childSize());
		_Operations::insertKey(curr, keyIdx, insertKey);
		Operations::insertChild(curr, keyIdx + 1, insertNode, curr->childSize());
		Operations::insertKey(curr, keyIdx, insertKey);

		if (curr->mLeaf) {

		@@ -102,52 +101,52 @@ class BLinkTree<KeyType, ValueType, Order, Devices::Host> {
		}
		}

		static inline void insertIntoFreeLeaf(BNode *curr, KeyType insertKey, ValueType value,
		BNode *insertNode) {
		static inline void insertIntoFreeLeaf(Node *curr, KeyType insertKey, ValueType value,
		Node *insertNode) {
		size_t keyIdx = upperBound(curr->mKeys, curr->mSize, insertKey);
		_Operations::insertChild(curr, keyIdx, insertNode, curr->childSize());
		_Operations::insertKey(curr, keyIdx, insertKey, value);
		Operations::insertChild(curr, keyIdx, insertNode, curr->childSize());
		Operations::insertKey(curr, keyIdx, insertKey, value);
		}

		static inline void increaseTreeHeight(BNode curr, KeyType insertKey, BNode insertNode,
		static inline void increaseTreeHeight(Node curr, KeyType insertKey, Node insertNode,
		Allocator &alloc) {

		BNode *leftNode = alloc.allocate();
		Node *leftNode = alloc.allocate();

		// copy the entire node to the element
		_Operations::init(leftNode, curr);
		Operations::init(leftNode, curr);
		for (size_t childKey = 0; childKey < curr->childSize(); ++childKey) {
		BNode *child = curr->mChildren[childKey];
		_Operations::appendChild(leftNode, child, childKey);
		Node *child = curr->mChildren[childKey];
		Operations::appendChild(leftNode, child, childKey);

		if (childKey < curr->mSize) {
		if (leftNode->mLeaf) {
		_Operations::appendKey(leftNode, curr->mKeys[childKey], curr->mValues[childKey]);
		Operations::appendKey(leftNode, curr->mKeys[childKey], curr->mValues[childKey]);
		} else {
		_Operations::appendKey(leftNode, curr->mKeys[childKey]);
		Operations::appendKey(leftNode, curr->mKeys[childKey]);
		}
		}
		}

		// this might cause some issues in synchronization
		_Operations::init(curr, false, nullptr, false);
		_Operations::appendChild(curr, leftNode, 0);
		_Operations::appendKey(curr, insertKey);
		_Operations::appendChild(curr, insertNode, 1);
		Operations::init(curr, false, nullptr, false);
		Operations::appendChild(curr, leftNode, 0);
		Operations::appendKey(curr, insertKey);
		Operations::appendChild(curr, insertNode, 1);
		}

		public:
		using Node = BNode;
		using Operations = _Operations;
		using NodeType = Node;
		using OperationsType = Operations;

		static inline BNode *init(Allocator &alloc) {
		BNode *root = alloc.allocate();
		_Operations::init(root, true, nullptr, false);
		static inline Node *init(Allocator &alloc) {
		Node *root = alloc.allocate();
		Operations::init(root, true, nullptr, false);
		return root;
		}

		static inline bool find(BNode *root, KeyType key, ValueType &result) {
		BNode *leaf = findLeaf(root, key);
		static inline bool find(Node *root, KeyType key, ValueType &result) {
		Node *leaf = findLeaf(root, key);
		if (leaf == nullptr)
		return false;

		@@ -165,10 +164,10 @@ public:
		return false;
		}

		static bool insert(BNode *root, KeyType key, ValueType value, Allocator &alloc, Latch &latch) {
		BNode prev = nullptr, curr = nullptr;
		static bool insert(Node *root, KeyType key, ValueType value, Allocator &alloc, Latch &latch) {
		Node prev = nullptr, curr = nullptr;
		do {
		BNode *tmpPrev = prev;
		Node *tmpPrev = prev;
		if (curr == nullptr) {
		curr = root;
		} else {
		@@ -186,7 +185,7 @@ public:
		size_t splitIdx = medianIdx + (!curr->mLeaf);

		KeyType medianKey = curr->mKeys[medianIdx];
		BNode *split = nullptr;
		Node *split = nullptr;

		split = splitNode(curr, alloc, splitIdx, medianIdx);

		@@ -204,7 +203,7 @@ public:
		return false;
		}

		BNode *tmp;
		Node *tmp;
		while (curr->mSibling != nullptr && (tmp = scan(curr, key)) == curr->mSibling) {
		curr = tmp;
		}
		@@ -213,8 +212,8 @@ public:
		return true;
		}

		static inline bool remove(BNode *root, KeyType key, Latch &latch) {
		BNode *cursor = findLeaf(root, key);
		static inline bool remove(Node *root, KeyType key, Latch &latch) {
		Node *cursor = findLeaf(root, key);

		if (cursor == nullptr)
		return true;
		@@ -239,8 +238,8 @@ public:
		return true;
		}

		_Operations::removeChild(cursor, keyIdx, cursor->childSize());
		_Operations::removeKey(cursor, keyIdx);
		Operations::removeChild(cursor, keyIdx, cursor->childSize());
		Operations::removeKey(cursor, keyIdx);

		latch.release(cursor);
		return true;