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;