fix(exp): Fix expt_0520 malloc failure

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

1 files changed, 50 insertions, 67 deletions

diff --git a/skiplistcustom.cuh b/skiplistcustom.cuh
index 54c4732..a16e80a 100644
--- a/skiplistcustom.cuh
+++ b/skiplistcustom.cuh
@@ -1,4 +1,5 @@
 #pragma once
+#include "sortlib.cuh"
 #ifndef LOCKFREE_SKIPLIST_CUH_
 #define LOCKFREE_SKIPLIST_CUH_
 
@@ -81,52 +82,6 @@ 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 {
@@ -138,9 +93,14 @@ public:
 
   unsigned int *pointerIndex;
 
-  key_type *randoms;
+  key_type *samples;
+
+  CustomSort searcher;
 
-  LockFreeSkipList(key_type *randoms) {
+  double slice_size;
+
+  LockFreeSkipList(key_type *samples, size_t sample_length)
+      : searcher(sample_length, sizeof(key_type)) {
     Node *h = new Node(0);
     // size_ = 0;
 #if __WORDSIZE == 64
@@ -162,8 +122,10 @@ public:
     cudaMalloc(&pointerIndex, sizeof(unsigned int));
     initDeviceVariable();
 
-    cudaMalloc(&randoms, sizeof(key_type) * SAMPLE_LENGTH);
-    cudaMemcpy(this->randoms, randoms, sizeof(key_type) * SAMPLE_LENGTH,
+    this->slice_size = 1.0 / sample_length;
+
+    cudaMalloc(&samples, sizeof(key_type) * sample_length);
+    cudaMemcpy(this->samples, samples, sizeof(key_type) * sample_length,
                cudaMemcpyHostToDevice);
   }
 
@@ -172,10 +134,33 @@ public:
     cudaMemcpy(pointerIndex, &x, sizeof(unsigned int), cudaMemcpyHostToDevice);
   }
 
-  __device__ bool find(key_type, Node **, Node **); // Helping method
+  ~LockFreeSkipList() { cudaFree(samples); }
+
+  __device__ unsigned static trailing_zeroes(size_t index) {
+    constexpr auto block_size = 2;
+    unsigned bits = 0;
+    unsigned x = index / block_size;
+
+    if (x) {
+      while (x % block_size == 0) {
+        ++bits;
+        x /= block_size;
+      }
+    }
+    return bits;
+  }
+
+  __device__ size_t calcLevel(key_type k) {
+    auto index =
+        this->searcher.sample_cdf_custom_version(this->samples, k) / slice_size;
+    auto level = trailing_zeroes(index);
+    return level;
+  }
+
+  __device__ bool find(key_type, Node **, Node **) const; // Helping method
   __device__ bool Add(key_type);
-  __device__ bool Delete(key_type);
-  __device__ bool Search(key_type);
+  //__device__ bool Delete(key_type);
+  __device__ bool Search(key_type) const;
 
 #ifdef MEASURE_ACCESS
   unsigned access_times = 0;
@@ -206,15 +191,13 @@ public:
 
 // Function for creating a new node when requested by an add operation
 
-__device__ Node *GetNewNode(Node **nodes, key_type key,
-                            unsigned int *pointerIndex,
-                            const key_type *randoms) {
+__device__ Node *GetNewNode(key_type key, unsigned int *pointerIndex,
+                            int topLevel) {
   key_type ind = atomicInc(pointerIndex, NUM_ITEMS);
   Node *n = nodes[ind];
   n->key = key;
-  n->topLevel = randoms[ind];
-  int i;
-  for (i = 0; i < n->topLevel + 1; i++) {
+  n->topLevel = topLevel;
+  for (int i = 0; i < n->topLevel + 1; i++) {
     n->SetRef(i, nullptr, false);
   }
   return n;
@@ -223,14 +206,14 @@ __device__ Node *GetNewNode(Node **nodes, key_type key,
 // Find the window holding key
 // On the way clean up logically deleted nodes (those with set marked bit)
 
-__device__ bool
-LockFreeSkipList::find(key_type key, Node **preds,
-                       Node **succs) { // preds and succs are arrays of pointers
+__device__ bool LockFreeSkipList::find(
+    key_type key, Node **preds,
+    Node **succs) const { // preds and succs are arrays of pointers
   int bottomLevel = 0;
   bool marked[] = {false};
   bool snip;
   Node *pred;
-  Node *curr = nullptr;
+  Node *curr;
   Node *succ;
   bool beenThereDoneThat;
   while (true) {
@@ -271,7 +254,7 @@ LockFreeSkipList::find(key_type key, Node **preds,
   }
 }
 
-__device__ bool LockFreeSkipList::Search(key_type key) {
+__device__ bool LockFreeSkipList::Search(key_type key) const {
   int bottomLevel = 0;
   bool marked = false;
   Node *pred = head;
@@ -306,7 +289,7 @@ __device__ bool LockFreeSkipList::Search(key_type key) {
   return (curr != nullptr && curr->key == key);
 }
 
-__device__ bool LockFreeSkipList::Delete(key_type key) {
+/*__device__ bool LockFreeSkipList::Delete(key_type key) {
   int bottomLevel = 0;
   Node *preds[MAX_LEVEL + 1];
   Node *succs[MAX_LEVEL + 1];
@@ -342,10 +325,10 @@ __device__ bool LockFreeSkipList::Delete(key_type key) {
       }
     }
   }
-}
+}*/
 
 __device__ bool LockFreeSkipList::Add(key_type key) {
-  Node *newNode = GetNewNode(nodes, key, this->pointerIndex, this->randoms);
+  Node *newNode = GetNewNode(key, this->pointerIndex, calcLevel(key));
   int topLevel = newNode->topLevel;
   int bottomLevel = 0;
   Node *preds[MAX_LEVEL + 1];