/*

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 <algorithm>
#include <cassert>
#include <cstdio>
#include <cstdlib>
#include <fstream>
#include <random>
#include <set>
#include <vector>

typedef unsigned long long key_type;

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

    for (key_type i; read_number < sample_length && !fin.eof();
         store_into_vector(i)) {
      fin >> i;
      read_number++;
    }

    if (random_number > 0) {
      read_number = random_number;
      for (key_type i; read_number > 0 && !fin.eof(); fin >> i)
        read_number--;
    }

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

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) {
  const 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) +
            ((double)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;
  }
};

__device__ Node **nodes;                  // Pool of pre-allocated nodes
__device__ unsigned int pointerIndex = 0; // Index into pool of free nodes
__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

// 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);
    while (true) {
      succ = curr->Get(level, &marked);
      while (marked) {
        curr = curr->GetReference(level);
        succ = curr->Get(level, &marked);
      }
      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];
    result[tid] = lockFreeSkipList->Search(item);
    assert(result[tid]);
  }
}

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

inline double calcSliceSize(size_t insertion_size) {
  return 1.0 / (double)insertion_size;
}

/// 計算 blocks 的大小，該大小和 NUM_THREADS
/// 相乘應能正好大於等於需要插入的大小
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;

  /// 因為該參數可選，所以需要進行判斷輸入的參數個數是否大於 4
  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);

  /// 宣告一個讀入 unsigned long long 的幫助類
  /// 將文件名和上面讀入的數值作為參數
  ReadHelper readHelper("normal_distribution.txt", sample_length,
                        insertion_length, total_row);

  /// 如果有指定需要隨機化讀入，則輸出跳過了多少的數才開始正式地讀取
  if (total_row) {
    printf("Skip: %ld ", readHelper.random_number);
    fflush(stdout);
  }

  /// 讀入文件
  readHelper.readFile();

  /// 宣告兩個 unsigned long long 的 vector ，用來存放樣例和插入的數值
  std::vector<key_type> _sample, _population;
  /// 調用讀入幫助類，將兩個數值分別賦值到上述宣告的 vector 中
  readHelper.split_into(_sample, _population);
  /// 重新排序
  rebuildSort(_sample);
  std::vector<key_type> _search(_population.begin(),
                                _population.begin() + search_length);

  // 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);
  /// 將需要查詢的物件複製到記憶體中
  cudaMemcpy(cudaOperatorItems, _population.data(),
             sizeof(key_type) * insertion_length, cudaMemcpyHostToDevice);
  /// 分配插入操作所需要的節點陣列
  Node **pointers = (Node **)new key_type[insertion_length];
  Node **Cpointers;

  // Allocate the pool of free nodes

  /// 給每個指標分配節點的記憶體
  for (int i = 0; i < insertion_length; i++) {
    cudaMalloc(&pointers[i], sizeof(Node));
  }
  /// 分配 GPU 處的節點記憶體
  cudaMalloc(&Cpointers, sizeof(Node *) * insertion_length);
  /// 將存儲節點的陣列複製到記憶體中
  cudaMemcpy(Cpointers, pointers, sizeof(Node *) * insertion_length,
             cudaMemcpyHostToDevice);

  // Allocate the skip list

  /// 宣告一個指向 GPU 記憶體的 Skip lists 的指標
  LockFreeSkipList *Clist;
  /// 在 CPU 中先把 Skip Lists 創建出來
  auto *list = new LockFreeSkipList(_sample.data(), _sample.size(),
                                    calcSliceSize(insertion_length));

  /// 分配 GPU 記憶體給 Skip Lists
  cudaMalloc(&Clist, sizeof(LockFreeSkipList));
  /// 将在 CPU 处建立的 Skip Lists 複製到 GPU 記憶體中
  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

  /// 計算 blocks 的大小，該大小和 NUM_THREADS
  /// 相乘應能正好大於等於需要插入的大小
  size_t blocks = calcBlocks(insertion_length);

  // Initialize the device memory
  /// 將需要的值賦值到 GPU 中的變數中
  init<<<1, 32>>>(Clist, Cpointers, insertion_length);
  cudaDeviceSynchronize();

  // Insertion to skiplist
  /// 該 GPU 語句啓動的 GPU 執行緒會將所需要的數據插入到 Skip lists 中
  kernelAdd<<<blocks, NUM_THREADS>>>(cudaOperatorItems, insertion_length);
  cudaDeviceSynchronize();

  // Re-allocate memory for search
  /// 將原來的記憶體釋放
  cudaFree(cudaOperatorItems);
  /// 分配新的記憶體大小給用於搜尋的數組
  cudaMalloc(&cudaOperatorItems, sizeof(key_type) * search_length);
  /// 將需要搜尋的數值複製到 GPU 的記憶體中
  cudaMemcpy(cudaOperatorItems, _search.data(),
             sizeof(key_type) * search_length, cudaMemcpyHostToDevice);

  /// 計算 blocks 的大小，該大小和 NUM_THREADS
  /// 相乘應能正好大於等於需要搜尋的大小
  blocks = calcBlocks(search_length);

  /// 創建用來計算時間的變數
  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start, nullptr);

  /// 執行 Kernel
  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);

  /// 列印出搜尋的長度以及所耗費的時間（單位為毫秒）
  printf("%lu: %lf\n", search_length, time);

  /// 分配用來存儲結果的記憶體
  auto *result = new key_type[search_length];

  /// 將存儲結果的陣列複製回 CPU 的記憶體中
  cudaMemcpy(result, cudaResult, sizeof(key_type) * search_length,
             cudaMemcpyDeviceToHost);

  /// 釋放記憶體
  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;
}