feat: Add skiplist cuda version

Signed-off-by: KunoiSayami <[email protected]>

14 files changed, 680 insertions, 24 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index ad9a488..8b6fac1 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -120,6 +120,8 @@ include(GNUInstallDirs)
 
 set(CMAKE_CUDA_STANDARD 14)
 find_package(CUDA REQUIRED)
+#find_package(Thrust REQUIRED CONFIG)
+#thrust_create_target(Thrust)
 
 add_library(leveldb "")
 target_sources(leveldb
@@ -146,6 +148,8 @@ target_sources(leveldb
     "db/memtable.h"
     "db/repair.cc"
     "db/skiplist.h"
+        "db/skiplist.cuh"
+    "db/skiplist.cu"
     "db/snapshot.h"
     "db/table_cache.cc"
     "db/table_cache.h"
@@ -174,8 +178,10 @@ target_sources(leveldb
     "table/table.cc"
     "table/two_level_iterator.cc"
     "table/two_level_iterator.h"
-    "util/arena.cu"
-    "util/arena.cuh"
+        "util/arena.cu"
+        "util/arena.cuh"
+        "util/arena.cc"
+        "util/arena.h"
     "util/bloom.cc"
     "util/cache.cc"
     "util/coding.cc"
@@ -193,6 +199,8 @@ target_sources(leveldb
     "util/no_destructor.h"
     "util/options.cc"
     "util/random.h"
+    "util/random.cuh"
+    "util/random.cu"
     "util/status.cc"
 
   # Only CMake 3.3+ supports PUBLIC sources in targets exported by "install".
@@ -213,7 +221,8 @@ target_sources(leveldb
     "${LEVELDB_PUBLIC_INCLUDE_DIR}/table.h"
     "${LEVELDB_PUBLIC_INCLUDE_DIR}/write_batch.h"
 )
-set_target_properties(leveldb PROPERTIES CUDA_ARCHITECTURES "35;50;72")
+set_target_properties(leveldb PROPERTIES CUDA_ARCHITECTURES "75")
+#target_link_libraries(leveldb Thrust)
 
 if (WIN32)
   target_sources(leveldb
@@ -291,6 +300,7 @@ add_executable(leveldbutil
   "db/leveldbutil.cc"
 )
 target_link_libraries(leveldbutil leveldb)
+set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Ldb")
 
 if(LEVELDB_BUILD_TESTS)
   enable_testing()
@@ -340,6 +350,7 @@ if(LEVELDB_BUILD_TESTS)
           LEVELDB_HAS_PORT_CONFIG_H=1
       )
     endif(NOT HAVE_CXX17_HAS_INCLUDE)
+    set_target_properties("${test_target_name}" PROPERTIES CUDA_ARCHITECTURES "75")
 
     add_test(NAME "${test_target_name}" COMMAND "${test_target_name}")
   endfunction(leveldb_test)
@@ -363,7 +374,7 @@ if(LEVELDB_BUILD_TESTS)
     leveldb_test("db/filename_test.cc")
     leveldb_test("db/log_test.cc")
     leveldb_test("db/recovery_test.cc")
-    leveldb_test("db/skiplist_test.cc")
+    leveldb_test("db/skiplist_test.cu")
     leveldb_test("db/version_edit_test.cc")
     leveldb_test("db/version_set_test.cc")
     leveldb_test("db/write_batch_test.cc")
@@ -373,7 +384,7 @@ if(LEVELDB_BUILD_TESTS)
     leveldb_test("table/filter_block_test.cc")
     leveldb_test("table/table_test.cc")
 
-    leveldb_test("util/arena_test.cu")
+    leveldb_test("util/arena_test.cc")
     leveldb_test("util/bloom_test.cc")
     leveldb_test("util/cache_test.cc")
     leveldb_test("util/coding_test.cc")
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.cu
index 402cb47..53f60f8 100644
--- a/db/skiplist_test.cc
+++ b/db/skiplist_test.cu
@@ -2,9 +2,8 @@
 // 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 "db/skiplist.cuh"
+//#include <atomic>
 #include <set>
 
 #include "leveldb/env.h"
@@ -13,7 +12,7 @@
 #include "port/thread_annotations.h"
 #include "util/arena.cuh"
 #include "util/hash.h"
-#include "util/random.h"
+#include "util/random.cuh"
 #include "util/testutil.h"
 
 #include "gtest/gtest.h"
@@ -34,7 +33,7 @@ struct Comparator {
   }
 };
 
-TEST(SkipTest, Empty) {
+/*TEST(SkipTest, Empty) {
   Arena arena;
   Comparator cmp;
   SkipList<Key, Comparator> list(cmp, &arena);
@@ -281,8 +280,61 @@ class ConcurrentTest {
       }
     }
   }
-};
+};*/
+
+
+__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;
 
@@ -366,6 +418,7 @@ TEST(SkipTest, Concurrent2) { RunConcurrent(2); }
 TEST(SkipTest, Concurrent3) { RunConcurrent(3); }
 TEST(SkipTest, Concurrent4) { RunConcurrent(4); }
 TEST(SkipTest, Concurrent5) { RunConcurrent(5); }
+*/
 
 }  // namespace leveldb
 
diff --git a/util/arena.cc b/util/arena.cc
new file mode 100644
index 0000000..41c6a0d
--- /dev/null
+++ b/util/arena.cc
@@ -0,0 +1,68 @@
+// 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 "util/arena.h"
+
+namespace leveldb {
+
+static const int kBlockSize = 4096;
+
+Arena::Arena()
+    : alloc_ptr_(nullptr), alloc_bytes_remaining_(0), memory_usage_(0) {}
+
+Arena::~Arena() {
+  for (size_t i = 0; i < blocks_.size(); i++) {
+    //cudaFree(blocks_[i]);
+    delete [] blocks_[i];
+  }
+}
+
+char* Arena::AllocateFallback(size_t bytes) {
+  if (bytes > kBlockSize / 4) {
+    // Object is more than a quarter of our block size.  Allocate it separately
+    // to avoid wasting too much space in leftover bytes.
+    char* result = AllocateNewBlock(bytes);
+    return result;
+  }
+
+  // We waste the remaining space in the current block.
+  alloc_ptr_ = AllocateNewBlock(kBlockSize);
+  alloc_bytes_remaining_ = kBlockSize;
+
+  char* result = alloc_ptr_;
+  alloc_ptr_ += bytes;
+  alloc_bytes_remaining_ -= bytes;
+  return result;
+}
+
+char* Arena::AllocateAligned(size_t bytes) {
+  const int align = (sizeof(void*) > 8) ? sizeof(void*) : 8;
+  static_assert((align & (align - 1)) == 0,
+                "Pointer size should be a power of 2");
+  size_t current_mod = reinterpret_cast<uintptr_t>(alloc_ptr_) & (align - 1);
+  size_t slop = (current_mod == 0 ? 0 : align - current_mod);
+  size_t needed = bytes + slop;
+  char* result;
+  if (needed <= alloc_bytes_remaining_) {
+    result = alloc_ptr_ + slop;
+    alloc_ptr_ += needed;
+    alloc_bytes_remaining_ -= needed;
+  } else {
+    // AllocateFallback always returned aligned memory
+    result = AllocateFallback(bytes);
+  }
+  assert((reinterpret_cast<uintptr_t>(result) & (align - 1)) == 0);
+  return result;
+}
+
+char* Arena::AllocateNewBlock(size_t block_bytes) {
+  char* result = new char[block_bytes];
+  //cudaMallocManaged((void **)&result, sizeof(char) * block_bytes);
+  blocks_.push_back(result);
+  memory_usage_.fetch_add(block_bytes + sizeof(char*),
+                          std::memory_order_relaxed);
+  return result;
+}
+
+}  // namespace leveldb
diff --git a/util/arena.cu b/util/arena.cu
index 09ec6e6..5075318 100644
--- a/util/arena.cu
+++ b/util/arena.cu
@@ -13,11 +13,12 @@ Arena::Arena()
 
 Arena::~Arena() {
   for (size_t i = 0; i < blocks_.size(); i++) {
-    cudaFree(blocks_[i]);
+    //cudaFree(blocks_[i]);
+    delete [] blocks_[i];
   }
 }
 
-char* Arena::AllocateFallback(size_t bytes) {
+__device__ char* Arena::AllocateFallback(size_t bytes) {
   if (bytes > kBlockSize / 4) {
     // Object is more than a quarter of our block size.  Allocate it separately
     // to avoid wasting too much space in leftover bytes.
@@ -35,7 +36,7 @@ char* Arena::AllocateFallback(size_t bytes) {
   return result;
 }
 
-char* Arena::AllocateAligned(size_t bytes) {
+__device__ char* Arena::AllocateAligned(size_t bytes) {
   const int align = (sizeof(void*) > 8) ? sizeof(void*) : 8;
   static_assert((align & (align - 1)) == 0,
                 "Pointer size should be a power of 2");
@@ -56,8 +57,8 @@ char* Arena::AllocateAligned(size_t bytes) {
 }
 
 char* Arena::AllocateNewBlock(size_t block_bytes) {
-  char* result = nullptr;
-  cudaMallocManaged((void **)&result, sizeof(char) * block_bytes);
+  char* result = new char[block_bytes];
+  //cudaMallocManaged((void **)&result, sizeof(char) * block_bytes);
   blocks_.push_back(result);
   memory_usage_.fetch_add(block_bytes + sizeof(char*),
                           std::memory_order_relaxed);
diff --git a/util/arena.cuh b/util/arena.cuh
index 68fc55d..618b426 100644
--- a/util/arena.cuh
+++ b/util/arena.cuh
@@ -11,6 +11,8 @@
 #include <cstdint>
 #include <vector>
 
+#include <cuda/atomic>
+
 namespace leveldb {
 
 class Arena {
@@ -26,30 +28,31 @@ class Arena {
   char* Allocate(size_t bytes);
 
   // Allocate memory with the normal alignment guarantees provided by malloc.
-  char* AllocateAligned(size_t bytes);
+  __device__ char* AllocateAligned(size_t bytes);
 
   // Returns an estimate of the total memory usage of data allocated
   // by the arena.
   size_t MemoryUsage() const {
-    return memory_usage_.load(std::memory_order_relaxed);
+    return memory_usage_.load(cuda::memory_order_relaxed);
   }
 
  private:
-  char* AllocateFallback(size_t bytes);
-  char* AllocateNewBlock(size_t block_bytes);
+  __device__ char* AllocateFallback(size_t bytes);
+  __device__ char* AllocateNewBlock(size_t block_bytes);
 
   // Allocation state
   char* alloc_ptr_;
   size_t alloc_bytes_remaining_;
 
   // Array of new[] allocated memory blocks
+  //thrust::host_vector<char *> blocks_;
   std::vector<char*> blocks_;
 
   // Total memory usage of the arena.
   //
   // TODO(costan): This member is accessed via atomics, but the others are
   //               accessed without any locking. Is this OK?
-  std::atomic<size_t> memory_usage_;
+  cuda::atomic<size_t> memory_usage_;
 };
 
 inline char* Arena::Allocate(size_t bytes) {
diff --git a/util/arena.h b/util/arena.h
new file mode 100644
index 0000000..68fc55d
--- /dev/null
+++ b/util/arena.h
@@ -0,0 +1,71 @@
+// 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_UTIL_ARENA_H_
+#define STORAGE_LEVELDB_UTIL_ARENA_H_
+
+#include <atomic>
+#include <cassert>
+#include <cstddef>
+#include <cstdint>
+#include <vector>
+
+namespace leveldb {
+
+class Arena {
+ public:
+  Arena();
+
+  Arena(const Arena&) = delete;
+  Arena& operator=(const Arena&) = delete;
+
+  ~Arena();
+
+  // Return a pointer to a newly allocated memory block of "bytes" bytes.
+  char* Allocate(size_t bytes);
+
+  // Allocate memory with the normal alignment guarantees provided by malloc.
+  char* AllocateAligned(size_t bytes);
+
+  // Returns an estimate of the total memory usage of data allocated
+  // by the arena.
+  size_t MemoryUsage() const {
+    return memory_usage_.load(std::memory_order_relaxed);
+  }
+
+ private:
+  char* AllocateFallback(size_t bytes);
+  char* AllocateNewBlock(size_t block_bytes);
+
+  // Allocation state
+  char* alloc_ptr_;
+  size_t alloc_bytes_remaining_;
+
+  // Array of new[] allocated memory blocks
+  std::vector<char*> blocks_;
+
+  // Total memory usage of the arena.
+  //
+  // TODO(costan): This member is accessed via atomics, but the others are
+  //               accessed without any locking. Is this OK?
+  std::atomic<size_t> memory_usage_;
+};
+
+inline char* Arena::Allocate(size_t bytes) {
+  // The semantics of what to return are a bit messy if we allow
+  // 0-byte allocations, so we disallow them here (we don't need
+  // them for our internal use).
+  assert(bytes > 0);
+  if (bytes <= alloc_bytes_remaining_) {
+    char* result = alloc_ptr_;
+    alloc_ptr_ += bytes;
+    alloc_bytes_remaining_ -= bytes;
+    return result;
+  }
+  return AllocateFallback(bytes);
+}
+
+}  // namespace leveldb
+
+#endif  // STORAGE_LEVELDB_UTIL_ARENA_H_
diff --git a/util/arena_test.cu b/util/arena_test.cc
index 9b87785..d2ead39 100644
--- a/util/arena_test.cu
+++ b/util/arena_test.cc
@@ -2,7 +2,7 @@
 // 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 "util/arena.cuh"
+#include "util/arena.h"
 #include "util/random.h"
 
 #include "gtest/gtest.h"
diff --git a/util/random.cu b/util/random.cu
new file mode 100644
index 0000000..715ab5a
--- /dev/null
+++ b/util/random.cu
@@ -0,0 +1 @@
+#include "random.cuh"
diff --git a/util/random.cuh b/util/random.cuh
new file mode 100644
index 0000000..38eafe2
--- /dev/null
+++ b/util/random.cuh
@@ -0,0 +1,63 @@
+// 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_UTIL_RANDOM_H_
+#define STORAGE_LEVELDB_UTIL_RANDOM_H_
+
+#include <cstdint>
+
+namespace leveldb {
+
+// A very simple random number generator.  Not especially good at
+// generating truly random bits, but good enough for our needs in this
+// package.
+class Random {
+ private:
+  uint32_t seed_;
+
+ public:
+  explicit Random(uint32_t s) : seed_(s & 0x7fffffffu) {
+    // Avoid bad seeds.
+    if (seed_ == 0 || seed_ == 2147483647L) {
+      seed_ = 1;
+    }
+  }
+  __device__ uint32_t Next() {
+    static const uint32_t M = 2147483647L;  // 2^31-1
+    static const uint64_t A = 16807;        // bits 14, 8, 7, 5, 2, 1, 0
+    // We are computing
+    //       seed_ = (seed_ * A) % M,    where M = 2^31-1
+    //
+    // seed_ must not be zero or M, or else all subsequent computed values
+    // will be zero or M respectively.  For all other values, seed_ will end
+    // up cycling through every number in [1,M-1]
+    uint64_t product = seed_ * A;
+
+    // Compute (product % M) using the fact that ((x << 31) % M) == x.
+    seed_ = static_cast<uint32_t>((product >> 31) + (product & M));
+    // The first reduction may overflow by 1 bit, so we may need to
+    // repeat.  mod == M is not possible; using > allows the faster
+    // sign-bit-based test.
+    if (seed_ > M) {
+      seed_ -= M;
+    }
+    return seed_;
+  }
+  // Returns a uniformly distributed value in the range [0..n-1]
+  // REQUIRES: n > 0
+  __device__ uint32_t Uniform(int n) { return Next() % n; }
+
+  // Randomly returns true ~"1/n" of the time, and false otherwise.
+  // REQUIRES: n > 0
+  __device__ bool OneIn(int n) { return (Next() % n) == 0; }
+
+  // Skewed: pick "base" uniformly from range [0,max_log] and then
+  // return "base" random bits.  The effect is to pick a number in the
+  // range [0,2^max_log-1] with exponential bias towards smaller numbers.
+  __device__ uint32_t Skewed(int max_log) { return Uniform(1 << Uniform(max_log + 1)); }
+};
+
+}  // namespace leveldb
+
+#endif  // STORAGE_LEVELDB_UTIL_RANDOM_H_
diff --git a/util/random.h b/util/random.h
index fe76ab4..d7cbdd9 100644
--- a/util/random.h
+++ b/util/random.h
@@ -58,6 +58,7 @@ class Random {
   uint32_t Skewed(int max_log) { return Uniform(1 << Uniform(max_log + 1)); }
 };
 
+
 }  // namespace leveldb
 
 #endif  // STORAGE_LEVELDB_UTIL_RANDOM_H_