From 01769e17e318deab2b9ffc2fc49464e50f865c15 Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Fri, 5 Nov 2021 16:34:40 +0800 Subject: chore: Refactor struct Signed-off-by: KunoiSayami --- CMakeLists.txt | 4 +- arena.cu | 68 ---------------------------------- arena.cuh | 68 ---------------------------------- headers/arena.cu | 69 ++++++++++++++++++++++++++++++++++ headers/arena.cuh | 68 ++++++++++++++++++++++++++++++++++ main.cu | 1 - tests/arena.cu | 108 ++++++++++++++++++++++++++++++++++++++++++++++++++++++ 7 files changed, 248 insertions(+), 138 deletions(-) delete mode 100644 arena.cu delete mode 100644 arena.cuh create mode 100644 headers/arena.cu create mode 100644 headers/arena.cuh create mode 100644 tests/arena.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index f1c2883..d004e93 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -3,7 +3,9 @@ project(cleveldb CUDA) set(CMAKE_CUDA_STANDARD 14) -add_executable(cleveldb main.cu arena.cu arena.cuh) +add_executable(cleveldb main.cu headers/arena.cu headers/arena.cuh) + +add_executable(arena_test tests/arena.cu headers/arena.cu headers/arena.cuh) set_target_properties(cleveldb PROPERTIES CUDA_SEPARABLE_COMPILATION ON) diff --git a/arena.cu b/arena.cu deleted file mode 100644 index 1bb0f91..0000000 --- a/arena.cu +++ /dev/null @@ -1,68 +0,0 @@ -// -// Created by user on 20/10/2021. -// - -#include "arena.cuh" - -namespace cleveldb { - -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++) { - //delete[] blocks_[i]; - cudaFree(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]; - char* result = nullptr; - 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 cleveldb \ No newline at end of file diff --git a/arena.cuh b/arena.cuh deleted file mode 100644 index 20da470..0000000 --- a/arena.cuh +++ /dev/null @@ -1,68 +0,0 @@ -// -// Created by user on 20/10/2021. -// - -#ifndef CLEVELDB_ARENA_CUH -#define CLEVELDB_ARENA_CUH - - -#include -#include -#include -#include -#include - -namespace cleveldb { - - -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); - - char* alloc_ptr_; - size_t alloc_bytes_remaining_; - - std::vector blocks_; - - 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 cleveldb - -#endif //CLEVELDB_ARENA_CUH diff --git a/headers/arena.cu b/headers/arena.cu new file mode 100644 index 0000000..7f5de7a --- /dev/null +++ b/headers/arena.cu @@ -0,0 +1,69 @@ +// +// Created by user on 20/10/2021. +// + +#include "arena.cuh" + +namespace cleveldb { + +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++) { + //delete[] blocks_[i]; + cudaFree(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]; + char* result = nullptr; + 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 cleveldb \ No newline at end of file diff --git a/headers/arena.cuh b/headers/arena.cuh new file mode 100644 index 0000000..20da470 --- /dev/null +++ b/headers/arena.cuh @@ -0,0 +1,68 @@ +// +// Created by user on 20/10/2021. +// + +#ifndef CLEVELDB_ARENA_CUH +#define CLEVELDB_ARENA_CUH + + +#include +#include +#include +#include +#include + +namespace cleveldb { + + +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); + + char* alloc_ptr_; + size_t alloc_bytes_remaining_; + + std::vector blocks_; + + 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 cleveldb + +#endif //CLEVELDB_ARENA_CUH diff --git a/main.cu b/main.cu index 8fbfc65..95d3da0 100644 --- a/main.cu +++ b/main.cu @@ -1,6 +1,5 @@ #include int main() { - std::cout << "Hello, World!" << std::endl; return 0; } diff --git a/tests/arena.cu b/tests/arena.cu new file mode 100644 index 0000000..f7a4ddc --- /dev/null +++ b/tests/arena.cu @@ -0,0 +1,108 @@ +#include + +#include "../headers/arena.cuh" + + +using namespace cleveldb; + +#include + + +// 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; + } + } + 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 + uint32_t Uniform(int n) { return Next() % n; } + + // Randomly returns true ~"1/n" of the time, and false otherwise. + // REQUIRES: n > 0 + 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. + uint32_t Skewed(int max_log) { return Uniform(1 << Uniform(max_log + 1)); } +}; + + + +int main() { + 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(arena.MemoryUsage() >= bytes); + if (i > N / 10) { + assert(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((int(p[b]) & 0xff) == (i % 256)); + } + } + return 0; +} -- cgit v1.3.1