dump: 20211129

Signed-off-by: KunoiSayami <[email protected]>

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