// 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);
}

MemTable::MemTable(const InternalKeyComparator& comparator)
    : comparator_(comparator), refs_(0) {

}

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_); }


__global__ void Add_(MemTable * mtb, size_t encoded_len,char * encode_data) {
  char* buf = mtb->arena_.Allocate(encoded_len);
  memcpy(buf, encode_data, 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 point  : point to host memory
  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;
  const size_t encoded_len = VarintLength(internal_key_size) + internal_key_size +
                             VarintLength(val_size) + 8;


    char * key_mem = this->arena_.Allocate( key_size);
    std::memcpy(key_mem, key.data(), key_size);

    //char * tag_mem = key_mem + key_size;
    //EncodeFixed64(tag_mem, (s << 8) | type);

    char * val_mem = this->arena_.Allocate(val_size);
    std::memcpy(val_mem, value.data(), val_size);

  char * insert_val = new char[encoded_len], *cuda_insert = nullptr;
  //  EncodeVarint32(insert_val, encoded_len);
  char * p = EncodeVarint32(insert_val + 8, internal_key_size);
  //EncodeFixed64(p, reinterpret_cast<uint64_t>(key_mem));
  memcpy(p, key_mem, key_size);
  p += key_size;
  EncodeFixed64(p, (s << 8) | type );
  p += 8;
  //EncodeFixed64(p, reinterpret_cast<uint64_t>(val_size));
  p = EncodeVarint32(p, val_size);
  EncodeFixed64(p, reinterpret_cast<uint64_t>(val_mem));

  assert(p + 8 == insert_val + encoded_len);

  cudaMalloc((void**)&cuda_insert, encoded_len);
  cudaMemcpy(cuda_insert, insert_val, encoded_len, cudaMemcpyHostToDevice);

  Add_<<<1, 1>>>(this, encoded_len, cuda_insert);
  cudaDeviceSynchronize();

  cudaFree(cuda_insert);
  delete[] insert_val;
}

__global__ void Get_(MemTable * met, char * memkey, char ** data, size_t* malloc_size) {
  *data = nullptr;
  auto iter = met->getIter();
  iter.Seek(memkey);
  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.
    size_t key_size;
    const char *entry = iter.key();
    const char * p = GetVarint32PtrCuda(entry, entry + 5,
                       reinterpret_cast<uint32_t*>(&key_size));
    p = GetVarint32PtrCuda(p + key_size, p + key_size + 5, nullptr);
    *malloc_size = (p - entry + 8);
    cudaMalloc((void**)*data, *malloc_size);
    memcpy(*data, entry, *malloc_size);
  }
}

/*
__global__ void Get_Phase2() {
    // Correct user key
    const uint64_t tag = DecodeFixed64Cuda(key_ptr + key_length - 8);
    switch (static_cast<ValueType>(tag & 0xff)) {
      case kTypeValue: {
        SizedString v = GetLengthPrefixedSliceCuda(key_ptr + key_length);
        *value=v;
        *ret = true;
        break;
      }
      case kTypeDeletion:
        //*s = Status::NotFound(Slice());
        *is_not_found = true;
        *ret = true;
        break;
    }
}
*/

bool MemTable::Get(const LookupKey& key, std::string* value, Status* s) {

  Slice memkey = key.memtable_key();
  char * cuda_mem_key = nullptr;
  cudaMalloc((void**)&cuda_mem_key, memkey.size());
  cudaMemcpy(cuda_mem_key, memkey.data(), memkey.size(), cudaMemcpyHostToDevice);

  char ** cuda_skiplist_key = nullptr;
  cudaMalloc((void**)&cuda_skiplist_key, sizeof(char *));

  size_t * cuda_malloc_size = nullptr;
  cudaMalloc((void**)&cuda_malloc_size, sizeof(cuda_malloc_size));

  Get_<<<1,1>>>(this, cuda_mem_key, cuda_skiplist_key, cuda_malloc_size);
  cudaDeviceSynchronize();
  cudaFree(cuda_mem_key);

  if (*cuda_skiplist_key == nullptr) {
    cudaFree(cuda_skiplist_key);
    cudaFree(cuda_malloc_size);
    return false;
  }
  auto * malloc_size = new size_t;
  cudaMemcpy(malloc_size, cuda_malloc_size, sizeof(size_t), cudaMemcpyDeviceToHost);

  char * entry = this->host_arena_.Allocate(*malloc_size);
  cudaMemcpy(entry, cuda_skiplist_key, *malloc_size, cudaMemcpyDeviceToHost);

  cudaFree(cuda_skiplist_key);
  cudaFree(cuda_malloc_size);
  delete malloc_size;

    uint32_t key_length;
    const char* key_ptr = GetVarint32Ptr(entry, entry + 5, &key_length);
    if (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: {
          // HOW CAN WE GET IT?
          Slice v = GetLengthPrefixedSlice(key_ptr + key_length);
          value->assign(v.data(), v.size());
          return true;
        }
        case kTypeDeletion:
          *s = Status::NotFound(Slice());
          return true;
      }
    }


/*  bool * cuda_return_value = nullptr, *return_value = new bool, *is_not_found = nullptr;
  SizedString * return_string_value = nullptr;
  cudaMalloc((void**)&cuda_return_value, sizeof(bool));
  cudaMalloc((void**)&is_not_found, sizeof(bool));
  cudaMallocManaged((void**)&return_string_value, sizeof(SizedString));

  Get_<<<1, 1>>>(this, cuda_return_value, cuda_mem_key, nullptr, is_not_found);
  cudaDeviceSynchronize();

  cudaMemcpy(return_value, cuda_return_value, sizeof(bool), cudaMemcpyDeviceToHost);

  if (*return_value) {
    if (*is_not_found) {
      *s = Status::NotFound(Slice());
    } else {
      auto * local_sized_string = new SizedString();
      cudaMemcpy(local_sized_string, return_string_value, sizeof(SizedString), cudaMemcpyDeviceToHost);

      char * local_string = new char[local_sized_string->length];
      cudaMemcpy(local_string, local_sized_string->data, local_sized_string->length, cudaMemcpyDeviceToHost);
      value->assign(local_string, local_sized_string->length);
      delete local_sized_string;
      delete [] local_string;
    }
  }

  bool rvalue = *return_value;
  delete return_value;
  cudaFree(is_not_found);
  cudaFree(cuda_return_value);
  cudaFree(cuda_mem_key);
  cudaFree(return_string_value);*/
    return false;
}

}  // namespace leveldb