feat: Add vaild_sort.cpp

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

5 files changed, 369 insertions, 6 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 7bc1eac..9d72a91 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -100,10 +100,22 @@ target_link_libraries(expt_0802 m stdc++)
 
 
 add_executable(expt_0804 expt_0804.cu)
-#set_target_properties(expt_0804 PROPERTIES LINKER_LANGUAGE CXX)
 target_link_libraries(expt_0804 m stdc++)
 
 set_target_properties(expt_0804 PROPERTIES
         CUDA_SEPARABLE_COMPILATION ON)
 set_target_properties(expt_0804 PROPERTIES CUDA_ARCHITECTURES "75")
 set_target_properties(expt_0804 PROPERTIES LINKER_LANGUAGE CUDA)
+
+
+add_executable(expt_0809 expt_0809.cu)
+target_link_libraries(expt_0809 m stdc++)
+
+set_target_properties(expt_0809 PROPERTIES
+        CUDA_SEPARABLE_COMPILATION ON)
+set_target_properties(expt_0809 PROPERTIES CUDA_ARCHITECTURES "75")
+set_target_properties(expt_0809 PROPERTIES LINKER_LANGUAGE CUDA)
+
+
+add_executable(valid_sort valid_sort_new.cpp)
+target_link_libraries(valid_sort m stdc++)
\ No newline at end of file
diff --git a/expt_0729.cu b/expt_0729.cu
index 2d85a49..8ba6aae 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 = 64, block_dim = 16;
+  dim3 grid_dim = 16, block_dim = 64;
 
   unsigned long step = test_size / (grid_dim.x * block_dim.x);
   printf("step: %lu\n", step);
diff --git a/expt_0804.cu b/expt_0804.cu
index 73d2c82..efb32f5 100644
--- a/expt_0804.cu
+++ b/expt_0804.cu
@@ -36,10 +36,6 @@ __device__ unsigned insert_value;
 __device__ unsigned int index_max, index_min;
 #endif
 
-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) {
   int step_limit = (int)log2f(SAMPLE_LENGTH);
diff --git a/expt_0809.cu b/expt_0809.cu
new file mode 100644
index 0000000..adc9d36
--- /dev/null
+++ b/expt_0809.cu
@@ -0,0 +1,206 @@
+#include <algorithm>
+#include <cassert>
+#include <cstdio>
+#include <cstring>
+#include <iostream>
+#include <type_traits>
+#include <vector>
+
+constexpr size_t length = 2097152;
+std::vector<unsigned long long> population_vector, sample_vector;
+
+unsigned long long max_value = 0, min_value = 0xfffffffff;
+
+inline void store_into_vector(unsigned long long value) {
+  if (max_value < value) {
+    max_value = value;
+  }
+  if (min_value > value) {
+    min_value = value;
+  }
+  population_vector.push_back(value);
+}
+
+typedef unsigned long long key_type;
+
+typedef unsigned long long *key_type_ptr;
+
+constexpr long SAMPLE_LENGTH = 1024;
+
+constexpr long TEST_SIZE = 1024;
+
+__device__ key_type *SampleItem, *PopulationItem;
+__device__ bool *cdf_result;
+__device__ unsigned insert_value;
+#ifdef TEST_BOUNDS
+__device__ unsigned int index_max, index_min;
+#endif
+
+__device__ const key_type *cudaBinarySearch(key_type *start, key_type *end,
+                                            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;
+    }
+    begin += ((mid_val - val) >> 63) & (mid + 1);
+    end -= (~((mid_val - val) >> 63)) & (mid + 1);
+    last_known_point = begin;
+  }
+  return last_known_point;
+}
+
+__device__ double sample_cdf(double x) {
+  auto it = cudaBinarySearch(SampleItem, SampleItem + SAMPLE_LENGTH + 2, x);
+  if (it == SampleItem + SAMPLE_LENGTH) {
+    return 1;
+  }
+  if (it == SampleItem) {
+    return 0;
+  }
+  auto it_prev = it - 1;
+  return (double(it_prev - SampleItem) +
+          (x - (double)*it_prev) / (double)(*it - *it_prev)) /
+         double(SAMPLE_LENGTH - 1);
+}
+
+__global__ void initStorage(unsigned scale, unsigned test_size) {
+  cudaFree(cdf_result);
+  cudaMalloc(&cdf_result, scale * test_size * sizeof(bool));
+  memset(cdf_result, 0, scale * test_size * sizeof(bool));
+  insert_value = 0;
+  // printf("init storage\n");
+#ifdef TEST_BOUNDS
+  index_max = 0;
+  index_min = 0x7fffffff;
+#endif
+}
+
+__global__ void init(key_type *sample_item, key_type *population_item) {
+  SampleItem = sample_item;
+  PopulationItem = population_item;
+  // cudaMalloc(&SampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long long));
+  //  cudaMalloc(&Storage, TEST_SIZE * sizeof(key_type));
+  cdf_result = nullptr;
+}
+
+__global__ void kernel(unsigned long step, const double slice_size,
+                       const long split_size) {
+  for (int i = 0; i < step; i++) {
+    auto tid = step * gridDim.x * blockDim.x + blockIdx.x * blockDim.x +
+               threadIdx.x + i;
+
+    // printf("%d\n", tid);
+    auto index = (int)(sample_cdf(PopulationItem[tid]) / slice_size);
+    if (cdf_result[index]) {
+      while (cdf_result[++index]) {
+        assert(index < split_size);
+      }
+    }
+    cdf_result[index] = true;
+    // atomicAdd(&insert_value, 1);
+
+#ifdef TEST_BOUNDS
+    while (true) {
+      unsigned tmp = index_min;
+      if (tmp < tid) {
+        break;
+      }
+      if (atomicCAS(&index_min, tmp, tid) == tmp) {
+        break;
+      }
+    }
+    while (true) {
+      unsigned tmp = index_max;
+      if (tmp > tid) {
+        break;
+      }
+      if (atomicCAS(&index_max, tmp, tid) == tmp) {
+        break;
+      }
+    }
+#endif
+  }
+}
+
+__global__ void print_function() {
+  // printf("%u\n", insert_value);
+#ifdef TEST_BOUNDS
+  printf("%u %u\n", index_min, index_max);
+#endif
+}
+
+int main(int argc, char const *argv[]) {
+  size_t test_size;
+  if (argc == 1) {
+    test_size = 1048576;
+  } else {
+    try {
+      test_size = std::stol(argv[1]);
+    } catch (...) {
+      return 1;
+    }
+  }
+
+  FILE *file = fopen("normal_distribution.txt", "r");
+  assert(file);
+  for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i))
+    ;
+  fclose(file);
+
+  assert(population_vector.size() == length);
+
+  sample_vector = std::vector<unsigned long long>(
+      population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH);
+  sample_vector.push_back(min_value);
+  sample_vector.push_back(max_value);
+
+  std::sort(sample_vector.begin(), sample_vector.end());
+
+  key_type *cudaSample = nullptr, *cudaPopulation = nullptr;
+  cudaMalloc(&cudaSample, sizeof(key_type) * (SAMPLE_LENGTH + 2));
+  cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE);
+  cudaMemcpy(cudaSample, &sample_vector[0],
+             sizeof(key_type) * sample_vector.size(), cudaMemcpyHostToDevice);
+  cudaMemcpy(cudaPopulation, &population_vector[SAMPLE_LENGTH],
+             sizeof(key_type) * TEST_SIZE, cudaMemcpyHostToDevice);
+  // memcpy(sample_heap, sample_vector.cbegin(),sizeof(key_type) *(SAMPLE_LENGTH
+  // + 2));
+
+  init<<<1, 1>>>(cudaSample, cudaPopulation);
+  cudaDeviceSynchronize();
+
+  dim3 grid_dim = 16, block_dim = 64;
+
+  unsigned long step = test_size / (grid_dim.x * block_dim.x);
+  // printf("(0809)step: %lu\n", step);
+  assert(!(test_size % (grid_dim.x * block_dim.x)));
+
+  for (int scale = 2; scale <= 8; scale++) {
+    const long split_size = test_size * scale;
+    const auto slice_size = 1.0 / (double)split_size;
+    printf("scale: %i ", scale);
+    initStorage<<<1, 1>>>(scale, test_size);
+    cudaDeviceSynchronize();
+    cudaEvent_t start, stop;
+    cudaEventCreate(&start);
+    cudaEventCreate(&stop);
+    cudaEventRecord(start, nullptr);
+    kernel<<<grid_dim, block_dim>>>(step, slice_size, split_size);
+    cudaDeviceSynchronize();
+    cudaEventRecord(stop, nullptr);
+    cudaEventSynchronize(stop);
+    float time;
+    cudaEventElapsedTime(&time, start, stop);
+    cudaEventDestroy(start);
+    cudaEventDestroy(stop);
+
+    printf("time: %lf\n", time);
+    print_function<<<1, 1>>>();
+    cudaDeviceSynchronize();
+  }
+  return 0;
+}
\ No newline at end of file
diff --git a/valid_sort_new.cpp b/valid_sort_new.cpp
new file mode 100644
index 0000000..91bdb6e
--- /dev/null
+++ b/valid_sort_new.cpp
@@ -0,0 +1,149 @@
+#include <algorithm>
+#include <cassert>
+#include <cmath>
+#include <iostream>
+#include <vector>
+
+typedef unsigned long long key_type;
+std::vector<unsigned long long> population_vector, sample_vector;
+std::vector<bool> result_storage;
+constexpr size_t SAMPLE_LENGTH = 1024, TEST_LENGTH = 1024;
+
+constexpr int MOVE_OFFSET = sizeof(key_type) * 8 - 1;
+
+size_t fast_log(size_t a) {
+  float t = a;
+  return (((*(int *)&t) >> 23) + 1) & 127;
+}
+
+size_t calculate_location(size_t index) {
+  size_t bit_low = SAMPLE_LENGTH >> fast_log(++index) >> 1;
+  return (!bit_low) ? SAMPLE_LENGTH
+                    : ((index << 1) | 1) * bit_low - SAMPLE_LENGTH;
+}
+
+size_t calculate_location_inverse(size_t index) {
+  ++index;
+  size_t low_bit = index & (-index);
+  return (low_bit == SAMPLE_LENGTH) ? SAMPLE_LENGTH
+                                    : ((SAMPLE_LENGTH + index) / low_bit) >> 1;
+}
+
+/*size_t calculate_location(size_t index) {
+  int step = (int)std::log2(++index) + 1;
+  auto location =
+      SAMPLE_LENGTH / (1 << step) * ((index & ((1 << (step - 1)) - 1)) * 2 + 1);
+  return location ? location : SAMPLE_LENGTH;
+}*/
+
+const key_type *cudaBinarySearch(key_type *const start, const key_type val) {
+  int step_limit = (int)fast_log(SAMPLE_LENGTH) - 1;
+  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;
+    }
+
+    // son = get_son_from_step(son, (*last_known_point > val));
+    son = son * 2 + (int)((*last_known_point - val) >> MOVE_OFFSET);
+    // if (son < 0) son = 0;
+    /*printf("%d %d %d\n", (1 << (i + 1)), son,
+           -(int)((*last_known_point - val) >> MOVE_OFFSET));*/
+    last_known_point = next_level_start + son;
+    assert(last_known_point < &*sample_vector.end());
+    // printf("%d\n", *last_known_point);
+  }
+  return last_known_point;
+}
+
+long double sample_cdf(long double x) {
+  auto it = cudaBinarySearch(sample_vector.data(), x);
+  assert(it < &*sample_vector.end());
+  if (it == (&*sample_vector.end() - 1)) {
+    return 1;
+  }
+  if (it == sample_vector.data()) {
+    return 0;
+  };
+  auto it_prev = sample_vector.data() +
+                 calculate_location_inverse(
+                     calculate_location(it - sample_vector.data()) - 2) -
+                 1;
+
+  // printf("------------ %ld %ld\n", it - &*sample_vector.begin(), it_prev -
+  // it);
+
+  auto test = ((long double)(it_prev - (&*sample_vector.begin())) +
+               (long double)(x - *it_prev) / (*it - *it_prev)) /
+              (long double)(sample_vector.size() - 1);
+  return test;
+}
+
+unsigned long long max_value = 0, min_value = 0xfffffffff;
+inline void store_into_vector(unsigned long long value) {
+  if (max_value < value) {
+    max_value = value;
+  }
+  if (min_value > value) {
+    min_value = value;
+  }
+  population_vector.push_back(value);
+}
+
+inline long double safe_ceil(long double value) {
+  auto c = std::ceil(value);
+  return c == 0 ? 1 : c;
+}
+
+int main() {
+  FILE *file = fopen("normal_distribution.txt", "r");
+  assert(file);
+  for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i))
+    ;
+  fclose(file);
+
+  std::vector<key_type> tmp(population_vector.begin(),
+                            population_vector.begin() + SAMPLE_LENGTH - 2);
+  tmp.push_back(max_value);
+  tmp.push_back(min_value);
+  std::sort(tmp.begin(), tmp.end());
+  sample_vector.resize(SAMPLE_LENGTH);
+
+  assert(&sample_vector[0] == &*sample_vector.begin());
+
+  for (int i = 0; i < SAMPLE_LENGTH; i++) {
+    sample_vector[calculate_location(i) - 1] = tmp[i];
+    auto location = calculate_location(i);
+    auto location2 = calculate_location_inverse(location - 1);
+    assert(location2 - 1 == i);
+    // printf("%zu %zu\n", location, location2);
+  }
+  // printf("--- %zu\n", calculate_location_inverse(1022));
+  for (int scale = 2; scale < 9; scale++) {
+    const long split_size = TEST_LENGTH * scale;
+    result_storage.clear();
+    result_storage.resize(scale * SAMPLE_LENGTH, false);
+    auto slice_size = 1.0 / (double)(split_size);
+
+    // printf("\nscale: %d\n", scale);
+    for (int i = 0; i < TEST_LENGTH; i++) {
+
+      auto index =
+          (int)safe_ceil(sample_cdf(population_vector[i + SAMPLE_LENGTH]) /
+                         slice_size) -
+          1;
+      // printf("index: %d\n", index);
+      while (result_storage[index]) {
+        index++;
+      }
+      result_storage[index] = true;
+    }
+    for (auto &&element : result_storage) {
+      printf("%d", element ? 1 : 0);
+    }
+    puts("");
+  }
+}
\ No newline at end of file