1 files changed, 26 insertions, 2 deletions

diff --git a/db/skiplist.cuh b/db/skiplist.cuh
index a6f7bb7..9f1cdaa 100644
--- a/db/skiplist.cuh
+++ b/db/skiplist.cuh
@@ -27,6 +27,7 @@
 //
 // ... prev vs. next pointer ordering ...
 
+#include <cstdio>
 #include <cassert>
 #include <cstdlib>
 #include <cuda/atomic>
@@ -184,6 +185,22 @@ struct SkipList<Key, Comparator>::Node {
     next_[n].store(x, cuda::memory_order_release);
   }
 
+  __device__ bool SetNextSafe(int n, Node *x, Node * origin) {
+    assert(n >= 0);
+    if (origin == nullptr) {
+      next_[n].store(x, cuda::memory_order_release);
+      return true;
+    }
+    printf("%d %p %p\n", n, x, origin);
+    bool ret = next_[n].compare_exchange_weak(origin, x, cuda::memory_order_acquire);
+    if (!ret) {
+      Node * rep = next_[n].load(cuda::memory_order_acquire);
+      printf("%p %p\n", rep, origin);
+      assert(rep != origin);
+    }
+    return ret;
+  }
+
   // No-barrier variants that can be safely used in a few locations.
   __device__ Node* NoBarrier_Next(int n) {
     assert(n >= 0);
@@ -379,12 +396,19 @@ __device__ void SkipList<Key, Comparator>::Insert(const Key& key) {
     max_height_.store(height, cuda::memory_order_relaxed);
   }
 
+  Node * original_next;
   x = NewNode(key, height);
   for (int i = 0; i < height; i++) {
     // NoBarrier_SetNext() suffices since we will add a barrier when
     // we publish a pointer to "x" in prev[i].
-    x->NoBarrier_SetNext(i, prev[i]->NoBarrier_Next(i));
-    prev[i]->SetNext(i, x);
+    while (true) {
+      original_next = prev[i]->NoBarrier_Next(i);
+      x->NoBarrier_SetNext(i, original_next);
+      if (prev[i]->SetNextSafe(i, x, original_next))
+        break;
+      printf("failed %p\n", original_next);
+      FindGreaterOrEqual(key, prev);
+    }
   }
 }