fix(exp): Fix expt_0520 malloc failure

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

5 files changed, 161 insertions, 365 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 883dbc4..118786c 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -222,7 +222,7 @@ set_target_properties(expt_0517 PROPERTIES
 set_target_properties(expt_0517 PROPERTIES CUDA_ARCHITECTURES "75")
 set_target_properties(expt_0517 PROPERTIES LINKER_LANGUAGE CUDA)
 
-add_executable(expt_0520 expt_0520.cu skiplistcustom.cuh)
+add_executable(expt_0520 expt_0520.cu skiplistcustom.cuh read_helper.h)
 target_link_libraries(expt_0520 m stdc++)
 
 set_target_properties(expt_0520 PROPERTIES
diff --git a/expt_0520.cu b/expt_0520.cu
index 839fef8..88815a7 100644
--- a/expt_0520.cu
+++ b/expt_0520.cu
@@ -64,12 +64,11 @@ Conference on Parallel and Distributed Systems, December 2012.
 
 // #include"cutil.h"		// Comment this if cutil.h is not available
 // #include "cuda_runtime.h"
-#include "skiplist.cuh"
-#include <algorithm>
+#include "read_helper.h"
+#include "skiplistcustom.cuh"
 #include <cassert>
 #include <cstdio>
 #include <cstdlib>
-#include <random>
 #include <set>
 
 #if __WORDSIZE == 64
@@ -116,11 +115,6 @@ constexpr size_t SAMPLE_SIZE = 1024;
 
 constexpr int block_size = STEP_SIZE;
 
-// Supported operations
-constexpr int ADD = 0;
-constexpr int DELETE = 1;
-constexpr int SEARCH = 2;
-
 typedef LL key_type;
 
 #ifdef RANDOM_TARGET
@@ -164,126 +158,31 @@ inline void __cudaCheckError(const char *file, const int line) {
 #endif
 }
 
-__device__ LockFreeSkipList *l; // The lock-free skip list
-
 __device__ key_type SampleStorage[SAMPLE_SIZE];
 
 // Kernel for initializing device memory
 
-__global__ void init(LockFreeSkipList *lockFreeSkipList, Node **n) {
-  nodes = n;
-  l = lockFreeSkipList;
-}
-
-__global__ void print() {
-  // For debugging
-  int tid = blockIdx.x * blockDim.x + threadIdx.x;
-  if (tid == 0) {
-    Node *p = l->head;
-    bool marked = false;
-    while (p != nullptr) {
-#if __WORDSIZE == 64
-      printf("%#llx, %u, marked=%u, address is %p : ", p->key, p->topLevel,
-             marked, p);
-#else
-      printf("%#x, %u, marked=%u, address is %p\n", p->key, p->topLevel, marked,
-             p);
-#endif
-      for (int i = 0; i < p->topLevel + 1; i++) {
-        printf(" %d ", (int)(p->GetReference(i)->key));
-      }
-      printf("\n");
-      p = p->Get(0, &marked);
-    }
-    printf("\n");
-  }
-}
+__global__ void init(Node **n) { nodes = n; }
 
 // The main kernel
 
-__global__ void kernel(LL *items, LL *op, LL *result) {
+__global__ void kernel(LockFreeSkipList *skipList, const key_type *population,
+                       size_t insertion_length) {
   // The array items holds the sequence of keys
   // The array op holds the sequence of operations
   // The array result, at the end, will hold the outcome of the operations
 
-  int tid, i;
-  for (i = 0; i < FACTOR;
+  for (int i = 0; i < FACTOR;
        i++) { // FACTOR is the number of operations per thread
-    tid = i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
-    if (tid >= NUM_ITEMS)
+    auto tid =
+        i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
+    if (tid >= insertion_length)
       return;
 
     // Grab the operation and the associated key and execute
-    LL item = items[tid];
-    if (op[tid] == ADD) {
-      result[tid] = l->Add(item);
-    }
-    if (op[tid] == DELETE) {
-      result[tid] = l->Delete(item);
-    }
-    if (op[tid] == SEARCH) {
-#ifdef MEASURE_TIME
-      unsigned long long start_time = clock64();
-#endif
-      result[tid] = l->Search(item);
-#ifdef MEASURE_TIME
-      unsigned long long end_time = clock64() - start_time;
-      if (l->spend_time[tid]) {
-        printf("conflict: %d\n", tid);
-      }
-      l->spend_time[tid] = (int)end_time;
-#endif
-    }
-  }
-}
-
-/*LL Randomlevel() {
-  LL v = 1;
-  double p = 0.5;
-  while (((rand() / (double)(RAND_MAX)) < p) && (v < MAX_LEVEL))
-    v++;
-  return v;
-}*/
-
-// Generate the level of a newly created node
-LL RandomLevel(std::mt19937 &randomEngine, double p) {
-  std::geometric_distribution<> distribution(p);
-  return std::min(MAX_LEVEL, (size_t)distribution(randomEngine));
-}
-
-std::vector<LL> storage;
-
-unsigned trailing_zeroes(size_t index) {
-  unsigned bits = 0;
-  LL x = index / block_size;
-
-  if (x) {
-    while (x % block_size == 0) {
-      ++bits;
-      x /= block_size;
-    }
+    key_type item = population[tid];
+    skipList->Add(item);
   }
-  return bits;
-}
-
-LL CustomLevel(LL value) {
-  auto left = std::lower_bound(storage.begin(), storage.end(), value);
-  auto right = std::upper_bound(storage.begin(), storage.end(), value);
-
-  if (right - left != 1) {
-    printf("%ld\n", right - left);
-  }
-  assert(right - left == 1);
-
-  auto index = left - storage.begin();
-
-  if (index % block_size == 0) {
-    auto level = trailing_zeroes(index) + 1;
-    // printf("%ld,%u\n", index, level);
-    return level;
-  }
-
-  return 1;
 }
 
 __global__ void print_function() {
@@ -292,155 +191,83 @@ __global__ void print_function() {
 #endif
 }
 
-std::vector<double> population;
-std::vector<double> sample;
-
-void initialize(const std::vector<double> &input_population,
-                const std::vector<double> &input_sample) {
-  population = input_population;
-  std::sort(population.begin(), population.end());
-  sample = input_sample;
-  std::sort(sample.begin(), sample.end());
-}
-
 int main(int argc, char **argv) {
   if (argc != 3) {
-    printf("Need two arguments: percent add ops and percent delete ops (e.g., "
+    printf("Need two arguments: sample_length add insertion_length delete ops "
+           "(e.g., "
            "30 50 for 30%% add and 50%% delete).\nAborting...\n");
     exit(1);
   }
 
-  // Extract operations ratio
-  long adds = strtol(argv[1], nullptr, 10);
-  long deletes = strtol(argv[2], nullptr, 10);
+  auto error = cudaGetLastError();
+  auto total_row = 0UL;
+  auto sample_length = strtol(argv[1], nullptr, 0);
+  auto insertion_length = strtol(argv[2], nullptr, 0);
+  printf("insert_length %ld, sample_length: %ld\n", insertion_length,
+         sample_length);
 
-  storage.reserve(NUM_ITEMS);
+  ReadHelper reader("normal_distribution.txt", sample_length, insertion_length);
+  printf("%d\n", __LINE__);
+  reader.readFile(total_row);
 
-  if (adds + deletes > 100) {
-    printf("Sum of add and delete percentages exceeds 100.\nAborting...\n");
-    exit(1);
-  }
-
-  // Allocate necessary arrays
-  LL *op = new LL[NUM_ITEMS];     //(LL *)malloc(sizeof(LL) * NUM_ITEMS);
-  LL *levels = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS);
-  LL *items = new LL[NUM_ITEMS];  //(LL *)malloc(sizeof(LL) * NUM_ITEMS);
-  LL *result = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS);
-  int i;
-
-  // NUM_ITEMS is the total number of operations to execute
-  // srand(0);
-
-  std::random_device randomDevice;
-  std::mt19937 randomEngine(randomDevice());
-  std::uniform_int_distribution<int> uniformIntDistributionArray(0,
-                                                                 NUM_ITEMS - 1);
-  // std::vector<LL> storage;
-
-  // std::normal_distribution<long double> normalDistribution{2147483647,
-  //                                                          2147483647};
-
-  for (i = 0; i < NUM_ITEMS; i++) {
-    items[i] = i + 3; // 10+rand()%KEYS;
-                      //  Keys associated with
-                      //  operations
-    storage.push_back(i + 3);
-    // auto key = (key_type)std::round(normalDistribution(randomEngine));
-    // items[i] = key;
-    // storage.push_back(key);
-  }
-
-  std::sort(storage.begin(), storage.end());
-
-#if 0
-  for (i = 0; i < NUM_ITEMS; i++) {
-    /*int first = rand() % NUM_ITEMS;
-    int second = rand() % NUM_ITEMS;*/
+  printf("%d\n", __LINE__);
+  key_type *cudaPopulation;
 
-    std::swap(items[uniformIntDistributionArray(randomEngine)],
-              items[uniformIntDistributionArray(randomEngine)]);
-    /*LL temp;
-    temp = items[first];
-    items[first] = items[second];
-    items[second] = temp;*/
-  }
-#endif
-
-  // Pre-generated levels of skip list nodes (relevant only if op[i] is add)
-  // srand(0);
-  for (i = 0; i < NUM_ITEMS; i++) {
-#ifdef RANDOM_HEIGHT
-    levels[i] = RandomLevel(1 / randomEngine) - 1; // 36/14
-#else
-    levels[i] = CustomLevel(items[i]) - 1; // 31/18
-#endif
-  }
-
-  // Populate the sequence of operations
-  for (i = 0; i < (NUM_ITEMS * adds) / 100; i++) {
-    op[i] = ADD;
-  }
-  for (; i < (NUM_ITEMS * (adds + deletes)) / 100; i++) {
-    op[i] = DELETE;
-  }
-  for (; i < NUM_ITEMS; i++) {
-    op[i] = SEARCH;
-  }
-
-  adds = (NUM_ITEMS * adds) / 100;
+  printf("%d\n", __LINE__);
+  cudaMalloc(&cudaPopulation, sizeof(key_type) * insertion_length);
+  std::vector<key_type> _sample, _population;
+  reader.split_into(_sample, _population);
+  cudaMemcpy(cudaPopulation, _population.data(),
+             sizeof(key_type) * insertion_length, cudaMemcpyHostToDevice);
 
+  printf("%d\n", __LINE__);
   // Allocate device memory
 
-  LL *Citems;
-  LL *Cop;
-  LL *Cresult;
-  // LL *Clevels;
-
-  cudaMalloc((void **)&Cresult, sizeof(LL) * NUM_ITEMS);
-  cudaMalloc((void **)&Citems, sizeof(LL) * NUM_ITEMS);
-  cudaMalloc((void **)&Cop, sizeof(LL) * NUM_ITEMS);
   // cudaMalloc((void **)&Clevels, sizeof(LL) * NUM_ITEMS);
   // 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 **)new LL[adds]; // malloc(sizeof(LL) * adds);
-  Node **Cpointers;
+  Node **pointers =
+      (Node **)new LL[insertion_length]; // malloc(sizeof(LL) * adds);
+  Node **cudaNodePointers;
+
+  printf("%d\n", __LINE__);
 
   // Allocate the pool of free nodes
 
-  for (i = 0; i < adds; i++) {
-    cudaMalloc((void **)&pointers[i], sizeof(Node));
+  for (int i = 0; i < insertion_length; i++) {
+    cudaMalloc(&pointers[i], sizeof(Node));
   }
-  cudaMalloc((void **)&Cpointers, sizeof(Node *) * adds);
-  cudaMemcpy(Cpointers, pointers, sizeof(Node *) * adds,
+  cudaMalloc((void **)&cudaNodePointers, sizeof(Node *) * insertion_length);
+  cudaMemcpy(cudaNodePointers, pointers, sizeof(Node *) * insertion_length,
              cudaMemcpyHostToDevice);
+  CudaCheckError();
 
+  printf("%d\n", __LINE__);
   // Allocate the skip list
 
   LockFreeSkipList *Clist;
-  auto *list = new LockFreeSkipList(levels);
+  auto *list = new LockFreeSkipList(_sample.data(), SAMPLE_SIZE);
 
   cudaMalloc((void **)&Clist, sizeof(LockFreeSkipList));
   cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice);
+  CudaCheckError();
   // Calculate the number of thread blocks
   // NUM_ITEMS = total number of operations to execute
   // NUM_THREADS = number of threads per block
   // FACTOR = number of operations per thread
 
-  int blocks = (NUM_ITEMS % (NUM_THREADS * FACTOR) == 0)
-                   ? NUM_ITEMS / (NUM_THREADS * FACTOR)
-                   : (NUM_ITEMS / (NUM_THREADS * FACTOR)) + 1;
+  size_t blocks = (insertion_length % (NUM_THREADS * FACTOR) == 0)
+                      ? insertion_length / (NUM_THREADS * FACTOR)
+                      : (insertion_length / (NUM_THREADS * FACTOR)) + 1;
 
   // Error checking code
-  cudaError_t error = cudaGetLastError();
   if (cudaSuccess != error) {
     printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
     exit(-1);
   }
 
   // Initialize the device memory
-  init<<<1, 32>>>(Clist, Cpointers);
+  init<<<1, 32>>>(cudaNodePointers);
   cudaDeviceSynchronize();
 
   // Launch main kernel
@@ -450,13 +277,8 @@ int main(int argc, char **argv) {
   cudaEventCreate(&stop);
   cudaEventRecord(start, nullptr);
 
-  kernel<<<blocks, NUM_THREADS>>>(Citems, Cop, Cresult);
+  kernel<<<blocks, NUM_THREADS>>>(Clist, cudaPopulation, insertion_length);
   CudaCheckError();
-  error = cudaGetLastError();
-  if (cudaSuccess != error) {
-    printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
-    // exit(-1);
-  }
   cudaDeviceSynchronize();
   cudaEventRecord(stop, nullptr);
   cudaEventSynchronize(stop);
@@ -471,58 +293,6 @@ int main(int argc, char **argv) {
 
   printf("%lu: %lf", NUM_ITEMS, time);
 
-  // Launch main kernel for query
-
-  // Populate the sequence of operations
-  for (i = 0; i < NUM_ITEMS; i++) {
-    op[i] = SEARCH;
-  }
-
-  LL *Cop2;
-  cudaMalloc((void **)&Cop2, sizeof(LL) * NUM_ITEMS);
-  cudaMemcpy(Cop2, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice);
-
-  cudaEventCreate(&start);
-  cudaEventCreate(&stop);
-  cudaEventRecord(start, nullptr);
-#ifdef MEASURE_TIME
-  kernel<<<NUM_ITEMS, 1>>>(Citems, Cop2, Cresult);
-#else
-  kernel<<<blocks, NUM_THREADS>>>(Citems, Cop2, Cresult);
-#endif
-  CudaCheckError();
-  error = cudaGetLastError();
-  if (cudaSuccess != error) {
-    printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
-    // exit(-1);
-  }
-  cudaDeviceSynchronize();
-  cudaEventRecord(stop, nullptr);
-  cudaEventSynchronize(stop);
-  cudaEventElapsedTime(&time, start, stop);
-  cudaEventDestroy(start);
-  cudaEventDestroy(stop);
-
-  // Print kernel execution time in milliseconds
-
-  printf(" %lf\n", time);
-
-  // Check for errors
-
-  error = cudaGetLastError();
-  if (cudaSuccess != error) {
-    printf("error1:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
-    exit(-1);
-  }
-
-  // Move results back to host memory
-
-  cudaMemcpy(result, Cresult, sizeof(LL) * NUM_ITEMS, cudaMemcpyDeviceToHost);
-
-  // Uncomment the following for debugging
-  // print<<<1,32>>>();
-  cudaDeviceSynchronize();
-
 #if (defined(MEASURE_TIME) || defined(MEASURE_ACCESS))
 
   print_function<<<1, 1>>>();
@@ -552,16 +322,12 @@ int main(int argc, char **argv) {
   //   printf("%d\n", element);
 #endif
 #endif
-  /*cudaFree(Clist);
-    cudaFree(Cop2);
-    cudaFree(Clevels);
-    cudaFree(Cop);
-    cudaFree(Citems);
-    cudaFree(Cresult);
-    free(pointers);
-    delete [] op;
-    delete [] levels;
-    delete [] items;
-    delete [] result;*/
+  cudaFree(Clist);
+  for (int i = 0; i < insertion_length; i++) {
+    cudaFree(pointers[i]);
+  }
+  cudaFree(cudaNodePointers);
+  delete list;
+  delete[] pointers;
   return 0;
 }
diff --git a/read_helper.h b/read_helper.h
index 311ebbb..292eaac 100644
--- a/read_helper.h
+++ b/read_helper.h
@@ -15,12 +15,18 @@
 #endif
 
 class ReadHelper {
-  unsigned long long max_value = 0, min_value = 0x7fffffffff;
+  typedef unsigned long long key_type;
+  key_type max_value = 0, min_value = 0x7fffffffff;
 
   char const *filename;
+  static constexpr size_t REVERSED_BLOCK = 256;
 
 public:
-  ReadHelper(char const *filename) : filename(filename) {}
+  size_t sample_length, needed_read_length;
+  ReadHelper(char const *filename, size_t sample_length,
+             size_t insertion_length)
+      : filename(filename), sample_length(sample_length),
+        needed_read_length(insertion_length + sample_length) {}
   unsigned long long maxValue() const { return max_value; }
   unsigned long long minValue() const { return min_value; }
 
@@ -35,6 +41,11 @@ public:
     population_vector.push_back(value);
   }
 
+  void finish_read() {
+    this->population_vector.push_back(this->max_value);
+    this->population_vector.push_back(this->min_value);
+  }
+
   static unsigned long randomRow(unsigned long max_value_) {
     std::random_device randomDevice;
     std::mt19937 mt19937(randomDevice());
@@ -42,20 +53,21 @@ public:
     return dst(mt19937);
   }
 
-  void readFile(long sample_length, unsigned long &total_row,
-                unsigned long max_number) {
+  void readFile(unsigned long &total_row) {
+    this->population_vector.clear();
     auto read_number = 0UL;
     FILE *file = fopen(filename, "r");
     assert(file);
     P_ERR("Reading sample");
     for (long long i;
-         read_number < max_number && fscanf(file, "%lld ", &i) != EOF;
+         read_number < needed_read_length && fscanf(file, "%lld ", &i) != EOF;
          store_into_vector(i))
       read_number++;
 
     if (total_row > 0) {
       P_ERR("\rReading skip");
-      read_number = randomRow(total_row - sample_length - max_number - 256);
+      read_number = randomRow(total_row - sample_length - needed_read_length -
+                              REVERSED_BLOCK);
       total_row = read_number;
       // fprintf(stderr, "Skip %lu\n", read_number);
       for (long long i; read_number > 0 && fscanf(file, "%lld ", &i) != EOF;)
@@ -65,12 +77,31 @@ public:
     P_ERR("\rReading population");
     read_number = 0;
 
-    auto remain = sample_length + 256;
+    auto remain = sample_length + REVERSED_BLOCK;
     for (long long i; read_number < remain && fscanf(file, "%lld ", &i) != EOF;
          store_into_vector(i))
       read_number++;
     fclose(file);
     P_ERR("\r");
+    finish_read();
+  }
+
+  void split_into(std::vector<key_type> &sample, std::vector<key_type> &p) {
+    sample.reserve(sample_length);
+    p.reserve(needed_read_length - sample_length);
+    memcpy(sample.data(), population_vector.data(),
+           sizeof(key_type) * sample_length);
+    memcpy(p.data(), population_vector.data() + sample_length,
+           sizeof(key_type) * (needed_read_length - sample_length));
+  }
+
+  void split_into(key_type *&sample, key_type *&p) {
+    sample = new key_type[sample_length];
+    p = new key_type[needed_read_length - sample_length];
+    memcpy(sample, population_vector.data(), sizeof(key_type) * sample_length);
+    memcpy(p, population_vector.data() + sample_length,
+           sizeof(key_type) * (needed_read_length - sample_length));
   }
 };
+
 #endif
\ No newline at end of file
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];
diff --git a/sortlib.cuh b/sortlib.cuh
index 5bacc11..ae0ea67 100644
--- a/sortlib.cuh
+++ b/sortlib.cuh
@@ -3,6 +3,7 @@
 
 #include <cassert>
 #include <cstdio>
+#include <vector>
 
 typedef unsigned long long key_type;
 
@@ -176,8 +177,23 @@ public:
     puts("You can remove this function if passed already");
 #endif
   }
+
+  __host__ __device__ size_t length() const { return this->LENGTH; }
 };
 
+template <typename T> void static rebuild(std::vector<T> original) {
+  auto sample_length = original.size();
+  auto sorter = CustomSort(sample_length);
+  auto tmp = new T[sample_length];
+
+  for (size_t i = 0; i < sample_length; i++) {
+    tmp[i] = original[sorter.calculate_index(i) - 1];
+  }
+  memcpy(original.data(), tmp, sizeof(T) * sample_length);
+
+  delete[] tmp;
+}
+
 __global__ void testCustomCalculation(size_t length) {
   CustomSort(length, sizeof(long) * 8).testCalculation();
 }