// 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