aboutsummaryrefslogtreecommitdiff
path: root/util
diff options
context:
space:
mode:
Diffstat (limited to 'util')
-rw-r--r--util/arena.cc68
-rw-r--r--util/arena.cu11
-rw-r--r--util/arena.cuh13
-rw-r--r--util/arena.h71
-rw-r--r--util/arena_test.cc (renamed from util/arena_test.cu)2
-rw-r--r--util/random.cu1
-rw-r--r--util/random.cuh63
-rw-r--r--util/random.h1
8 files changed, 219 insertions, 11 deletions
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<uintptr_t>(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<uintptr_t>(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 <cstdint>
#include <vector>
+#include <cuda/atomic>
+
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<char *> blocks_;
std::vector<char*> 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<size_t> memory_usage_;
+ cuda::atomic<size_t> 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 <atomic>
+#include <cassert>
+#include <cstddef>
+#include <cstdint>
+#include <vector>
+
+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<char*> 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<size_t> 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.cu b/util/arena_test.cc
index 9b87785..d2ead39 100644
--- a/util/arena_test.cu
+++ b/util/arena_test.cc
@@ -2,7 +2,7 @@
// 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/arena.h"
#include "util/random.h"
#include "gtest/gtest.h"
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 <cstdint>
+
+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<uint32_t>((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_
generated by cgit v1.3.1 (git 2.54.0) at 2026-09-04 07:56:54 +0000