From a66cea8f77aeafa8523d28ea63a159826cef38a3 Mon Sep 17 00:00:00 2001
From: KunoiSayami <i@leanhe.dev>
Date: Fri, 18 Feb 2022 11:07:36 +0800
Subject: feat(skiplist): Temporary use spin lock

Signed-off-by: KunoiSayami <i@leanhe.dev>
---
 db/skiplist.cuh | 14 ++++++++------
 1 file changed, 8 insertions(+), 6 deletions(-)

(limited to 'db/skiplist.cuh')

diff --git a/db/skiplist.cuh b/db/skiplist.cuh
index 6f0cc6c..1fd7091 100644
--- a/db/skiplist.cuh
+++ b/db/skiplist.cuh
@@ -228,11 +228,10 @@ struct SkipList<Key, Comparator>::Node {
   __device__ bool SetNextSafe(int n, Node *x) {
     assert(n >= 0);
     Node * origin = next_[n].load(cuda::memory_order_acquire);
-    //assert(origin != x);
-    if (origin == nullptr) {
-      next_[n].store(x);
-      x->NoBarrier_SetNext(n, origin);
-      return true;
+    assert(origin != x);
+    if (origin == nullptr && next_[n].compare_exchange_weak(origin, x)) {
+        x->NoBarrier_SetNext(n, origin);
+        return true;
     }
     if (origin->key < x->key) {
       return false;
@@ -414,7 +413,7 @@ __device__ SkipList<Key, Comparator>::SkipList(Comparator cmp, Arena* arena)
       head_(NewNode(0 /* any key will do */, kMaxHeight)),
       max_height_(1),
       rnd_(0xdeadbeef),
-      arena_lock_(),
+      arena_lock_(true),
       find_lock_() {
   for (int i = 0; i < kMaxHeight; i++) {
     head_->SetNext(i, nullptr);
@@ -431,6 +430,8 @@ __device__ void SkipList<Key, Comparator>::Insert(const Key& key) {
   // Our data structure does not allow duplicate insertion
   assert(x == nullptr || !Equal(key, x->key));
 
+  // TODO: Remove spin lock in feature
+  this->find_lock_.lock();
 
   int height = RandomHeight();
   if (height > GetMaxHeight()) {
@@ -462,6 +463,7 @@ __device__ void SkipList<Key, Comparator>::Insert(const Key& key) {
     /*x->NoBarrier_SetNext(i, prev[i]->NoBarrier_Next(i));
     prev[i]->SetNext(i, x);*/
   }
+  this->find_lock_.unlock();
 }
 
 template <typename Key, class Comparator>
-- 
cgit v1.3.1