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

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


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


class CudaSlice;

// A Comparator object provides a total order across slices that are
// used as keys in an sstable or a database.  A Comparator implementation
// must be thread-safe since leveldb may invoke its methods concurrently
// from multiple threads.
class LEVELDB_EXPORT CudaComparator {
 public:
  virtual ~CudaComparator();

  virtual  __device__ int Compare(const CudaSlice& a, const CudaSlice& b) const = 0;

  virtual __device__ __host__ const char* Name() const = 0;

  virtual __device__ void FindShortestSeparator(const CudaSlice* start,
                                     const CudaSlice& limit) const = 0;

  virtual __device__ void FindShortSuccessor(CudaSlice* key) const = 0;
};

// Return a builtin comparator that uses lexicographic byte-wise
// ordering.  The result remains the property of this module and
// must not be deleted.
LEVELDB_EXPORT const Comparator* BytewiseComparator();


// A comparator for internal keys that uses a specified comparator for
// the user key portion and breaks ties by decreasing sequence number.
class CudaInternalKeyComparator : public CudaComparator {
 private:
  const CudaComparator* user_comparator_;

 public:
  explicit CudaInternalKeyComparator(const CudaComparator* c) : user_comparator_(c) {}
  __device__ __host__ const char* Name() const override;
  __device__ int Compare(const CudaSlice& a, const CudaSlice& b) const override;
  __device__ void FindShortestSeparator(const CudaSlice* start,
                             const CudaSlice& limit) const override;
  __device__ void FindShortSuccessor(CudaSlice* key) const override;

  __device__ __host__ const CudaComparator* user_comparator() const { return user_comparator_; }

  __device__ int Compare(const InternalKey& a, const InternalKey& b) const;
};


class LEVELDB_EXPORT CudaSlice {
 public:
  // Create an empty slice.
  CudaSlice() : data_(""), size_(0) {}

  // Create a slice that refers to d[0,n-1].
  CudaSlice(const char* d, size_t n) : data_(d), size_(n) {}

  // Create a slice that refers to the contents of "s"
  //CudaSlice(const std::string& s) : data_(s.data()), size_(s.size()) {}

  // Create a slice that refers to s[0,strlen(s)-1]
  CudaSlice(const char* s) : data_(s), size_(strlen(s)) {}

  // Intentionally copyable.
  CudaSlice(const CudaSlice&) = default;
  CudaSlice& operator=(const CudaSlice&) = default;

  // Return a pointer to the beginning of the referenced data
  __device__ const char* data() const { return data_; }

  // Return the length (in bytes) of the referenced data
  __device__ __host__ size_t size() const { return size_; }

  // Return true iff the length of the referenced data is zero
  __device__ __host__ bool empty() const { return size_ == 0; }

  // Return the ith byte in the referenced data.
  // REQUIRES: n < size()
  /*__device__ char operator[](size_t n) const {
    assert(n < size());
    return data_[n];
  }*/
  __device__ char get_index(size_t n) const {
    assert(n < size());
    return data_[n];
  }

  // Change this slice to refer to an empty array
  __device__ __host__ void clear() {
    data_ = nullptr;
    size_ = 0;
  }

  // Drop the first "n" bytes from this slice.
  void remove_prefix(size_t n) {
    assert(n <= size());
    data_ += n;
    size_ -= n;
  }

  // Return a string that contains the copy of the referenced data.
  //std::string ToString() const { return std::string(data_, size_); }

  // Three-way comparison.  Returns value:
  //   <  0 iff "*this" <  "b",
  //   == 0 iff "*this" == "b",
  //   >  0 iff "*this" >  "b"
  __device__ int compare(const CudaSlice& b) const;

  // Return true iff "x" is a prefix of "*this"
  bool starts_with(const CudaSlice& x) const {
    return ((size_ >= x.size_) && (memcmp(data_, x.data_, x.size_) == 0));
  }

 private:
  const char* data_;
  size_t size_;
};


template<typename T>
__device__ T cudaMin(const T & a, const T &b) {
  return a>b?a:b;
}

CudaComparator::~CudaComparator() = default;

namespace {
class BytewiseComparatorImpl : public CudaComparator {
 public:
  BytewiseComparatorImpl() = default;



  __device__ __host__ const char* Name() const override { return "leveldb.BytewiseComparator"; }

  __device__ int Compare(const CudaSlice& a, const CudaSlice& b) const override {
    return a.compare(b);
  }

  __device__ void FindShortestSeparator(const CudaSlice * start,
                             const CudaSlice& limit) const override {
    // Find length of common prefix
    size_t min_length = std::min(start->size(), limit.size());
    size_t diff_index = 0;
    while ((diff_index < min_length) &&
           ((*start)[diff_index] == limit[diff_index])) {
      diff_index++;
    }

    if (diff_index >= min_length) {
      // Do not shorten if one string is a prefix of the other
    } else {
      uint8_t diff_byte = static_cast<uint8_t>((*start)[diff_index]);
      if (diff_byte < static_cast<uint8_t>(0xff) &&
          diff_byte + 1 < static_cast<uint8_t>(limit[diff_index])) {
        (*start)[diff_index]++;
        start->resize(diff_index + 1);
        assert(Compare(*start, limit) < 0);
      }
    }
  }

  void FindShortSuccessor(CudaSlice* key) const override {
    // Find first character that can be incremented
    size_t n = key->size();
    for (size_t i = 0; i < n; i++) {
      const uint8_t byte = (*key)[i];
      if (byte != static_cast<uint8_t>(0xff)) {
        (*key)[i] = byte + 1;
        key->resize(i + 1);
        return;
      }
    }
    // *key is a run of 0xffs.  Leave it alone.
  }
};
}  // namespace

const Comparator* BytewiseComparator() {
  static NoDestructor<BytewiseComparatorImpl> singleton;
  return singleton.get();
}

template<typename T>
__device__ int cudaMemcmp(const T * a, const T * b, size_t min_length) {
  for (size_t i = 0; i < min_length ; i++) {
    if (a[i] == b[i])
      continue;
    return a[i] < b[i] ? -1 : 1;
  }
  return 0;
}

/*__device__ inline bool operator==(const CudaSlice& x, const CudaSlice& y) {
  return ((x.size() == y.size()) &&
          (cudaMemcmp(x.data(), y.data(), x.size()) == 0));
}*/

/*inline bool operator!=(const CudaSlice& x, const CudaSlice& y) { return !(x == y); }*/



__device__ inline int CudaSlice::compare(const CudaSlice& b) const {
  const size_t min_len = (size_ < b.size_) ? size_ : b.size_;
  int r = cudaMemcmp(data_, b.data_, min_len);
  if (r == 0) {
    if (size_ < b.size_)
      r = -1;
    else if (size_ > b.size_)
      r = +1;
  }
  return r;
}



class MemTableCuda {
 public:
  // MemTables are reference counted.  The initial reference count
  // is zero and the caller must call Ref() at least once.
  __device__ explicit MemTableCuda(const InternalKeyComparator& comparator);

  MemTableCuda(const MemTableCuda&) = delete;
  MemTableCuda& operator=(const MemTableCuda&) = 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.
  __device__ void Add(SequenceNumber seq, ValueType type, const CudaSlice& 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.
  __device__ bool Get(const LookupKey& key, std::string* value, Status* s);

 private:
  friend class MemTableCudaIterator;
  friend class MemTableCudaBackwardIterator;
  friend __global__ void MemTableCudaAdd_(MemTableCuda *, size_t, char *);
  friend __global__ void MemTableCudaGet_(MemTableCuda *, char *, char **, size_t* malloc_size);

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

  typedef SkipList<const char*, KeyComparator> Table;

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

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

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

}  // namespace leveldb

#endif  // STORAGE_LEVELDB_DB_MEMTABLE_H_