From d3362a1d4c8af292365e16d466d3f57bb0af326f Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Wed, 15 Jun 2022 21:55:30 +0800 Subject: feat: Add MEASURE_ACCESS and MEASURE_TIME macro Signed-off-by: KunoiSayami --- main.cu | 80 +++++++++++++++++++++++++++++++++++++++++++++++++++++++---------- 1 file changed, 68 insertions(+), 12 deletions(-) (limited to 'main.cu') diff --git a/main.cu b/main.cu index b1239d7..fa5520b 100644 --- a/main.cu +++ b/main.cu @@ -84,9 +84,16 @@ typedef unsigned int LL; #define STEP_SIZE 2 #endif +//#define MEASURE_TIME +//#define MEASURE_ACCESS + +#if (defined(MEASURE_ACCESS) && defined(MEASURE_TIME)) +#error "Shouldn't define MEASURE_TIME and MEASURE_ACCESS at the same time" +#endif + // Maximum level of a node in the skip list //#define MAX_LEVEL 32 -constexpr size_t MAX_LEVEL = 32; +constexpr size_t MAX_LEVEL = 16; // Number of threads per block //#define NUM_THREADS 512 @@ -99,7 +106,7 @@ constexpr size_t FACTOR = 1; // should change this to dynamic next time constexpr size_t KEY_INDEX_SIZE = 32; -volatile int block_size = STEP_SIZE; +constexpr int block_size = STEP_SIZE; // Supported operations constexpr int ADD = 0; @@ -257,7 +264,6 @@ class LockFreeSkipList { public: Node *head; Node *tail; - unsigned size_ = 0; MemMap key_map; LockFreeSkipList() { Node *h = new Node(0); @@ -282,6 +288,25 @@ public: __device__ bool Add(LL); __device__ bool Delete(LL); __device__ bool Search(LL); + +#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 _count = 1) {} +#endif + +#ifdef MEASURE_TIME + unsigned round = 0; + __device__ void increaseRoundCount(unsigned count = 1) { + atomicAdd(&this->round, count); + } + __device__ unsigned getRoundCount() const { return this->round; } +#endif }; __device__ Node **nodes; // Pool of pre-allocated nodes @@ -373,13 +398,24 @@ __device__ bool LockFreeSkipList::Search(LL key) { Node *curr = nullptr; Node *succ; int level; +#ifdef MEASURE_TIME +#endif 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; @@ -389,6 +425,9 @@ __device__ bool LockFreeSkipList::Search(LL key) { } } } +#ifdef MEASURE_TIME + atomicAdd(&this->round, 1); +#endif return (curr != nullptr && curr->key == key); } @@ -420,7 +459,7 @@ __device__ bool LockFreeSkipList::Delete(LL key) { if (iMarkedIt) { find(key, preds, succs); // size_ -= 1; - atomicDec(&size_, 1); + // atomicDec(&size_, 1); return true; } else if (marked[0]) { return false; @@ -475,7 +514,7 @@ __device__ bool LockFreeSkipList::Add(LL key) { } // size_ += 1; // this->key_map.insert(MapNode(ll, newNode)); - atomicAdd(&size_, 1); + // atomicAdd(&size_, 1); return true; } } @@ -542,7 +581,7 @@ __global__ void kernel(LL *items, LL *op, LL *result) { }*/ // Generate the level of a newly created node -LL Randomlevel(std::mt19937 &randomEngine, double p) { +LL RandomLevel(std::mt19937 &randomEngine, double p) { std::geometric_distribution<> distribution(p); return std::min(MAX_LEVEL, (size_t)distribution(randomEngine)); } @@ -582,7 +621,14 @@ LL CustomLevel(LL value) { return 1; } -__global__ void print_function() { printf("size: %u\n", l->size_); } +__global__ void print_function() { +#ifdef MEASURE_ACCESS + printf("count: %u\n", l->getAccessCount()); +#endif +#ifdef MEASURE_TIME + printf("round: %u\n", l->getRoundCount()); +#endif +} int main(int argc, char **argv) { if (argc != 3) { @@ -618,15 +664,22 @@ int main(int argc, char **argv) { NUM_ITEMS - 1); // std::vector storage; + // std::normal_distribution normalDistribution{2147483647, + // 2147483647}; + for (i = 0; i < NUM_ITEMS; i++) { items[i] = i + 3; // 10+rand()%KEYS; - // Keys associated with - // operations + // Keys associated with + // operations storage.push_back(i + 3); + // auto key = (key_type)std::round(normalDistribution(randomEngine)); + // items[i] = key; + // storage.push_back(key); } std::sort(storage.begin(), storage.end()); +#if 0 for (i = 0; i < NUM_ITEMS; i++) { /*int first = rand() % NUM_ITEMS; int second = rand() % NUM_ITEMS;*/ @@ -638,12 +691,13 @@ int main(int argc, char **argv) { items[first] = items[second]; items[second] = temp;*/ } +#endif // Pre-generated levels of skip list nodes (relevant only if op[i] is add) // srand(0); for (i = 0; i < NUM_ITEMS; i++) { #ifdef RANDOM_HEIGHT - levels[i] = Randomlevel(1 / randomEngine) - 1; // 36/14 + levels[i] = RandomLevel(1 / randomEngine) - 1; // 36/14 #else levels[i] = CustomLevel(items[i]) - 1; // 31/18 #endif @@ -790,11 +844,13 @@ int main(int argc, char **argv) { // Uncomment the following for debugging // print<<<1,32>>>(); - // cudaDeviceSynchronize(); + cudaDeviceSynchronize(); - // print_function<<<1, 1>>>(); +#if (defined(MEASURE_TIME) || defined(MEASURE_ACCESS)) + print_function<<<1, 1>>>(); cudaDeviceSynchronize(); +#endif /*cudaFree(Clist); cudaFree(Cop2); cudaFree(Clevels); -- cgit v1.3.1