feat: Finish implement cuda memtable

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

9 files changed, 273 insertions, 135 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index aa46c4b..9ea26f4 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -181,6 +181,8 @@ target_sources(leveldb
     "table/two_level_iterator.h"
         "util/arena.cu"
         "util/arena.cuh"
+        "util/arena.h"
+        "util/arena.cc"
     "util/bloom.cc"
     "util/cache.cc"
     "util/coding.cu"
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
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_;
 };
 
diff --git a/db/skiplist.cuh b/db/skiplist.cuh
index 125a790..2d4d92b 100644
--- a/db/skiplist.cuh
+++ b/db/skiplist.cuh
@@ -74,7 +74,7 @@ class CudaSpinLock {
   }
 };
 
-class Arena;
+class ArenaCuda;
 
 template <typename Key, class Comparator>
 class SkipList {
@@ -86,7 +86,7 @@ class SkipList {
   // and will allocate memory using "*arena".  Objects allocated in the arena
   // must remain allocated for the lifetime of the skiplist object.
   //explicit SkipList(Comparator cmp, Arena* arena);
-  __device__ explicit SkipList(Comparator cmp, Arena* arena);
+  __device__ explicit SkipList(Comparator cmp, ArenaCuda* arena);
 
   SkipList(const SkipList&) = delete;
   SkipList& operator=(const SkipList&) = delete;
@@ -168,7 +168,7 @@ class SkipList {
 
   // Immutable after construction
   Comparator const compare_;
-  Arena* const arena_;  // Arena used for allocations of nodes
+  ArenaCuda* const arena_;  // Arena used for allocations of nodes
 
   Node* const head_;
 
@@ -407,7 +407,7 @@ __device__ typename SkipList<Key, Comparator>::Node* SkipList<Key, Comparator>::
 }
 
 template <typename Key, class Comparator>
-__device__ SkipList<Key, Comparator>::SkipList(Comparator cmp, Arena* arena)
+__device__ SkipList<Key, Comparator>::SkipList(Comparator cmp, ArenaCuda* arena)
     : compare_(cmp),
       arena_(arena),
       head_(NewNode(0 /* any key will do */, kMaxHeight)),
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index e5a1256..f77fc60 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -441,13 +441,13 @@ __global__ void testKeysIsEqualLists(SkipList<Key, Comparator> * skiplist, const
   assert(!iter.Valid());
 }
 
-__global__ void initSkipList(Arena ** pArena, SkipList<Key, Comparator> ** pSkipList) {
+__global__ void initSkipList(ArenaCuda ** pArena, SkipList<Key, Comparator> ** pSkipList) {
   Comparator cmp;
-  *pArena = new Arena();
+  *pArena = new ArenaCuda();
   *pSkipList = new SkipList<Key, Comparator>(cmp, *pArena);
 }
 
-__global__ void freeSkipList(Arena *** pArena, SkipList<Key, Comparator> *** pSkipList) {
+__global__ void freeSkipList(ArenaCuda *** pArena, SkipList<Key, Comparator> *** pSkipList) {
     cudaFree(**pArena);
     cudaFree(**pSkipList);
     cudaFree(*pArena);
@@ -456,7 +456,7 @@ __global__ void freeSkipList(Arena *** pArena, SkipList<Key, Comparator> *** pSk
 
 TEST(SkipTest, TestInitSkiplist) {
 
-  Arena ** pArena;
+  ArenaCuda ** pArena;
   SkipList<Key, Comparator> ** pSkipList;
 
   cudaMallocManaged((void**)&pArena, sizeof(void*));
@@ -495,7 +495,7 @@ TEST(SkipTest, TestSingleCudaInsert) {
     keys[i] = .Next();
   }*/
   Key * keys = new Key[SKIPLIST_TEST_SIZE], * sorted_keys = new Key[SKIPLIST_TEST_SIZE];
-  Arena ** pArena;
+  ArenaCuda ** pArena;
   SkipList<Key, Comparator> ** skipList;
   std::set<Key> k;
 
@@ -561,7 +561,7 @@ TEST(SkipTest, TestSingleCudaInsert) {
 
 TEST(SkipTest, TestMultiThreadInsert) {
   Key * keys = new Key[SKIPLIST_TEST_SIZE], * sorted_keys = new Key[SKIPLIST_TEST_SIZE];
-  Arena ** pArena;
+  ArenaCuda ** pArena;
   SkipList<Key, Comparator> ** pSkipList;
   std::set<Key> k;
   Key * device_keys = nullptr;
diff --git a/util/arena.cu b/util/arena.cu
index e6c8c04..e91d652 100644
--- a/util/arena.cu
+++ b/util/arena.cu
@@ -8,11 +8,11 @@ namespace leveldb {
 
 static const int kBlockSize = 4096;
 
-__device__ __host__ Arena::Arena()
+__device__ __host__ ArenaCuda::ArenaCuda()
     : alloc_ptr_(nullptr), alloc_bytes_remaining_(0), memory_usage_(0),
       head_(nullptr), blocks_(nullptr) {}
 
-__host__ __device__ Arena::~Arena() {
+__host__ __device__ ArenaCuda::~ArenaCuda() {
   ArenaNode * current = this->head_;
   while (current != nullptr) {
     ArenaNode * next = current->next;
@@ -22,7 +22,7 @@ __host__ __device__ Arena::~Arena() {
   }
 }
 
-__device__ char* Arena::AllocateFallback(size_t bytes) {
+__device__ char* ArenaCuda::AllocateFallback(size_t bytes) {
   if (bytes > kBlockSize / 4) {
     // Object is more than a quarter of our block size.  Allocate it separately
     // to avoid wasting too much space in leftover bytes.
@@ -40,7 +40,7 @@ __device__ char* Arena::AllocateFallback(size_t bytes) {
   return result;
 }
 
-__device__ char* Arena::AllocateAligned(size_t bytes) {
+__device__ char* ArenaCuda::AllocateAligned(size_t bytes) {
   const int align = (sizeof(void*) > 8) ? sizeof(void*) : 8;
   static_assert((align & (align - 1)) == 0,
                 "Pointer size should be a power of 2");
@@ -60,7 +60,7 @@ __device__ char* Arena::AllocateAligned(size_t bytes) {
   return result;
 }
 
-__device__ char* Arena::AllocateNewBlock(size_t block_bytes) {
+__device__ char* ArenaCuda::AllocateNewBlock(size_t block_bytes) {
   // Allocate new target
   char* result = nullptr;
   // Allocate storage target
diff --git a/util/arena.cuh b/util/arena.cuh
index e9abd31..a92547f 100644
--- a/util/arena.cuh
+++ b/util/arena.cuh
@@ -2,8 +2,8 @@
 // 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_UTIL_ARENA_H_
-#define STORAGE_LEVELDB_UTIL_ARENA_H_
+#ifndef STORAGE_LEVELDB_UTIL_ARENA_CUDA_H_
+#define STORAGE_LEVELDB_UTIL_ARENA_CUDA_H_
 
 #include <atomic>
 #include <cassert>
@@ -15,14 +15,14 @@
 
 namespace leveldb {
 
-class Arena {
+class ArenaCuda {
  public:
-  explicit __device__ __host__ Arena();
+  explicit __device__ __host__ ArenaCuda();
 
-  Arena(const Arena&) = delete;
-  Arena& operator=(const Arena&) = delete;
+  ArenaCuda(const ArenaCuda&) = delete;
+  ArenaCuda& operator=(const ArenaCuda&) = delete;
 
-  __host__ __device__ ~Arena();
+  __host__ __device__ ~ArenaCuda();
 
   // Return a pointer to a newly allocated memory block of "bytes" bytes.
   __device__ char* Allocate(size_t bytes);
@@ -64,7 +64,7 @@ class Arena {
   cuda::atomic<size_t> memory_usage_;
 };
 
-__device__ inline char* Arena::Allocate(size_t bytes) {
+__device__ inline char* ArenaCuda::Allocate(size_t bytes) {
   // The semantics of what to return are a bit messy if we allow
   // 0-byte allocations, so we disallow them here (we don't need
   // them for our internal use).
@@ -80,4 +80,4 @@ __device__ inline char* Arena::Allocate(size_t bytes) {
 
 }  // namespace leveldb
 
-#endif  // STORAGE_LEVELDB_UTIL_ARENA_H_
+#endif  // STORAGE_LEVELDB_UTIL_ARENA_CUDA_H_
diff --git a/util/coding.cu b/util/coding.cu
index 63ddf28..ee93ddb 100644
--- a/util/coding.cu
+++ b/util/coding.cu
@@ -75,7 +75,7 @@ void PutLengthPrefixedSlice(std::string* dst, const Slice& value) {
   dst->append(value.data(), value.size());
 }
 
-int VarintLength(uint64_t v) {
+__device__ __host__ int VarintLength(uint64_t v) {
   int len = 1;
   while (v >= 128) {
     v >>= 7;
diff --git a/util/coding.cuh b/util/coding.cuh
index f89177b..dbe5564 100644
--- a/util/coding.cuh
+++ b/util/coding.cuh
@@ -40,7 +40,7 @@ const char* GetVarint32Ptr(const char* p, const char* limit, uint32_t* v);
 const char* GetVarint64Ptr(const char* p, const char* limit, uint64_t* v);
 
 // Returns the length of the varint32 or varint64 encoding of "v"
-int VarintLength(uint64_t v);
+__device__ __host__ int VarintLength(uint64_t v);
 
 // Lower-level versions of Put... that write directly into a character buffer
 // and return a pointer just past the last byte written.