feat: Add MEASURE_ACCESS and MEASURE_TIME macro

Signed-off-by: KunoiSayami <[email protected]>

1 files changed, 68 insertions, 12 deletions

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<LL> storage;
 
+  // std::normal_distribution<long double> 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);