2023-07-01 00:51

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

1 files changed, 988 insertions, 0 deletions

diff --git a/publish_0630.cu b/publish_0630.cu
new file mode 100644
index 0000000..1befc23
--- /dev/null
+++ b/publish_0630.cu
@@ -0,0 +1,988 @@
+/*
+
+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 SHAkey_type 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 <cassert>
+#include <cstdio>
+#include <cstdlib>
+#include <random>
+#include <set>
+
+typedef unsigned long long key_type;
+
+#include "read_helper.h"
+
+// Maximum level of a node in the skip list
+constexpr size_t MAX_LEVEL = 32;
+
+// Number of threads per block
+constexpr size_t NUM_THREADS = 512;
+
+constexpr size_t FACTOR = 1;
+
+#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
+}
+
+template <typename key_type> class ReadHelper_ {
+  // typedef unsigned long long key_type;
+
+  key_type max_value = std::numeric_limits<key_type>::min(),
+           min_value = std::numeric_limits<key_type>::max();
+
+  inline void store_into_vector(key_type value) {
+    if (max_value < value) {
+      max_value = value;
+    }
+    if (min_value > value) {
+      min_value = value;
+    }
+    population_vector.push_back(value);
+  }
+
+  char const *filename;
+  static constexpr size_t REVERSED_BLOCK = 256;
+
+  auto genRandomRow(size_t total_row) const {
+    if (total_row != 0) {
+      assert(total_row > (population_length + sample_length + REVERSED_BLOCK));
+      return randomRow(total_row - population_length - sample_length -
+                       REVERSED_BLOCK);
+    }
+    return 0UL;
+  }
+
+public:
+  const size_t sample_length, population_length, random_number;
+  ReadHelper_(char const *filename, size_t sample_length,
+              size_t population_length, size_t total_row = 0)
+      : filename(filename), sample_length(sample_length),
+        population_length(population_length),
+        random_number(genRandomRow(total_row)) {}
+  key_type maxValue() const { return max_value; }
+  key_type minValue() const { return min_value; }
+
+  std::vector<key_type> population_vector;
+
+  static unsigned long randomRow(unsigned long max_value_) {
+    std::random_device randomDevice;
+    std::mt19937 mt19937(randomDevice());
+    std::uniform_int_distribution<std::mt19937::result_type> dst(0, max_value_);
+    return dst(mt19937);
+  }
+
+  bool readFile() {
+    this->population_vector.clear();
+    auto read_number = 0UL;
+    std::ifstream fin(filename);
+
+    if (!fin.is_open()) {
+      return false;
+    }
+
+    P_ERR("Reading sample");
+    for (key_type i; read_number < sample_length && !fin.eof();
+         store_into_vector(i)) {
+      fin >> i;
+      read_number++;
+    }
+
+    if (random_number > 0) {
+      P_ERR("\rReading skip");
+      read_number = random_number;
+      for (key_type i; read_number > 0 && !fin.eof(); fin >> i)
+        read_number--;
+    }
+
+    P_ERR("\rReading population");
+    read_number = 0;
+
+    auto remain = population_length + REVERSED_BLOCK;
+    for (key_type i; read_number < remain && !fin.eof(); store_into_vector(i)) {
+      fin >> i;
+      read_number++;
+    }
+
+    fin.close();
+    P_ERR("\r");
+    return true;
+  }
+
+  void split_into(std::vector<key_type> &sample, std::vector<key_type> &p) {
+    sample.resize(sample_length - 2);
+    p.resize(population_length);
+    // printf("%zu\n", needed_read_length - sample_length);
+    memcpy(sample.data(), population_vector.data(),
+           sizeof(key_type) * (sample_length - 2));
+    sample.push_back(this->max_value);
+    sample.push_back(this->min_value);
+    memcpy(p.data(), population_vector.data() + sample_length,
+           sizeof(key_type) * population_length);
+  }
+
+  void split_into(key_type *&sample, key_type *&p) {
+    sample = new key_type[sample_length];
+    p = new key_type[population_length];
+    memcpy(sample, population_vector.data(),
+           sizeof(key_type) * (sample_length - 2));
+    sample[sample_length - 2] = this->max_value;
+    sample[sample_length - 1] = this->min_value;
+    memcpy(p, population_vector.data() + sample_length,
+           sizeof(key_type) * (population_length));
+  }
+
+  size_t size() const { return this->population_vector.size(); }
+};
+
+typedef ReadHelper_<unsigned long long> ReadHelper;
+
+#include <algorithm>
+#include <cassert>
+#include <cstdio>
+#include <vector>
+
+typedef unsigned long long key_type;
+
+class CustomSort {
+public:
+  explicit __device__ __host__ CustomSort(size_t length, int move_offset)
+      : LENGTH(length), MOVE_OFFSET(move_offset - 1),
+        STEP_LIMIT(fast_log(LENGTH)) {}
+  const size_t LENGTH;
+
+  /// MOVE_OFFSET means bit to select branch
+  const int MOVE_OFFSET;
+  const unsigned int STEP_LIMIT;
+
+  __device__ __host__ static size_t fast_log(size_t a) {
+#ifdef __CUDA_ARCH__
+    return (size_t)log2((double)a);
+#else
+    return (size_t)std::log2(a);
+#endif
+  }
+
+  __device__ __host__ size_t calculate_index(size_t rank) const {
+    size_t bit_low = (LENGTH + 1) >> fast_log(++rank) >> 1;
+    return (((rank << 1) | 1) * bit_low - LENGTH - 1);
+  }
+
+  __device__ __host__ size_t calculate_rank(size_t index) const {
+    index++;
+    size_t low_bit = index & (-index);
+    return ((LENGTH + index) / low_bit) >> 1;
+  }
+
+  __device__ __host__ const key_type *binary_search(key_type *const start,
+                                                    const key_type val) const {
+
+    key_type *last_known_point = start;
+    auto son = 0UL;
+
+    for (int i = 0; i < STEP_LIMIT; i++) {
+      const auto next_level_start = start + (1 << (i + 1)) - 1;
+
+      if (*last_known_point == val) {
+        return last_known_point;
+      }
+      auto branch_selector = ((*last_known_point - val) >> MOVE_OFFSET);
+      son = son * 2 + branch_selector;
+      last_known_point = next_level_start + son;
+    }
+    return last_known_point;
+  }
+
+  /// Should be correct version
+  __device__ __host__ double sample_cdf_custom_version(key_type *start,
+                                                       key_type x) const {
+    auto it = this->binary_search(start, x);
+    auto prev_real_location = calculate_rank(it - start) - 1;
+
+    if (prev_real_location == this->LENGTH) {
+      return 1;
+    }
+
+    if (prev_real_location == 0) {
+      return 0;
+    }
+
+    auto it_prev = start + calculate_index(prev_real_location - 1) - 1;
+
+    return ((double)prev_real_location +
+            (double)(x - *it_prev) / (double)(*it - *it_prev)) /
+           (double)(this->LENGTH - 1);
+  }
+
+  __host__ __device__ size_t length() const { return this->LENGTH; }
+};
+
+template <typename T> void rebuild(std::vector<T> &original) {
+  auto sample_length = original.size();
+  auto sorter = CustomSort(sample_length, sizeof(T) * 8);
+  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;
+}
+
+template <typename T> void rebuildSort(std::vector<T> &original) {
+  std::sort(original.begin(), original.end());
+  rebuild(original);
+}
+
+class FactorySort {
+
+public:
+  __device__ __host__ static const key_type *
+  cudaBinarySearch(key_type *start, const key_type *end, const key_type val) {
+    auto begin = start;
+    key_type *last_known_point = begin;
+    assert(begin < end);
+    while (begin <= end) {
+      auto mid = begin + (end - begin) / 2;
+      auto mid_val = *mid;
+      if (val == mid_val) {
+        return mid;
+      } else if (val > mid_val) {
+        begin = mid + 1;
+      } else {
+        end = mid - 1;
+      }
+      last_known_point = begin;
+    }
+    return last_known_point;
+  }
+
+  __device__ __host__ double static sample_cdf(key_type *begin,
+                                               unsigned long length,
+                                               key_type x) {
+    // printf("%f\n", x);
+    auto end = begin + length;
+    auto it = cudaBinarySearch(begin, end, x);
+    if (it == end) {
+      return 1;
+    }
+    if (it == begin) {
+      return 0;
+    }
+    auto it_prev = it - 1;
+    return (double(it_prev - begin) +
+            (x - (double)*it_prev) / (double)(*it - *it_prev)) /
+           double(length - 1);
+  }
+};
+
+// class Node;
+
+// Definition of generic node class
+
+class
+#ifndef _MSC_VER
+    __attribute__((aligned(16)))
+#else
+    __declspec(align(16))
+#endif
+    Node {
+public:
+  int topLevel;                   // Level of the node
+  key_type key;                   // Key value
+  key_type next[MAX_LEVEL + 1]{}; // Array of next links
+
+  // Create a next field from a reference and mark bit
+  static __device__ __host__ key_type CreateRef(Node *ref, bool mark) {
+    auto val = (key_type)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) {
+    key_type 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) {
+    key_type oldVal = (key_type)expectedRef | oldMark;
+    key_type newVal = (key_type)newRef | newMark;
+    key_type *ref = &(next[index]);
+    key_type oldValOut = atomicCAS(ref, oldVal, newVal);
+    if (oldValOut == oldVal)
+      return true;
+    return false;
+  }
+
+  // Constructor for sentinel nodes
+  explicit Node(key_type k) {
+    key = k;
+    topLevel = MAX_LEVEL;
+    int i;
+    for (i = 0; i < MAX_LEVEL + 1; i++) {
+      next[i] = CreateRef((Node *)nullptr, false);
+    }
+  }
+};
+
+// Definition of lock-free skip list
+
+class LockFreeSkipList {
+
+  key_type *sample = nullptr;
+  size_t sampleLength;
+
+  CustomSort customSort;
+
+  double scaleSize;
+
+public:
+  Node *head = nullptr;
+  Node *tail = nullptr;
+  LockFreeSkipList(key_type *_sample, size_t sample_length, double scale_size)
+      : sampleLength(sample_length),
+        customSort(sample_length, sizeof(key_type) * 8), scaleSize(scale_size) {
+    Node *h = new Node(0);
+    // size_ = 0;
+    Node *t = new Node(std::numeric_limits<key_type>::max() - 1);
+    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(key_type, Node **, Node **); // Helping method
+  __device__ bool Add(key_type);
+  __device__ bool Delete(key_type);
+  __device__ bool Search(key_type);
+  //~LockFreeSkipList(){cudaFree()}
+
+  __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;
+  }
+
+  static __device__ unsigned trailing_zeroes(size_t index) {
+    constexpr auto block_size = 2;
+    unsigned bits = 0;
+    key_type x = index / block_size;
+
+    if (x) {
+      while (x % block_size == 0) {
+        ++bits;
+        x /= block_size;
+      }
+    }
+    return bits;
+  }
+
+  __device__ unsigned calcIndex(key_type key) {
+    auto result = customSort.sample_cdf_custom_version(this->sample, key);
+    auto cdf_index = result / scaleSize;
+    auto index = trailing_zeroes((size_t)cdf_index);
+    return index;
+  }
+
+#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 = 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__ key_type *randoms; // Array storing the levels of the nodes in the
+// free
+// pool
+__device__ unsigned int NODE_LIMIT;
+
+// Function for creating a new node when requested by an add operation
+
+__device__ Node *GetNewNode(key_type key, size_t topLevel) {
+  key_type ind = atomicInc(&pointerIndex, NODE_LIMIT);
+  Node *n = nodes[ind];
+  n->key = key;
+  // n->topLevel = randoms[ind];
+  n->topLevel = (int)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__ key_type KeyIndex[KEY_INDEX_SIZE];
+
+//__device__ key_type SampleStorage[SAMPLE_SIZE];
+
+// Kernel for initializing device memory
+
+__global__ void init(LockFreeSkipList *l1, Node **n,
+                     unsigned int insertion_limit) {
+  // randoms = rands;
+  nodes = n;
+  lockFreeSkipList = l1;
+  NODE_LIMIT = insertion_limit;
+}
+
+// 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
+  int bottomLevel = 0;
+  bool marked[] = {false};
+  bool snip;
+  Node *pred;
+  Node *curr;
+  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(key_type 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(key_type 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(key_type key) {
+  Node *newNode = GetNewNode(key, calcIndex(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;
+    }
+  }
+}
+
+// The main kernel
+
+__global__ void kernel(const key_type *items, size_t search_length,
+                       key_type *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
+    key_type item = items[tid];
+#ifdef MEASURE_TIME
+    unsigned long long start_time = clock64();
+#endif
+    result[tid] = lockFreeSkipList->Search(item);
+    assert(result[tid]);
+#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(key_type *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]);
+  }
+}
+
+// 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 double calcSliceSize(size_t insertion_size) {
+  return 1.0 / (double)insertion_size;
+}
+
+inline auto 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("Usage %s <sample> <search> <insertion> [file length]\n", argv[0]);
+    exit(1);
+  }
+
+  auto sample_length = strtol(argv[1], nullptr, 10);
+  auto search_length = strtol(argv[2], nullptr, 10);
+  auto insertion_length = strtol(argv[3], nullptr, 10);
+  auto total_row = 0L;
+
+  if (argc > 4) {
+    total_row = strtol(argv[4], nullptr, 10);
+  }
+
+  if (insertion_length < search_length) {
+    printf("Search should smaller than insertion\n");
+  }
+
+  printf("Sample: %ld, Insertion: %ld, Search: %ld ", sample_length,
+         insertion_length, search_length);
+  fflush(stdout);
+
+  ReadHelper readHelper("normal_distribution.txt", sample_length,
+                        insertion_length, total_row);
+
+  if (total_row) {
+    printf("Skip: %ld ", readHelper.random_number);
+    fflush(stdout);
+  }
+
+  readHelper.readFile();
+  std::vector<key_type> _sample, _population;
+  readHelper.split_into(_sample, _population);
+  /*printf("%lu, Create search vector: %ld\n", _population.size(),
+         _population.end() - (_population.begin() + insertion_length));*/
+  rebuildSort(_sample);
+  std::vector<key_type> _search(_population.begin(),
+                                _population.begin() + search_length);
+  //_population.resize(insertion_length);
+
+  // Allocate necessary arrays
+  // key_type *op = new key_type[NUM_ITEMS];     //(key_type
+  // *)malloc(sizeof(key_type) * NUM_ITEMS); key_type *levels = new
+  // key_type[NUM_ITEMS]; //(key_type *)malloc(sizeof(key_type) * NUM_ITEMS);
+  // key_type *items = new key_type[NUM_ITEMS];  //(key_type
+  // *)malloc(sizeof(key_type) * NUM_ITEMS);
+  auto *result =
+      new key_type[search_length]; //(key_type *)malloc(sizeof(key_type) *
+                                   // NUM_ITEMS);
+
+  // Allocate device memory
+
+  key_type *cudaOperatorItems;
+  // key_type *Cop;
+  key_type *cudaResult;
+
+  cudaMalloc(&cudaResult, sizeof(key_type) * search_length);
+  cudaMalloc(&cudaOperatorItems, sizeof(key_type) * insertion_length);
+  // cudaMalloc(&Cop, sizeof(key_type) * NUM_ITEMS);
+  // cudaMemcpy(Clevels, levels, sizeof(key_type) * NUM_ITEMS,
+  // cudaMemcpyHostToDevice);
+  cudaMemcpy(cudaOperatorItems, _population.data(),
+             sizeof(key_type) * insertion_length, cudaMemcpyHostToDevice);
+  // cudaMemcpy(Cop, op, sizeof(key_type) * NUM_ITEMS, cudaMemcpyHostToDevice);
+  Node **pointers =
+      (Node *
+           *)new key_type[insertion_length]; // malloc(sizeof(key_type) * 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(),
+                                    calcSliceSize(insertion_length));
+
+  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, insertion_length);
+  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
+  blocks = calcBlocks(search_length);
+
+  cudaEvent_t start, stop;
+  cudaEventCreate(&start);
+  cudaEventCreate(&stop);
+  cudaEventRecord(start, nullptr);
+
+  kernel<<<blocks, NUM_THREADS>>>(cudaOperatorItems, search_length, cudaResult);
+  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\n", search_length, time);
+  // Check for errors
+
+  // Move results back to host memory
+
+  cudaMemcpy(result, cudaResult, sizeof(key_type) * search_length,
+             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(cudaResult);
+  cudaFree(cudaOperatorItems);
+  cudaFree(Cpointers);
+  for (int i = 0; i < insertion_length; i++) {
+    cudaFree(pointers[i]);
+  }
+  cudaFree(pointers);
+  delete[] result;
+  delete list;
+  return 0;
+}