feat(exp): Add expt_0525.cu

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

5 files changed, 896 insertions, 39 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 118786c..4553b8c 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -222,6 +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 read_helper.h)
 target_link_libraries(expt_0520 m stdc++)
 
@@ -230,6 +231,11 @@ set_target_properties(expt_0520 PROPERTIES
 set_target_properties(expt_0520 PROPERTIES CUDA_ARCHITECTURES "75")
 set_target_properties(expt_0520 PROPERTIES LINKER_LANGUAGE CUDA)
 
+add_executable(expt_0525 expt_0525.cu read_helper.h)
+target_link_libraries(expt_0525 m stdc++)
 
-
+set_target_properties(expt_0525 PROPERTIES
+        CUDA_SEPARABLE_COMPILATION ON)
+set_target_properties(expt_0525 PROPERTIES CUDA_ARCHITECTURES "75")
+set_target_properties(expt_0525 PROPERTIES LINKER_LANGUAGE CUDA)
 
diff --git a/expt_0520.cu b/expt_0520.cu
index 88815a7..d31dc78 100644
--- a/expt_0520.cu
+++ b/expt_0520.cu
@@ -111,18 +111,12 @@ constexpr size_t FACTOR = 1;
 
 // should change this to dynamic next time
 // constexpr size_t KEY_INDEX_SIZE = 32;
-constexpr size_t SAMPLE_SIZE = 1024;
+// constexpr size_t SAMPLE_SIZE = 1023;
 
 constexpr int block_size = STEP_SIZE;
 
 typedef LL key_type;
 
-#ifdef RANDOM_TARGET
-constexpr const char *TARGET_STRING = "RANDOM";
-#else
-constexpr const char *TARGET_STRING = "PERFECT";
-#endif
-
 #define CUDA_ERROR_CHECK
 
 #define CudaSafeCall(err) __cudaSafeCall(err, __FILE__, __LINE__)
@@ -158,30 +152,38 @@ inline void __cudaCheckError(const char *file, const int line) {
 #endif
 }
 
-__device__ key_type SampleStorage[SAMPLE_SIZE];
+//__device__ key_type SampleStorage[SAMPLE_SIZE];
+__device__ LockFreeSkipList *cudaGlobalSkipList;
 
 // Kernel for initializing device memory
 
-__global__ void init(Node **n) { nodes = n; }
+__global__ void init(LockFreeSkipList *list, Node **n) {
+  cudaGlobalSkipList = list;
+  nodes = n;
+}
+
+__global__ void testFunctions() { cudaGlobalSkipList->testSample(); }
 
 // The main kernel
 
-__global__ void kernel(LockFreeSkipList *skipList, const key_type *population,
-                       size_t insertion_length) {
+__global__ void kernel(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
 
   for (int i = 0; i < FACTOR;
        i++) { // FACTOR is the number of operations per thread
-    auto tid =
-        i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
+    // auto tid = i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x +
+    // threadIdx.x;
+    auto tid = FACTOR * (blockIdx.x * blockDim.x + threadIdx.x) + i;
+    // printf("%lu\n", tid);
     if (tid >= insertion_length)
       return;
 
     // Grab the operation and the associated key and execute
     key_type item = population[tid];
-    skipList->Add(item);
+    // printf("%llu\n", population[tid]);
+    cudaGlobalSkipList->Add(item);
   }
 }
 
@@ -207,20 +209,16 @@ int main(int argc, char **argv) {
          sample_length);
 
   ReadHelper reader("normal_distribution.txt", sample_length, insertion_length);
-  printf("%d\n", __LINE__);
   reader.readFile(total_row);
 
-  printf("%d\n", __LINE__);
   key_type *cudaPopulation;
 
-  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
 
   // cudaMalloc((void **)&Clevels, sizeof(LL) * NUM_ITEMS);
@@ -230,8 +228,6 @@ int main(int argc, char **argv) {
       (Node **)new LL[insertion_length]; // malloc(sizeof(LL) * adds);
   Node **cudaNodePointers;
 
-  printf("%d\n", __LINE__);
-
   // Allocate the pool of free nodes
 
   for (int i = 0; i < insertion_length; i++) {
@@ -242,14 +238,15 @@ int main(int argc, char **argv) {
              cudaMemcpyHostToDevice);
   CudaCheckError();
 
-  printf("%d\n", __LINE__);
   // Allocate the skip list
 
-  LockFreeSkipList *Clist;
-  auto *list = new LockFreeSkipList(_sample.data(), SAMPLE_SIZE);
+  // LockFreeSkipList *cudaLockFreeSkipList;
+  auto *list = new LockFreeSkipList(_sample.data(), sample_length);
 
-  cudaMalloc((void **)&Clist, sizeof(LockFreeSkipList));
-  cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice);
+  LockFreeSkipList *cudaSkipList = nullptr;
+  cudaMalloc(&cudaSkipList, sizeof(LockFreeSkipList));
+  cudaMemcpy(cudaSkipList, list, sizeof(LockFreeSkipList),
+             cudaMemcpyHostToDevice);
   CudaCheckError();
   // Calculate the number of thread blocks
   // NUM_ITEMS = total number of operations to execute
@@ -267,7 +264,10 @@ int main(int argc, char **argv) {
   }
 
   // Initialize the device memory
-  init<<<1, 32>>>(cudaNodePointers);
+  init<<<1, 32>>>(cudaSkipList, cudaNodePointers);
+  cudaDeviceSynchronize();
+
+  testFunctions<<<1, 1>>>();
   cudaDeviceSynchronize();
 
   // Launch main kernel
@@ -277,7 +277,7 @@ int main(int argc, char **argv) {
   cudaEventCreate(&stop);
   cudaEventRecord(start, nullptr);
 
-  kernel<<<blocks, NUM_THREADS>>>(Clist, cudaPopulation, insertion_length);
+  kernel<<<blocks, NUM_THREADS>>>(cudaPopulation, insertion_length);
   CudaCheckError();
   cudaDeviceSynchronize();
   cudaEventRecord(stop, nullptr);
@@ -289,8 +289,6 @@ int main(int argc, char **argv) {
 
   // Print kernel execution time in milliseconds
 
-  printf("%s %d ", TARGET_STRING, block_size);
-
   printf("%lu: %lf", NUM_ITEMS, time);
 
 #if (defined(MEASURE_TIME) || defined(MEASURE_ACCESS))
@@ -322,7 +320,7 @@ int main(int argc, char **argv) {
   //   printf("%d\n", element);
 #endif
 #endif
-  cudaFree(Clist);
+  // cudaFree(cudaLockFreeSkipList);
   for (int i = 0; i < insertion_length; i++) {
     cudaFree(pointers[i]);
   }
diff --git a/expt_0525.cu b/expt_0525.cu
new file mode 100644
index 0000000..af024bc
--- /dev/null
+++ b/expt_0525.cu
@@ -0,0 +1,837 @@
+/*
+
+Copyright 2012-2013 Indian Institute of Technology Kanpur. All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions are met:
+
+1. Redistributions of source code must retain the above copyright notice, this
+list of conditions, and the following disclaimer.
+2. Redistributions in binary form must reproduce the above copyright notice,
+this list of conditions, and the following disclaimer in the documentation
+and/or other materials provided with the distribution.
+
+THIS SOFTWARE IS PROVIDED BY INDIAN INSTITUTE OF TECHNOLOGY KANPUR ``AS IS''
+AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+DISCLAIMED. IN NO EVENT SHALL INDIAN INSTITUTE OF TECHNOLOGY KANPUR OR
+THE CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT
+OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+POSSIBILITY OF SUCH DAMAGE.
+
+The views and conclusions contained in the software and documentation are
+those of the authors and should not be interpreted as representing official
+policies, either expressed or implied, of Indian Institute of Technology Kanpur.
+
+*/
+
+/**********************************************************************************
+
+ Lock-free skip list for CUDA; tested for CUDA 4.2 on 32-bit Ubuntu 10.10 and
+64-bit Ubuntu 12.04. Developed at IIT Kanpur.
+
+ Inputs: Percentage of add and delete operations (e.g., 30 50 for 30% add and
+50% delete) Output: Prints the total time (in milliseconds) to execute the the
+sequence of operations
+
+ Compilation flags: -O3 -arch sm_20 -I ~/NVIDIA_GPU_Computing_SDK/C/common/inc/
+-DNUM_ITEMS=num_ops -DFACTOR=num_ops_per_thread -DKEYS=num_keys
+
+ NUM_ITEMS is the total number of operations (mix of add, delete, search) to
+execute.
+
+ FACTOR is the number of operations per thread.
+
+ KEYS is the number of integer keys assumed in the range [10, 9+KEYS].
+ The paper cited below states that the key range is [0, KEYS-1]. However, we
+have shifted the range by +10 so that the head sentinel key (the minimum key)
+can be chosen as zero. Any positive shift other than +10 would also work.
+
+ The include path ~/NVIDIA_GPU_Computing_SDK/C/common/inc/ is needed for
+cutil.h.
+
+ Related work:
+
+ Prabhakar Misra and Mainak Chaudhuri. Performance Evaluation of Concurrent
+Lock-free Data Structures on GPUs. In Proceedings of the 18th IEEE International
+Conference on Parallel and Distributed Systems, December 2012.
+
+***************************************************************************************/
+
+// #include"cutil.h"		// Comment this if cutil.h is not available
+// #include "cuda_runtime.h"
+#include "read_helper.h"
+#include "sortlib.cuh"
+#include <algorithm>
+#include <cassert>
+#include <cstdio>
+#include <cstdlib>
+#include <random>
+#include <set>
+
+#if __WORDSIZE == 64
+typedef unsigned long long LL;
+#else
+typedef unsigned int LL;
+#endif
+
+#ifndef BUILD_SIZE
+#define BUILD_SIZE 1048576
+#endif
+
+#ifndef STEP_SIZE
+#define STEP_SIZE 2
+#endif
+
+// #define MEASURE_TIME
+//  #define MEASURE_ACCESS
+
+#if (defined(MEASURE_ACCESS) && defined(MEASURE_TIME))
+#error "Shouldn't define MEASURE_TIME and MEASURE_ACCESS at the same time"
+#endif
+
+#ifdef MEASURE_TIME
+#undef BUILD_SIZE
+#define BUILD_SIZE 1024
+#endif
+
+// Maximum level of a node in the skip list
+// #define MAX_LEVEL 32
+constexpr size_t MAX_LEVEL = 16;
+
+// Number of threads per block
+// #define NUM_THREADS 512
+constexpr size_t NUM_THREADS = 512;
+
+constexpr size_t NUM_ITEMS = BUILD_SIZE;
+// constexpr size_t KEYS = 1048576;
+constexpr size_t FACTOR = 1;
+
+// should change this to dynamic next time
+// constexpr size_t KEY_INDEX_SIZE = 32;
+// constexpr size_t SAMPLE_SIZE = 1023;
+
+// 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
+constexpr const char *TARGET_STRING = "RANDOM";
+#else
+// constexpr const char *TARGET_STRING = "PERFECT";
+#endif
+
+#define CUDA_ERROR_CHECK
+
+#define CudaSafeCall(err) __cudaSafeCall(err, __FILE__, __LINE__)
+#define CudaCheckError() __cudaCheckError(__FILE__, __LINE__)
+
+inline void cudaSafeCall_(cudaError err, const char *file, const int line) {
+#ifdef CUDA_ERROR_CHECK
+  if (cudaSuccess != err) {
+    fprintf(stderr, "cudaSafeCall() failed at %s:%i : %s\n", file, line,
+            cudaGetErrorString(err));
+    exit(-1);
+  }
+#endif
+}
+
+inline void __cudaCheckError(const char *file, const int line) {
+#ifdef CUDA_ERROR_CHECK
+  cudaError err = cudaGetLastError();
+  if (cudaSuccess != err) {
+    fprintf(stderr, "cudaCheckError() failed at %s:%i : %s\n", file, line,
+            cudaGetErrorString(err));
+    exit(-1);
+  }
+
+  // More careful checking. However, this will affect performance.
+  // Comment away if needed.
+  err = cudaDeviceSynchronize();
+  if (cudaSuccess != err) {
+    fprintf(stderr, "cudaCheckError() with sync failed at %s:%i : %s\n", file,
+            line, cudaGetErrorString(err));
+    exit(-1);
+  }
+#endif
+}
+
+class Node;
+
+// Definition of generic node class
+
+class __attribute__((aligned(16))) Node {
+public:
+  int topLevel;             // Level of the node
+  LL key;                   // Key value
+  LL next[MAX_LEVEL + 1]{}; // Array of next links
+
+  // Create a next field from a reference and mark bit
+  static __device__ __host__ LL CreateRef(Node *ref, bool mark) {
+    LL val = (LL)ref;
+    val = val | mark;
+    return val;
+  }
+
+  __device__ __host__ void SetRef(int index, Node *ref, bool mark) {
+    next[index] = CreateRef(ref, mark);
+  }
+
+  // Extract the reference from a next field
+  __device__ Node *GetReference(int index) {
+    LL ref = next[index];
+    return (Node *)((ref >> 1) << 1);
+  }
+
+  // Extract the reference and mark bit from a next field
+  __device__ Node *Get(int index, bool *marked) {
+    marked[0] = next[index] % 2;
+    return (Node *)((next[index] >> 1) << 1);
+  }
+
+  // CompareAndSet wrapper
+  __device__ bool CompareAndSet(int index, Node *expectedRef, Node *newRef,
+                                bool oldMark, bool newMark) {
+    LL oldVal = (LL)expectedRef | oldMark;
+    LL newVal = (LL)newRef | newMark;
+    LL *ref = &(next[index]);
+    LL oldValOut = atomicCAS(ref, oldVal, newVal);
+    if (oldValOut == oldVal)
+      return true;
+    return false;
+  }
+
+  // Constructor for sentinel nodes
+  explicit Node(LL k) {
+    key = k;
+    topLevel = MAX_LEVEL;
+    int i;
+    for (i = 0; i < MAX_LEVEL + 1; i++) {
+      next[i] = CreateRef((Node *)nullptr, false);
+    }
+  }
+};
+
+/*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 {
+
+  key_type *sample;
+  size_t sampleLength;
+
+  CustomSort customSort;
+
+public:
+  Node *head;
+  Node *tail;
+  LockFreeSkipList(key_type *_sample, size_t sample_length)
+      : sampleLength(sample_length),
+        customSort(sample_length, sizeof(key_type) * 8) {
+    Node *h = new Node(0);
+    // size_ = 0;
+#if __WORDSIZE == 64
+    Node *t = new Node(std::numeric_limits<key_type>::max() - 1);
+#else
+    Node *t = new Node((LL)0xffffffff);
+#endif
+    cudaMalloc(&head, sizeof(Node));
+
+    cudaMalloc(&tail, sizeof(Node));
+    for (auto i = 0; i < h->topLevel + 1; i++) {
+      h->SetRef(i, tail, false);
+    }
+    cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice);
+
+    cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice);
+
+    cudaMalloc(&this->sample, sizeof(key_type) * sampleLength);
+    cudaMemcpy(this->sample, _sample, sizeof(key_type) * sampleLength,
+               cudaMemcpyHostToDevice);
+  }
+  __device__ bool find(LL, Node **, Node **); // Helping method
+  __device__ bool Add(LL);
+  __device__ bool Delete(LL);
+  __device__ bool Search(LL);
+
+  __device__ size_t searchIndex(key_type key) {
+    auto result = customSort.binary_search(this->sample, key);
+    auto index = customSort.calculate_index(result - this->sample);
+    return index;
+  }
+
+  __device__ unsigned trailing_zeroes(size_t index) {
+    constexpr auto block_size = 2;
+    unsigned bits = 0;
+    LL x = index / block_size;
+
+    if (x) {
+      while (x % block_size == 0) {
+        ++bits;
+        x /= block_size;
+      }
+    }
+    return bits;
+  }
+
+#ifdef MEASURE_ACCESS
+  unsigned access_times = 0;
+
+  __device__ unsigned getAccessCount() const { return this->access_times; }
+  __device__ void increaseAccessCount(unsigned count = 1) {
+    atomicAdd(&this->access_times, count);
+  }
+#else
+  __device__ void increaseAccessCount(unsigned _count = 1) {}
+#endif
+
+#ifdef MEASURE_TIME
+  unsigned round = 0;
+  __device__ void increaseRoundCount(unsigned count = 1) {
+    atomicAdd(&this->round, count);
+  }
+  int spend_time[NUM_ITEMS]{0};
+  unsigned long long total_time = 0;
+  __device__ unsigned getRoundCount() const { return this->round; }
+#endif
+};
+
+__device__ Node **nodes;                  // Pool of pre-allocated nodes
+__device__ unsigned int pointerIndex = 0; // Index into pool of free nodes
+__device__ LL
+    *randoms; // Array storing the levels of the nodes in the free pool
+
+// Function for creating a new node when requested by an add operation
+
+__device__ Node *GetNewNode(LL key, size_t topLevel) {
+  LL ind = atomicInc(&pointerIndex, NUM_ITEMS);
+  Node *n = nodes[ind];
+  n->key = key;
+  // n->topLevel = randoms[ind];
+  n->topLevel = topLevel;
+  int i;
+  for (i = 0; i < n->topLevel + 1; i++) {
+    n->SetRef(i, nullptr, false);
+  }
+  return n;
+}
+
+__device__ LockFreeSkipList *lockFreeSkipList; // The lock-free skip list
+
+//__device__ LL KeyIndex[KEY_INDEX_SIZE];
+
+//__device__ key_type SampleStorage[SAMPLE_SIZE];
+
+// Kernel for initializing device memory
+
+__global__ void init(LockFreeSkipList *l1, Node **n, LL *rands) {
+  randoms = rands;
+  nodes = n;
+  lockFreeSkipList = l1;
+}
+
+// Find the window holding key
+// On the way clean up logically deleted nodes (those with set marked bit)
+
+__device__ bool
+LockFreeSkipList::find(LL key, Node **preds,
+                       Node **succs) { // preds and succs are arrays of pointers
+  int bottomLevel = 0;
+  bool marked[] = {false};
+  bool snip;
+  Node *pred;
+  Node *curr = nullptr;
+  Node *succ;
+  bool beenThereDoneThat;
+  while (true) {
+    beenThereDoneThat = false;
+    pred = head;
+    int level;
+    for (level = MAX_LEVEL; level >= bottomLevel; level--) {
+      curr = pred->GetReference(level);
+      while (true) {
+        succ = curr->Get(level, marked);
+        while (marked[0]) {
+          snip = pred->CompareAndSet(level, curr, succ, false, false);
+          beenThereDoneThat = true;
+          if (!snip)
+            break;
+          curr = pred->GetReference(level);
+          succ = curr->Get(level, marked);
+          beenThereDoneThat = false;
+          // printf("find key is %d \n",(int)key);
+        }
+        if (beenThereDoneThat)
+          break;
+        if (curr->key <= key) {
+          pred = curr;
+          curr = succ;
+        } else {
+          break;
+        }
+      }
+      if (beenThereDoneThat)
+        break;
+      preds[level] = pred;
+      succs[level] = curr;
+    }
+    if (beenThereDoneThat)
+      continue;
+    return ((curr->key == key));
+  }
+}
+
+__device__ bool LockFreeSkipList::Search(LL key) {
+  int bottomLevel = 0;
+  bool marked = false;
+  Node *pred = head;
+  Node *curr = nullptr;
+  Node *succ;
+  int level;
+  for (level = MAX_LEVEL; level >= bottomLevel; level--) {
+    curr = pred->GetReference(level);
+#ifdef MEASURE_ACCESS
+    this->increaseAccessCount();
+#endif
+    while (true) {
+      succ = curr->Get(level, &marked);
+#ifdef MEASURE_ACCESS
+      this->increaseAccessCount();
+#endif
+      while (marked) {
+        curr = curr->GetReference(level);
+        succ = curr->Get(level, &marked);
+#ifdef MEASURE_ACCESS
+        this->increaseAccessCount(2);
+#endif
+      }
+      if (curr->key < key) {
+        pred = curr;
+        curr = succ;
+      } else {
+        break;
+      }
+    }
+  }
+  return (curr != nullptr && curr->key == key);
+}
+
+__device__ bool LockFreeSkipList::Delete(LL key) {
+  int bottomLevel = 0;
+  Node *preds[MAX_LEVEL + 1];
+  Node *succs[MAX_LEVEL + 1];
+  Node *succ;
+  bool marked[] = {false};
+  while (true) {
+    bool found = find(key, preds, succs);
+    if (!found) {
+      return false;
+    } else {
+      Node *nodeToDelete = succs[bottomLevel];
+      int level;
+      for (level = nodeToDelete->topLevel; level >= bottomLevel + 1; level--) {
+        succ = nodeToDelete->Get(level, marked);
+        while (!marked[0]) {
+          nodeToDelete->CompareAndSet(level, succ, succ, false, true);
+          succ = nodeToDelete->Get(level, marked);
+        }
+      }
+      succ = nodeToDelete->Get(bottomLevel, marked);
+      while (true) {
+        bool iMarkedIt =
+            nodeToDelete->CompareAndSet(bottomLevel, succ, succ, false, true);
+        succ = succs[bottomLevel]->Get(bottomLevel, marked);
+        if (iMarkedIt) {
+          find(key, preds, succs);
+          // size_ -= 1;
+          // atomicDec(&size_, 1);
+          return true;
+        } else if (marked[0]) {
+          return false;
+        }
+      }
+    }
+  }
+}
+
+__device__ bool LockFreeSkipList::Add(LL key) {
+  Node *newNode =
+      GetNewNode(key, this->trailing_zeroes(this->searchIndex(key)));
+  int topLevel = newNode->topLevel;
+  int bottomLevel = 0;
+  Node *preds[MAX_LEVEL + 1];
+  Node *succs[MAX_LEVEL + 1];
+  int level;
+  while (true) {
+    bool found = find(key, preds, succs);
+    if (found) {
+      return false;
+    } else {
+      Node *pred;
+      Node *succ;
+      for (level = bottomLevel; level <= topLevel; level++) {
+        succ = succs[level];
+        newNode->SetRef(level, succ, false);
+      }
+      pred = preds[bottomLevel];
+      succ = succs[bottomLevel];
+      bool t;
+      // printf("--- key is %d pred is %d succ is %d level is %d
+      // \n",(int)key,(int)pred->key,(int)succ->key,0);
+      t = pred->CompareAndSet(bottomLevel, succ, newNode, false, false);
+      if (!t) {
+        continue;
+      }
+      for (level = bottomLevel + 1; level <= topLevel; level++) {
+
+        while (true) {
+          pred = preds[level];
+          succ = succs[level];
+          newNode->SetRef(level, succ, false);
+          // printf("-- key is %d pred is %d succ is %d level is %d
+          // \n",(int)key,(int)pred->key,(int)succ->key,(int)level);
+          if (pred->CompareAndSet(level, succ, newNode, false, false)) {
+            break;
+          }
+          // printf("key is %d pred is %d succ is %d level is %d
+          // \n",(int)key,(int)pred->key,(int)succ->key,(int)level);
+          find(key, preds, succs);
+        }
+      }
+      // size_ += 1;
+      // this->key_map.insert(MapNode(ll, newNode));
+      // atomicAdd(&size_, 1);
+      return true;
+    }
+  }
+}
+
+__global__ void print() {
+  // For debugging
+  int tid = blockIdx.x * blockDim.x + threadIdx.x;
+  if (tid == 0) {
+    Node *p = lockFreeSkipList->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");
+  }
+}
+
+// The main kernel
+
+__global__ void kernel(const LL *items, size_t search_length, LL *result) {
+  // 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
+
+  for (int i = 0; i < FACTOR;
+       i++) { // FACTOR is the number of operations per thread
+    auto tid =
+        i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
+    if (tid >= search_length)
+      return;
+
+    // Grab the operation and the associated key and execute
+    LL item = items[tid];
+#ifdef MEASURE_TIME
+    unsigned long long start_time = clock64();
+#endif
+    result[tid] = lockFreeSkipList->Search(item);
+#ifdef MEASURE_TIME
+    unsigned long long end_time = clock64() - start_time;
+    if (lockFreeSkipList->spend_time[tid]) {
+      printf("conflict: %d\n", tid);
+    }
+    lockFreeSkipList->spend_time[tid] = (int)end_time;
+#endif
+  }
+}
+
+__global__ void kernelAdd(LL *item, size_t insertion_length) {
+
+  for (int i = 0; i < FACTOR;
+       i++) { // FACTOR is the number of operations per thread
+    auto tid =
+        i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
+
+    if (tid >= insertion_length)
+      return;
+    lockFreeSkipList->Add(item[tid]);
+  }
+}
+
+/*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
+
+__global__ void print_function() {
+#ifdef MEASURE_ACCESS
+  printf("count: %u\n", l->getAccessCount());
+#endif
+}
+
+/*__global__ void copy_function(int *spend_time) {
+  memcpy(spend_time, lockFreeSkipList->spend_time, sizeof(int) * NUM_ITEMS);
+}*/
+
+inline size_t calcBlocks(size_t input) {
+  return (input % (NUM_THREADS * FACTOR) == 0)
+             ? input / (NUM_THREADS * FACTOR)
+             : (input / (NUM_THREADS * FACTOR)) + 1;
+}
+
+int main(int argc, char **argv) {
+  if (argc != 4) {
+    printf("Need two arguments: percent add ops and percent delete ops (e.g., "
+           "30 50 for 30%% add and 50%% delete).\nAborting...\n");
+    exit(1);
+  }
+
+  auto sample_length = strtol(argv[2], nullptr, 10);
+  auto insertion_length = strtol(argv[2], nullptr, 10);
+  auto search_length = strtol(argv[2], nullptr, 10);
+  auto total_row = 0UL;
+
+  ReadHelper readHelper("normal_distribution.txt", sample_length,
+                        insertion_length + search_length);
+  readHelper.readFile(total_row);
+  std::vector<key_type> _sample, _population;
+  readHelper.split_into(_sample, _population);
+  printf("Create search vector: %ld\n",
+         _population.end() - _population.begin() + insertion_length);
+  std::vector<key_type> _search(_population.begin() + insertion_length,
+                                _population.end());
+  _population.resize(insertion_length);
+
+  // 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);
+
+  // Allocate device memory
+
+  LL *cudaOperatorItems;
+  // LL *Cop;
+  LL *Cresult;
+  LL *Clevels;
+
+  cudaMalloc(&Cresult, sizeof(LL) * NUM_ITEMS);
+  cudaMalloc(&cudaOperatorItems, sizeof(LL) * insertion_length);
+  // cudaMalloc(&Cop, sizeof(LL) * NUM_ITEMS);
+  cudaMalloc(&Clevels, sizeof(LL) * NUM_ITEMS);
+  // cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS,
+  // cudaMemcpyHostToDevice);
+  cudaMemcpy(cudaOperatorItems, _population.data(),
+             sizeof(key_type) * insertion_length, cudaMemcpyHostToDevice);
+  // cudaMemcpy(Cop, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice);
+  Node **pointers =
+      (Node **)new LL[insertion_length]; // malloc(sizeof(LL) * adds);
+  Node **Cpointers;
+
+  // Allocate the pool of free nodes
+
+  for (int i = 0; i < insertion_length; i++) {
+    cudaMalloc(&pointers[i], sizeof(Node));
+  }
+  cudaMalloc(&Cpointers, sizeof(Node *) * insertion_length);
+  cudaMemcpy(Cpointers, pointers, sizeof(Node *) * insertion_length,
+             cudaMemcpyHostToDevice);
+
+  // Allocate the skip list
+
+  LockFreeSkipList *Clist;
+  auto *list = new LockFreeSkipList(_sample.data(), _sample.size());
+
+  cudaMalloc(&Clist, sizeof(LockFreeSkipList));
+  cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice);
+  // 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
+
+  size_t blocks = calcBlocks(insertion_length);
+
+  CudaCheckError();
+
+  // Initialize the device memory
+  init<<<1, 32>>>(Clist, Cpointers, Clevels);
+  cudaDeviceSynchronize();
+
+  // Insertion to skiplist
+  kernelAdd<<<blocks, NUM_THREADS>>>(cudaOperatorItems, insertion_length);
+  cudaDeviceSynchronize();
+
+  // Re-allocate memory for search
+  cudaFree(cudaOperatorItems);
+  cudaMalloc(&cudaOperatorItems, sizeof(key_type) * search_length);
+  cudaMemcpy(cudaOperatorItems, _search.data(),
+             sizeof(key_type) * search_length, cudaMemcpyHostToDevice);
+
+  // Launch main kernel
+
+  cudaEvent_t start, stop;
+  cudaEventCreate(&start);
+  cudaEventCreate(&stop);
+  cudaEventRecord(start, nullptr);
+
+  blocks = calcBlocks(search_length);
+  kernel<<<blocks, NUM_THREADS>>>(cudaOperatorItems, search_length, Cresult);
+  CudaCheckError();
+  cudaDeviceSynchronize();
+  cudaEventRecord(stop, nullptr);
+  cudaEventSynchronize(stop);
+  float time;
+  cudaEventElapsedTime(&time, start, stop);
+  cudaEventDestroy(start);
+  cudaEventDestroy(stop);
+
+  // Print kernel execution time in milliseconds
+
+  printf("%lu: %lf", NUM_ITEMS, time);
+
+  // LL *Cop2;
+  //  cudaMalloc((void **)&Cop2, sizeof(LL) * NUM_ITEMS);
+  //  cudaMemcpy(Cop2, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice);
+
+  cudaEventCreate(&start);
+  cudaEventCreate(&stop);
+  cudaEventRecord(start, nullptr);
+  kernel<<<blocks, NUM_THREADS>>>(cudaOperatorItems, search_length, Cresult);
+  CudaCheckError();
+  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
+
+  // 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>>>();
+
+  cudaDeviceSynchronize();
+#ifdef MEASURE_TIME
+  {
+    int *cuda_tmp = nullptr;
+    cudaMalloc(&cuda_tmp, sizeof(int) * NUM_ITEMS);
+    copy_function<<<1, 1>>>(cuda_tmp);
+    cudaDeviceSynchronize();
+    CudaCheckError();
+    FILE *file = fopen("spend_time.txt", "w");
+    int *tmp = new int[NUM_ITEMS];
+    cudaMemcpy(tmp, cuda_tmp, sizeof(int) * NUM_ITEMS, cudaMemcpyDeviceToHost);
+    // memcpy(tmp, SpendTime, sizeof(int) * NUM_ITEMS);
+    for (int i = 0; i < NUM_ITEMS; i++) {
+      if (tmp[i] == 0)
+        break;
+      fprintf(file, "%d\n", tmp[i]);
+    }
+    // printf("%d\n", i);
+    delete[] tmp;
+    fclose(file);
+  }
+  // for (auto element : SpendTimeVec)
+  //   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;*/
+  return 0;
+}
diff --git a/skiplistcustom.cuh b/skiplistcustom.cuh
index a16e80a..80ecc19 100644
--- a/skiplistcustom.cuh
+++ b/skiplistcustom.cuh
@@ -1,5 +1,6 @@
 #pragma once
 #include "sortlib.cuh"
+#include <limits>
 #ifndef LOCKFREE_SKIPLIST_CUH_
 #define LOCKFREE_SKIPLIST_CUH_
 
@@ -85,7 +86,7 @@ public:
 // Definition of lock-free skip list
 
 class LockFreeSkipList {
-  static constexpr size_t SAMPLE_LENGTH = 1024;
+  // static constexpr size_t SAMPLE_LENGTH = 1023;
 
 public:
   Node *head;
@@ -99,20 +100,21 @@ public:
 
   double slice_size;
 
+  size_t sampleLength;
+
   LockFreeSkipList(key_type *samples, size_t sample_length)
-      : searcher(sample_length, sizeof(key_type)) {
+      : searcher(sample_length, sizeof(key_type)), sampleLength(sample_length) {
     Node *h = new Node(0);
     // size_ = 0;
 #if __WORDSIZE == 64
-    Node *t = new Node((key_type)NUM_ITEMS + 10);
+    Node *t = new Node(std::numeric_limits<key_type>::max());
 #else
     Node *t = new Node(0xffffffffULL);
 #endif
     cudaMalloc(&head, sizeof(Node));
 
     cudaMalloc(&tail, sizeof(Node));
-    int i;
-    for (i = 0; i < h->topLevel + 1; i++) {
+    for (int i = 0; i < h->topLevel + 1; i++) {
       h->SetRef(i, tail, false);
     }
     cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice);
@@ -123,8 +125,10 @@ public:
     initDeviceVariable();
 
     this->slice_size = 1.0 / sample_length;
+    // printf("%lf\n", this->slice_size);
 
-    cudaMalloc(&samples, sizeof(key_type) * sample_length);
+    printf("%ld\n", sample_length);
+    cudaMalloc(&this->samples, sizeof(key_type) * sample_length);
     cudaMemcpy(this->samples, samples, sizeof(key_type) * sample_length,
                cudaMemcpyHostToDevice);
   }
@@ -150,9 +154,11 @@ public:
     return bits;
   }
 
-  __device__ size_t calcLevel(key_type k) {
+  __device__ size_t calcLevel(key_type k) const {
+    // printf("Samples\n");
     auto index =
         this->searcher.sample_cdf_custom_version(this->samples, k) / slice_size;
+    // printf("Samples finish\n");
     auto level = trailing_zeroes(index);
     return level;
   }
@@ -162,6 +168,12 @@ public:
   //__device__ bool Delete(key_type);
   __device__ bool Search(key_type) const;
 
+  __device__ void testSample() {
+    for (int i = 0; i < sampleLength; i++) {
+      printf("%lld\n", this->samples[i]);
+    }
+  }
+
 #ifdef MEASURE_ACCESS
   unsigned access_times = 0;
 
@@ -193,6 +205,7 @@ public:
 
 __device__ Node *GetNewNode(key_type key, unsigned int *pointerIndex,
                             int topLevel) {
+  printf("%d\n", topLevel);
   key_type ind = atomicInc(pointerIndex, NUM_ITEMS);
   Node *n = nodes[ind];
   n->key = key;
@@ -328,7 +341,9 @@ __device__ bool LockFreeSkipList::Search(key_type key) const {
 }*/
 
 __device__ bool LockFreeSkipList::Add(key_type key) {
-  Node *newNode = GetNewNode(key, this->pointerIndex, calcLevel(key));
+  auto level1 = calcLevel(key) - 1;
+  printf("%lu\n", level1);
+  Node *newNode = GetNewNode(key, this->pointerIndex, level1);
   int topLevel = newNode->topLevel;
   int bottomLevel = 0;
   Node *preds[MAX_LEVEL + 1];
diff --git a/sortlib.cuh b/sortlib.cuh
index ae0ea67..4c339ca 100644
--- a/sortlib.cuh
+++ b/sortlib.cuh
@@ -55,6 +55,7 @@ public:
                                                     const key_type val) const {
 
     // int step_limit = (int)fast_log(LENGTH);
+    // printf("Binary search\n");
     key_type *last_known_point = start;
     auto son = 0UL;