From b338a2083d68b9a8d1e3eb17538535aa050426df Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Mon, 13 Dec 2021 01:36:45 +0800 Subject: test(skiplist): Add insert and check test Signed-off-by: KunoiSayami --- CMakeLists.txt | 4 +- db/skiplist.cuh | 24 +++--- db/skiplist_test.cu | 219 ++++++++++++++++++++++++++++++++++------------------ util/arena.cu | 2 +- 4 files changed, 160 insertions(+), 89 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 29b83e9..f807626 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -220,7 +220,7 @@ target_sources(leveldb "${LEVELDB_PUBLIC_INCLUDE_DIR}/write_batch.h" ) set_target_properties(leveldb PROPERTIES CUDA_ARCHITECTURES "75") -#set_target_properties(leveldb PROPERTIES CUDA_SEPARABLE_COMPILATION ON) +set_target_properties(leveldb PROPERTIES CUDA_SEPARABLE_COMPILATION ON) set(CMAKE_CUDA_SEPARABLE_COMPILATION ON) #target_link_libraries(leveldb Thrust) @@ -352,7 +352,7 @@ if(LEVELDB_BUILD_TESTS) ) endif(NOT HAVE_CXX17_HAS_INCLUDE) set_target_properties("${test_target_name}" PROPERTIES CUDA_ARCHITECTURES "75") - #set_target_properties("${test_target_name}" PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + set_target_properties("${test_target_name}" PROPERTIES CUDA_SEPARABLE_COMPILATION ON) add_test(NAME "${test_target_name}" COMMAND "${test_target_name}") endfunction(leveldb_test) 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::Node* SkipList:: } template -inline SkipList::Iterator::Iterator(const SkipList* list) { +__device__ __host__ inline SkipList::Iterator::Iterator(const SkipList* list) { list_ = list; node_ = nullptr; } template -inline bool SkipList::Iterator::Valid() const { +__device__ __host__ inline bool SkipList::Iterator::Valid() const { return node_ != nullptr; } template -inline const Key& SkipList::Iterator::key() const { +__device__ inline const Key& SkipList::Iterator::key() const { assert(Valid()); return node_->key; } @@ -248,12 +248,12 @@ __device__ inline void SkipList::Iterator::Seek(const Key& targ } template -inline void SkipList::Iterator::SeekToFirst() { +__device__ inline void SkipList::Iterator::SeekToFirst() { node_ = list_->head_->Next(0); } template -inline void SkipList::Iterator::SeekToLast() { +__device__ inline void SkipList::Iterator::SeekToLast() { node_ = list_->FindLast(); if (node_ == list_->head_) { node_ = nullptr; @@ -390,7 +390,7 @@ __device__ void SkipList::Insert(const Key& key) { } template -bool SkipList::Contains(const Key& key) const { +__device__ bool SkipList::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 keys; - Arena arena; - Comparator cmp; - SkipList 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::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::Iterator iter(&list); - iter.Seek(i); - - // Compare against model iterator - std::set::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::Iterator iter(&list); - iter.SeekToLast(); - - // Compare against model iterator - for (std::set::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 * skipList, Random *de } } +struct Node { + Key num; + Node* next; + +}; + +template +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 +__device__ void ASSERT_EQ_dev(T a, U b) { + assert(a == b); +} + +__global__ void insert_and_lookup(SkipList * skipList) { + + const int N = 2000; + const int R = 5000; + Random rnd(1000); + + MemorySet keys; + Arena arena; + Comparator cmp; + SkipList 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::Iterator iter(&list); + iter.Seek(i); + + // Compare against model iterator + std::set::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::Iterator iter(&list); + iter.SeekToLast(); + + // Compare against model iterator + for (std::set::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 ** pSkipList) { Comparator cmp; *pArena = new Arena(); @@ -324,6 +392,9 @@ __host__ void host_insert_test() { cudaDeviceSynchronize(); + insert_and_lookup<<>>(skipList); + cudaDeviceSynchronize(); + } diff --git a/util/arena.cu b/util/arena.cu index a9556e5..82dbfe6 100644 --- a/util/arena.cu +++ b/util/arena.cu @@ -12,7 +12,7 @@ __device__ Arena::Arena() : alloc_ptr_(nullptr), alloc_bytes_remaining_(0), memory_usage_(0), head_(nullptr), blocks_(nullptr) {} -Arena::~Arena() { +__device__ Arena::~Arena() { ArenaNode * current = this->head_; while (current != nullptr) { ArenaNode * next = current->next; -- cgit v1.3.1