2 files changed, 50 insertions, 7 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);
+    }
   }
 }
 
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index a48bd06..8d924f4 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -404,6 +404,25 @@ class CudaSpinLock {
   }
 };
 
+__global__ void testCudaAtomic() {
+  cuda::atomic<void *> node;
+  void * test_point = reinterpret_cast<void*>(0xdeadbeef);
+  void * rep_point = reinterpret_cast<void*>(0xdead);
+  node.store(rep_point);
+  assert(node.compare_exchange_weak(rep_point, test_point, cuda::memory_order_acquire));
+  void * null_ptr = nullptr;
+  node.store(nullptr);
+  assert(node.compare_exchange_weak(null_ptr, test_point, cuda::memory_order_acquire));
+}
+
+TEST(SkipTest, TestCudaAtomic) {
+  testCudaAtomic<<<1, 1>>>();
+  cudaDeviceSynchronize();
+  int last_error = cudaGetLastError();
+  if (last_error != cudaSuccess) {
+    printf("Got error: %d\n", last_error);
+  }
+}
 
 __global__ void testLock() {
   CudaSpinLock lock;
@@ -421,16 +440,16 @@ TEST(SkipTest, TestLock) {
 
 __global__ void testParallel(SkipList<Key, Comparator> * skipList, Key * keys, CudaSpinLock * lock) {
   unsigned int start = threadIdx.x;
-  //printf("start: %u\n", start);
-  lock->lock();
+  printf("start: %u\n", start);
+  //lock->lock();
   //printf("start insert: %u\n", start);
   for (unsigned i = start * TEST_STEP; i < (start + 1) * TEST_STEP; i++) {
-    //printf("%u %02u %lu\n", start, i, keys[i]);
+    printf("%u %02u %lu\n", start, i, keys[i]);
     //printf("key: %lu\n", keys[i]);
     skipList->Insert(keys[i]);
   }
-  lock->unlock();
-  //printf("done: %u\n", start);
+  //lock->unlock();
+  printf("done: %u\n", start);
 }
 
 __global__ void testSingle(SkipList<Key, Comparator>* skipList, Key * keys) {