aboutsummaryrefslogtreecommitdiff
path: root/db/memtable.cuh
diff options
context:
space:
mode:
Diffstat (limited to 'db/memtable.cuh')
-rw-r--r--db/memtable.cuh325
1 files changed, 288 insertions, 37 deletions
diff --git a/db/memtable.cuh b/db/memtable.cuh
index 766eb03..6f0b2cd 100644
--- a/db/memtable.cuh
+++ b/db/memtable.cuh
@@ -18,36 +18,6 @@ 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;
@@ -73,13 +43,6 @@ struct SizedString {
};
-__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
@@ -154,6 +117,294 @@ class MemTable {
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_