/*

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);
    }
  }
};

/*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 {
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);
        }
      }
      // 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 = 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);
}

std::vector<double> population;
std::vector<double> sample;

void initialize(const std::vector<double> &input_population,
                const std::vector<double> &input_sample) {
  population = input_population;
  std::sort(population.begin(), population.end());
  sample = input_sample;
  std::sort(sample.begin(), sample.end());
}

__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;
}