feat: Add MemMap class

Signed-off-by: KunoiSayami <[email protected]>

1 files changed, 53 insertions, 1 deletions

diff --git a/main.cu b/main.cu
index fe801d4..bef698c 100644
--- a/main.cu
+++ b/main.cu
@@ -88,6 +88,9 @@ constexpr size_t NUM_ITEMS = 1048576;
 // constexpr size_t KEYS = 1048576;
 constexpr size_t FACTOR = 1;
 
+// should change this to dynamic next time
+constexpr size_t KEY_INDEX_SIZE = 32;
+
 // Supported operations
 constexpr int ADD = 0;
 constexpr int DELETE = 1;
@@ -184,6 +187,52 @@ public:
   }
 };
 
+struct MapNode {
+  LL key;
+  // Node *point[MAX_LEVEL + 1];
+  Node *point;
+};
+
+class MemMap {
+public:
+  size_t size;
+  size_t real_size;
+  MapNode *store;
+
+  MemMap() : size(0), store(nullptr), real_size(0) {}
+
+  __device__ bool insert(MapNode node) {
+    bool need_extend = this->size + 1 > this->real_size;
+    if (need_extend) {
+      bool need_copy = this->real_size == 0;
+      if (!need_copy) {
+        this->real_size += 1;
+      }
+      this->real_size *= 2;
+      MapNode *old = this->store;
+      this->store = new MapNode[this->real_size];
+      if (need_copy) {
+        memcpy(this->store, old, this->real_size * sizeof(MapNode *));
+      }
+      delete[] old;
+    }
+    // need sort after insert
+    this->store[size] = node;
+    this->size += 1;
+  }
+
+  __device__ Node *search(LL key) {
+    for (int offset = 0; offset < this->size; offset++) {
+      if (this->store[offset].key >= key) {
+        return this->store[offset].point;
+      }
+    }
+    return nullptr;
+  }
+
+  __device__ ~MemMap() { delete[] store; }
+};
+
 // Definition of lock-free skip list
 
 class LockFreeSkipList {
@@ -191,6 +240,7 @@ public:
   Node *head;
   Node *tail;
   unsigned size_ = 0;
+  MemMap key_map;
   LockFreeSkipList() {
     Node *h = new Node(0);
     // size_ = 0;
@@ -237,6 +287,8 @@ __device__ Node *GetNewNode(LL key) {
 
 __device__ LockFreeSkipList *l; // The lock-free skip list
 
+__device__ LL KeyIndex[KEY_INDEX_SIZE];
+
 // Kernel for initializing device memory
 
 __global__ void init(LockFreeSkipList *l1, Node **n, LL *rands) {
@@ -559,7 +611,7 @@ int main(int argc, char **argv) {
   cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice);
   cudaMemcpy(Citems, items, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice);
   cudaMemcpy(Cop, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice);
-  Node **pointers = (Node **)malloc(sizeof(LL) * adds);
+  Node **pointers = (Node **)new LL[adds]; // malloc(sizeof(LL) * adds);
   Node **Cpointers;
 
   // Allocate the pool of free nodes