diff options
Diffstat (limited to 'publish_0630.cu')
| -rw-r--r-- | publish_0630.cu | 63 |
1 files changed, 42 insertions, 21 deletions
diff --git a/publish_0630.cu b/publish_0630.cu index 424aa7c..583b46f 100644 --- a/publish_0630.cu +++ b/publish_0630.cu @@ -127,9 +127,9 @@ template <typename key_type> class ReadHelper_ { key_type max_value = std::numeric_limits<key_type>::min(), min_value = std::numeric_limits<key_type>::max(); - /// 將值存儲到 vector 的幫助函數 - /// 該函數可以幫助最將最大值和最小值更新 - /// 因為其資料需要存在上下界，所以該函數在存入的時候可以更新上下界的值 + /// 將值存儲到 vector 的幫助函式 + /// 該函式可以幫助最將最大值和最小值更新 + /// 因為其資料需要存在上下界，所以該函式在存入的時候可以更新上下界的值 inline void store_into_vector(key_type value) { if (max_value < value) { max_value = value; @@ -172,7 +172,7 @@ public: /// 用於存放讀入的數值 std::vector<key_type> population_vector; - /// 生成 [0, max_value_] 之間隨機整數的函數 + /// 生成 [0, max_value_] 之間隨機整數的函式 static unsigned long randomRow(unsigned long max_value_) { std::random_device randomDevice; /// 用於生成隨機整數的發生器 @@ -204,7 +204,7 @@ public: read_number++; } - /// 如果在讀入時需要跳過數值，則執行下面的函數 + /// 如果在讀入時需要跳過數值，則執行下面的函式 if (random_number > 0) { /// 將隨機數賦值到本地變數中 read_number = random_number; @@ -264,7 +264,7 @@ public: /// 用來存儲二元樹的上限 const unsigned int STEP_LIMIT; - /// log2 函數，用於和 CPU 使用的 LOG 做區分 + /// log2 函式，用於和 CPU 使用的 LOG 做區分 __device__ __host__ static size_t fast_log(size_t a) { #ifdef __CUDA_ARCH__ return (size_t)log2((double)a); @@ -409,8 +409,8 @@ public: } // CompareAndSet wrapper - /// 該函數是 atomicCAS 的包裝函數，用來在有 mark - /// 時讓函數可以按照預期的效果正常工作 + /// 該函式是 atomicCAS 的包裝函式，用來在有 mark + /// 時讓函式可以按照預期的效果正常工作 /// CAS 是 CompareAndSet /// 的簡寫，該原子操作用來確保設置值時不會被另一個操作篡改 該 atomicCAS /// 操作也是無鎖(Lock-free)的核心 @@ -518,8 +518,11 @@ public: /// 包裝的函式，用於計算 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; } @@ -534,7 +537,7 @@ __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]; @@ -570,8 +573,8 @@ __global__ void init(LockFreeSkipList *l1, Node **n, __device__ bool LockFreeSkipList::find(key_type key, Node **preds, Node **succs) { // preds and succs are arrays of pointers - int bottomLevel = 0; - bool marked[] = {false}; + const int bottomLevel = 0; + bool marked[1]{}; bool snip; Node *pred; Node *curr; @@ -580,14 +583,17 @@ LockFreeSkipList::find(key_type key, Node **preds, while (true) { beenThereDoneThat = false; pred = head; - 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[0]) { snip = pred->CompareAndSet(level, curr, succ, false, false); beenThereDoneThat = true; + /// 如果操作失敗，則表明該值已經被變更 + /// 由前項跳出後重新搜尋 if (!snip) break; curr = pred->GetReference(level); @@ -642,33 +648,45 @@ __device__ bool LockFreeSkipList::Search(key_type key) { } __device__ bool LockFreeSkipList::Add(key_type key) { + /// 宣告一個指標來存放需要插入的節點 + /// 參數中會先呼叫計算高度的函式算出高度 Node *newNode = GetNewNode(key, calcLevel(key)); + /// 宣告一個變數來存儲節點的最高高度 int topLevel = newNode->topLevel; - int bottomLevel = 0; + /// 宣告一個變數來表示最低的層級 + constexpr 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++) { + /// 從最高層向下設置下一層的節點 + for (int 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) { + /// 先從最底層插入該節點 + bool set_success = + pred->CompareAndSet(bottomLevel, succ, newNode, false, false); + /// 如果失敗，則重新開始這個過程 + if (!set_success) { continue; } - for (level = bottomLevel + 1; level <= topLevel; level++) { + /// 插入成功後，從底層向上插入節點 + for (int level = bottomLevel + 1; level <= topLevel; level++) { while (true) { pred = preds[level]; @@ -676,17 +694,20 @@ __device__ bool LockFreeSkipList::Add(key_type key) { 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; } } @@ -785,7 +806,7 @@ int main(int argc, char **argv) { /// 宣告兩個 unsigned long long 的 vector ，用來存放樣例和插入的數值 std::vector<key_type> _sample, _population; - /// 調用讀入幫助類，將兩個數值分別賦值到上述宣告的 vector 中 + /// 呼叫讀入幫助類，將兩個數值分別賦值到上述宣告的 vector 中 readHelper.split_into(_sample, _population); /// 重新排序 rebuildSort(_sample); |