2023-07-15 03:58HEADmaster

Signed-off-by: KunoiSayami <[email protected]>

1 files changed, 41 insertions, 35 deletions

diff --git a/publish_0630.cu b/publish_0630.cu
index c0a7ced..ea3e05f 100644
--- a/publish_0630.cu
+++ b/publish_0630.cu
@@ -73,6 +73,7 @@ Conference on Parallel and Distributed Systems, December 2012.
 #include <vector>
 
 typedef unsigned long long key_type;
+typedef unsigned long long pKey;
 
 // Maximum level of a node in the skip list
 constexpr size_t MAX_LEVEL = 32;
@@ -282,7 +283,7 @@ public:
 
   __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;
@@ -369,13 +370,13 @@ public:
   int topLevel;
   /// 該變數用於存放該節點的值
   key_type key;
-  /// 定義一個執行下一個節點的指標陣列
-  key_type next[MAX_LEVEL + 1]{};
+  /// 定義一個指向下一個節點的指標陣列
+  pKey next[MAX_LEVEL + 1]{};
 
   /// 用來創建指向下一個節點的指標
   /// 如果 mark 被標記為 1 則為該節點已經被刪除
-  static __device__ __host__ key_type CreateRef(Node *ref, bool mark) {
-    auto val = (key_type)ref;
+  static __device__ __host__ pKey CreateRef(Node *ref, bool mark) {
+    auto val = (pKey)ref;
     /// 如果 mark 為 false （會被轉換為 0），則指標值不變
     /// 如果 mark 為 true （會被轉換為 1），則指標的二進制個位數變為 1
     val = val | mark;
@@ -387,9 +388,9 @@ public:
     next[index] = CreateRef(ref, mark);
   }
 
-  /// 用來獲得指標指向的下一個值的實際指標
+  /// 用來取得指標指向的下一個值的實際指標
   __device__ Node *GetReference(int index) {
-    key_type ref = next[index];
+    pKey ref = next[index];
     return (Node *)((ref >> 1) << 1);
   }
 
@@ -407,13 +408,13 @@ public:
   __device__ bool CompareAndSet(int index, Node *expectedRef, Node *newRef,
                                 bool oldMark, bool newMark) {
     /// 宣告一個變數來存儲舊的值
-    key_type oldVal = (key_type)expectedRef | oldMark;
+    pKey oldVal = (pKey)expectedRef | oldMark;
     /// 宣告一個變數來存儲新的值
-    key_type newVal = (key_type)newRef | newMark;
+    pKey newVal = (pKey)newRef | newMark;
     /// 宣告一個指標來存儲需要替換的目標位址
-    key_type *ref = &(next[index]);
+    pKey *ref = &(next[index]);
     /// 進行 atomicCAS 操作，同時判斷其是否正確地交換
-    key_type oldValOut = atomicCAS(ref, oldVal, newVal);
+    pKey oldValOut = atomicCAS(ref, oldVal, newVal);
     /// 如果輸出值為舊值，則操作成功，否則失敗
     if (oldValOut == oldVal)
       return true;
@@ -426,8 +427,9 @@ public:
     key = k;
     /// 預設最高層為 MAX_LEVEL
     topLevel = MAX_LEVEL;
-    for (int i = 0; i < MAX_LEVEL + 1; i++) {
-      next[i] = CreateRef(nullptr, false);
+    /// 遍歷所有指標，將指標初始化為空指標
+    for (unsigned long long &element : next) {
+      element = CreateRef(nullptr, false);
     }
   }
 };
@@ -483,7 +485,7 @@ public:
   }
   /// 該函式用於在插入的時候查詢目標節點是否被標記為刪除
   /// 搜尋函式不需要用到該函式
-  __device__ bool find(key_type, Node **, Node **); // Helping method
+  __device__ bool find(key_type, Node **, Node **);
   __device__ bool Add(key_type);
   __device__ bool Search(key_type);
 
@@ -556,13 +558,12 @@ __device__ bool LockFreeSkipList::find(key_type key, Node **preds,
                                        Node **succs) {
   const int bottomLevel = 0;
   bool marked[1]{};
-  bool snip;
+  /// 宣告三個指標分別用來指向前繼和後繼與目前搜尋的節點
   Node *pred;
   Node *curr;
   Node *succ;
-  bool beenThereDoneThat;
   while (true) {
-    beenThereDoneThat = false;
+    bool beenThereDoneThat = false;
     pred = head;
     /// 從最上層到最下層搜尋
     for (int level = MAX_LEVEL; level >= bottomLevel; level--) {
@@ -571,7 +572,7 @@ __device__ bool LockFreeSkipList::find(key_type key, Node **preds,
         succ = curr->Get(level, marked);
         /// 如被標記為已刪除，則繼續搜尋
         while (marked[0]) {
-          snip = pred->CompareAndSet(level, curr, succ, false, false);
+          bool snip = pred->CompareAndSet(level, curr, succ, false, false);
           beenThereDoneThat = true;
           /// 如果操作失敗，則表明該值已經被變更
           /// 由前項跳出後重新搜尋
@@ -580,8 +581,8 @@ __device__ bool LockFreeSkipList::find(key_type key, Node **preds,
           curr = pred->GetReference(level);
           succ = curr->Get(level, marked);
           beenThereDoneThat = false;
-          // printf("find key is %d \n",(int)key);
         }
+        /// 如果已經把key清理過，則跳出迴圈該層迴圈
         if (beenThereDoneThat)
           break;
         if (curr->key <= key) {
@@ -602,29 +603,39 @@ __device__ bool LockFreeSkipList::find(key_type key, Node **preds,
   }
 }
 
+/// 用於搜尋的函式
 __device__ bool LockFreeSkipList::Search(key_type key) {
-  int bottomLevel = 0;
+  const int bottomLevel = 0;
+  /// 源於存儲該節點是否被標記刪除
   bool marked = false;
+  /// 宣告一系列指標來存儲必須的節點
   Node *pred = head;
   Node *curr = nullptr;
   Node *succ;
-  int level;
-  for (level = MAX_LEVEL; level >= bottomLevel; level--) {
+  /// 從最上層開始搜尋
+  for (int 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);
 }
 
@@ -641,7 +652,7 @@ __device__ bool LockFreeSkipList::Add(key_type key) {
   Node *succs[MAX_LEVEL + 1];
   while (true) {
     /// 先嘗試搜尋是否已經在資料結構中可以找到這個值
-    bool found = find(key, preds, succs);
+    bool found = Search(key);
     if (found) {
       /// 如果找到，則返回未能插入
       return false;
@@ -657,8 +668,6 @@ __device__ bool LockFreeSkipList::Add(key_type key) {
       /// 將指標分別設定為當前位置的前繼和後繼節點
       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);
@@ -673,21 +682,14 @@ __device__ bool LockFreeSkipList::Add(key_type key) {
           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;
     }
@@ -696,15 +698,20 @@ __device__ bool LockFreeSkipList::Add(key_type key) {
 
 __global__ void kernel(const key_type *items, size_t search_length,
                        key_type *result) {
-
+  /// 執行緒所需操作數量的迴圈
   for (int i = 0; i < FACTOR; i++) {
+    /// 計算出當前執行緒所需要運算的值
     auto tid =
         i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
+    /// 確保需要計算的值不會超過搜尋數目的上限
     if (tid >= search_length)
       return;
 
+    /// 宣告一個變數存放取出的值
     key_type item = items[tid];
+    /// 呼叫 Skip Lists 的搜尋函式搜尋取出的值
     result[tid] = lockFreeSkipList->Search(item);
+    /// 確保返回的值為 true 即搜尋成功
     assert(result[tid]);
   }
 }
@@ -732,7 +739,7 @@ inline auto calcBlocks(size_t input) {
 int main(int argc, char **argv) {
   if (argc < 4) {
     printf("Usage %s <sample> <search> <insertion> [file length]\n", argv[0]);
-    exit(1);
+    return 1;
   }
 
   /// 從命令行中讀取樣例的大小
@@ -789,7 +796,6 @@ int main(int argc, char **argv) {
 
   /// 宣告存儲需要操作的陣列指標
   key_type *cudaOperatorItems;
-  // key_type *Cop;
   /// 宣告存儲結果陣列的指標
   key_type *cudaResult;