diff options
Diffstat (limited to 'db/skiplist.cuh')
| -rw-r--r-- | db/skiplist.cuh | 14 |
1 files changed, 8 insertions, 6 deletions
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> |