diff options
| -rw-r--r-- | db/skiplist.cuh | 16 | ||||
| -rw-r--r-- | db/skiplist_test.cu | 34 |
2 files changed, 40 insertions, 10 deletions
diff --git a/db/skiplist.cuh b/db/skiplist.cuh index 8fe864a..a285b8e 100644 --- a/db/skiplist.cuh +++ b/db/skiplist.cuh @@ -77,11 +77,11 @@ class SkipList { // Advances to the next position. // REQUIRES: Valid() - void Next(); + __device__ void Next(); // Advances to the previous position. // REQUIRES: Valid() - void Prev(); + __device__ void Prev(); // Advance to the first entry with a key >= target __device__ void Seek(const Key& target); @@ -123,11 +123,11 @@ class SkipList { // Return the latest node with a key < key. // Return head_ if there is no such node. - Node* FindLessThan(const Key& key) const; + __device__ Node* FindLessThan(const Key& key) const; // Return the last node in the list. // Return head_ if list is empty. - Node* FindLast() const; + __device__ Node* FindLast() const; // Immutable after construction Comparator const compare_; @@ -226,13 +226,13 @@ __device__ inline const Key& SkipList<Key, Comparator>::Iterator::key() const { } template <typename Key, class Comparator> -inline void SkipList<Key, Comparator>::Iterator::Next() { +__device__ inline void SkipList<Key, Comparator>::Iterator::Next() { assert(Valid()); node_ = node_->Next(0); } template <typename Key, class Comparator> -inline void SkipList<Key, Comparator>::Iterator::Prev() { +__device__ inline void SkipList<Key, Comparator>::Iterator::Prev() { // Instead of using explicit "prev" links, we just search for the // last node that falls before key. assert(Valid()); @@ -304,7 +304,7 @@ __device__ SkipList<Key, Comparator>::FindGreaterOrEqual(const Key& key, template <typename Key, class Comparator> typename SkipList<Key, Comparator>::Node* -SkipList<Key, Comparator>::FindLessThan(const Key& key) const { +__device__ SkipList<Key, Comparator>::FindLessThan(const Key& key) const { Node* x = head_; int level = GetMaxHeight() - 1; while (true) { @@ -324,7 +324,7 @@ SkipList<Key, Comparator>::FindLessThan(const Key& key) const { } template <typename Key, class Comparator> -typename SkipList<Key, Comparator>::Node* SkipList<Key, Comparator>::FindLast() +__device__ typename SkipList<Key, Comparator>::Node* SkipList<Key, Comparator>::FindLast() const { Node* x = head_; int level = GetMaxHeight() - 1; diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu index 2b42b48..0022c84 100644 --- a/db/skiplist_test.cu +++ b/db/skiplist_test.cu @@ -2,11 +2,13 @@ // Use of this source code is governed by a BSD-style license that can be // found in the LICENSE file. See the AUTHORS file for names of contributors. + #include "db/skiplist.cuh" //#include <atomic> #include <set> #include "leveldb/env.h" +#include "cassert" #include "port/port.h" #include "port/thread_annotations.h" @@ -387,10 +389,10 @@ __host__ void host_insert_test() { init<<<gridSize, blockSize>>>(&pArena, &skipList); cudaDeviceSynchronize(); - insert_skiplist<<<gridSize, blockSize>>>(skipList, device_rnd); + //insert_skiplist<<<gridSize, blockSize>>>(skipList, device_rnd); - cudaDeviceSynchronize(); + //cudaDeviceSynchronize(); insert_and_lookup<<<gridSize, blockSize>>>(skipList); cudaDeviceSynchronize(); @@ -402,6 +404,34 @@ TEST(SkipTest, TestCudaInsert) { host_insert_test(); } +#define cudaAssert(condition) \ + if (!(condition)){ printf("Assertion %s failed!\n", #condition); asm("trap;"); } + +__global__ void test_empty() { + Arena arena; + Comparator cmp; + SkipList<Key, Comparator> list(cmp, &arena); + assert(!list.Contains(10)); + SkipList<Key, Comparator>::Iterator iter(&list); + assert(!iter.Valid()); + iter.SeekToFirst(); + assert(!iter.Valid()); + iter.Seek(100); + assert(!iter.Valid()); + iter.SeekToLast(); + assert(!iter.Valid()); +} + +__global__ void test_fail() { + std::printf("show some message"); + assert(0); +} + +TEST(SkipTest, Empty) { + test_fail<<<1,1>>>(); + cudaDeviceSynchronize(); +} + /* // Needed when building in C++11 mode. constexpr uint32_t ConcurrentTest::K; |