aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2021-11-29 10:09:43 +0800
committerKunoiSayami <[email protected]>2021-11-29 10:09:43 +0800
commit8126602f094fa6644ee109d7f8b49fa0a02b5451 (patch)
treec6d67ce671bb10f89f9e4ba9681731500584f86a
parent0e18a4800b0737dc8f1d3720bc79bf000529ddef (diff)
dump: 20211129
Signed-off-by: KunoiSayami <[email protected]>
-rw-r--r--CMakeLists.txt52
-rw-r--r--db/skiplist_test.cu26
-rw-r--r--util/arena.cuh2
-rw-r--r--util/random.cuh10
4 files changed, 70 insertions, 20 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 8b6fac1..b9a4ca1 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -180,8 +180,6 @@ target_sources(leveldb
"table/two_level_iterator.h"
"util/arena.cu"
"util/arena.cuh"
- "util/arena.cc"
- "util/arena.h"
"util/bloom.cc"
"util/cache.cc"
"util/coding.cc"
@@ -300,7 +298,7 @@ add_executable(leveldbutil
"db/leveldbutil.cc"
)
target_link_libraries(leveldbutil leveldb)
-set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Ldb")
+#set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -gencode=arch=compute_70,code=[compute_70,lto_70]")
if(LEVELDB_BUILD_TESTS)
enable_testing()
@@ -355,6 +353,52 @@ if(LEVELDB_BUILD_TESTS)
add_test(NAME "${test_target_name}" COMMAND "${test_target_name}")
endfunction(leveldb_test)
+ function(leveldb_test_arena test_file)
+ get_filename_component(test_target_name "${test_file}" NAME_WE)
+
+ cuda_add_library(arena_cuda "")
+ target_sources(arena_cuda
+ PRIVATE
+ "util/arena.cu"
+ )
+ set_target_properties(arena_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+ target_compile_options(arena_cuda PRIVATE
+ $<$<COMPILE_LANGUAGE:CUDA>: -dc>)
+
+#[[ add_library(arena_cuda_link "")
+ target_sources(arena_cuda_link
+ PRIVATE
+ ${CMAKE_CURRENT_SOURCE_DIR}/libs/arena.o
+ )
+ target_compile_options(arena_cuda_link PRIVATE
+ $<$<COMPILE_LANGUAGE:CUDA>: -dlink>)]]
+
+ add_executable("${test_target_name}" "")
+ target_sources("${test_target_name}"
+ PRIVATE
+ "${PROJECT_BINARY_DIR}/${LEVELDB_PORT_CONFIG_DIR}/port_config.h"
+ "util/testutil.cc"
+ "util/testutil.h"
+
+ "${test_file}"
+ )
+ set_target_properties(arena_cuda PROPERTIES CUDA_ARCHITECTURES "75")
+ target_link_libraries("${test_target_name}" leveldb gmock gtest benchmark)
+ target_compile_definitions("${test_target_name}"
+ PRIVATE
+ ${LEVELDB_PLATFORM_NAME}=1
+ )
+ if (NOT HAVE_CXX17_HAS_INCLUDE)
+ target_compile_definitions("${test_target_name}"
+ PRIVATE
+ LEVELDB_HAS_PORT_CONFIG_H=1
+ )
+ endif(NOT HAVE_CXX17_HAS_INCLUDE)
+ set_target_properties("${test_target_name}" PROPERTIES CUDA_ARCHITECTURES "75")
+
+ add_test(NAME "${test_target_name}" COMMAND "${test_target_name}")
+ endfunction(leveldb_test_arena)
+
leveldb_test("db/c_test.c")
leveldb_test("db/fault_injection_test.cc")
@@ -374,7 +418,7 @@ if(LEVELDB_BUILD_TESTS)
leveldb_test("db/filename_test.cc")
leveldb_test("db/log_test.cc")
leveldb_test("db/recovery_test.cc")
- leveldb_test("db/skiplist_test.cu")
+ leveldb_test_arena("db/skiplist_test.cu")
leveldb_test("db/version_edit_test.cc")
leveldb_test("db/version_set_test.cc")
leveldb_test("db/write_batch_test.cc")
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index 9b7ec7e..e83af79 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -288,30 +288,36 @@ __device__ void update_list(SkipList<Key, Comparator> * l, Key key) {
}
__global__ void insert_skiplist(int random_seed) {
- Arena arena;
+ //unsigned row = threadIdx.x + blockIdx.x * blockDim.x;
+ /*for (int i = 0; i < 1000; i++ ) {
+ update_list(&skipList, device_rnd->Next());
+ }*/
+}
+
+__global__ void init(Arena ** pArena, SkipList<Key, Comparator> ** pSkipList) {
Comparator cmp;
- auto * skipList = new SkipList<Key, Comparator>(cmp, &arena);
- auto * device_rnd = new Random(random_seed);
- unsigned row = threadIdx.x + blockIdx.x * blockDim.x;
- for (int i = 0; i < 1000; i++ ) {
- update_list(skipList, device_rnd->Next());
- }
+ *pArena = new Arena();
+ *pSkipList = new SkipList<Key, Comparator>(cmp, reinterpret_cast<Arena*>(&*pArena));
}
__host__ void host_insert_test() {
- Key * keys;
+ //Key * keys;
//SkipList<Key, Comparator> list(cmp, &arena);
/*
for (int i = 0; i < 1000; i++) {
keys[i] = .Next();
}*/
+ Arena * pArena;
+ SkipList<Key, Comparator> * skipList;
+ auto * device_rnd = new Random(test::RandomSeed());
- dim3 blockSize(32, 1);
+ dim3 blockSize(1, 1);
- dim3 gridSize(32, 1);
+ dim3 gridSize(1, 1);
+ init<<<gridSize, blockSize>>>(&pArena, &skipList);
insert_skiplist<<<gridSize, blockSize>>>(test::RandomSeed());
cudaDeviceSynchronize();
diff --git a/util/arena.cuh b/util/arena.cuh
index 950f286..f9ccb1e 100644
--- a/util/arena.cuh
+++ b/util/arena.cuh
@@ -17,7 +17,7 @@ namespace leveldb {
class Arena {
public:
- __device__ Arena();
+ explicit __device__ Arena();
Arena(const Arena&) = delete;
Arena& operator=(const Arena&) = delete;
diff --git a/util/random.cuh b/util/random.cuh
index 95845f1..64eaadf 100644
--- a/util/random.cuh
+++ b/util/random.cuh
@@ -17,13 +17,13 @@ class Random {
uint32_t seed_;
public:
- __device__ explicit Random(uint32_t s) : seed_(s & 0x7fffffffu) {
+ __host__ __device__ explicit Random(uint32_t s) : seed_(s & 0x7fffffffu) {
// Avoid bad seeds.
if (seed_ == 0 || seed_ == 2147483647L) {
seed_ = 1;
}
}
- __device__ uint32_t Next() {
+ __host__ __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
@@ -46,16 +46,16 @@ class Random {
}
// Returns a uniformly distributed value in the range [0..n-1]
// REQUIRES: n > 0
- __device__ uint32_t Uniform(int n) { return Next() % n; }
+ __host__ __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; }
+ __host__ __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)); }
+ __host__ __device__ uint32_t Skewed(int max_log) { return Uniform(1 << Uniform(max_log + 1)); }
};
} // namespace leveldb