feat: Finish implement cuda memtable

Signed-off-by: KunoiSayami <[email protected]>

1 files changed, 169 insertions, 102 deletions

diff --git a/db/memtable.cu b/db/memtable.cu
index c4c3a75..8b1fab5 100644
--- a/db/memtable.cu
+++ b/db/memtable.cu
@@ -18,8 +18,10 @@ static Slice GetLengthPrefixedSlice(const char* data) {
   return Slice(p, len);
 }
 
-__device__ MemTable::MemTable(const InternalKeyComparator& comparator)
-    : comparator_(comparator), refs_(0), table_(comparator_, &arena_) {}
+MemTable::MemTable(const InternalKeyComparator& comparator)
+    : comparator_(comparator), refs_(0) {
+
+}
 
 MemTable::~MemTable() { assert(refs_ == 0); }
 
@@ -73,131 +75,196 @@ class MemTableIterator : public Iterator {
 
 Iterator* MemTable::NewIterator() { return new MemTableIterator(&table_); }
 
-__device__ char* EncodeVarint32Device(char* dst, uint32_t v) {
-  // Operate on characters as unsigneds
-  uint8_t* ptr = reinterpret_cast<uint8_t*>(dst);
-  static const int B = 128;
-  if (v < (1 << 7)) {
-    *(ptr++) = v;
-  } else if (v < (1 << 14)) {
-    *(ptr++) = v | B;
-    *(ptr++) = v >> 7;
-  } else if (v < (1 << 21)) {
-    *(ptr++) = v | B;
-    *(ptr++) = (v >> 7) | B;
-    *(ptr++) = v >> 14;
-  } else if (v < (1 << 28)) {
-    *(ptr++) = v | B;
-    *(ptr++) = (v >> 7) | B;
-    *(ptr++) = (v >> 14) | B;
-    *(ptr++) = v >> 21;
-  } else {
-    *(ptr++) = v | B;
-    *(ptr++) = (v >> 7) | B;
-    *(ptr++) = (v >> 14) | B;
-    *(ptr++) = (v >> 21) | B;
-    *(ptr++) = v >> 28;
-  }
-  return reinterpret_cast<char*>(ptr);
-}
-
-__device__ inline void EncodeFixed64Device(char* dst, uint64_t value) {
-  uint8_t* const buffer = reinterpret_cast<uint8_t*>(dst);
 
-  // Recent clang and gcc optimize this to a single mov / str instruction.
-  buffer[0] = static_cast<uint8_t>(value);
-  buffer[1] = static_cast<uint8_t>(value >> 8);
-  buffer[2] = static_cast<uint8_t>(value >> 16);
-  buffer[3] = static_cast<uint8_t>(value >> 24);
-  buffer[4] = static_cast<uint8_t>(value >> 32);
-  buffer[5] = static_cast<uint8_t>(value >> 40);
-  buffer[6] = static_cast<uint8_t>(value >> 48);
-  buffer[7] = static_cast<uint8_t>(value >> 56);
-}
-
-__global__ void Add_(MemTable * mtb, SequenceNumber s, ValueType type,
-                     size_t encoded_len, size_t internal_key_size,
-                     char * key_data, size_t key_size, char * value_data, size_t val_size) {
+__global__ void Add_(MemTable * mtb, size_t encoded_len,char * encode_data) {
   char* buf = mtb->arena_.Allocate(encoded_len);
-  char* p = EncodeVarint32Device(buf, internal_key_size);
-  memcpy(p, key_data, key_size);
-  //std::memcpy(p, key_data, key_size);
-  p += key_size;
-  EncodeFixed64Device(p, (s << 8) | type);
-  p += 8;
-  p = EncodeVarint32Device(p, val_size);
-  //std::memcpy(p, value.data(), val_size);
-  memcpy(p, value_data, val_size);
-  assert(p + val_size == buf + encoded_len);
+  memcpy(buf, encode_data, encoded_len);
   mtb->table_.Insert(buf);
 }
 
+
 void MemTable::Add(SequenceNumber s, ValueType type, const Slice& key,
                    const Slice& value) {
   // Format of an entry is concatenation of:
   //  key_size     : varint32 of internal_key.size()
-  //  key bytes    : char[internal_key.size()]
+  //  key point    : point to host memory
   //  tag          : uint64((sequence << 8) | type)
   //  value_size   : varint32 of value.size()
-  //  value bytes  : char[value.size()]
+  //  value point  : point to host memory
   size_t key_size = key.size();
   size_t val_size = value.size();
   size_t internal_key_size = key_size + 8;
-  const size_t encoded_len = VarintLength(internal_key_size) +
-                             internal_key_size + VarintLength(val_size) +
-                             val_size;
-  char * key_mem = nullptr;
-  cudaMalloc((void**)&key_mem, key_size);
-  cudaMemcpy(key_mem, key.data(), key_size, cudaMemcpyHostToDevice);
-  char * value_mem = nullptr;
-  cudaMalloc((void**)&value_mem, val_size);
-  cudaMemcpy(value_mem, value.data(), val_size, cudaMemcpyHostToDevice);
+  //const size_t encoded_len = VarintLength(internal_key_size) +
+  //                           internal_key_size + VarintLength(val_size) +
+  //                           val_size;
+  const size_t encoded_len = VarintLength(internal_key_size) + 8 +
+                             VarintLength(val_size) + 8;
+
+
+    char * key_mem = this->arena_.Allocate( key_size);
+    std::memcpy(key_mem, key.data(), key_size);
 
-  Add_<<<1, 1>>>(this, s, type, encoded_len, internal_key_size, key_mem, key_size, value_mem, val_size);
+    //char * tag_mem = key_mem + key_size;
+    //EncodeFixed64(tag_mem, (s << 8) | type);
+
+    char * val_mem = this->arena_.Allocate(val_size);
+    std::memcpy(val_mem, value.data(), val_size);
+
+  char * insert_val = new char[encoded_len], *cuda_insert = nullptr;
+  //  EncodeVarint32(insert_val, encoded_len);
+  char * p = EncodeVarint32(insert_val + 8, internal_key_size);
+  EncodeFixed64(p, reinterpret_cast<uint64_t>(key_mem));
+  p += 8;
+  EncodeFixed64(p, (s << 8) | type );
+  p += 8;
+  //EncodeFixed64(p, reinterpret_cast<uint64_t>(val_size));
+  p = EncodeVarint32(p, val_size);
+  EncodeFixed64(p, reinterpret_cast<uint64_t>(val_mem));
+
+  assert(p + 8 == insert_val + encoded_len);
+
+  cudaMalloc((void**)&cuda_insert, encoded_len);
+  cudaMemcpy(cuda_insert, insert_val, encoded_len, cudaMemcpyHostToDevice);
+
+  Add_<<<1, 1>>>(this, encoded_len, cuda_insert);
   cudaDeviceSynchronize();
 
-  cudaFree(key_mem);
-  cudaFree(value_mem);
+  cudaFree(cuda_insert);
+  delete[] insert_val;
 }
 
-__global__ void Get_(MemTable * met, bool * ret) {
+__global__ void Get_(MemTable * met, char * memkey, char ** data, size_t* malloc_size) {
+  *data = nullptr;
+  auto iter = met->getIter();
+  iter.Seek(memkey);
+  if (iter.Valid()) {
+    // entry format is:
+    //    klength  varint32
+    //    userkey  char[klength]
+    //    tag      uint64
+    //    vlength  varint32
+    //    value    char[vlength]
+    // Check that it belongs to same user key.  We do not check the
+    // sequence number since the Seek() call above should have skipped
+    // all entries with overly large sequence numbers.
+    const char *entry = iter.key();
+    const char * p = GetVarint32PtrCuda(entry, entry + 5,
+                       reinterpret_cast<uint32_t*>(malloc_size));
+    p = GetVarint32PtrCuda(p + 8, p + 13, nullptr);
+    *malloc_size = p - entry;
+    cudaMalloc((void**)*data, *malloc_size);
+    memcpy(*data, entry, *malloc_size);
+  }
+}
 
-    Slice memkey = key.memtable_key();
-    met->Table::Iterator iter(&met->table_);
-    iter.Seek(memkey.data());
-    if (iter.Valid()) {
-      // entry format is:
-      //    klength  varint32
-      //    userkey  char[klength]
-      //    tag      uint64
-      //    vlength  varint32
-      //    value    char[vlength]
-      // Check that it belongs to same user key.  We do not check the
-      // sequence number since the Seek() call above should have skipped
-      // all entries with overly large sequence numbers.
-      const char* entry = iter.key();
-      uint32_t key_length;
-      const char* key_ptr = GetVarint32Ptr(entry, entry + 5, &key_length);
-      if (met->comparator_.comparator.user_comparator()->Compare(
-              Slice(key_ptr, key_length - 8), key.user_key()) == 0) {
-        // Correct user key
-        const uint64_t tag = DecodeFixed64(key_ptr + key_length - 8);
-        switch (static_cast<ValueType>(tag & 0xff)) {
-          case kTypeValue: {
-            Slice v = GetLengthPrefixedSlice(key_ptr + key_length);
-            value->assign(v.data(), v.size());
-            return true;
-          }
-          case kTypeDeletion:
-            *s = Status::NotFound(Slice());
-            return true;
-        }
+/*
+__global__ void Get_Phase2() {
+    // Correct user key
+    const uint64_t tag = DecodeFixed64Cuda(key_ptr + key_length - 8);
+    switch (static_cast<ValueType>(tag & 0xff)) {
+      case kTypeValue: {
+        SizedString v = GetLengthPrefixedSliceCuda(key_ptr + key_length);
+        *value=v;
+        *ret = true;
+        break;
       }
+      case kTypeDeletion:
+        //*s = Status::NotFound(Slice());
+        *is_not_found = true;
+        *ret = true;
+        break;
     }
 }
+*/
 
 bool MemTable::Get(const LookupKey& key, std::string* value, Status* s) {
-  return false;
+
+  Slice memkey = key.memtable_key();
+  char * cuda_mem_key = nullptr;
+  cudaMalloc((void**)&cuda_mem_key, memkey.size());
+  cudaMemcpy(cuda_mem_key, memkey.data(), memkey.size(), cudaMemcpyHostToDevice);
+
+  char ** cuda_skiplist_key = nullptr;
+  cudaMalloc((void**)&cuda_skiplist_key, sizeof(char *));
+
+  size_t * cuda_malloc_size = nullptr;
+  cudaMalloc((void**)&cuda_malloc_size, sizeof(cuda_malloc_size));
+
+  Get_<<<1,1>>>(this, cuda_mem_key, cuda_skiplist_key, cuda_malloc_size);
+  cudaDeviceSynchronize();
+
+  if (*cuda_skiplist_key == nullptr) {
+    cudaFree(cuda_mem_key);
+    cudaFree(cuda_skiplist_key);
+    cudaFree(cuda_malloc_size);
+    return false;
+  }
+  auto * malloc_size = new size_t;
+  cudaMemcpy(malloc_size, cuda_malloc_size, sizeof(size_t), cudaMemcpyDeviceToHost);
+
+  char * entry = this->host_arena_.Allocate(*malloc_size);
+  cudaMemcpy(entry, cuda_skiplist_key, *malloc_size, cudaMemcpyDeviceToHost);
+
+  cudaFree(cuda_mem_key);
+  cudaFree(cuda_skiplist_key);
+  cudaFree(cuda_malloc_size);
+  delete malloc_size;
+
+    uint32_t key_length;
+    const char* key_ptr_ptr = GetVarint32Ptr(entry, entry + 5, &key_length);
+    const char * key_ptr =
+        reinterpret_cast<const char*>(DecodeFixed64(key_ptr_ptr));
+    if (comparator_.comparator.user_comparator()->Compare(
+            Slice(key_ptr, key_length - 8), key.user_key()) == 0) {
+      // Correct user key
+      const uint64_t tag = DecodeFixed64(key_ptr_ptr + key_length - 8);
+      switch (static_cast<ValueType>(tag & 0xff)) {
+        case kTypeValue: {
+          // HOW CAN WE GET IT?
+          Slice v = GetLengthPrefixedSlice(key_ptr_ptr + key_length);
+          value->assign(v.data(), v.size());
+          return true;
+        }
+        case kTypeDeletion:
+          *s = Status::NotFound(Slice());
+          return true;
+      }
+    }
+
+
+/*  bool * cuda_return_value = nullptr, *return_value = new bool, *is_not_found = nullptr;
+  SizedString * return_string_value = nullptr;
+  cudaMalloc((void**)&cuda_return_value, sizeof(bool));
+  cudaMalloc((void**)&is_not_found, sizeof(bool));
+  cudaMallocManaged((void**)&return_string_value, sizeof(SizedString));
+
+  Get_<<<1, 1>>>(this, cuda_return_value, cuda_mem_key, nullptr, is_not_found);
+  cudaDeviceSynchronize();
+
+  cudaMemcpy(return_value, cuda_return_value, sizeof(bool), cudaMemcpyDeviceToHost);
+
+  if (*return_value) {
+    if (*is_not_found) {
+      *s = Status::NotFound(Slice());
+    } else {
+      auto * local_sized_string = new SizedString();
+      cudaMemcpy(local_sized_string, return_string_value, sizeof(SizedString), cudaMemcpyDeviceToHost);
+
+      char * local_string = new char[local_sized_string->length];
+      cudaMemcpy(local_string, local_sized_string->data, local_sized_string->length, cudaMemcpyDeviceToHost);
+      value->assign(local_string, local_sized_string->length);
+      delete local_sized_string;
+      delete [] local_string;
+    }
+  }
+
+  bool rvalue = *return_value;
+  delete return_value;
+  cudaFree(is_not_found);
+  cudaFree(cuda_return_value);
+  cudaFree(cuda_mem_key);
+  cudaFree(return_string_value);*/
+    return false;
 }
 
 }  // namespace leveldb