diff options
| -rw-r--r-- | publish_0630.cu | 192 |
1 files changed, 59 insertions, 133 deletions
diff --git a/publish_0630.cu b/publish_0630.cu index 1befc23..107b024 100644 --- a/publish_0630.cu +++ b/publish_0630.cu @@ -66,16 +66,17 @@ 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; -#include "read_helper.h" - // Maximum level of a node in the skip list constexpr size_t MAX_LEVEL = 32; @@ -175,7 +176,6 @@ public: return false; } - P_ERR("Reading sample"); for (key_type i; read_number < sample_length && !fin.eof(); store_into_vector(i)) { fin >> i; @@ -183,13 +183,11 @@ public: } if (random_number > 0) { - P_ERR("\rReading skip"); read_number = random_number; for (key_type i; read_number > 0 && !fin.eof(); fin >> i) read_number--; } - P_ERR("\rReading population"); read_number = 0; auto remain = population_length + REVERSED_BLOCK; @@ -199,7 +197,6 @@ public: } fin.close(); - P_ERR("\r"); return true; } @@ -231,11 +228,6 @@ public: typedef ReadHelper_<unsigned long long> ReadHelper; -#include <algorithm> -#include <cassert> -#include <cstdio> -#include <vector> - typedef unsigned long long key_type; class CustomSort { @@ -312,7 +304,7 @@ public: }; template <typename T> void rebuild(std::vector<T> &original) { - auto sample_length = original.size(); + const auto sample_length = original.size(); auto sorter = CustomSort(sample_length, sizeof(T) * 8); auto tmp = new T[sample_length]; @@ -366,7 +358,7 @@ public: } auto it_prev = it - 1; return (double(it_prev - begin) + - (x - (double)*it_prev) / (double)(*it - *it_prev)) / + ((double)x - (double)*it_prev) / (double)(*it - *it_prev)) / double(length - 1); } }; @@ -499,34 +491,10 @@ public: auto index = trailing_zeroes((size_t)cdf_index); return index; } - -#ifdef MEASURE_ACCESS - unsigned access_times = 0; - - __device__ unsigned getAccessCount() const { return this->access_times; } - __device__ void increaseAccessCount(unsigned count = 1) { - atomicAdd(&this->access_times, count); - } -#else - __device__ void increaseAccessCount(unsigned = 1) {} -#endif - -#ifdef MEASURE_TIME - unsigned round = 0; - __device__ void increaseRoundCount(unsigned count = 1) { - atomicAdd(&this->round, count); - } - int spend_time[NUM_ITEMS]{0}; - unsigned long long total_time = 0; - __device__ unsigned getRoundCount() const { return this->round; } -#endif }; __device__ Node **nodes; // Pool of pre-allocated nodes __device__ unsigned int pointerIndex = 0; // Index into pool of free nodes -//__device__ key_type *randoms; // Array storing the levels of the nodes in the -// free -// pool __device__ unsigned int NODE_LIMIT; // Function for creating a new node when requested by an add operation @@ -546,10 +514,6 @@ __device__ Node *GetNewNode(key_type key, size_t topLevel) { __device__ LockFreeSkipList *lockFreeSkipList; // The lock-free skip list -//__device__ key_type KeyIndex[KEY_INDEX_SIZE]; - -//__device__ key_type SampleStorage[SAMPLE_SIZE]; - // Kernel for initializing device memory __global__ void init(LockFreeSkipList *l1, Node **n, @@ -620,20 +584,11 @@ __device__ bool LockFreeSkipList::Search(key_type key) { int level; for (level = MAX_LEVEL; level >= bottomLevel; level--) { curr = pred->GetReference(level); -#ifdef MEASURE_ACCESS - this->increaseAccessCount(); -#endif while (true) { succ = curr->Get(level, &marked); -#ifdef MEASURE_ACCESS - this->increaseAccessCount(); -#endif while (marked) { curr = curr->GetReference(level); succ = curr->Get(level, &marked); -#ifdef MEASURE_ACCESS - this->increaseAccessCount(2); -#endif } if (curr->key < key) { pred = curr; @@ -752,18 +707,8 @@ __global__ void kernel(const key_type *items, size_t search_length, // Grab the operation and the associated key and execute key_type item = items[tid]; -#ifdef MEASURE_TIME - unsigned long long start_time = clock64(); -#endif result[tid] = lockFreeSkipList->Search(item); assert(result[tid]); -#ifdef MEASURE_TIME - unsigned long long end_time = clock64() - start_time; - if (lockFreeSkipList->spend_time[tid]) { - printf("conflict: %d\n", tid); - } - lockFreeSkipList->spend_time[tid] = (int)end_time; -#endif } } @@ -780,22 +725,12 @@ __global__ void kernelAdd(key_type *item, size_t insertion_length) { } } -// Generate the level of a newly created node - -__global__ void print_function() { -#ifdef MEASURE_ACCESS - printf("count: %u\n", l->getAccessCount()); -#endif -} - -/*__global__ void copy_function(int *spend_time) { - memcpy(spend_time, lockFreeSkipList->spend_time, sizeof(int) * NUM_ITEMS); -}*/ - inline double calcSliceSize(size_t insertion_size) { return 1.0 / (double)insertion_size; } +/// 計算 blocks 的大小，該大小和 NUM_THREADS +/// 相乘應能正好大於等於需要插入的大小 inline auto calcBlocks(size_t input) { return (input % (NUM_THREADS * FACTOR) == 0) ? input / (NUM_THREADS * FACTOR) @@ -808,121 +743,145 @@ int main(int argc, char **argv) { 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); - /*printf("%lu, Create search vector: %ld\n", _population.size(), - _population.end() - (_population.begin() + insertion_length));*/ + /// 重新排序 rebuildSort(_sample); std::vector<key_type> _search(_population.begin(), _population.begin() + search_length); - //_population.resize(insertion_length); - - // Allocate necessary arrays - // key_type *op = new key_type[NUM_ITEMS]; //(key_type - // *)malloc(sizeof(key_type) * NUM_ITEMS); key_type *levels = new - // key_type[NUM_ITEMS]; //(key_type *)malloc(sizeof(key_type) * NUM_ITEMS); - // key_type *items = new key_type[NUM_ITEMS]; //(key_type - // *)malloc(sizeof(key_type) * NUM_ITEMS); - auto *result = - new key_type[search_length]; //(key_type *)malloc(sizeof(key_type) * - // NUM_ITEMS); // 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); - // cudaMalloc(&Cop, sizeof(key_type) * NUM_ITEMS); - // cudaMemcpy(Clevels, levels, sizeof(key_type) * NUM_ITEMS, - // cudaMemcpyHostToDevice); + /// 將需要查詢的物件複製到記憶體中 cudaMemcpy(cudaOperatorItems, _population.data(), sizeof(key_type) * insertion_length, cudaMemcpyHostToDevice); - // cudaMemcpy(Cop, op, sizeof(key_type) * NUM_ITEMS, cudaMemcpyHostToDevice); - Node **pointers = - (Node * - *)new key_type[insertion_length]; // malloc(sizeof(key_type) * adds); + /// 分配插入操作所需要的節點陣列 + 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(), calcSliceSize(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); - CudaCheckError(); - // 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); - // Launch main kernel + /// 計算 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; @@ -930,50 +889,17 @@ int main(int argc, char **argv) { cudaEventDestroy(start); cudaEventDestroy(stop); - // Print kernel execution time in milliseconds - + /// 列印出搜尋的長度以及所耗費的時間（單位為毫秒） printf("%lu: %lf\n", search_length, time); - // Check for errors - // Move results back to host memory + /// 分配用來存儲結果的記憶體 + auto *result = new key_type[search_length]; + /// 將存儲結果的陣列複製回 CPU 的記憶體中 cudaMemcpy(result, cudaResult, sizeof(key_type) * search_length, cudaMemcpyDeviceToHost); - // Uncomment the following for debugging - // print<<<1,32>>>(); - cudaDeviceSynchronize(); - -#if (defined(MEASURE_TIME) || defined(MEASURE_ACCESS)) - - print_function<<<1, 1>>>(); - - cudaDeviceSynchronize(); -#ifdef MEASURE_TIME - { - int *cuda_tmp = nullptr; - cudaMalloc(&cuda_tmp, sizeof(int) * NUM_ITEMS); - copy_function<<<1, 1>>>(cuda_tmp); - cudaDeviceSynchronize(); - CudaCheckError(); - FILE *file = fopen("spend_time.txt", "w"); - int *tmp = new int[NUM_ITEMS]; - cudaMemcpy(tmp, cuda_tmp, sizeof(int) * NUM_ITEMS, cudaMemcpyDeviceToHost); - // memcpy(tmp, SpendTime, sizeof(int) * NUM_ITEMS); - for (int i = 0; i < NUM_ITEMS; i++) { - if (tmp[i] == 0) - break; - fprintf(file, "%d\n", tmp[i]); - } - // printf("%d\n", i); - delete[] tmp; - fclose(file); - } - // for (auto element : SpendTimeVec) - // printf("%d\n", element); -#endif -#endif - + /// 釋放記憶體 cudaFree(Clist); cudaFree(cudaResult); cudaFree(cudaOperatorItems); |