aboutsummaryrefslogtreecommitdiff
path: root/db/skiplist.cuh
diff options
context:
space:
mode:
Diffstat (limited to 'db/skiplist.cuh')
-rw-r--r--db/skiplist.cuh28
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);
+ }
}
}