feat: Optimize binary search

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

3 files changed, 7 insertions, 10 deletions

diff --git a/expt_0729.cu b/expt_0729.cu
index 31aac7f..2d85a49 100644
--- a/expt_0729.cu
+++ b/expt_0729.cu
@@ -174,7 +174,7 @@ int main(int argc, char const *argv[]) {
   init<<<1, 1>>>(cudaSample, cudaPopulation);
   cudaDeviceSynchronize();
 
-  dim3 grid_dim = 2, block_dim = 512;
+  dim3 grid_dim = 64, block_dim = 16;
 
   unsigned long step = test_size / (grid_dim.x * block_dim.x);
   printf("step: %lu\n", step);
diff --git a/expt_0802.cpp b/expt_0802.cpp
index b1a906a..7d520e2 100644
--- a/expt_0802.cpp
+++ b/expt_0802.cpp
@@ -25,10 +25,6 @@ size_t calculate_location(size_t index) {
 }
 typedef int key_type;
 
-int get_son_from_step(int point, bool negative) {
-  return negative ? point * 2 : point * 2 + 1;
-}
-
 const key_type *cudaBinarySearch(key_type *const start, const key_type *end,
                                  const key_type val) {
   int step_limit = (int)std::log2(LENGTH) + 1;
@@ -40,7 +36,8 @@ const key_type *cudaBinarySearch(key_type *const start, const key_type *end,
     if (*last_known_point == val) {
       return last_known_point;
     }
-    son = get_son_from_step(son, (*last_known_point > val));
+    // son = get_son_from_step(son, (*last_known_point > val));
+    son = son * 2 + -((*last_known_point - val) >> 31);
     last_known_point = next_level_start + son;
     printf("%d\n", *last_known_point);
   }
@@ -73,6 +70,6 @@ int main() {
     assert(calculate_location(i) == result[i]);
   }
   puts("");
-  cudaBinarySearch(&result[0], &*result.end(), 1024);
+  cudaBinarySearch(&result[0], &*result.end(), 928);
   return 0;
 }
diff --git a/expt_0804.cu b/expt_0804.cu
index 3a599ef..73d2c82 100644
--- a/expt_0804.cu
+++ b/expt_0804.cu
@@ -51,7 +51,7 @@ __device__ const key_type *cudaBinarySearch(key_type *start,
     if (*last_known_point == val) {
       return last_known_point;
     }
-    son = get_son_from_step(son, (*last_known_point > val));
+    son = son * 2 + -((*last_known_point - val) >> 63);
     last_known_point = next_level_start + son;
   }
   return last_known_point;
@@ -191,10 +191,10 @@ int main(int argc, char const *argv[]) {
   initCuda<<<1, 1>>>(cudaSample, cudaPopulation);
   cudaDeviceSynchronize();
 
-  dim3 grid_dim = 2, block_dim = 512;
+  dim3 grid_dim = 64, block_dim = 16;
 
   unsigned long step = test_size / (grid_dim.x * block_dim.x);
-  printf("step: %lu\n", step);
+  printf("(new) step: %lu\n", step);
   assert(!(test_size % (grid_dim.x * block_dim.x)));
 
   for (int scale = 2; scale <= 8; scale++) {