diff options
| -rw-r--r-- | expt_0707.cpp | 8 | ||||
| -rw-r--r-- | publish_0630.cu | 83 |
2 files changed, 29 insertions, 62 deletions
diff --git a/expt_0707.cpp b/expt_0707.cpp index b3e51ff..fb327fd 100644 --- a/expt_0707.cpp +++ b/expt_0707.cpp @@ -23,8 +23,9 @@ PRINTF_BYTE_TO_BINARY_INT32((i) >> 32), PRINTF_BYTE_TO_BINARY_INT32(i) /* --- end macros --- */ -template <typename T> void print_type_max(T _) { - std::cout << typeid(T).name() << "" +template <typename T> void print_type_max(T) { + std::cout << typeid(T).name() << "" << std::numeric_limits<size_t>::max() + << std::endl; } int main() { @@ -36,5 +37,6 @@ int main() { "= " PRINTF_BINARY_PATTERN_INT64 "\n", last_known_point, val, PRINTF_BYTE_TO_BINARY_INT64(last_known_point - val)); - std::cout << std::numeric_limits<size_t>::max() << std::endl; + + print_type_max((size_t)114514); }
\ No newline at end of file diff --git a/publish_0630.cu b/publish_0630.cu index 9cdd7d3..c0a7ced 100644 --- a/publish_0630.cu +++ b/publish_0630.cu @@ -243,7 +243,6 @@ public: size_t size() const { return this->population_vector.size(); } }; -/// typedef ReadHelper_<unsigned long long> ReadHelper; class CustomSort { @@ -291,19 +290,20 @@ public: /// 利用迴圈一直搜尋到完美最後二元樹的一層 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; + + /// 計算出下層的層級 + const auto next_level_start = start + (1 << (i + 1)) - 1; + /// 將最後的指標更新 last_known_point = next_level_start + son; } @@ -339,41 +339,39 @@ public: }; /// 該函式將 vector 得的值重組成二元搜尋樹可用的結構 -template <typename T> void rebuild(std::vector<T> &original) { +template <typename T> void rebuildSort(std::vector<T> &original) { + /// 將資料重新排序 + std::sort(original.begin(), original.end()); /// 先定義一個長度的變數 const auto sample_length = original.size(); + /// 宣告一個 CustomSort 類用來排序 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 key; /// 定義一個執行下一個節點的指標陣列 - key_type next[MAX_LEVEL + 1]{}; // Array of next links + key_type next[MAX_LEVEL + 1]{}; - // Create a next field from a reference and mark bit /// 用來創建指向下一個節點的指標 /// 如果 mark 被標記為 1 則為該節點已經被刪除 static __device__ __host__ key_type CreateRef(Node *ref, bool mark) { @@ -389,21 +387,18 @@ public: 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 @@ -425,20 +420,18 @@ public: return false; } - // Constructor for sentinel nodes /// 用來新建節點的建構子 explicit Node(key_type k) { + /// 將key存儲到自身 key = k; + /// 預設最高層為 MAX_LEVEL topLevel = MAX_LEVEL; - int i; - for (i = 0; i < MAX_LEVEL + 1; i++) { + for (int i = 0; i < MAX_LEVEL + 1; i++) { next[i] = CreateRef(nullptr, false); } } }; -// Definition of lock-free skip list - /// 定義 Skip Lists 的類 class LockFreeSkipList { @@ -548,26 +541,19 @@ __device__ Node *GetNewNode(key_type key, size_t topLevel) { } /// 存放在 GPU 中的 Skip Lists 的指標 -__device__ LockFreeSkipList *lockFreeSkipList; // The lock-free skip list - -// Kernel for initializing device memory +__device__ LockFreeSkipList *lockFreeSkipList; /// 將值賦予到全域變數中 /// 分別為 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 +__device__ bool LockFreeSkipList::find(key_type key, Node **preds, + Node **succs) { const int bottomLevel = 0; bool marked[1]{}; bool snip; @@ -708,22 +694,15 @@ __device__ bool LockFreeSkipList::Add(key_type key) { } } -// 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 + 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; - // Grab the operation and the associated key and execute key_type item = items[tid]; result[tid] = lockFreeSkipList->Search(item); assert(result[tid]); @@ -732,8 +711,7 @@ __global__ void kernel(const key_type *items, size_t search_length, __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 + for (int i = 0; i < FACTOR; i++) { auto tid = i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x; @@ -809,8 +787,6 @@ int main(int argc, char **argv) { std::vector<key_type> _search(_population.begin(), _population.begin() + search_length); - // Allocate device memory - /// 宣告存儲需要操作的陣列指標 key_type *cudaOperatorItems; // key_type *Cop; @@ -828,8 +804,6 @@ int main(int argc, char **argv) { 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)); @@ -840,8 +814,6 @@ int main(int argc, char **argv) { cudaMemcpy(Cpointers, pointers, sizeof(Node *) * insertion_length, cudaMemcpyHostToDevice); - // Allocate the skip list - /// 宣告一個指向 GPU 記憶體的 Skip lists 的指標 LockFreeSkipList *Clist; /// 在 CPU 中先把 Skip Lists 創建出來 @@ -852,26 +824,19 @@ int main(int argc, char **argv) { 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); /// 分配新的記憶體大小給用於搜尋的陣列 |