1 files changed, 0 insertions, 893 deletions

diff --git a/work.cu b/work.cu
deleted file mode 100644
index 46bdfb4..0000000
--- a/work.cu
+++ /dev/null
@@ -1,893 +0,0 @@
-/*
-
-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 <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 = 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
-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
-  __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
-  Node(LL 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 {
-public:
-  Node *head;
-  Node *tail;
-  LockFreeSkipList() {
-    Node *h = new Node(0);
-    // size_ = 0;
-#if __WORDSIZE == 64
-    Node *t = new Node((LL)NUM_ITEMS + 10);
-#else
-    Node *t = new Node((LL)0xffffffff);
-#endif
-    cudaMalloc((void **)&head, sizeof(Node));
-
-    cudaMalloc((void **)&tail, sizeof(Node));
-    int i;
-    for (i = 0; i < h->topLevel + 1; i++) {
-      h->SetRef(i, tail, false);
-    }
-    cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice);
-
-    cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice);
-  }
-  __device__ bool find(LL, Node **, Node **); // Helping method
-  __device__ bool Add(LL);
-  __device__ bool Delete(LL);
-  __device__ bool Search(LL);
-
-#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) {
-  LL 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->SetRef(i, nullptr, false);
-  }
-  return n;
-}
-
-__device__ LockFreeSkipList *l; // 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;
-  l = 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);
-  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);
-        }
-      }
-      return true;
-    }
-  }
-}
-
-__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");
-  }
-}
-
-// The main kernel
-
-__global__ void kernel(LL *items, LL *op, 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
-
-  int tid, i;
-  for (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)
-      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;
-    }
-  }
-  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() {
-#ifdef MEASURE_ACCESS
-  printf("count: %u\n", l->getAccessCount());
-#endif
-}
-
-__global__ void copy_function(int *spend_time) {
-  memcpy(spend_time, l->spend_time, sizeof(int) * NUM_ITEMS);
-}
-
-__device__ key_type *cudaBinarySearch(key_type *start, key_type *end,
-                                      key_type val) {
-  auto begin = start;
-  key_type *last_known_point = nullptr;
-  while (begin < end) {
-    auto mid = (end - begin) / 2;
-    auto mid_val = *(start + mid);
-    if (val == mid_val) {
-      return start + mid;
-    } else if (val > mid_val) {
-      begin = begin + mid + 1;
-    } else {
-      end = end - mid - 1;
-    }
-    last_known_point = begin;
-  }
-  return last_known_point;
-}
-
-__device__ long double sample_cdf(long double x) {
-  auto it = cudaBinarySearch(SampleStorage, SampleStorage + SAMPLE_SIZE, x);
-  if (it == SampleStorage + SAMPLE_SIZE) {
-    return 1;
-  }
-  if (it == SampleStorage) {
-    return 0;
-  }
-  auto it_prev = it - 1;
-  return (double(it_prev - SampleStorage) +
-          (x - (long double)*it_prev) / (long double)(*it - *it_prev)) /
-         double(SAMPLE_SIZE - 1);
-}
-
-int main(int argc, char **argv) {
-  if (argc != 3) {
-    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);
-  }
-
-  // Extract operations ratio
-  long adds = strtol(argv[1], nullptr, 10);
-  long deletes = strtol(argv[2], nullptr, 10);
-
-  storage.reserve(NUM_ITEMS);
-
-  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;*/
-
-    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;
-
-  // 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;
-
-  // Allocate the pool of free nodes
-
-  for (i = 0; i < adds; i++) {
-    cudaMalloc((void **)&pointers[i], sizeof(Node));
-  }
-  cudaMalloc((void **)&Cpointers, sizeof(Node *) * adds);
-  cudaMemcpy(Cpointers, pointers, sizeof(Node *) * adds,
-             cudaMemcpyHostToDevice);
-
-  // Allocate the skip list
-
-  LockFreeSkipList *Clist;
-  auto *list = new LockFreeSkipList();
-
-  cudaMalloc((void **)&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
-
-  int blocks = (NUM_ITEMS % (NUM_THREADS * FACTOR) == 0)
-                   ? NUM_ITEMS / (NUM_THREADS * FACTOR)
-                   : (NUM_ITEMS / (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, Clevels);
-  cudaDeviceSynchronize();
-
-  // Launch main kernel
-
-  cudaEvent_t start, stop;
-  cudaEventCreate(&start);
-  cudaEventCreate(&stop);
-  cudaEventRecord(start, nullptr);
-
-  kernel<<<blocks, NUM_THREADS>>>(Citems, Cop, Cresult);
-  CudaCheckError();
-  error = cudaGetLastError();
-  if (cudaSuccess != error) {
-    printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
-    // exit(-1);
-  }
-  cudaDeviceSynchronize();
-  cudaEventRecord(stop, nullptr);
-  cudaEventSynchronize(stop);
-  float time;
-  cudaEventElapsedTime(&time, start, stop);
-  cudaEventDestroy(start);
-  cudaEventDestroy(stop);
-
-  // Print kernel execution time in milliseconds
-
-  printf("%s %d ", TARGET_STRING, block_size);
-
-  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>>>();
-
-  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 (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;
-}