/*

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 "sortlib.cuh"
#include <cassert>
#include <cstdio>
#include <cstdlib>
#include <random>
#include <set>

typedef unsigned long long LL;

// #define MEASURE_TIME
//  #define MEASURE_ACCESS

#define READ_NO_OUTPUT
#include "read_helper.h"

#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 = 32;

// 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;

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 = 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(LL, Node **, Node **); // Helping method
  __device__ bool Add(LL);
  __device__ bool Delete(LL);
  __device__ bool Search(LL);
  //~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;
    LL 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__ LL *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(LL key, size_t topLevel) {
  LL 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__ LL 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(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;
  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, 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;
    }
  }
}

__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);
    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(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]);
  }
}

// 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 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("Usage %s [sample] [search] [insertion]\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);

  if (insertion_length < search_length) {
    printf("Search should smaller than insertion\n");
  }

  printf("Sample: %ld, Insertion: %ld, Search: %ld\n", sample_length,
         insertion_length, search_length);

  ReadHelper readHelper("normal_distribution.txt", sample_length,
                        insertion_length);
  readHelper.readFile(nullptr);
  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
  // 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[search_length]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS);

  // Allocate device memory

  LL *cudaOperatorItems;
  // LL *Cop;
  LL *cudaResult;

  cudaMalloc(&cudaResult, sizeof(key_type) * search_length);
  cudaMalloc(&cudaOperatorItems, sizeof(key_type) * insertion_length);
  // cudaMalloc(&Cop, 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(),
                                    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

  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start, nullptr);

  blocks = calcBlocks(search_length);
  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(LL) * 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(Cop2);
    cudaFree(Clevels);
    cudaFree(Cop);
    cudaFree(Citems);
    cudaFree(Cresult);
    free(pointers);
    delete [] op;
    delete [] levels;
    delete [] items;
    delete [] result;*/
  return 0;
}