3 files changed, 294 insertions, 1 deletions
diff --git a/db/memtable.cu b/db/memtable.cu
new file mode 100644
index 0000000..c4c3a75
--- /dev/null
+++ b/db/memtable.cu
@@ -0,0 +1,203 @@
+// 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/memtable.cuh"
+#include "db/dbformat.h"
+#include "leveldb/comparator.h"
+#include "leveldb/env.h"
+#include "leveldb/iterator.h"
+#include "util/coding.cuh"
+
+namespace leveldb {
+
+static Slice GetLengthPrefixedSlice(const char* data) {
+  uint32_t len;
+  const char* p = data;
+  p = GetVarint32Ptr(p, p + 5, &len);  // +5: we assume "p" is not corrupted
+  return Slice(p, len);
+}
+
+__device__ MemTable::MemTable(const InternalKeyComparator& comparator)
+    : comparator_(comparator), refs_(0), table_(comparator_, &arena_) {}
+
+MemTable::~MemTable() { assert(refs_ == 0); }
+
+size_t MemTable::ApproximateMemoryUsage() { return arena_.MemoryUsage(); }
+
+int MemTable::KeyComparator::operator()(const char* aptr,
+                                        const char* bptr) const {
+  // Internal keys are encoded as length-prefixed strings.
+  Slice a = GetLengthPrefixedSlice(aptr);
+  Slice b = GetLengthPrefixedSlice(bptr);
+  return comparator.Compare(a, b);
+}
+
+// Encode a suitable internal key target for "target" and return it.
+// Uses *scratch as scratch space, and the returned pointer will point
+// into this scratch space.
+static const char* EncodeKey(std::string* scratch, const Slice& target) {
+  scratch->clear();
+  PutVarint32(scratch, target.size());
+  scratch->append(target.data(), target.size());
+  return scratch->data();
+}
+
+class MemTableIterator : public Iterator {
+ public:
+  explicit MemTableIterator(MemTable::Table* table) : iter_(table) {}
+
+  MemTableIterator(const MemTableIterator&) = delete;
+  MemTableIterator& operator=(const MemTableIterator&) = delete;
+
+  ~MemTableIterator() override = default;
+
+  bool Valid() const override { return iter_.Valid(); }
+  void Seek(const Slice& k) override { iter_.Seek(EncodeKey(&tmp_, k)); }
+  void SeekToFirst() override { iter_.SeekToFirst(); }
+  void SeekToLast() override { iter_.SeekToLast(); }
+  void Next() override { iter_.Next(); }
+  void Prev() override { iter_.Prev(); }
+  Slice key() const override { return GetLengthPrefixedSlice(iter_.key()); }
+  Slice value() const override {
+    Slice key_slice = GetLengthPrefixedSlice(iter_.key());
+    return GetLengthPrefixedSlice(key_slice.data() + key_slice.size());
+  }
+
+  Status status() const override { return Status::OK(); }
+
+ private:
+  MemTable::Table::Iterator iter_;
+  std::string tmp_;  // For passing to EncodeKey
+};
+
+Iterator* MemTable::NewIterator() { return new MemTableIterator(&table_); }
+
+__device__ char* EncodeVarint32Device(char* dst, uint32_t v) {
+  // Operate on characters as unsigneds
+  uint8_t* ptr = reinterpret_cast<uint8_t*>(dst);
+  static const int B = 128;
+  if (v < (1 << 7)) {
+    *(ptr++) = v;
+  } else if (v < (1 << 14)) {
+    *(ptr++) = v | B;
+    *(ptr++) = v >> 7;
+  } else if (v < (1 << 21)) {
+    *(ptr++) = v | B;
+    *(ptr++) = (v >> 7) | B;
+    *(ptr++) = v >> 14;
+  } else if (v < (1 << 28)) {
+    *(ptr++) = v | B;
+    *(ptr++) = (v >> 7) | B;
+    *(ptr++) = (v >> 14) | B;
+    *(ptr++) = v >> 21;
+  } else {
+    *(ptr++) = v | B;
+    *(ptr++) = (v >> 7) | B;
+    *(ptr++) = (v >> 14) | B;
+    *(ptr++) = (v >> 21) | B;
+    *(ptr++) = v >> 28;
+  }
+  return reinterpret_cast<char*>(ptr);
+}
+
+__device__ inline void EncodeFixed64Device(char* dst, uint64_t value) {
+  uint8_t* const buffer = reinterpret_cast<uint8_t*>(dst);
+
+  // Recent clang and gcc optimize this to a single mov / str instruction.
+  buffer[0] = static_cast<uint8_t>(value);
+  buffer[1] = static_cast<uint8_t>(value >> 8);
+  buffer[2] = static_cast<uint8_t>(value >> 16);
+  buffer[3] = static_cast<uint8_t>(value >> 24);
+  buffer[4] = static_cast<uint8_t>(value >> 32);
+  buffer[5] = static_cast<uint8_t>(value >> 40);
+  buffer[6] = static_cast<uint8_t>(value >> 48);
+  buffer[7] = static_cast<uint8_t>(value >> 56);
+}
+
+__global__ void Add_(MemTable * mtb, SequenceNumber s, ValueType type,
+                     size_t encoded_len, size_t internal_key_size,
+                     char * key_data, size_t key_size, char * value_data, size_t val_size) {
+  char* buf = mtb->arena_.Allocate(encoded_len);
+  char* p = EncodeVarint32Device(buf, internal_key_size);
+  memcpy(p, key_data, key_size);
+  //std::memcpy(p, key_data, key_size);
+  p += key_size;
+  EncodeFixed64Device(p, (s << 8) | type);
+  p += 8;
+  p = EncodeVarint32Device(p, val_size);
+  //std::memcpy(p, value.data(), val_size);
+  memcpy(p, value_data, val_size);
+  assert(p + val_size == buf + encoded_len);
+  mtb->table_.Insert(buf);
+}
+
+void MemTable::Add(SequenceNumber s, ValueType type, const Slice& key,
+                   const Slice& value) {
+  // Format of an entry is concatenation of:
+  //  key_size     : varint32 of internal_key.size()
+  //  key bytes    : char[internal_key.size()]
+  //  tag          : uint64((sequence << 8) | type)
+  //  value_size   : varint32 of value.size()
+  //  value bytes  : char[value.size()]
+  size_t key_size = key.size();
+  size_t val_size = value.size();
+  size_t internal_key_size = key_size + 8;
+  const size_t encoded_len = VarintLength(internal_key_size) +
+                             internal_key_size + VarintLength(val_size) +
+                             val_size;
+  char * key_mem = nullptr;
+  cudaMalloc((void**)&key_mem, key_size);
+  cudaMemcpy(key_mem, key.data(), key_size, cudaMemcpyHostToDevice);
+  char * value_mem = nullptr;
+  cudaMalloc((void**)&value_mem, val_size);
+  cudaMemcpy(value_mem, value.data(), val_size, cudaMemcpyHostToDevice);
+
+  Add_<<<1, 1>>>(this, s, type, encoded_len, internal_key_size, key_mem, key_size, value_mem, val_size);
+  cudaDeviceSynchronize();
+
+  cudaFree(key_mem);
+  cudaFree(value_mem);
+}
+
+__global__ void Get_(MemTable * met, bool * ret) {
+
+    Slice memkey = key.memtable_key();
+    met->Table::Iterator iter(&met->table_);
+    iter.Seek(memkey.data());
+    if (iter.Valid()) {
+      // entry format is:
+      //    klength  varint32
+      //    userkey  char[klength]
+      //    tag      uint64
+      //    vlength  varint32
+      //    value    char[vlength]
+      // Check that it belongs to same user key.  We do not check the
+      // sequence number since the Seek() call above should have skipped
+      // all entries with overly large sequence numbers.
+      const char* entry = iter.key();
+      uint32_t key_length;
+      const char* key_ptr = GetVarint32Ptr(entry, entry + 5, &key_length);
+      if (met->comparator_.comparator.user_comparator()->Compare(
+              Slice(key_ptr, key_length - 8), key.user_key()) == 0) {
+        // Correct user key
+        const uint64_t tag = DecodeFixed64(key_ptr + key_length - 8);
+        switch (static_cast<ValueType>(tag & 0xff)) {
+          case kTypeValue: {
+            Slice v = GetLengthPrefixedSlice(key_ptr + key_length);
+            value->assign(v.data(), v.size());
+            return true;
+          }
+          case kTypeDeletion:
+            *s = Status::NotFound(Slice());
+            return true;
+        }
+      }
+    }
+}
+
+bool MemTable::Get(const LookupKey& key, std::string* value, Status* s) {
+  return false;
+}
+
+}  // namespace leveldb
diff --git a/db/memtable.cuh b/db/memtable.cuh
new file mode 100644
index 0000000..ac3c2b6
--- /dev/null
+++ b/db/memtable.cuh
@@ -0,0 +1,90 @@
+// 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_MEMTABLE_H_
+#define STORAGE_LEVELDB_DB_MEMTABLE_H_
+
+#include <string>
+
+#include "db/dbformat.h"
+#include "db/skiplist.cuh"
+#include "leveldb/db.h"
+#include "util/arena.cuh"
+
+namespace leveldb {
+
+class InternalKeyComparator;
+class MemTableIterator;
+
+class MemTable {
+ public:
+  // MemTables are reference counted.  The initial reference count
+  // is zero and the caller must call Ref() at least once.
+  __device__ explicit MemTable(const InternalKeyComparator& comparator);
+
+  MemTable(const MemTable&) = delete;
+  MemTable& operator=(const MemTable&) = delete;
+
+  // Increase reference count.
+  void Ref() { ++refs_; }
+
+  // Drop reference count.  Delete if no more references exist.
+  void Unref() {
+    --refs_;
+    assert(refs_ >= 0);
+    if (refs_ <= 0) {
+      delete this;
+    }
+  }
+
+  // Returns an estimate of the number of bytes of data in use by this
+  // data structure. It is safe to call when MemTable is being modified.
+  size_t ApproximateMemoryUsage();
+
+  // Return an iterator that yields the contents of the memtable.
+  //
+  // The caller must ensure that the underlying MemTable remains live
+  // while the returned iterator is live.  The keys returned by this
+  // iterator are internal keys encoded by AppendInternalKey in the
+  // db/format.{h,cc} module.
+  Iterator* NewIterator();
+
+  // Add an entry into memtable that maps key to value at the
+  // specified sequence number and with the specified type.
+  // Typically value will be empty if type==kTypeDeletion.
+  void Add(SequenceNumber seq, ValueType type, const Slice& key,
+           const Slice& value);
+
+  // If memtable contains a value for key, store it in *value and return true.
+  // If memtable contains a deletion for key, store a NotFound() error
+  // in *status and return true.
+  // Else, return false.
+  bool Get(const LookupKey& key, std::string* value, Status* s);
+
+ private:
+  friend class MemTableIterator;
+  friend class MemTableBackwardIterator;
+  friend __global__ void Add_(MemTable *, SequenceNumber, ValueType, size_t, size_t, char*, size_t, char *, size_t);
+  friend __global__ void Get_(MemTable *, bool * ret);
+
+  struct KeyComparator {
+    const InternalKeyComparator comparator;
+    __device__ explicit KeyComparator(const InternalKeyComparator& c) : comparator(c) {}
+    __device__ int operator()(const char* a, const char* b) const;
+    __device__ ~KeyComparator() = default;
+  };
+
+  typedef SkipList<const char*, KeyComparator> Table;
+
+  ~MemTable();  // Private since only Unref() should be used to delete it
+
+  KeyComparator comparator_;
+  int refs_;
+  Arena arena_;
+  Table table_;
+};
+
+}  // namespace leveldb
+
+#endif  // STORAGE_LEVELDB_DB_MEMTABLE_H_
diff --git a/db/skiplist.cuh b/db/skiplist.cuh
index 1fd7091..125a790 100644
--- a/db/skiplist.cuh
+++ b/db/skiplist.cuh
@@ -218,7 +218,7 @@ struct SkipList<Key, Comparator>::Node {
     // version of the returned Node.
     return next_[n].load(cuda::memory_order_acquire);
   }
-  __device__ void SetNext(int n, Node* x) {
+  __device__ __host__ 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.