From 0406eca917270ebc4d05b20de4eecc45645d60f0 Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Mon, 15 Nov 2021 11:57:39 +0800 Subject: feat: Add skiplist cuda version Signed-off-by: KunoiSayami --- util/arena.cc | 68 +++++++++++++++++++++++++++++++++++++++++++++++++++ util/arena.cu | 11 +++++---- util/arena.cuh | 13 ++++++---- util/arena.h | 71 ++++++++++++++++++++++++++++++++++++++++++++++++++++++ util/arena_test.cc | 66 ++++++++++++++++++++++++++++++++++++++++++++++++++ util/arena_test.cu | 66 -------------------------------------------------- util/random.cu | 1 + util/random.cuh | 63 ++++++++++++++++++++++++++++++++++++++++++++++++ util/random.h | 1 + 9 files changed, 284 insertions(+), 76 deletions(-) create mode 100644 util/arena.cc create mode 100644 util/arena.h create mode 100644 util/arena_test.cc delete mode 100644 util/arena_test.cu create mode 100644 util/random.cu create mode 100644 util/random.cuh (limited to 'util') diff --git a/util/arena.cc b/util/arena.cc new file mode 100644 index 0000000..41c6a0d --- /dev/null +++ b/util/arena.cc @@ -0,0 +1,68 @@ +// 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. + +#include "util/arena.h" + +namespace leveldb { + +static const int kBlockSize = 4096; + +Arena::Arena() + : alloc_ptr_(nullptr), alloc_bytes_remaining_(0), memory_usage_(0) {} + +Arena::~Arena() { + for (size_t i = 0; i < blocks_.size(); i++) { + //cudaFree(blocks_[i]); + delete [] blocks_[i]; + } +} + +char* Arena::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. + char* result = AllocateNewBlock(bytes); + return result; + } + + // We waste the remaining space in the current block. + alloc_ptr_ = AllocateNewBlock(kBlockSize); + alloc_bytes_remaining_ = kBlockSize; + + char* result = alloc_ptr_; + alloc_ptr_ += bytes; + alloc_bytes_remaining_ -= bytes; + return result; +} + +char* Arena::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"); + size_t current_mod = reinterpret_cast(alloc_ptr_) & (align - 1); + size_t slop = (current_mod == 0 ? 0 : align - current_mod); + size_t needed = bytes + slop; + char* result; + if (needed <= alloc_bytes_remaining_) { + result = alloc_ptr_ + slop; + alloc_ptr_ += needed; + alloc_bytes_remaining_ -= needed; + } else { + // AllocateFallback always returned aligned memory + result = AllocateFallback(bytes); + } + assert((reinterpret_cast(result) & (align - 1)) == 0); + return result; +} + +char* Arena::AllocateNewBlock(size_t block_bytes) { + char* result = new char[block_bytes]; + //cudaMallocManaged((void **)&result, sizeof(char) * block_bytes); + blocks_.push_back(result); + memory_usage_.fetch_add(block_bytes + sizeof(char*), + std::memory_order_relaxed); + return result; +} + +} // namespace leveldb diff --git a/util/arena.cu b/util/arena.cu index 09ec6e6..5075318 100644 --- a/util/arena.cu +++ b/util/arena.cu @@ -13,11 +13,12 @@ Arena::Arena() Arena::~Arena() { for (size_t i = 0; i < blocks_.size(); i++) { - cudaFree(blocks_[i]); + //cudaFree(blocks_[i]); + delete [] blocks_[i]; } } -char* Arena::AllocateFallback(size_t bytes) { +__device__ char* Arena::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. @@ -35,7 +36,7 @@ char* Arena::AllocateFallback(size_t bytes) { return result; } -char* Arena::AllocateAligned(size_t bytes) { +__device__ char* Arena::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"); @@ -56,8 +57,8 @@ char* Arena::AllocateAligned(size_t bytes) { } char* Arena::AllocateNewBlock(size_t block_bytes) { - char* result = nullptr; - cudaMallocManaged((void **)&result, sizeof(char) * block_bytes); + char* result = new char[block_bytes]; + //cudaMallocManaged((void **)&result, sizeof(char) * block_bytes); blocks_.push_back(result); memory_usage_.fetch_add(block_bytes + sizeof(char*), std::memory_order_relaxed); diff --git a/util/arena.cuh b/util/arena.cuh index 68fc55d..618b426 100644 --- a/util/arena.cuh +++ b/util/arena.cuh @@ -11,6 +11,8 @@ #include #include +#include + namespace leveldb { class Arena { @@ -26,30 +28,31 @@ class Arena { char* Allocate(size_t bytes); // Allocate memory with the normal alignment guarantees provided by malloc. - char* AllocateAligned(size_t bytes); + __device__ char* AllocateAligned(size_t bytes); // Returns an estimate of the total memory usage of data allocated // by the arena. size_t MemoryUsage() const { - return memory_usage_.load(std::memory_order_relaxed); + return memory_usage_.load(cuda::memory_order_relaxed); } private: - char* AllocateFallback(size_t bytes); - char* AllocateNewBlock(size_t block_bytes); + __device__ char* AllocateFallback(size_t bytes); + __device__ char* AllocateNewBlock(size_t block_bytes); // Allocation state char* alloc_ptr_; size_t alloc_bytes_remaining_; // Array of new[] allocated memory blocks + //thrust::host_vector blocks_; std::vector blocks_; // Total memory usage of the arena. // // TODO(costan): This member is accessed via atomics, but the others are // accessed without any locking. Is this OK? - std::atomic memory_usage_; + cuda::atomic memory_usage_; }; inline char* Arena::Allocate(size_t bytes) { diff --git a/util/arena.h b/util/arena.h new file mode 100644 index 0000000..68fc55d --- /dev/null +++ b/util/arena.h @@ -0,0 +1,71 @@ +// 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_UTIL_ARENA_H_ +#define STORAGE_LEVELDB_UTIL_ARENA_H_ + +#include +#include +#include +#include +#include + +namespace leveldb { + +class Arena { + public: + Arena(); + + Arena(const Arena&) = delete; + Arena& operator=(const Arena&) = delete; + + ~Arena(); + + // Return a pointer to a newly allocated memory block of "bytes" bytes. + char* Allocate(size_t bytes); + + // Allocate memory with the normal alignment guarantees provided by malloc. + char* AllocateAligned(size_t bytes); + + // Returns an estimate of the total memory usage of data allocated + // by the arena. + size_t MemoryUsage() const { + return memory_usage_.load(std::memory_order_relaxed); + } + + private: + char* AllocateFallback(size_t bytes); + char* AllocateNewBlock(size_t block_bytes); + + // Allocation state + char* alloc_ptr_; + size_t alloc_bytes_remaining_; + + // Array of new[] allocated memory blocks + std::vector blocks_; + + // Total memory usage of the arena. + // + // TODO(costan): This member is accessed via atomics, but the others are + // accessed without any locking. Is this OK? + std::atomic memory_usage_; +}; + +inline char* Arena::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). + assert(bytes > 0); + if (bytes <= alloc_bytes_remaining_) { + char* result = alloc_ptr_; + alloc_ptr_ += bytes; + alloc_bytes_remaining_ -= bytes; + return result; + } + return AllocateFallback(bytes); +} + +} // namespace leveldb + +#endif // STORAGE_LEVELDB_UTIL_ARENA_H_ diff --git a/util/arena_test.cc b/util/arena_test.cc new file mode 100644 index 0000000..d2ead39 --- /dev/null +++ b/util/arena_test.cc @@ -0,0 +1,66 @@ +// 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. + +#include "util/arena.h" +#include "util/random.h" + +#include "gtest/gtest.h" + +namespace leveldb { + +TEST(ArenaTest, Empty) { Arena arena; } + +TEST(ArenaTest, Simple) { + std::vector> allocated; + Arena arena; + const int N = 100000; + size_t bytes = 0; + Random rnd(301); + for (int i = 0; i < N; i++) { + size_t s; + if (i % (N / 10) == 0) { + s = i; + } else { + s = rnd.OneIn(4000) + ? rnd.Uniform(6000) + : (rnd.OneIn(10) ? rnd.Uniform(100) : rnd.Uniform(20)); + } + if (s == 0) { + // Our arena disallows size 0 allocations. + s = 1; + } + char* r; + if (rnd.OneIn(10)) { + r = arena.AllocateAligned(s); + } else { + r = arena.Allocate(s); + } + + for (size_t b = 0; b < s; b++) { + // Fill the "i"th allocation with a known bit pattern + r[b] = i % 256; + } + bytes += s; + allocated.push_back(std::make_pair(s, r)); + ASSERT_GE(arena.MemoryUsage(), bytes); + if (i > N / 10) { + ASSERT_LE(arena.MemoryUsage(), bytes * 1.10); + } + } + for (size_t i = 0; i < allocated.size(); i++) { + size_t num_bytes = allocated[i].first; + const char* p = allocated[i].second; + for (size_t b = 0; b < num_bytes; b++) { + // Check the "i"th allocation for the known bit pattern + ASSERT_EQ(int(p[b]) & 0xff, i % 256); + } + } +} + +} // namespace leveldb + +int main(int argc, char** argv) { + testing::InitGoogleTest(&argc, argv); + return RUN_ALL_TESTS(); +} diff --git a/util/arena_test.cu b/util/arena_test.cu deleted file mode 100644 index 9b87785..0000000 --- a/util/arena_test.cu +++ /dev/null @@ -1,66 +0,0 @@ -// 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. - -#include "util/arena.cuh" -#include "util/random.h" - -#include "gtest/gtest.h" - -namespace leveldb { - -TEST(ArenaTest, Empty) { Arena arena; } - -TEST(ArenaTest, Simple) { - std::vector> allocated; - Arena arena; - const int N = 100000; - size_t bytes = 0; - Random rnd(301); - for (int i = 0; i < N; i++) { - size_t s; - if (i % (N / 10) == 0) { - s = i; - } else { - s = rnd.OneIn(4000) - ? rnd.Uniform(6000) - : (rnd.OneIn(10) ? rnd.Uniform(100) : rnd.Uniform(20)); - } - if (s == 0) { - // Our arena disallows size 0 allocations. - s = 1; - } - char* r; - if (rnd.OneIn(10)) { - r = arena.AllocateAligned(s); - } else { - r = arena.Allocate(s); - } - - for (size_t b = 0; b < s; b++) { - // Fill the "i"th allocation with a known bit pattern - r[b] = i % 256; - } - bytes += s; - allocated.push_back(std::make_pair(s, r)); - ASSERT_GE(arena.MemoryUsage(), bytes); - if (i > N / 10) { - ASSERT_LE(arena.MemoryUsage(), bytes * 1.10); - } - } - for (size_t i = 0; i < allocated.size(); i++) { - size_t num_bytes = allocated[i].first; - const char* p = allocated[i].second; - for (size_t b = 0; b < num_bytes; b++) { - // Check the "i"th allocation for the known bit pattern - ASSERT_EQ(int(p[b]) & 0xff, i % 256); - } - } -} - -} // namespace leveldb - -int main(int argc, char** argv) { - testing::InitGoogleTest(&argc, argv); - return RUN_ALL_TESTS(); -} diff --git a/util/random.cu b/util/random.cu new file mode 100644 index 0000000..715ab5a --- /dev/null +++ b/util/random.cu @@ -0,0 +1 @@ +#include "random.cuh" diff --git a/util/random.cuh b/util/random.cuh new file mode 100644 index 0000000..38eafe2 --- /dev/null +++ b/util/random.cuh @@ -0,0 +1,63 @@ +// 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_UTIL_RANDOM_H_ +#define STORAGE_LEVELDB_UTIL_RANDOM_H_ + +#include + +namespace leveldb { + +// A very simple random number generator. Not especially good at +// generating truly random bits, but good enough for our needs in this +// package. +class Random { + private: + uint32_t seed_; + + public: + explicit Random(uint32_t s) : seed_(s & 0x7fffffffu) { + // Avoid bad seeds. + if (seed_ == 0 || seed_ == 2147483647L) { + seed_ = 1; + } + } + __device__ uint32_t Next() { + static const uint32_t M = 2147483647L; // 2^31-1 + static const uint64_t A = 16807; // bits 14, 8, 7, 5, 2, 1, 0 + // We are computing + // seed_ = (seed_ * A) % M, where M = 2^31-1 + // + // seed_ must not be zero or M, or else all subsequent computed values + // will be zero or M respectively. For all other values, seed_ will end + // up cycling through every number in [1,M-1] + uint64_t product = seed_ * A; + + // Compute (product % M) using the fact that ((x << 31) % M) == x. + seed_ = static_cast((product >> 31) + (product & M)); + // The first reduction may overflow by 1 bit, so we may need to + // repeat. mod == M is not possible; using > allows the faster + // sign-bit-based test. + if (seed_ > M) { + seed_ -= M; + } + return seed_; + } + // Returns a uniformly distributed value in the range [0..n-1] + // REQUIRES: n > 0 + __device__ uint32_t Uniform(int n) { return Next() % n; } + + // Randomly returns true ~"1/n" of the time, and false otherwise. + // REQUIRES: n > 0 + __device__ bool OneIn(int n) { return (Next() % n) == 0; } + + // Skewed: pick "base" uniformly from range [0,max_log] and then + // return "base" random bits. The effect is to pick a number in the + // range [0,2^max_log-1] with exponential bias towards smaller numbers. + __device__ uint32_t Skewed(int max_log) { return Uniform(1 << Uniform(max_log + 1)); } +}; + +} // namespace leveldb + +#endif // STORAGE_LEVELDB_UTIL_RANDOM_H_ diff --git a/util/random.h b/util/random.h index fe76ab4..d7cbdd9 100644 --- a/util/random.h +++ b/util/random.h @@ -58,6 +58,7 @@ class Random { uint32_t Skewed(int max_log) { return Uniform(1 << Uniform(max_log + 1)); } }; + } // namespace leveldb #endif // STORAGE_LEVELDB_UTIL_RANDOM_H_ -- cgit v1.3.1