1 files changed, 145 insertions, 74 deletions

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();
+
 }