Loading implementation/src/BLinkTree/BLinkNode.hpp +4 −1 Original line number Diff line number Diff line Loading @@ -5,7 +5,10 @@ #include <cstddef> #include <stdint.h> template <typename KeyType, typename ValueType, size_t Order> struct BLinkNode { template <typename _KeyType, typename _ValueType, size_t Order> struct BLinkNode { using KeyType = _KeyType; using ValueType = _ValueType; BLinkNode *volatile mSibling; KeyType mHighKey; Loading implementation/src/BLinkTree/BLinkTreeCuda.cu +9 −1 Original line number Diff line number Diff line Loading @@ -366,7 +366,15 @@ public: uint16_t size = cursor->mSize; uint16_t keyIdx = lowerBoundWarp(cursor->mKeys, size, key, warp); if (keyIdx >= size || cursor->mKeys[keyIdx] != key) { if (!(keyIdx < cursor->mSize)) { releaseLatchWarp(cursor, latch, warp); return true; } KeyType cursorKey = cursor->mKeys[keyIdx]; // cursorKey != key if (cursorKey < key || key < cursorKey) { releaseLatchWarp(cursor, latch, warp); return true; } Loading implementation/src/BLinkTree/BLinkTreeHost.hpp +10 −1 Original line number Diff line number Diff line Loading @@ -225,7 +225,16 @@ public: // remove item from keys and remove the respective node size_t keyIdx = lowerBound(cursor->mKeys, cursor->mSize, key); if (keyIdx == cursor->mSize || cursor->mKeys[keyIdx] != key) { if (keyIdx == cursor->mSize) { latch.release(cursor); return true; } KeyType cursorKey = cursor->mKeys[keyIdx]; // cursorKey != key if (cursorKey < key || key < cursorKey) { latch.release(cursor); return true; } Loading implementation/src/BLinkTree/BNodeOperations/BNodeOperationsCuda.cu +2 −2 Original line number Diff line number Diff line Loading @@ -119,8 +119,8 @@ struct BLinkOperations<KeyType, ValueType, Order, TNL::Devices::Cuda> { __device__ static void removeKey(Node *volatile node, size_t index, Warp &warp) { auto rank = warp.thread_rank(); if (rank < Order) { KeyType thrKey = 0; ValueType thrVal = 0; KeyType thrKey; ValueType thrVal; if (rank >= index) { if (rank + 1 < Order) { Loading implementation/src/BPlusTree/BNodeOperations/BNodeOperationsCuda.cu +2 −2 Original line number Diff line number Diff line Loading @@ -129,8 +129,8 @@ struct BNodeOperations<KeyType, ValueType, Order, TNL::Devices::Cuda> { __device__ static void removeKey(Node *volatile node, size_t index, Warp &warp) { auto rank = warp.thread_rank(); if (rank < Order) { KeyType thrKey = 0; ValueType thrVal = 0; KeyType thrKey; ValueType thrVal; if (rank >= index) { if (rank + 1 < Order) { Loading Loading
implementation/src/BLinkTree/BLinkNode.hpp +4 −1 Original line number Diff line number Diff line Loading @@ -5,7 +5,10 @@ #include <cstddef> #include <stdint.h> template <typename KeyType, typename ValueType, size_t Order> struct BLinkNode { template <typename _KeyType, typename _ValueType, size_t Order> struct BLinkNode { using KeyType = _KeyType; using ValueType = _ValueType; BLinkNode *volatile mSibling; KeyType mHighKey; Loading
implementation/src/BLinkTree/BLinkTreeCuda.cu +9 −1 Original line number Diff line number Diff line Loading @@ -366,7 +366,15 @@ public: uint16_t size = cursor->mSize; uint16_t keyIdx = lowerBoundWarp(cursor->mKeys, size, key, warp); if (keyIdx >= size || cursor->mKeys[keyIdx] != key) { if (!(keyIdx < cursor->mSize)) { releaseLatchWarp(cursor, latch, warp); return true; } KeyType cursorKey = cursor->mKeys[keyIdx]; // cursorKey != key if (cursorKey < key || key < cursorKey) { releaseLatchWarp(cursor, latch, warp); return true; } Loading
implementation/src/BLinkTree/BLinkTreeHost.hpp +10 −1 Original line number Diff line number Diff line Loading @@ -225,7 +225,16 @@ public: // remove item from keys and remove the respective node size_t keyIdx = lowerBound(cursor->mKeys, cursor->mSize, key); if (keyIdx == cursor->mSize || cursor->mKeys[keyIdx] != key) { if (keyIdx == cursor->mSize) { latch.release(cursor); return true; } KeyType cursorKey = cursor->mKeys[keyIdx]; // cursorKey != key if (cursorKey < key || key < cursorKey) { latch.release(cursor); return true; } Loading
implementation/src/BLinkTree/BNodeOperations/BNodeOperationsCuda.cu +2 −2 Original line number Diff line number Diff line Loading @@ -119,8 +119,8 @@ struct BLinkOperations<KeyType, ValueType, Order, TNL::Devices::Cuda> { __device__ static void removeKey(Node *volatile node, size_t index, Warp &warp) { auto rank = warp.thread_rank(); if (rank < Order) { KeyType thrKey = 0; ValueType thrVal = 0; KeyType thrKey; ValueType thrVal; if (rank >= index) { if (rank + 1 < Order) { Loading
implementation/src/BPlusTree/BNodeOperations/BNodeOperationsCuda.cu +2 −2 Original line number Diff line number Diff line Loading @@ -129,8 +129,8 @@ struct BNodeOperations<KeyType, ValueType, Order, TNL::Devices::Cuda> { __device__ static void removeKey(Node *volatile node, size_t index, Warp &warp) { auto rank = warp.thread_rank(); if (rank < Order) { KeyType thrKey = 0; ValueType thrVal = 0; KeyType thrKey; ValueType thrVal; if (rank >= index) { if (rank + 1 < Order) { Loading