1 files changed, 203 insertions, 0 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