From b642c58f9d8750804b95f6ab9ccbfbee1f5a671d Mon Sep 17 00:00:00 2001
From: KunoiSayami <i@leanhe.dev>
Date: Fri, 17 Jun 2022 00:29:54 +0800
Subject: feat: Finish implement measure time function * feat: Add
 normal_distribution program

Signed-off-by: KunoiSayami <i@leanhe.dev>
---
 CMakeLists.txt          |  7 ++++--
 main.cu                 | 57 ++++++++++++++++++++++++++++++++++++++++++++++---
 normal_distribution.cpp | 17 +++++++++++++++
 3 files changed, 76 insertions(+), 5 deletions(-)
 create mode 100644 normal_distribution.cpp

diff --git a/CMakeLists.txt b/CMakeLists.txt
index f816aed..e249fba 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1,6 +1,5 @@
 cmake_minimum_required(VERSION 3.22)
-project(lockfree CUDA)
-
+project(lockfree CUDA CXX)
 set(CMAKE_CUDA_STANDARD 14)
 
 #set(GXX_INCLUDE /usr/lib/gcc/x86_64-pc-linux-gnu/11.3.0/include)
@@ -19,6 +18,7 @@ target_link_libraries(lockfree m stdc++)
 
 set_target_properties(lockfree PROPERTIES
         CUDA_SEPARABLE_COMPILATION ON)
+set_target_properties(lockfree PROPERTIES CUDA_ARCHITECTURES "75")
 
 function(build_different_target is_random build_size step_size)
 
@@ -62,3 +62,6 @@ build_different_target(true 1048576 2)
 build_different_target(true 1048576 8)
 
 
+add_executable(normal_distribution normal_distribution.cpp)
+set_target_properties(normal_distribution PROPERTIES LINKER_LANGUAGE CXX)
+
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();
diff --git a/normal_distribution.cpp b/normal_distribution.cpp
new file mode 100644
index 0000000..c9cf701
--- /dev/null
+++ b/normal_distribution.cpp
@@ -0,0 +1,17 @@
+#include <algorithm>
+#include <cstdio>
+#include <random>
+int main() {
+
+  std::random_device randomDevice;
+  std::mt19937 randomEngine(randomDevice());
+  std::normal_distribution<long double> normalDistribution(2147483647,
+                                                           2147483647);
+
+  for (int i = 0; i < 1048576; i++) {
+    printf("%llu\n",
+           (unsigned long long)std::round(normalDistribution(randomEngine)));
+  }
+
+  return 0;
+}
\ No newline at end of file
-- 
cgit v1.3.1