feat(exp): Add expt_0525.cu

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

1 files changed, 23 insertions, 8 deletions

diff --git a/skiplistcustom.cuh b/skiplistcustom.cuh
index a16e80a..80ecc19 100644
--- a/skiplistcustom.cuh
+++ b/skiplistcustom.cuh
@@ -1,5 +1,6 @@
 #pragma once
 #include "sortlib.cuh"
+#include <limits>
 #ifndef LOCKFREE_SKIPLIST_CUH_
 #define LOCKFREE_SKIPLIST_CUH_
 
@@ -85,7 +86,7 @@ public:
 // Definition of lock-free skip list
 
 class LockFreeSkipList {
-  static constexpr size_t SAMPLE_LENGTH = 1024;
+  // static constexpr size_t SAMPLE_LENGTH = 1023;
 
 public:
   Node *head;
@@ -99,20 +100,21 @@ public:
 
   double slice_size;
 
+  size_t sampleLength;
+
   LockFreeSkipList(key_type *samples, size_t sample_length)
-      : searcher(sample_length, sizeof(key_type)) {
+      : searcher(sample_length, sizeof(key_type)), sampleLength(sample_length) {
     Node *h = new Node(0);
     // size_ = 0;
 #if __WORDSIZE == 64
-    Node *t = new Node((key_type)NUM_ITEMS + 10);
+    Node *t = new Node(std::numeric_limits<key_type>::max());
 #else
     Node *t = new Node(0xffffffffULL);
 #endif
     cudaMalloc(&head, sizeof(Node));
 
     cudaMalloc(&tail, sizeof(Node));
-    int i;
-    for (i = 0; i < h->topLevel + 1; i++) {
+    for (int i = 0; i < h->topLevel + 1; i++) {
       h->SetRef(i, tail, false);
     }
     cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice);
@@ -123,8 +125,10 @@ public:
     initDeviceVariable();
 
     this->slice_size = 1.0 / sample_length;
+    // printf("%lf\n", this->slice_size);
 
-    cudaMalloc(&samples, sizeof(key_type) * sample_length);
+    printf("%ld\n", sample_length);
+    cudaMalloc(&this->samples, sizeof(key_type) * sample_length);
     cudaMemcpy(this->samples, samples, sizeof(key_type) * sample_length,
                cudaMemcpyHostToDevice);
   }
@@ -150,9 +154,11 @@ public:
     return bits;
   }
 
-  __device__ size_t calcLevel(key_type k) {
+  __device__ size_t calcLevel(key_type k) const {
+    // printf("Samples\n");
     auto index =
         this->searcher.sample_cdf_custom_version(this->samples, k) / slice_size;
+    // printf("Samples finish\n");
     auto level = trailing_zeroes(index);
     return level;
   }
@@ -162,6 +168,12 @@ public:
   //__device__ bool Delete(key_type);
   __device__ bool Search(key_type) const;
 
+  __device__ void testSample() {
+    for (int i = 0; i < sampleLength; i++) {
+      printf("%lld\n", this->samples[i]);
+    }
+  }
+
 #ifdef MEASURE_ACCESS
   unsigned access_times = 0;
 
@@ -193,6 +205,7 @@ public:
 
 __device__ Node *GetNewNode(key_type key, unsigned int *pointerIndex,
                             int topLevel) {
+  printf("%d\n", topLevel);
   key_type ind = atomicInc(pointerIndex, NUM_ITEMS);
   Node *n = nodes[ind];
   n->key = key;
@@ -328,7 +341,9 @@ __device__ bool LockFreeSkipList::Search(key_type key) const {
 }*/
 
 __device__ bool LockFreeSkipList::Add(key_type key) {
-  Node *newNode = GetNewNode(key, this->pointerIndex, calcLevel(key));
+  auto level1 = calcLevel(key) - 1;
+  printf("%lu\n", level1);
+  Node *newNode = GetNewNode(key, this->pointerIndex, level1);
   int topLevel = newNode->topLevel;
   int bottomLevel = 0;
   Node *preds[MAX_LEVEL + 1];