1 files changed, 24 insertions, 15 deletions
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);