diff options
Diffstat (limited to 'db/memtable.cu')
| -rw-r--r-- | db/memtable.cu | 271 |
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 |