From 9c0ce2f5570dedd5ad68c7e8fc222928046c5908 Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Sun, 28 May 2023 23:20:31 +0800 Subject: fix(exp): Fix search action Signed-off-by: KunoiSayami --- expt_0525.cu | 63 ++++++++++++++++++++++++++++++++++++------------------------ sortlib.cuh | 8 +++++++- 2 files changed, 45 insertions(+), 26 deletions(-) diff --git a/expt_0525.cu b/expt_0525.cu index b4e4da5..cff4f73 100644 --- a/expt_0525.cu +++ b/expt_0525.cu @@ -65,20 +65,21 @@ Conference on Parallel and Distributed Systems, December 2012. // #include"cutil.h" // Comment this if cutil.h is not available // #include "cuda_runtime.h" #include "sortlib.cuh" +#include #include #include #include #include #include -#define READ_NO_OUTPUT -#include "read_helper.h" - typedef unsigned long long LL; // #define MEASURE_TIME // #define MEASURE_ACCESS +#define READ_NO_OUTPUT +#include "read_helper.h" + #if (defined(MEASURE_ACCESS) && defined(MEASURE_TIME)) #error "Shouldn't define MEASURE_TIME and MEASURE_ACCESS at the same time" #endif @@ -260,12 +261,14 @@ class LockFreeSkipList { CustomSort customSort; + double scaleSize; + public: Node *head = nullptr; - Node *tail = nullptr; - LockFreeSkipList(key_type *_sample, size_t sample_length) + Node *tail; + LockFreeSkipList(key_type *_sample, size_t sample_length, double scale_size) : sampleLength(sample_length), - customSort(sample_length, sizeof(key_type) * 8) { + customSort(sample_length, sizeof(key_type) * 8), scaleSize(scale_size) { Node *h = new Node(0); // size_ = 0; Node *t = new Node(std::numeric_limits::max() - 1); @@ -309,6 +312,13 @@ 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; + } + #ifdef MEASURE_ACCESS unsigned access_times = 0; @@ -588,6 +598,7 @@ __global__ void kernel(const LL *items, size_t search_length, LL *result) { 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]) { @@ -611,14 +622,6 @@ __global__ void kernelAdd(LL *item, size_t insertion_length) { } } -/*LL Randomlevel() { - LL v = 1; - double p = 0.5; - while (((rand() / (double)(RAND_MAX)) < p) && (v < MAX_LEVEL)) - v++; - return v; -}*/ - // Generate the level of a newly created node __global__ void print_function() { @@ -631,6 +634,10 @@ __global__ void print_function() { memcpy(spend_time, lockFreeSkipList->spend_time, sizeof(int) * NUM_ITEMS); }*/ +inline double calcSliceSize(size_t insertion_size) { + return 1.0 / (double)insertion_size; +} + inline size_t calcBlocks(size_t input) { return (input % (NUM_THREADS * FACTOR) == 0) ? input / (NUM_THREADS * FACTOR) @@ -639,27 +646,32 @@ inline size_t calcBlocks(size_t input) { int main(int argc, char **argv) { if (argc != 4) { - printf("Usage %s [sample] [insertion] [search]\n", argv[0]); + printf("Usage %s [sample] [search] [insertion]\n", argv[0]); exit(1); } auto sample_length = strtol(argv[1], nullptr, 10); - auto insertion_length = strtol(argv[2], nullptr, 10); - auto search_length = strtol(argv[3], nullptr, 10); + auto search_length = strtol(argv[2], nullptr, 10); + auto insertion_length = strtol(argv[3], nullptr, 10); + + if (insertion_length < sample_length) { + printf("Search should smaller than insertion\n"); + } + printf("Sample: %ld, Insertion: %ld, Search: %ld\n", sample_length, insertion_length, search_length); ReadHelper readHelper("normal_distribution.txt", sample_length, - insertion_length + search_length); + insertion_length); readHelper.readFile(nullptr); std::vector _sample, _population; readHelper.split_into(_sample, _population); /*printf("%lu, Create search vector: %ld\n", _population.size(), _population.end() - (_population.begin() + insertion_length));*/ - rebuild(_sample); - std::vector _search(_population.begin() + insertion_length, - _population.end()); - _population.resize(insertion_length); + rebuildSort(_sample); + std::vector _search(_population.begin(), + _population.begin() + search_length); + //_population.resize(insertion_length); // Allocate necessary arrays // LL *op = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); @@ -673,8 +685,8 @@ int main(int argc, char **argv) { // LL *Cop; LL *cudaResult; - cudaMalloc(&cudaResult, sizeof(LL) * search_length); - cudaMalloc(&cudaOperatorItems, sizeof(LL) * insertion_length); + cudaMalloc(&cudaResult, sizeof(key_type) * search_length); + cudaMalloc(&cudaOperatorItems, sizeof(key_type) * insertion_length); // cudaMalloc(&Cop, sizeof(LL) * NUM_ITEMS); // cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS, // cudaMemcpyHostToDevice); @@ -697,7 +709,8 @@ int main(int argc, char **argv) { // Allocate the skip list LockFreeSkipList *Clist; - auto *list = new LockFreeSkipList(_sample.data(), _sample.size()); + auto *list = new LockFreeSkipList(_sample.data(), _sample.size(), + calcSliceSize(insertion_length)); cudaMalloc(&Clist, sizeof(LockFreeSkipList)); cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice); diff --git a/sortlib.cuh b/sortlib.cuh index 6b19d3e..a7d3e2d 100644 --- a/sortlib.cuh +++ b/sortlib.cuh @@ -1,6 +1,7 @@ #ifndef LOCKFREE_SORTLIB_CUH #define LOCKFREE_SORTLIB_CUH +#include #include #include #include @@ -182,7 +183,7 @@ public: __host__ __device__ size_t length() const { return this->LENGTH; } }; -template void rebuild(std::vector original) { +template void rebuild(std::vector &original) { auto sample_length = original.size(); auto sorter = CustomSort(sample_length, sizeof(T) * 8); auto tmp = new T[sample_length]; @@ -195,6 +196,11 @@ template void rebuild(std::vector original) { delete[] tmp; } +template void rebuildSort(std::vector &original) { + std::sort(original.begin(), original.end()); + rebuild(original); +} + __global__ void testCustomCalculation(size_t length) { CustomSort(length, sizeof(long) * 8).testCalculation(); } -- cgit v1.3.1