From 5e5d17b0f5d97be4c11a96a2ad2c8e10a6127832 Mon Sep 17 00:00:00 2001
From: KunoiSayami <i@leanhe.dev>
Date: Sun, 7 Aug 2022 03:06:29 +0800
Subject: feat: Implement expt_0804

Signed-off-by: KunoiSayami <i@leanhe.dev>
---
 expt_0802.cpp | 15 ---------------
 expt_0804.cu  | 39 ++++++++++++++++++++++++---------------
 2 files changed, 24 insertions(+), 30 deletions(-)

diff --git a/expt_0802.cpp b/expt_0802.cpp
index fbd6c23..b1a906a 100644
--- a/expt_0802.cpp
+++ b/expt_0802.cpp
@@ -25,8 +25,6 @@ size_t calculate_location(size_t index) {
 }
 typedef int key_type;
 
-inline double safeStep(double step) { return step < 0 ? 0 : step; }
-
 int get_son_from_step(int point, bool negative) {
   return negative ? point * 2 : point * 2 + 1;
 }
@@ -34,31 +32,18 @@ int get_son_from_step(int point, bool negative) {
 const key_type *cudaBinarySearch(key_type *const start, const key_type *end,
                                  const key_type val) {
   int step_limit = (int)std::log2(LENGTH) + 1;
-  const auto length = (end - start);
   key_type *last_known_point = start;
   auto son = 0;
 
   for (int i = 0; i < step_limit; i++) {
     const auto next_level_start = start + (1 << (i + 1)) - 1;
-    /*const auto current_time = (1.0 / (1 << (i + 1)));
-    printf("%d ", *level_start);
-    printf("%d %d\n", *last_known_point, (1 << i));
     if (*last_known_point == val) {
-      puts("find");
-      return last_known_point;
-    }
-    sum = safeStep(sum +
-                   ((*last_known_point > val) ? -current_time : current_time));
-    last_known_point = level_start + (int)(sum * (double)(1 << i));*/
-    if (*last_known_point == val) {
-      puts("found");
       return last_known_point;
     }
     son = get_son_from_step(son, (*last_known_point > val));
     last_known_point = next_level_start + son;
     printf("%d\n", *last_known_point);
   }
-  puts("");
   return last_known_point;
 }
 
diff --git a/expt_0804.cu b/expt_0804.cu
index 47c269e..3a599ef 100644
--- a/expt_0804.cu
+++ b/expt_0804.cu
@@ -36,21 +36,23 @@ __device__ unsigned insert_value;
 __device__ unsigned int index_max, index_min;
 #endif
 
-__device__ const key_type *cudaBinarySearch(key_type *start, key_type *end,
+inline __device__ int get_son_from_step(int point, bool negative) {
+  return negative ? point * 2 : point * 2 + 1;
+}
+
+__device__ const key_type *cudaBinarySearch(key_type *start,
                                             const key_type val) {
-  auto begin = start;
-  key_type *last_known_point = nullptr;
-  while (begin < end) {
-    auto mid = (end - begin) / 2;
-    auto mid_val = *(start + mid);
-    if (val == mid_val) {
-      return start + mid;
-    } else if (val > mid_val) {
-      begin = begin + mid + 1;
-    } else {
-      end = end - mid - 1;
+  int step_limit = (int)log2f(SAMPLE_LENGTH);
+  key_type *last_known_point = start;
+  auto son = 0;
+
+  for (int i = 0; i < step_limit; i++) {
+    const auto next_level_start = start + (1 << (i + 1)) - 1;
+    if (*last_known_point == val) {
+      return last_known_point;
     }
-    last_known_point = begin;
+    son = get_son_from_step(son, (*last_known_point > val));
+    last_known_point = next_level_start + son;
   }
   return last_known_point;
 }
@@ -63,7 +65,7 @@ __device__ __host__ size_t calculate_location(size_t index) {
 }
 
 __device__ double sample_cdf(double x) {
-  auto it = cudaBinarySearch(SampleItem, SampleItem + SAMPLE_LENGTH + 2, x);
+  auto it = cudaBinarySearch(SampleItem, x);
   if (it == SampleItem + SAMPLE_LENGTH) {
     return 1;
   }
@@ -91,6 +93,8 @@ __global__ void initStorage(unsigned scale, unsigned test_size) {
 __global__ void initCuda(key_type *sample_item, key_type *population_item) {
   SampleItem = sample_item;
   PopulationItem = population_item;
+  /*for (int i = 0; i < SAMPLE_LENGTH; i++) {
+  }*/
   // cudaMalloc(&SampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long long));
   //  cudaMalloc(&Storage, TEST_SIZE * sizeof(key_type));
   cdf_result = nullptr;
@@ -169,10 +173,15 @@ int main(int argc, char const *argv[]) {
 
   std::sort(sample_vector.begin(), sample_vector.end());
 
+  sample_vector_into_cuda.resize(SAMPLE_LENGTH);
+  for (int i = 0; i < SAMPLE_LENGTH; i++) {
+    sample_vector_into_cuda[i] = sample_vector[calculate_location(i) - 1];
+  }
+
   key_type *cudaSample = nullptr, *cudaPopulation = nullptr;
   cudaMalloc(&cudaSample, sizeof(key_type) * (SAMPLE_LENGTH));
   cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE);
-  cudaMemcpy(cudaSample, &sample_vector[0],
+  cudaMemcpy(cudaSample, &sample_vector_into_cuda[0],
              sizeof(key_type) * sample_vector.size(), cudaMemcpyHostToDevice);
   cudaMemcpy(cudaPopulation, &population_vector[SAMPLE_LENGTH],
              sizeof(key_type) * TEST_SIZE, cudaMemcpyHostToDevice);
-- 
cgit v1.3.1