From 706cae9f460fe698857f47aaa42eb517ea2cb09f Mon Sep 17 00:00:00 2001
From: KunoiSayami <i@leanhe.dev>
Date: Thu, 7 Apr 2022 01:32:17 +0800
Subject: feat: Finish implement cuda memtable

Signed-off-by: KunoiSayami <i@leanhe.dev>
---
 db/memtable.cuh | 83 ++++++++++++++++++++++++++++++++++++++++++++++++++++-----
 1 file changed, 76 insertions(+), 7 deletions(-)

(limited to 'db/memtable.cuh')

diff --git a/db/memtable.cuh b/db/memtable.cuh
index ac3c2b6..766eb03 100644
--- a/db/memtable.cuh
+++ b/db/memtable.cuh
@@ -11,17 +11,80 @@
 #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.
-  __device__ explicit MemTable(const InternalKeyComparator& comparator);
+  explicit MemTable(const InternalKeyComparator& comparator);
 
   MemTable(const MemTable&) = delete;
   MemTable& operator=(const MemTable&) = delete;
@@ -65,23 +128,29 @@ class MemTable {
  private:
   friend class MemTableIterator;
   friend class MemTableBackwardIterator;
-  friend __global__ void Add_(MemTable *, SequenceNumber, ValueType, size_t, size_t, char*, size_t, char *, size_t);
-  friend __global__ void Get_(MemTable *, bool * ret);
+  friend __global__ void Add_(MemTable *, size_t, char *);
+  friend __global__ void Get_(MemTable *, char *, char **, size_t* malloc_size);
 
   struct KeyComparator {
     const InternalKeyComparator comparator;
-    __device__ explicit KeyComparator(const InternalKeyComparator& c) : comparator(c) {}
-    __device__ int operator()(const char* a, const char* b) const;
-    __device__ ~KeyComparator() = default;
+    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 arena_;
+  Arena host_arena_;
+  ArenaCuda arena_;
   Table table_;
 };
 
-- 
cgit v1.3.1