// 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"
#include "util/arena.h"

namespace leveldb {

class InternalKeyComparator;
class MemTableIterator;

__device__ const char* GetVarint32PtrFallbackCuda(const char* p, const char* limit,
                                                  uint32_t* value) {
  uint32_t result = 0;
  for (uint32_t shift = 0; shift <= 28 && p < limit; shift += 7) {
    uint32_t byte = *(reinterpret_cast<const uint8_t*>(p));
    p++;
    if (byte & 128) {
      // More bytes are present
      result |= ((byte & 127) << shift);
    } else {
      result |= (byte << shift);
      *value = result;
      return reinterpret_cast<const char*>(p);
    }
  }
  return nullptr;
}

__device__ inline const char* GetVarint32PtrCuda(const char* p, const char* limit,
                                                 uint32_t* value) {
  if (p < limit) {
    uint32_t result = *(reinterpret_cast<const uint8_t*>(p));
    if ((result & 128) == 0) {
      *value = result;
      return p + 1;
    }
  }
  return GetVarint32PtrFallbackCuda(p, limit, value);
}

struct SizedString {
  char * data;
  size_t length;

  __host__ explicit SizedString() {
    this->data = nullptr;
    this->length = 0;
  }

  __device__ explicit SizedString(const char* p, size_t len) {
    this->data = nullptr;
    cudaMalloc((void**)&this->data, len);
    memcpy(this->data, p, len);
    this->length = len;
  }

  __host__ __device__ SizedString(const SizedString & other) {
    this->data = other.data;
    this->length = other.length;
  }

  __host__ __device__ SizedString & operator=(const SizedString& other) = default;
};


__device__ static SizedString GetLengthPrefixedSliceCuda(const char* data) {
  uint32_t len;
  const char* p = data;
  p = GetVarint32PtrCuda(p, p + 5, &len);  // +5: we assume "p" is not corrupted
  return SizedString(p, len);
}

class MemTable {
 public:
  // MemTables are reference counted.  The initial reference count
  // is zero and the caller must call Ref() at least once.
  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 *, size_t, char *);
  friend __global__ void Get_(MemTable *, char *, char **, size_t* malloc_size);

  struct KeyComparator {
    const InternalKeyComparator comparator;
    explicit KeyComparator(const InternalKeyComparator& c) : comparator(c) {}
    int operator()(const char* a, const char* b) const;
    ~KeyComparator() = default;
  };

  typedef SkipList<const char*, KeyComparator> Table;

  __device__  Table::Iterator getIter() {
    Table::Iterator iter(&this->table_);
    return iter;
  }

  ~MemTable();  // Private since only Unref() should be used to delete it

  KeyComparator comparator_;
  int refs_;
  Arena host_arena_;
  ArenaCuda arena_;
  Table table_;
};

}  // namespace leveldb

#endif  // STORAGE_LEVELDB_DB_MEMTABLE_H_