/*

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

  /// 將值存儲到 vector 的幫助函式
  /// 該函式可以幫助最將最大值和最小值更新
  /// 因為其資料需要存在上下界，所以該函式在存入的時候可以更新上下界的值
  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;

  /// 用於生成讀入時要跳過的數的數量
  auto genRandomRow(size_t total_row) const {
    /// 檢查是否滿足需要跳過行數的條件
    if (total_row != 0) {
      /// 檢查總行數是否少於需要讀入的數量相加
      assert(total_row > (population_length + sample_length));

      return randomRow(total_row - population_length - sample_length);
    }
    /// 否則返回 0 則會不跳過
    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)) {}

  /// 用於存放讀入的數值
  std::vector<key_type> population_vector;

  /// 生成 [0, max_value_] 之間隨機整數的函式
  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;
    /// 用 ifstream 類打開文件，準備進行讀入
    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--;
    }

    /// 重設已讀取數值為 0
    read_number = 0;

    /// 將剩下的資料存入 vector 中
    auto remain = population_length;
    for (key_type i; read_number < remain && !fin.eof(); store_into_vector(i)) {
      fin >> i;
      read_number++;
    }

    /// 關閉文件
    fin.close();
    return true;
  }

  /// 將讀入后的數值根據大小放置到指定的vector中
  void split_into(std::vector<key_type> &sample, std::vector<key_type> &p) {
    /// 把樣例 vector 的大小設為需要的大小
    sample.resize(sample_length - 2);
    /// 把插入的 vector 的大小設為需要的大小
    p.resize(population_length);
    /// 將資料複製到指定的樣例 vector 中
    memcpy(sample.data(), population_vector.data(),
           sizeof(key_type) * (sample_length - 2));
    /// 將極值放進樣例中
    sample.push_back(this->max_value);
    sample.push_back(this->min_value);
    /// 將資料複製到指定的插入 vector 中
    memcpy(p.data(), 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;

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;

  /// log2 函式，用於和 CPU 使用的 LOG 做區分
  __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++) {

      /// 判斷在該層是否指標已經與其相等
      if (*last_known_point == val) {
        return last_known_point;
      }

      /// 計算出該層的層級
      const auto next_level_start = start + (1 << (i + 1)) - 1;

      /// 為了避免分支分歧，我們用位運算來代替 if-else 結構
      /// 如果兩數相減 > 0 會使該變數為 1 ,否則為 0
      auto branch_selector = ((*last_known_point - val) >> MOVE_OFFSET);
      /// 將上一層節點乘以2即可得到下一層指標節點的偏移位置
      son = son * 2 + branch_selector;
      /// 將最後的指標更新
      last_known_point = next_level_start + son;
    }
    /// 返回最後訪問的位置，即最接近搜尋元素的位置
    return last_known_point;
  }

  /// 用於計算 CDF 的函式
  __device__ __host__ double cdf(key_type *start, key_type x) const {
    /// 首先通過搜尋得到最接近這個數的位置
    auto it = this->binary_search(start, x);
    /// 通過逆向演算法算出該數的實際排名
    auto prev_real_location = calculate_rank(it - start) - 1;

    /// 如果算出的排名是末尾元素，則返回 1
    if (prev_real_location == this->LENGTH) {
      return 1;
    }

    /// 如果算出的排名是頭元素，則返回 0
    if (prev_real_location == 0) {
      return 0;
    }

    /// 有過逆向演算法算出該數排名前一位的實際位置
    auto it_prev = start + calculate_index(prev_real_location - 1) - 1;

    /// 通過公式返回 CDF 值
    return ((double)prev_real_location +
            (double)(x - *it_prev) / (double)(*it - *it_prev)) /
           (double)(this->LENGTH - 1);
  }
};

/// 該函式將 vector 得的值重組成二元搜尋樹可用的結構
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 Node;

// Definition of generic node class

/// 基本節點類的定義，該節點對齊 16 字節
class __attribute__((aligned(16))) Node {
public:
  /// 定義存放該節點級別的變數
  int topLevel;
  /// 該變數用於存放該節點的值
  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
  /// 用來創建指向下一個節點的指標
  /// 如果 mark 被標記為 1 則為該節點已經被刪除
  static __device__ __host__ key_type CreateRef(Node *ref, bool mark) {
    auto val = (key_type)ref;
    /// 如果 mark 為 false （會被轉換為 0），則指標值不變
    /// 如果 mark 為 true （會被轉換為 1），則指標的二進制個位數變為 1
    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
  /// 該函式是 atomicCAS 的包裝函式，用來在有 mark
  /// 時讓函式可以按照預期的效果正常工作
  /// CAS 是 CompareAndSet
  /// 的簡寫，該原子操作用來確保設置值時不會被另一個操作篡改 該 atomicCAS
  /// 操作也是無鎖(Lock-free)的核心
  __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]);
    /// 進行 atomicCAS 操作，同時判斷其是否正確地交換
    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(nullptr, false);
    }
  }
};

// Definition of lock-free skip list

/// 定義 Skip Lists 的類
class LockFreeSkipList {

  /// 該指標指向存儲樣例的陣列
  key_type *sample = nullptr;

  /// 用於存儲樣例長度的變數
  size_t sampleLength;

  /// 用於查詢的改進的二元樹搜尋類
  CustomSort customSort;

  /// 存儲總共需要插入的大小
  /// 用於做資料大小放縮
  /// 因為 CDF 演算後是 0-1 之間的值，所以我們需要
  double insertSize;

public:
  Node *head = nullptr;
  Node *tail = nullptr;
  LockFreeSkipList(key_type *_sample, size_t sample_length, size_t insertSize)
      : sampleLength(sample_length),
        /// 初始化二元搜索樹類
        customSort(sample_length, sizeof(key_type) * 8),
        insertSize((double)insertSize) {
    /// 宣告一個頭節點
    Node *h = new Node(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);
    }

    /// 將頭和尾指標分別複製到 GPU 記憶體中
    cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice);
    cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice);

    /// 分配用於存放樣例的記憶體
    cudaMalloc(&this->sample, sizeof(key_type) * sampleLength);
    /// 將樣例複製到 GPU 的記憶體中
    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 Search(key_type);

  /// 用來計算出在完美二元樹條件下該節點位置應有的高度
  /// 實際上是計算該位置的末尾在二進制下有多少個 0
  /// 末尾 0 的個數則為該節點的高度
  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;
  }

  /// 包裝的函式，用於計算 CDF 以及高度
  __device__ unsigned calcLevel(key_type key) {
    /// 宣告一個變數來存儲 CDF 的值
    auto result = customSort.cdf(this->sample, key);
    /// 計算 CDF 縮放後的實際位置
    auto cdf_index = result * insertSize;
    /// 宣告一個變數用來存儲從位置計算出的高度
    auto level = trailing_zeroes((size_t)cdf_index);
    return level;
  }
};

/// 用於存儲節點的 GPU 指標
__device__ Node **nodes;
/// 用於使用中標示節點的使用情況的指標
__device__ unsigned int pointerIndex = 0;
/// 用於在 GPU 中標識節點池的使用情況
__device__ unsigned int NODE_LIMIT;

/// 在Skip Lists 中生成新的節點，高度為 topLevel 所標識的高度
__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 = (int)topLevel;
  /// 利用迴圈初始化所有高的的指標
  for (int i = 0; i < n->topLevel + 1; i++) {
    n->SetRef(i, nullptr, false);
  }
  return n;
}

/// 存放在 GPU 中的 Skip Lists 的指標
__device__ LockFreeSkipList *lockFreeSkipList; // The lock-free skip list

// Kernel for initializing device memory

/// 將值賦予到全域變數中
/// 分別為 Skip Lists 的指標，預分配的節點，以及可插入個數的最大數值
__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
  const int bottomLevel = 0;
  bool marked[1]{};
  bool snip;
  Node *pred;
  Node *curr;
  Node *succ;
  bool beenThereDoneThat;
  while (true) {
    beenThereDoneThat = false;
    pred = head;
    /// 從最上層到最下層搜尋
    for (int 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::Add(key_type key) {
  /// 宣告一個指標來存放需要插入的節點
  /// 參數中會先呼叫計算高度的函式算出高度
  Node *newNode = GetNewNode(key, calcLevel(key));
  /// 宣告一個變數來存儲節點的最高高度
  int topLevel = newNode->topLevel;
  /// 宣告一個變數來表示最低的層級
  constexpr int bottomLevel = 0;
  /// 宣告兩個指標陣列分別用來指向前繼和後繼
  Node *preds[MAX_LEVEL + 1];
  Node *succs[MAX_LEVEL + 1];
  while (true) {
    /// 先嘗試搜尋是否已經在資料結構中可以找到這個值
    bool found = find(key, preds, succs);
    if (found) {
      /// 如果找到，則返回未能插入
      return false;
    } else {
      /// 宣告兩個指標來存儲前繼節點和後繼節點
      Node *pred;
      Node *succ;
      /// 從最高層向下設置下一層的節點
      for (int level = bottomLevel; level <= topLevel; level++) {
        succ = succs[level];
        newNode->SetRef(level, succ, false);
      }
      /// 將指標分別設定為當前位置的前繼和後繼節點
      pred = preds[bottomLevel];
      succ = succs[bottomLevel];
      // printf("--- key is %d pred is %d succ is %d level is %d
      // \n",(int)key,(int)pred->key,(int)succ->key,0);
      /// 先從最底層插入該節點
      bool set_success =
          pred->CompareAndSet(bottomLevel, succ, newNode, false, false);
      /// 如果失敗，則重新開始這個過程
      if (!set_success) {
        continue;
      }
      /// 插入成功後，從底層向上插入節點
      for (int 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);
      /// 插入成功，返回 true
      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]);
  }
}

/// 計算 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);
  /// 將搜尋的 vector 從插入陣列的部分獨立出來
  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(), 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;
}