feat: Finish implement measure time function

* feat: Add normal_distribution program

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

1 files changed, 54 insertions, 3 deletions

diff --git a/main.cu b/main.cu
index fa5520b..524b6d3 100644
--- a/main.cu
+++ b/main.cu
@@ -84,13 +84,19 @@ typedef unsigned int LL;
 #define STEP_SIZE 2
 #endif
 
-//#define MEASURE_TIME
+#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
 
+#ifdef MEASURE_TIME
+#include <cuda/atomic>
+#undef BUILD_SIZE
+#define BUILD_SIZE 1024
+#endif
+
 // Maximum level of a node in the skip list
 //#define MAX_LEVEL 32
 constexpr size_t MAX_LEVEL = 16;
@@ -258,6 +264,43 @@ public:
   __device__ ~MemMap() { delete[] store; }
 };
 
+class CudaSpinLock {
+  static constexpr int UNLOCKED = 0;
+  static constexpr int LOCKED = 1;
+
+  cuda::atomic<int> m_value;
+  bool isFake;
+
+public:
+  __device__ __host__ explicit CudaSpinLock()
+      : m_value(UNLOCKED), isFake(false) {}
+
+  __device__ __host__ explicit CudaSpinLock(bool fake)
+      : m_value(UNLOCKED), isFake(fake) {}
+
+  __device__ void lock() {
+    if (!isFake) {
+      while (true) {
+        int expected = UNLOCKED;
+        // this->m_value.wait(LOCKED);
+        if (this->m_value.compare_exchange_weak(expected, LOCKED))
+          break;
+      }
+    }
+  }
+
+  __device__ void unlock() {
+    if (!isFake) {
+      m_value.store(UNLOCKED);
+    }
+  }
+
+  __device__ bool isLock() {
+    // printf("%d\n", this->m_value.load());
+    return !isFake && this->m_value.load() == LOCKED;
+  }
+};
+
 // Definition of lock-free skip list
 
 class LockFreeSkipList {
@@ -302,9 +345,11 @@ public:
 
 #ifdef MEASURE_TIME
   unsigned round = 0;
+  CudaSpinLock spinLock;
   __device__ void increaseRoundCount(unsigned count = 1) {
     atomicAdd(&this->round, count);
   }
+  unsigned long long total_time = 0;
   __device__ unsigned getRoundCount() const { return this->round; }
 #endif
 };
@@ -399,6 +444,8 @@ __device__ bool LockFreeSkipList::Search(LL key) {
   Node *succ;
   int level;
 #ifdef MEASURE_TIME
+  this->spinLock.lock();
+  auto start_time = clock64();
 #endif
   for (level = MAX_LEVEL; level >= bottomLevel; level--) {
     curr = pred->GetReference(level);
@@ -426,7 +473,11 @@ __device__ bool LockFreeSkipList::Search(LL key) {
     }
   }
 #ifdef MEASURE_TIME
+  this->spinLock.unlock();
+  unsigned long long end_time = clock64() - start_time;
+  // printf("%lld\n", end_time);
   atomicAdd(&this->round, 1);
+  atomicAdd(&this->total_time, end_time);
 #endif
   return (curr != nullptr && curr->key == key);
 }
@@ -626,7 +677,8 @@ __global__ void print_function() {
   printf("count: %u\n", l->getAccessCount());
 #endif
 #ifdef MEASURE_TIME
-  printf("round: %u\n", l->getRoundCount());
+  printf("round: %u time: %llu avg: %.2lf\n", l->getRoundCount(), l->total_time,
+         l->total_time / (double)l->getRoundCount());
 #endif
 }
 
@@ -811,7 +863,6 @@ int main(int argc, char **argv) {
   cudaEventCreate(&start);
   cudaEventCreate(&stop);
   cudaEventRecord(start, nullptr);
-
   kernel<<<blocks, NUM_THREADS>>>(Citems, Cop2, Cresult);
   CudaCheckError();
   error = cudaGetLastError();