diff options
Diffstat (limited to 'db')
| -rw-r--r-- | db/skiplist.cuh | 24 | ||||
| -rw-r--r-- | db/skiplist_test.cu | 219 |
2 files changed, 157 insertions, 86 deletions
diff --git a/db/skiplist.cuh b/db/skiplist.cuh index 4fcd220..8fe864a 100644 --- a/db/skiplist.cuh +++ b/db/skiplist.cuh @@ -59,21 +59,21 @@ class SkipList { __device__ void Insert(const Key& key); // Returns true iff an entry that compares equal to key is in the list. - bool Contains(const Key& key) const; + __device__ bool Contains(const Key& key) const; // Iteration over the contents of a skip list class Iterator { public: // Initialize an iterator over the specified list. // The returned iterator is not valid. - explicit Iterator(const SkipList* list); + __device__ __host__ explicit Iterator(const SkipList* list); // Returns true iff the iterator is positioned at a valid node. - bool Valid() const; + __device__ __host__ bool Valid() const; // Returns the key at the current position. // REQUIRES: Valid() - const Key& key() const; + __device__ const Key& key() const; // Advances to the next position. // REQUIRES: Valid() @@ -88,11 +88,11 @@ class SkipList { // Position at the first entry in list. // Final state of iterator is Valid() iff list is not empty. - void SeekToFirst(); + __device__ void SeekToFirst(); // Position at the last entry in list. // Final state of iterator is Valid() iff list is not empty. - void SeekToLast(); + __device__ void SeekToLast(); private: const SkipList* list_; @@ -209,18 +209,18 @@ __device__ typename SkipList<Key, Comparator>::Node* SkipList<Key, Comparator>:: } template <typename Key, class Comparator> -inline SkipList<Key, Comparator>::Iterator::Iterator(const SkipList* list) { +__device__ __host__ inline SkipList<Key, Comparator>::Iterator::Iterator(const SkipList* list) { list_ = list; node_ = nullptr; } template <typename Key, class Comparator> -inline bool SkipList<Key, Comparator>::Iterator::Valid() const { +__device__ __host__ inline bool SkipList<Key, Comparator>::Iterator::Valid() const { return node_ != nullptr; } template <typename Key, class Comparator> -inline const Key& SkipList<Key, Comparator>::Iterator::key() const { +__device__ inline const Key& SkipList<Key, Comparator>::Iterator::key() const { assert(Valid()); return node_->key; } @@ -248,12 +248,12 @@ __device__ inline void SkipList<Key, Comparator>::Iterator::Seek(const Key& targ } template <typename Key, class Comparator> -inline void SkipList<Key, Comparator>::Iterator::SeekToFirst() { +__device__ inline void SkipList<Key, Comparator>::Iterator::SeekToFirst() { node_ = list_->head_->Next(0); } template <typename Key, class Comparator> -inline void SkipList<Key, Comparator>::Iterator::SeekToLast() { +__device__ inline void SkipList<Key, Comparator>::Iterator::SeekToLast() { node_ = list_->FindLast(); if (node_ == list_->head_) { node_ = nullptr; @@ -390,7 +390,7 @@ __device__ void SkipList<Key, Comparator>::Insert(const Key& key) { } template <typename Key, class Comparator> -bool SkipList<Key, Comparator>::Contains(const Key& key) const { +__device__ bool SkipList<Key, Comparator>::Contains(const Key& key) const { Node* x = FindGreaterOrEqual(key, nullptr); if (x != nullptr && Equal(key, x->key)) { return true; diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu index d18f11f..2b42b48 100644 --- a/db/skiplist_test.cu +++ b/db/skiplist_test.cu @@ -50,80 +50,6 @@ struct Comparator { } TEST(SkipTest, InsertAndLookup) { - const int N = 2000; - const int R = 5000; - Random rnd(1000); - std::set<Key> keys; - Arena arena; - Comparator cmp; - SkipList<Key, Comparator> list(cmp, &arena); - for (int i = 0; i < N; i++) { - Key key = rnd.Next() % R; - if (keys.insert(key).second) { - list.Insert(key); - } - } - - for (int i = 0; i < R; i++) { - if (list.Contains(i)) { - ASSERT_EQ(keys.count(i), 1); - } else { - ASSERT_EQ(keys.count(i), 0); - } - } - - // Simple iterator tests - { - SkipList<Key, Comparator>::Iterator iter(&list); - ASSERT_TRUE(!iter.Valid()); - - iter.Seek(0); - ASSERT_TRUE(iter.Valid()); - ASSERT_EQ(*(keys.begin()), iter.key()); - - iter.SeekToFirst(); - ASSERT_TRUE(iter.Valid()); - ASSERT_EQ(*(keys.begin()), iter.key()); - - iter.SeekToLast(); - ASSERT_TRUE(iter.Valid()); - ASSERT_EQ(*(keys.rbegin()), iter.key()); - } - - // Forward iteration test - for (int i = 0; i < R; i++) { - SkipList<Key, Comparator>::Iterator iter(&list); - iter.Seek(i); - - // Compare against model iterator - std::set<Key>::iterator model_iter = keys.lower_bound(i); - for (int j = 0; j < 3; j++) { - if (model_iter == keys.end()) { - ASSERT_TRUE(!iter.Valid()); - break; - } else { - ASSERT_TRUE(iter.Valid()); - ASSERT_EQ(*model_iter, iter.key()); - ++model_iter; - iter.Next(); - } - } - } - - // Backward iteration test - { - SkipList<Key, Comparator>::Iterator iter(&list); - iter.SeekToLast(); - - // Compare against model iterator - for (std::set<Key>::reverse_iterator model_iter = keys.rbegin(); - model_iter != keys.rend(); ++model_iter) { - ASSERT_TRUE(iter.Valid()); - ASSERT_EQ(*model_iter, iter.key()); - iter.Prev(); - } - ASSERT_TRUE(!iter.Valid()); - } } // We want to make sure that with a single writer and multiple @@ -294,6 +220,148 @@ __global__ void insert_skiplist(SkipList<Key, Comparator> * skipList, Random *de } } +struct Node { + Key num; + Node* next; + +}; + +template<typename Key> +class MemorySet { + public: + explicit __device__ MemorySet(): first(nullptr), current(nullptr) {}; + + __device__ ~MemorySet() { + Node * crt = first, * prev; + while (crt != nullptr) { + prev = crt; + crt = crt->next; + delete prev; + } + first = nullptr; + } + + __device__ bool insert(const Key & k) { + if (first == nullptr) { + first = new Node; + first->num = k; + first->next = nullptr; + current = first; + return true; + } + if (this->find(k)) { + return false; + } + current->next = new Node; + current = current->next; + current->num = k; + current->next = nullptr; + return true; + } + + __device__ bool find(const Key & value) { + Node * crt = first; + while (crt != nullptr) { + if (crt->num == value) { + return true; + } + crt = crt->next; + } + return false; + } + + __device__ Node * get_first() { + return this->first; + } + + __device__ size_t count(const Key & k) { + if (this->find(k)) { + return 1; + } + return 0; + } + + private: + size_t total; + Node * first; + Node * current; +}; + +template<typename T, typename U> +__device__ void ASSERT_EQ_dev(T a, U b) { + assert(a == b); +} + +__global__ void insert_and_lookup(SkipList<Key, Comparator> * skipList) { + + const int N = 2000; + const int R = 5000; + Random rnd(1000); + + MemorySet<Key> keys; + Arena arena; + Comparator cmp; + SkipList<Key, Comparator> list(cmp, &arena); + for (int i = 0; i < N; i++) { + Key key = rnd.Next() % R; + if (keys.insert(key)) { + list.Insert(key); + } + } + + for (int i = 0; i < R; i++) { + if (list.Contains(i)) { + ASSERT_EQ_dev(keys.count(i), 1); + } else { + ASSERT_EQ_dev(keys.count(i), 0); + } + } + + Node * cur = keys.get_first(); + while (cur != nullptr) { + assert(list.Contains(cur->num)); + cur = cur->next; + } +/* + + // Forward iteration test + for (int i = 0; i < R; i++) { + SkipList<Key, Comparator>::Iterator iter(&list); + iter.Seek(i); + + // Compare against model iterator + std::set<Key>::iterator model_iter = keys.lower_bound(i); + for (int j = 0; j < 3; j++) { + if (model_iter == keys.end()) { + ASSERT_TRUE(!iter.Valid()); + break; + } else { + ASSERT_TRUE(iter.Valid()); + ASSERT_EQ(*model_iter, iter.key()); + ++model_iter; + iter.Next(); + } + } + } + + // Backward iteration test + { + SkipList<Key, Comparator>::Iterator iter(&list); + iter.SeekToLast(); + + // Compare against model iterator + for (std::set<Key>::reverse_iterator model_iter = keys.rbegin(); + model_iter != keys.rend(); ++model_iter) { + ASSERT_TRUE(iter.Valid()); + ASSERT_EQ(*model_iter, iter.key()); + iter.Prev(); + } + ASSERT_TRUE(!iter.Valid()); + } +*/ + +} + __global__ void init(Arena ** pArena, SkipList<Key, Comparator> ** pSkipList) { Comparator cmp; *pArena = new Arena(); @@ -324,6 +392,9 @@ __host__ void host_insert_test() { cudaDeviceSynchronize(); + insert_and_lookup<<<gridSize, blockSize>>>(skipList); + cudaDeviceSynchronize(); + } |