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