From 0406eca917270ebc4d05b20de4eecc45645d60f0 Mon Sep 17 00:00:00 2001
From: KunoiSayami <i@leanhe.dev>
Date: Mon, 15 Nov 2021 11:57:39 +0800
Subject: feat: Add skiplist cuda version

Signed-off-by: KunoiSayami <i@leanhe.dev>
---
 db/memtable.h       |   2 +-
 db/skiplist.cu      |   1 +
 db/skiplist.cuh     | 383 ++++++++++++++++++++++++++++++++++++++++++++++
 db/skiplist.h       |   2 +-
 db/skiplist_test.cc | 375 ---------------------------------------------
 db/skiplist_test.cu | 428 ++++++++++++++++++++++++++++++++++++++++++++++++++++
 6 files changed, 814 insertions(+), 377 deletions(-)
 create mode 100644 db/skiplist.cu
 create mode 100644 db/skiplist.cuh
 delete mode 100644 db/skiplist_test.cc
 create mode 100644 db/skiplist_test.cu

(limited to 'db')

diff --git a/db/memtable.h b/db/memtable.h
index 3d3444b..1099504 100644
--- a/db/memtable.h
+++ b/db/memtable.h
@@ -11,7 +11,7 @@
 
 #include "leveldb/db.h"
 
-#include "util/arena.cuh"
+#include "util/arena.h"
 
 namespace leveldb {
 
diff --git a/db/skiplist.cu b/db/skiplist.cu
new file mode 100644
index 0000000..c6517e2
--- /dev/null
+++ b/db/skiplist.cu
@@ -0,0 +1 @@
+#include "db/skiplist.cuh"
diff --git a/db/skiplist.cuh b/db/skiplist.cuh
new file mode 100644
index 0000000..5310d9c
--- /dev/null
+++ b/db/skiplist.cuh
@@ -0,0 +1,383 @@
+// Copyright (c) 2011 The LevelDB Authors. All rights reserved.
+// 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.
+
+#ifndef STORAGE_LEVELDB_DB_SKIPLIST_H_
+#define STORAGE_LEVELDB_DB_SKIPLIST_H_
+
+// Thread safety
+// -------------
+//
+// Writes require external synchronization, most likely a mutex.
+// Reads require a guarantee that the SkipList will not be destroyed
+// while the read is in progress.  Apart from that, reads progress
+// without any internal locking or synchronization.
+//
+// Invariants:
+//
+// (1) Allocated nodes are never deleted until the SkipList is
+// destroyed.  This is trivially guaranteed by the code since we
+// never delete any skip list nodes.
+//
+// (2) The contents of a Node except for the next/prev pointers are
+// immutable after the Node has been linked into the SkipList.
+// Only Insert() modifies the list, and it is careful to initialize
+// a node and use release-stores to publish the nodes in one or
+// more lists.
+//
+// ... prev vs. next pointer ordering ...
+
+#include <atomic>
+#include <cassert>
+#include <cstdlib>
+#include <cuda/atomic>
+
+#include "util/arena.cuh"
+#include "util/random.cuh"
+
+namespace leveldb {
+
+class Arena;
+
+template <typename Key, class Comparator>
+class SkipList {
+ private:
+  struct Node;
+
+ public:
+  // Create a new SkipList object that will use "cmp" for comparing keys,
+  // and will allocate memory using "*arena".  Objects allocated in the arena
+  // must remain allocated for the lifetime of the skiplist object.
+  explicit SkipList(Comparator cmp, Arena* arena);
+
+  SkipList(const SkipList&) = delete;
+  SkipList& operator=(const SkipList&) = delete;
+
+  // Insert key into the list.
+  // REQUIRES: nothing that compares equal to key is currently in the list.
+  __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;
+
+  // 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);
+
+    // Returns true iff the iterator is positioned at a valid node.
+    bool Valid() const;
+
+    // Returns the key at the current position.
+    // REQUIRES: Valid()
+    const Key& key() const;
+
+    // Advances to the next position.
+    // REQUIRES: Valid()
+    void Next();
+
+    // Advances to the previous position.
+    // REQUIRES: Valid()
+    void Prev();
+
+    // Advance to the first entry with a key >= target
+    __device__ void Seek(const Key& target);
+
+    // Position at the first entry in list.
+    // Final state of iterator is Valid() iff list is not empty.
+    void SeekToFirst();
+
+    // Position at the last entry in list.
+    // Final state of iterator is Valid() iff list is not empty.
+    void SeekToLast();
+
+   private:
+    const SkipList* list_;
+    Node* node_;
+    // Intentionally copyable
+  };
+
+ private:
+  enum { kMaxHeight = 12 };
+
+  __device__ inline int GetMaxHeight() const {
+    return max_height_.load(cuda::memory_order_relaxed);
+  }
+
+  __device__ Node* NewNode(const Key& key, int height);
+  __device__ int RandomHeight();
+  __device__ bool Equal(const Key& a, const Key& b) const { return (compare_(a, b) == 0); }
+
+  // Return true if key is greater than the data stored in "n"
+  bool KeyIsAfterNode(const Key& key, Node* n) const;
+
+  // Return the earliest node that comes at or after key.
+  // Return nullptr if there is no such node.
+  //
+  // If prev is non-null, fills prev[level] with pointer to previous
+  // node at "level" for every level in [0..max_height_-1].
+  __device__ Node* FindGreaterOrEqual(const Key& key, Node** prev) const;
+
+  // Return the latest node with a key < key.
+  // Return head_ if there is no such node.
+  Node* FindLessThan(const Key& key) const;
+
+  // Return the last node in the list.
+  // Return head_ if list is empty.
+  Node* FindLast() const;
+
+  // Immutable after construction
+  Comparator const compare_;
+  Arena* const arena_;  // Arena used for allocations of nodes
+
+  Node* const head_;
+
+  // Modified only by Insert().  Read racily by readers, but stale
+  // values are ok.
+  cuda::atomic<int> max_height_; // Height of the entire list
+
+  // Read/written only by Insert().
+  Random rnd_;
+};
+
+// Implementation details follow
+template <typename Key, class Comparator>
+struct SkipList<Key, Comparator>::Node {
+  explicit Node(const Key& k) : key(k) {}
+
+  Key const key;
+
+  // Accessors/mutators for links.  Wrapped in methods so we can
+  // add the appropriate barriers as necessary.
+  Node* Next(int n) {
+    assert(n >= 0);
+    // Use an 'acquire load' so that we observe a fully initialized
+    // version of the returned Node.
+    return next_[n].load(std::memory_order_acquire);
+  }
+  __device__ void SetNext(int n, Node* x) {
+    assert(n >= 0);
+    // Use a 'release store' so that anybody who reads through this
+    // pointer observes a fully initialized version of the inserted node.
+    next_[n].store(x, std::memory_order_release);
+  }
+
+  // No-barrier variants that can be safely used in a few locations.
+  __device__ Node* NoBarrier_Next(int n) {
+    assert(n >= 0);
+    return next_[n].load(std::memory_order_relaxed);
+  }
+  __device__ void NoBarrier_SetNext(int n, Node* x) {
+    assert(n >= 0);
+    next_[n].store(x, std::memory_order_relaxed);
+  }
+
+ private:
+  // Array of length equal to the node height.  next_[0] is lowest level link.
+  std::atomic<Node*> next_[1];
+};
+
+template <typename Key, class Comparator>
+__device__ typename SkipList<Key, Comparator>::Node* SkipList<Key, Comparator>::NewNode(
+    const Key& key, int height) {
+  char* const node_memory = arena_->AllocateAligned(
+      sizeof(Node) + sizeof(std::atomic<Node*>) * (height - 1));
+  return new (node_memory) Node(key);
+}
+
+template <typename Key, class Comparator>
+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 {
+  return node_ != nullptr;
+}
+
+template <typename Key, class Comparator>
+inline const Key& SkipList<Key, Comparator>::Iterator::key() const {
+  assert(Valid());
+  return node_->key;
+}
+
+template <typename Key, class Comparator>
+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() {
+  // Instead of using explicit "prev" links, we just search for the
+  // last node that falls before key.
+  assert(Valid());
+  node_ = list_->FindLessThan(node_->key);
+  if (node_ == list_->head_) {
+    node_ = nullptr;
+  }
+}
+
+template <typename Key, class Comparator>
+__device__ inline void SkipList<Key, Comparator>::Iterator::Seek(const Key& target) {
+  node_ = list_->FindGreaterOrEqual(target, nullptr);
+}
+
+template <typename Key, class Comparator>
+inline void SkipList<Key, Comparator>::Iterator::SeekToFirst() {
+  node_ = list_->head_->Next(0);
+}
+
+template <typename Key, class Comparator>
+inline void SkipList<Key, Comparator>::Iterator::SeekToLast() {
+  node_ = list_->FindLast();
+  if (node_ == list_->head_) {
+    node_ = nullptr;
+  }
+}
+
+template <typename Key, class Comparator>
+__device__ int SkipList<Key, Comparator>::RandomHeight() {
+  // Increase height with probability 1 in kBranching
+  static const unsigned int kBranching = 4;
+  int height = 1;
+  while (height < kMaxHeight && ((rnd_.Next() % kBranching) == 0)) {
+    height++;
+  }
+  assert(height > 0);
+  assert(height <= kMaxHeight);
+  return height;
+}
+
+template <typename Key, class Comparator>
+bool SkipList<Key, Comparator>::KeyIsAfterNode(const Key& key, Node* n) const {
+  // null n is considered infinite
+  return (n != nullptr) && (compare_(n->key, key) < 0);
+}
+
+template <typename Key, class Comparator>
+typename SkipList<Key, Comparator>::Node*
+__device__ SkipList<Key, Comparator>::FindGreaterOrEqual(const Key& key,
+                                              Node** prev) const {
+  Node* x = head_;
+  int level = GetMaxHeight() - 1;
+  while (true) {
+    Node* next = x->Next(level);
+    if (KeyIsAfterNode(key, next)) {
+      // Keep searching in this list
+      x = next;
+    } else {
+      if (prev != nullptr) prev[level] = x;
+      if (level == 0) {
+        return next;
+      } else {
+        // Switch to next list
+        level--;
+      }
+    }
+  }
+}
+
+template <typename Key, class Comparator>
+typename SkipList<Key, Comparator>::Node*
+SkipList<Key, Comparator>::FindLessThan(const Key& key) const {
+  Node* x = head_;
+  int level = GetMaxHeight() - 1;
+  while (true) {
+    assert(x == head_ || compare_(x->key, key) < 0);
+    Node* next = x->Next(level);
+    if (next == nullptr || compare_(next->key, key) >= 0) {
+      if (level == 0) {
+        return x;
+      } else {
+        // Switch to next list
+        level--;
+      }
+    } else {
+      x = next;
+    }
+  }
+}
+
+template <typename Key, class Comparator>
+typename SkipList<Key, Comparator>::Node* SkipList<Key, Comparator>::FindLast()
+    const {
+  Node* x = head_;
+  int level = GetMaxHeight() - 1;
+  while (true) {
+    Node* next = x->Next(level);
+    if (next == nullptr) {
+      if (level == 0) {
+        return x;
+      } else {
+        // Switch to next list
+        level--;
+      }
+    } else {
+      x = next;
+    }
+  }
+}
+
+template <typename Key, class Comparator>
+SkipList<Key, Comparator>::SkipList(Comparator cmp, Arena* arena)
+    : compare_(cmp),
+      arena_(arena),
+      head_(NewNode(0 /* any key will do */, kMaxHeight)),
+      max_height_(1),
+      rnd_(0xdeadbeef) {
+  for (int i = 0; i < kMaxHeight; i++) {
+    head_->SetNext(i, nullptr);
+  }
+}
+
+template <typename Key, class Comparator>
+__device__ void SkipList<Key, Comparator>::Insert(const Key& key) {
+  // TODO(opt): We can use a barrier-free variant of FindGreaterOrEqual()
+  // here since Insert() is externally synchronized.
+  Node* prev[kMaxHeight];
+  Node* x = FindGreaterOrEqual(key, prev);
+
+  // Our data structure does not allow duplicate insertion
+  assert(x == nullptr || !Equal(key, x->key));
+
+  int height = RandomHeight();
+  if (height > GetMaxHeight()) {
+    for (int i = GetMaxHeight(); i < height; i++) {
+      prev[i] = head_;
+    }
+    // It is ok to mutate max_height_ without any synchronization
+    // with concurrent readers.  A concurrent reader that observes
+    // the new value of max_height_ will see either the old value of
+    // new level pointers from head_ (nullptr), or a new value set in
+    // the loop below.  In the former case the reader will
+    // immediately drop to the next level since nullptr sorts after all
+    // keys.  In the latter case the reader will use the new node.
+    max_height_.store(height, cuda::memory_order_relaxed);
+  }
+
+  x = NewNode(key, height);
+  for (int i = 0; i < height; i++) {
+    // NoBarrier_SetNext() suffices since we will add a barrier when
+    // we publish a pointer to "x" in prev[i].
+    x->NoBarrier_SetNext(i, prev[i]->NoBarrier_Next(i));
+    prev[i]->SetNext(i, x);
+  }
+}
+
+template <typename Key, class Comparator>
+bool SkipList<Key, Comparator>::Contains(const Key& key) const {
+  Node* x = FindGreaterOrEqual(key, nullptr);
+  if (x != nullptr && Equal(key, x->key)) {
+    return true;
+  } else {
+    return false;
+  }
+}
+
+}  // namespace leveldb
+
+#endif  // STORAGE_LEVELDB_DB_SKIPLIST_H_
diff --git a/db/skiplist.h b/db/skiplist.h
index 13df778..a59b45b 100644
--- a/db/skiplist.h
+++ b/db/skiplist.h
@@ -31,7 +31,7 @@
 #include <cassert>
 #include <cstdlib>
 
-#include "util/arena.cuh"
+#include "util/arena.h"
 #include "util/random.h"
 
 namespace leveldb {
diff --git a/db/skiplist_test.cc b/db/skiplist_test.cc
deleted file mode 100644
index 402cb47..0000000
--- a/db/skiplist_test.cc
+++ /dev/null
@@ -1,375 +0,0 @@
-// Copyright (c) 2011 The LevelDB Authors. All rights reserved.
-// 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.h"
-
-#include <atomic>
-#include <set>
-
-#include "leveldb/env.h"
-
-#include "port/port.h"
-#include "port/thread_annotations.h"
-#include "util/arena.cuh"
-#include "util/hash.h"
-#include "util/random.h"
-#include "util/testutil.h"
-
-#include "gtest/gtest.h"
-
-namespace leveldb {
-
-typedef uint64_t Key;
-
-struct Comparator {
-  int operator()(const Key& a, const Key& b) const {
-    if (a < b) {
-      return -1;
-    } else if (a > b) {
-      return +1;
-    } else {
-      return 0;
-    }
-  }
-};
-
-TEST(SkipTest, Empty) {
-  Arena arena;
-  Comparator cmp;
-  SkipList<Key, Comparator> list(cmp, &arena);
-  ASSERT_TRUE(!list.Contains(10));
-
-  SkipList<Key, Comparator>::Iterator iter(&list);
-  ASSERT_TRUE(!iter.Valid());
-  iter.SeekToFirst();
-  ASSERT_TRUE(!iter.Valid());
-  iter.Seek(100);
-  ASSERT_TRUE(!iter.Valid());
-  iter.SeekToLast();
-  ASSERT_TRUE(!iter.Valid());
-}
-
-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
-// concurrent readers (with no synchronization other than when a
-// reader's iterator is created), the reader always observes all the
-// data that was present in the skip list when the iterator was
-// constructed.  Because insertions are happening concurrently, we may
-// also observe new values that were inserted since the iterator was
-// constructed, but we should never miss any values that were present
-// at iterator construction time.
-//
-// We generate multi-part keys:
-//     <key,gen,hash>
-// where:
-//     key is in range [0..K-1]
-//     gen is a generation number for key
-//     hash is hash(key,gen)
-//
-// The insertion code picks a random key, sets gen to be 1 + the last
-// generation number inserted for that key, and sets hash to Hash(key,gen).
-//
-// At the beginning of a read, we snapshot the last inserted
-// generation number for each key.  We then iterate, including random
-// calls to Next() and Seek().  For every key we encounter, we
-// check that it is either expected given the initial snapshot or has
-// been concurrently added since the iterator started.
-class ConcurrentTest {
- private:
-  static constexpr uint32_t K = 4;
-
-  static uint64_t key(Key key) { return (key >> 40); }
-  static uint64_t gen(Key key) { return (key >> 8) & 0xffffffffu; }
-  static uint64_t hash(Key key) { return key & 0xff; }
-
-  static uint64_t HashNumbers(uint64_t k, uint64_t g) {
-    uint64_t data[2] = {k, g};
-    return Hash(reinterpret_cast<char*>(data), sizeof(data), 0);
-  }
-
-  static Key MakeKey(uint64_t k, uint64_t g) {
-    static_assert(sizeof(Key) == sizeof(uint64_t), "");
-    assert(k <= K);  // We sometimes pass K to seek to the end of the skiplist
-    assert(g <= 0xffffffffu);
-    return ((k << 40) | (g << 8) | (HashNumbers(k, g) & 0xff));
-  }
-
-  static bool IsValidKey(Key k) {
-    return hash(k) == (HashNumbers(key(k), gen(k)) & 0xff);
-  }
-
-  static Key RandomTarget(Random* rnd) {
-    switch (rnd->Next() % 10) {
-      case 0:
-        // Seek to beginning
-        return MakeKey(0, 0);
-      case 1:
-        // Seek to end
-        return MakeKey(K, 0);
-      default:
-        // Seek to middle
-        return MakeKey(rnd->Next() % K, 0);
-    }
-  }
-
-  // Per-key generation
-  struct State {
-    std::atomic<int> generation[K];
-    void Set(int k, int v) {
-      generation[k].store(v, std::memory_order_release);
-    }
-    int Get(int k) { return generation[k].load(std::memory_order_acquire); }
-
-    State() {
-      for (int k = 0; k < K; k++) {
-        Set(k, 0);
-      }
-    }
-  };
-
-  // Current state of the test
-  State current_;
-
-  Arena arena_;
-
-  // SkipList is not protected by mu_.  We just use a single writer
-  // thread to modify it.
-  SkipList<Key, Comparator> list_;
-
- public:
-  ConcurrentTest() : list_(Comparator(), &arena_) {}
-
-  // REQUIRES: External synchronization
-  void WriteStep(Random* rnd) {
-    const uint32_t k = rnd->Next() % K;
-    const intptr_t g = current_.Get(k) + 1;
-    const Key key = MakeKey(k, g);
-    list_.Insert(key);
-    current_.Set(k, g);
-  }
-
-  void ReadStep(Random* rnd) {
-    // Remember the initial committed state of the skiplist.
-    State initial_state;
-    for (int k = 0; k < K; k++) {
-      initial_state.Set(k, current_.Get(k));
-    }
-
-    Key pos = RandomTarget(rnd);
-    SkipList<Key, Comparator>::Iterator iter(&list_);
-    iter.Seek(pos);
-    while (true) {
-      Key current;
-      if (!iter.Valid()) {
-        current = MakeKey(K, 0);
-      } else {
-        current = iter.key();
-        ASSERT_TRUE(IsValidKey(current)) << current;
-      }
-      ASSERT_LE(pos, current) << "should not go backwards";
-
-      // Verify that everything in [pos,current) was not present in
-      // initial_state.
-      while (pos < current) {
-        ASSERT_LT(key(pos), K) << pos;
-
-        // Note that generation 0 is never inserted, so it is ok if
-        // <*,0,*> is missing.
-        ASSERT_TRUE((gen(pos) == 0) ||
-                    (gen(pos) > static_cast<Key>(initial_state.Get(key(pos)))))
-            << "key: " << key(pos) << "; gen: " << gen(pos)
-            << "; initgen: " << initial_state.Get(key(pos));
-
-        // Advance to next key in the valid key space
-        if (key(pos) < key(current)) {
-          pos = MakeKey(key(pos) + 1, 0);
-        } else {
-          pos = MakeKey(key(pos), gen(pos) + 1);
-        }
-      }
-
-      if (!iter.Valid()) {
-        break;
-      }
-
-      if (rnd->Next() % 2) {
-        iter.Next();
-        pos = MakeKey(key(pos), gen(pos) + 1);
-      } else {
-        Key new_target = RandomTarget(rnd);
-        if (new_target > pos) {
-          pos = new_target;
-          iter.Seek(new_target);
-        }
-      }
-    }
-  }
-};
-
-// Needed when building in C++11 mode.
-constexpr uint32_t ConcurrentTest::K;
-
-// Simple test that does single-threaded testing of the ConcurrentTest
-// scaffolding.
-TEST(SkipTest, ConcurrentWithoutThreads) {
-  ConcurrentTest test;
-  Random rnd(test::RandomSeed());
-  for (int i = 0; i < 10000; i++) {
-    test.ReadStep(&rnd);
-    test.WriteStep(&rnd);
-  }
-}
-
-class TestState {
- public:
-  ConcurrentTest t_;
-  int seed_;
-  std::atomic<bool> quit_flag_;
-
-  enum ReaderState { STARTING, RUNNING, DONE };
-
-  explicit TestState(int s)
-      : seed_(s), quit_flag_(false), state_(STARTING), state_cv_(&mu_) {}
-
-  void Wait(ReaderState s) LOCKS_EXCLUDED(mu_) {
-    mu_.Lock();
-    while (state_ != s) {
-      state_cv_.Wait();
-    }
-    mu_.Unlock();
-  }
-
-  void Change(ReaderState s) LOCKS_EXCLUDED(mu_) {
-    mu_.Lock();
-    state_ = s;
-    state_cv_.Signal();
-    mu_.Unlock();
-  }
-
- private:
-  port::Mutex mu_;
-  ReaderState state_ GUARDED_BY(mu_);
-  port::CondVar state_cv_ GUARDED_BY(mu_);
-};
-
-static void ConcurrentReader(void* arg) {
-  TestState* state = reinterpret_cast<TestState*>(arg);
-  Random rnd(state->seed_);
-  int64_t reads = 0;
-  state->Change(TestState::RUNNING);
-  while (!state->quit_flag_.load(std::memory_order_acquire)) {
-    state->t_.ReadStep(&rnd);
-    ++reads;
-  }
-  state->Change(TestState::DONE);
-}
-
-static void RunConcurrent(int run) {
-  const int seed = test::RandomSeed() + (run * 100);
-  Random rnd(seed);
-  const int N = 1000;
-  const int kSize = 1000;
-  for (int i = 0; i < N; i++) {
-    if ((i % 100) == 0) {
-      std::fprintf(stderr, "Run %d of %d\n", i, N);
-    }
-    TestState state(seed + 1);
-    Env::Default()->Schedule(ConcurrentReader, &state);
-    state.Wait(TestState::RUNNING);
-    for (int i = 0; i < kSize; i++) {
-      state.t_.WriteStep(&rnd);
-    }
-    state.quit_flag_.store(true, std::memory_order_release);
-    state.Wait(TestState::DONE);
-  }
-}
-
-TEST(SkipTest, Concurrent1) { RunConcurrent(1); }
-TEST(SkipTest, Concurrent2) { RunConcurrent(2); }
-TEST(SkipTest, Concurrent3) { RunConcurrent(3); }
-TEST(SkipTest, Concurrent4) { RunConcurrent(4); }
-TEST(SkipTest, Concurrent5) { RunConcurrent(5); }
-
-}  // namespace leveldb
-
-int main(int argc, char** argv) {
-  testing::InitGoogleTest(&argc, argv);
-  return RUN_ALL_TESTS();
-}
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
new file mode 100644
index 0000000..53f60f8
--- /dev/null
+++ b/db/skiplist_test.cu
@@ -0,0 +1,428 @@
+// Copyright (c) 2011 The LevelDB Authors. All rights reserved.
+// 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 "port/port.h"
+#include "port/thread_annotations.h"
+#include "util/arena.cuh"
+#include "util/hash.h"
+#include "util/random.cuh"
+#include "util/testutil.h"
+
+#include "gtest/gtest.h"
+
+namespace leveldb {
+
+typedef uint64_t Key;
+
+struct Comparator {
+  int operator()(const Key& a, const Key& b) const {
+    if (a < b) {
+      return -1;
+    } else if (a > b) {
+      return +1;
+    } else {
+      return 0;
+    }
+  }
+};
+
+/*TEST(SkipTest, Empty) {
+  Arena arena;
+  Comparator cmp;
+  SkipList<Key, Comparator> list(cmp, &arena);
+  ASSERT_TRUE(!list.Contains(10));
+
+  SkipList<Key, Comparator>::Iterator iter(&list);
+  ASSERT_TRUE(!iter.Valid());
+  iter.SeekToFirst();
+  ASSERT_TRUE(!iter.Valid());
+  iter.Seek(100);
+  ASSERT_TRUE(!iter.Valid());
+  iter.SeekToLast();
+  ASSERT_TRUE(!iter.Valid());
+}
+
+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
+// concurrent readers (with no synchronization other than when a
+// reader's iterator is created), the reader always observes all the
+// data that was present in the skip list when the iterator was
+// constructed.  Because insertions are happening concurrently, we may
+// also observe new values that were inserted since the iterator was
+// constructed, but we should never miss any values that were present
+// at iterator construction time.
+//
+// We generate multi-part keys:
+//     <key,gen,hash>
+// where:
+//     key is in range [0..K-1]
+//     gen is a generation number for key
+//     hash is hash(key,gen)
+//
+// The insertion code picks a random key, sets gen to be 1 + the last
+// generation number inserted for that key, and sets hash to Hash(key,gen).
+//
+// At the beginning of a read, we snapshot the last inserted
+// generation number for each key.  We then iterate, including random
+// calls to Next() and Seek().  For every key we encounter, we
+// check that it is either expected given the initial snapshot or has
+// been concurrently added since the iterator started.
+class ConcurrentTest {
+ private:
+  static constexpr uint32_t K = 4;
+
+  static uint64_t key(Key key) { return (key >> 40); }
+  static uint64_t gen(Key key) { return (key >> 8) & 0xffffffffu; }
+  static uint64_t hash(Key key) { return key & 0xff; }
+
+  static uint64_t HashNumbers(uint64_t k, uint64_t g) {
+    uint64_t data[2] = {k, g};
+    return Hash(reinterpret_cast<char*>(data), sizeof(data), 0);
+  }
+
+  static Key MakeKey(uint64_t k, uint64_t g) {
+    static_assert(sizeof(Key) == sizeof(uint64_t), "");
+    assert(k <= K);  // We sometimes pass K to seek to the end of the skiplist
+    assert(g <= 0xffffffffu);
+    return ((k << 40) | (g << 8) | (HashNumbers(k, g) & 0xff));
+  }
+
+  static bool IsValidKey(Key k) {
+    return hash(k) == (HashNumbers(key(k), gen(k)) & 0xff);
+  }
+
+  static Key RandomTarget(Random* rnd) {
+    switch (rnd->Next() % 10) {
+      case 0:
+        // Seek to beginning
+        return MakeKey(0, 0);
+      case 1:
+        // Seek to end
+        return MakeKey(K, 0);
+      default:
+        // Seek to middle
+        return MakeKey(rnd->Next() % K, 0);
+    }
+  }
+
+  // Per-key generation
+  struct State {
+    std::atomic<int> generation[K];
+    void Set(int k, int v) {
+      generation[k].store(v, std::memory_order_release);
+    }
+    int Get(int k) { return generation[k].load(std::memory_order_acquire); }
+
+    State() {
+      for (int k = 0; k < K; k++) {
+        Set(k, 0);
+      }
+    }
+  };
+
+  // Current state of the test
+  State current_;
+
+  Arena arena_;
+
+  // SkipList is not protected by mu_.  We just use a single writer
+  // thread to modify it.
+  SkipList<Key, Comparator> list_;
+
+ public:
+  ConcurrentTest() : list_(Comparator(), &arena_) {}
+
+  // REQUIRES: External synchronization
+  void WriteStep(Random* rnd) {
+    const uint32_t k = rnd->Next() % K;
+    const intptr_t g = current_.Get(k) + 1;
+    const Key key = MakeKey(k, g);
+    list_.Insert(key);
+    current_.Set(k, g);
+  }
+
+  void ReadStep(Random* rnd) {
+    // Remember the initial committed state of the skiplist.
+    State initial_state;
+    for (int k = 0; k < K; k++) {
+      initial_state.Set(k, current_.Get(k));
+    }
+
+    Key pos = RandomTarget(rnd);
+    SkipList<Key, Comparator>::Iterator iter(&list_);
+    iter.Seek(pos);
+    while (true) {
+      Key current;
+      if (!iter.Valid()) {
+        current = MakeKey(K, 0);
+      } else {
+        current = iter.key();
+        ASSERT_TRUE(IsValidKey(current)) << current;
+      }
+      ASSERT_LE(pos, current) << "should not go backwards";
+
+      // Verify that everything in [pos,current) was not present in
+      // initial_state.
+      while (pos < current) {
+        ASSERT_LT(key(pos), K) << pos;
+
+        // Note that generation 0 is never inserted, so it is ok if
+        // <*,0,*> is missing.
+        ASSERT_TRUE((gen(pos) == 0) ||
+                    (gen(pos) > static_cast<Key>(initial_state.Get(key(pos)))))
+            << "key: " << key(pos) << "; gen: " << gen(pos)
+            << "; initgen: " << initial_state.Get(key(pos));
+
+        // Advance to next key in the valid key space
+        if (key(pos) < key(current)) {
+          pos = MakeKey(key(pos) + 1, 0);
+        } else {
+          pos = MakeKey(key(pos), gen(pos) + 1);
+        }
+      }
+
+      if (!iter.Valid()) {
+        break;
+      }
+
+      if (rnd->Next() % 2) {
+        iter.Next();
+        pos = MakeKey(key(pos), gen(pos) + 1);
+      } else {
+        Key new_target = RandomTarget(rnd);
+        if (new_target > pos) {
+          pos = new_target;
+          iter.Seek(new_target);
+        }
+      }
+    }
+  }
+};*/
+
+
+__device__ void update_list(SkipList<Key, Comparator> * l, Key key) {
+  l->Insert(key);
+}
+
+__global__ void insert_skiplist(SkipList<Key, Comparator> * l, Random * rnd) {
+  unsigned row = threadIdx.x + blockIdx.x * blockDim.x;
+  for (int i = 0; i < 1000; i++ ) {
+    update_list(l, rnd->Next());
+  }
+}
+
+
+
+__host__ void host_insert_test() {
+  constexpr size_t sz_size = 1024;
+  Arena arena;
+  Comparator cmp;
+  Key * keys;
+  //SkipList<Key, Comparator> list(cmp, &arena);
+
+  SkipList<Key, Comparator> * ptr_list;
+  Random * device_rnd;
+
+  //cudaMallocManaged((void**)&keys, sizeof(Key) * sz_size);
+  cudaMallocManaged((void**)&device_rnd, sizeof(Random));
+  cudaMallocManaged((void**)&ptr_list, sizeof(SkipList<Key, Comparator>));
+
+  ptr_list = new SkipList<Key, Comparator>(cmp, &arena);
+  device_rnd = new Random(test::RandomSeed());
+  /*
+  for (int i = 0; i < 1000; i++) {
+    keys[i] = .Next();
+  }*/
+
+  dim3 blockSize(32, 1);
+
+  dim3 gridSize(32, 1);
+
+  insert_skiplist<<<gridSize, blockSize>>>(ptr_list, device_rnd);
+
+  cudaDeviceSynchronize();
+
+  cudaFree(&device_rnd);
+  cudaFree(ptr_list);
+}
+
+
+TEST(SkipTest, TestCudaInsert) {
+  host_insert_test();
+}
+
+/*
+// Needed when building in C++11 mode.
+constexpr uint32_t ConcurrentTest::K;
+
+// Simple test that does single-threaded testing of the ConcurrentTest
+// scaffolding.
+TEST(SkipTest, ConcurrentWithoutThreads) {
+  ConcurrentTest test;
+  Random rnd(test::RandomSeed());
+  for (int i = 0; i < 10000; i++) {
+    test.ReadStep(&rnd);
+    test.WriteStep(&rnd);
+  }
+}
+
+class TestState {
+ public:
+  ConcurrentTest t_;
+  int seed_;
+  std::atomic<bool> quit_flag_;
+
+  enum ReaderState { STARTING, RUNNING, DONE };
+
+  explicit TestState(int s)
+      : seed_(s), quit_flag_(false), state_(STARTING), state_cv_(&mu_) {}
+
+  void Wait(ReaderState s) LOCKS_EXCLUDED(mu_) {
+    mu_.Lock();
+    while (state_ != s) {
+      state_cv_.Wait();
+    }
+    mu_.Unlock();
+  }
+
+  void Change(ReaderState s) LOCKS_EXCLUDED(mu_) {
+    mu_.Lock();
+    state_ = s;
+    state_cv_.Signal();
+    mu_.Unlock();
+  }
+
+ private:
+  port::Mutex mu_;
+  ReaderState state_ GUARDED_BY(mu_);
+  port::CondVar state_cv_ GUARDED_BY(mu_);
+};
+
+static void ConcurrentReader(void* arg) {
+  TestState* state = reinterpret_cast<TestState*>(arg);
+  Random rnd(state->seed_);
+  int64_t reads = 0;
+  state->Change(TestState::RUNNING);
+  while (!state->quit_flag_.load(std::memory_order_acquire)) {
+    state->t_.ReadStep(&rnd);
+    ++reads;
+  }
+  state->Change(TestState::DONE);
+}
+
+static void RunConcurrent(int run) {
+  const int seed = test::RandomSeed() + (run * 100);
+  Random rnd(seed);
+  const int N = 1000;
+  const int kSize = 1000;
+  for (int i = 0; i < N; i++) {
+    if ((i % 100) == 0) {
+      std::fprintf(stderr, "Run %d of %d\n", i, N);
+    }
+    TestState state(seed + 1);
+    Env::Default()->Schedule(ConcurrentReader, &state);
+    state.Wait(TestState::RUNNING);
+    for (int i = 0; i < kSize; i++) {
+      state.t_.WriteStep(&rnd);
+    }
+    state.quit_flag_.store(true, std::memory_order_release);
+    state.Wait(TestState::DONE);
+  }
+}
+
+TEST(SkipTest, Concurrent1) { RunConcurrent(1); }
+TEST(SkipTest, Concurrent2) { RunConcurrent(2); }
+TEST(SkipTest, Concurrent3) { RunConcurrent(3); }
+TEST(SkipTest, Concurrent4) { RunConcurrent(4); }
+TEST(SkipTest, Concurrent5) { RunConcurrent(5); }
+*/
+
+}  // namespace leveldb
+
+int main(int argc, char** argv) {
+  testing::InitGoogleTest(&argc, argv);
+  return RUN_ALL_TESTS();
+}
-- 
cgit v1.3.1