diff options
Diffstat (limited to 'publish_0630.cu')
| -rw-r--r-- | publish_0630.cu | 297 |
1 files changed, 149 insertions, 148 deletions
diff --git a/publish_0630.cu b/publish_0630.cu index 107b024..ce61d25 100644 --- a/publish_0630.cu +++ b/publish_0630.cu @@ -12,7 +12,7 @@ this list of conditions, and the following disclaimer. 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 +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 @@ -120,12 +120,18 @@ inline void __cudaCheckError(const char *file, const int line) { #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; @@ -136,111 +142,131 @@ template <typename key_type> class ReadHelper_ { 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); } + /// 否則返回 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)) {} - key_type maxValue() const { return max_value; } - key_type minValue() const { return min_value; } + /// 用於存放讀入的數值 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 + REVERSED_BLOCK; 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); - // printf("%zu\n", needed_read_length - sample_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); } - 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; + /// log2 函數，用於和 CPU 使用的 LOG 做區分 __device__ __host__ static size_t fast_log(size_t a) { #ifdef __CUDA_ARCH__ return (size_t)log2((double)a); @@ -249,11 +275,13 @@ public: #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); @@ -262,49 +290,66 @@ 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; + /// 利用迴圈一直搜尋到完美最後二元樹的一層 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; } + + /// 計算出該層的層級 + 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; } - /// Should be correct version - __device__ __host__ double sample_cdf_custom_version(key_type *start, - key_type x) const { + /// 用於計算 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); } - - __host__ __device__ size_t length() const { return this->LENGTH; } }; +/// 該函式將 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]; @@ -321,156 +366,144 @@ template <typename T> void rebuildSort(std::vector<T> &original) { 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 { +/// 基本節點類的定義，該節點對齊 16 字節 +class __attribute__((aligned(16))) Node { public: - int topLevel; // Level of the node - key_type key; // Key value + /// 定義存放該節點級別的變數 + 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((Node *)nullptr, false); + next[i] = CreateRef(nullptr, false); } } }; // Definition of lock-free skip list +/// 定義 Skip Lists 的類 class LockFreeSkipList { + /// 該指標指向存儲樣例的陣列 key_type *sample = nullptr; + + /// 用於存儲樣例長度的變數 size_t sampleLength; + /// 用於查詢的改進的二元樹搜尋類 CustomSort customSort; - double scaleSize; + /// 存儲總共需要插入的大小 + /// 用於做資料大小放縮 + /// 因為 CDF 演算後是 0-1 之間的值，所以我們需要 + double insertSize; public: Node *head = nullptr; Node *tail = nullptr; - LockFreeSkipList(key_type *_sample, size_t sample_length, double scale_size) + LockFreeSkipList(key_type *_sample, size_t sample_length, size_t insertSize) : sampleLength(sample_length), - customSort(sample_length, sizeof(key_type) * 8), scaleSize(scale_size) { + /// 初始化二元搜索樹類 + customSort(sample_length, sizeof(key_type) * 8), + insertSize((double)insertSize) { + /// 宣告一個頭節點 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); + /// 將頭和尾指標分別複製到 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 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; - } + /// 用來計算出在完美二元樹條件下該節點位置應有的高度 + /// 實際上是計算該位置的末尾在二進制下有多少個 0 + /// 末尾 0 的個數則為該節點的高度 static __device__ unsigned trailing_zeroes(size_t index) { constexpr auto block_size = 2; unsigned bits = 0; @@ -485,37 +518,46 @@ public: 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; + /// 包裝的函式，用於計算 CDF 以及高度 + __device__ unsigned calcLevel(key_type key) { + auto result = customSort.cdf(this->sample, key); + auto cdf_index = result * insertSize; + auto level = trailing_zeroes((size_t)cdf_index); + return level; } }; -__device__ Node **nodes; // Pool of pre-allocated nodes -__device__ unsigned int pointerIndex = 0; // Index into pool of free nodes +/// 用於存儲節點的 GPU 指標 +__device__ Node **nodes; +/// 用於使用中標示節點的使用情況的指標 +__device__ unsigned int pointerIndex = 0; +/// 用於在 GPU 中標識節點池的使用情況 __device__ unsigned int NODE_LIMIT; -// Function for creating a new node when requested by an add operation - +/// 在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 = randoms[ind]; + /// 賦值該指標的最大值 n->topLevel = (int)topLevel; - int i; - for (i = 0; i < n->topLevel + 1; i++) { + /// 利用迴圈初始化所有高的的指標 + 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; @@ -601,46 +643,8 @@ __device__ bool LockFreeSkipList::Search(key_type key) { 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)); + Node *newNode = GetNewNode(key, calcLevel(key)); int topLevel = newNode->topLevel; int bottomLevel = 0; Node *preds[MAX_LEVEL + 1]; @@ -725,10 +729,6 @@ __global__ void kernelAdd(key_type *item, size_t insertion_length) { } } -inline double calcSliceSize(size_t insertion_size) { - return 1.0 / (double)insertion_size; -} - /// 計算 blocks 的大小，該大小和 NUM_THREADS /// 相乘應能正好大於等於需要插入的大小 inline auto calcBlocks(size_t input) { @@ -791,6 +791,7 @@ int main(int argc, char **argv) { readHelper.split_into(_sample, _population); /// 重新排序 rebuildSort(_sample); + /// 將搜尋的 vector 從插入陣列的部分獨立出來 std::vector<key_type> _search(_population.begin(), _population.begin() + search_length); @@ -830,8 +831,8 @@ int main(int argc, char **argv) { /// 宣告一個指向 GPU 記憶體的 Skip lists 的指標 LockFreeSkipList *Clist; /// 在 CPU 中先把 Skip Lists 創建出來 - auto *list = new LockFreeSkipList(_sample.data(), _sample.size(), - calcSliceSize(insertion_length)); + auto *list = + new LockFreeSkipList(_sample.data(), _sample.size(), insertion_length); /// 分配 GPU 記憶體給 Skip Lists cudaMalloc(&Clist, sizeof(LockFreeSkipList)); @@ -859,7 +860,7 @@ int main(int argc, char **argv) { // Re-allocate memory for search /// 將原來的記憶體釋放 cudaFree(cudaOperatorItems); - /// 分配新的記憶體大小給用於搜尋的數組 + /// 分配新的記憶體大小給用於搜尋的陣列 cudaMalloc(&cudaOperatorItems, sizeof(key_type) * search_length); /// 將需要搜尋的數值複製到 GPU 的記憶體中 cudaMemcpy(cudaOperatorItems, _search.data(), |