aboutsummaryrefslogtreecommitdiff
path: root/db/memtable.cuh
blob: 6f0b2cdcffbef649f43e40e72abe49d7ec59af80 (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
// 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_