feat: Write search cost to file

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

1 files changed, 53 insertions, 56 deletions

diff --git a/main.cu b/main.cu
index 524b6d3..bc42b11 100644
--- a/main.cu
+++ b/main.cu
@@ -218,7 +218,7 @@ public:
   }
 };
 
-struct MapNode {
+/*struct MapNode {
   LL key;
   // Node *point[MAX_LEVEL + 1];
   Node *point;
@@ -262,44 +262,7 @@ 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
 
@@ -307,7 +270,6 @@ class LockFreeSkipList {
 public:
   Node *head;
   Node *tail;
-  MemMap key_map;
   LockFreeSkipList() {
     Node *h = new Node(0);
     // size_ = 0;
@@ -345,10 +307,10 @@ public:
 
 #ifdef MEASURE_TIME
   unsigned round = 0;
-  CudaSpinLock spinLock;
   __device__ void increaseRoundCount(unsigned count = 1) {
     atomicAdd(&this->round, count);
   }
+  int spend_time[NUM_ITEMS];
   unsigned long long total_time = 0;
   __device__ unsigned getRoundCount() const { return this->round; }
 #endif
@@ -376,6 +338,9 @@ __device__ Node *GetNewNode(LL key) {
 __device__ LockFreeSkipList *l; // The lock-free skip list
 
 __device__ LL KeyIndex[KEY_INDEX_SIZE];
+#ifdef MEASURE_TIME
+__device__ int *SpendTime;
+#endif
 
 // Kernel for initializing device memory
 
@@ -443,10 +408,6 @@ __device__ bool LockFreeSkipList::Search(LL key) {
   Node *curr = nullptr;
   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);
 #ifdef MEASURE_ACCESS
@@ -472,13 +433,6 @@ __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);
 }
 
@@ -618,7 +572,19 @@ __global__ void kernel(LL *items, LL *op, LL *result) {
       result[tid] = l->Delete(item);
     }
     if (op[tid] == SEARCH) {
+#ifdef MEASURE_TIME
+      unsigned long long start_time = clock64();
+#endif
       result[tid] = l->Search(item);
+#ifdef MEASURE_TIME
+      unsigned long long end_time = clock64() - start_time;
+      // SpendTime[tid] = (int)end_time;
+      if (l->spend_time[tid]) {
+        printf("%d\n", tid);
+      }
+      l->spend_time[tid] = (int)end_time;
+      // printf("%d %d\n", SpendTime[tid], tid);
+#endif
     }
   }
 }
@@ -638,6 +604,9 @@ LL RandomLevel(std::mt19937 &randomEngine, double p) {
 }
 
 std::vector<LL> storage;
+#ifdef MEASURE_TIME
+std::vector<int> SpendTimeVec;
+#endif
 
 unsigned trailing_zeroes(size_t index) {
   unsigned bits = 0;
@@ -676,10 +645,10 @@ __global__ void print_function() {
 #ifdef MEASURE_ACCESS
   printf("count: %u\n", l->getAccessCount());
 #endif
-#ifdef MEASURE_TIME
-  printf("round: %u time: %llu avg: %.2lf\n", l->getRoundCount(), l->total_time,
-         l->total_time / (double)l->getRoundCount());
-#endif
+}
+
+__global__ void copy_function(int *spend_time) {
+  memcpy(spend_time, l->spend_time, sizeof(int) * NUM_THREADS);
 }
 
 int main(int argc, char **argv) {
@@ -863,7 +832,11 @@ int main(int argc, char **argv) {
   cudaEventCreate(&start);
   cudaEventCreate(&stop);
   cudaEventRecord(start, nullptr);
+#ifdef MEASURE_TIME
+  kernel<<<NUM_THREADS, 1>>>(Citems, Cop2, Cresult);
+#else
   kernel<<<blocks, NUM_THREADS>>>(Citems, Cop2, Cresult);
+#endif
   CudaCheckError();
   error = cudaGetLastError();
   if (cudaSuccess != error) {
@@ -898,9 +871,33 @@ int main(int argc, char **argv) {
   cudaDeviceSynchronize();
 
 #if (defined(MEASURE_TIME) || defined(MEASURE_ACCESS))
+
   print_function<<<1, 1>>>();
 
   cudaDeviceSynchronize();
+#ifdef MEASURE_TIME
+  {
+    int *cuda_tmp = nullptr;
+    cudaMalloc(&cuda_tmp, sizeof(int) * NUM_ITEMS);
+    copy_function<<<1, 1>>>(cuda_tmp);
+    cudaDeviceSynchronize();
+    CudaCheckError();
+    FILE *file = fopen("spend_time.txt", "w");
+    int *tmp = new int[NUM_ITEMS];
+    cudaMemcpy(tmp, cuda_tmp, sizeof(int) * NUM_ITEMS, cudaMemcpyDeviceToHost);
+    // memcpy(tmp, SpendTime, sizeof(int) * NUM_ITEMS);
+    for (i = 0; i < NUM_ITEMS; i++) {
+      if (tmp[i] == 0)
+        break;
+      fprintf(file, "%d\n", tmp[i]);
+    }
+    // printf("%d\n", i);
+    delete[] tmp;
+    fclose(file);
+  }
+  // for (auto element : SpendTimeVec)
+  //   printf("%d\n", element);
+#endif
 #endif
   /*cudaFree(Clist);
     cudaFree(Cop2);