2023-04-27 15:59

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

2 files changed, 445 insertions, 0 deletions

diff --git a/expt_0406.cu b/expt_0406.cu
new file mode 100644
index 0000000..fd7915d
--- /dev/null
+++ b/expt_0406.cu
@@ -0,0 +1,220 @@
+#include <algorithm>
+#include <cassert>
+#include <cstdio>
+#include <iostream>
+#include <vector>
+
+constexpr size_t length = 2097152;
+std::vector<unsigned long long> population_vector, sample_vector,
+    sample_vector_into_cuda;
+
+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;
+  } else 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 = 32'768;
+
+__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;
+  assert(begin < end);
+  while (begin < end) {
+    auto mid = (end - begin) / 2;
+    auto mid_val = *(begin + mid);
+    if (val == mid_val) {
+      return begin + mid;
+    } else if (val > mid_val) {
+      begin = begin + mid + 1;
+    } else {
+      end = end - 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("initCuda storage\n");
+#ifdef TEST_BOUNDS
+  index_max = 0;
+  index_min = 0x7fffffff;
+#endif
+}
+
+__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;
+}
+
+__global__ void kernel(unsigned long step, const double slice_size) {
+  for (int i = 0; i < step; i++) {
+    auto tid = step * (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);
+
+  puts("Read success");
+
+  sample_vector = std::vector<unsigned long long>(
+      population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH - 2);
+  sample_vector.push_back(min_value);
+  sample_vector.push_back(max_value);
+
+  printf("Sample vector length: %zu\n", sample_vector.size());
+
+  std::sort(sample_vector.begin(), sample_vector.end());
+
+  sample_vector_into_cuda.resize(SAMPLE_LENGTH);
+  assert(sample_vector.size() == 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));
+  cudaMemcpy(cudaSample, sample_vector.data(), sizeof(key_type) * SAMPLE_LENGTH,
+             cudaMemcpyHostToDevice);
+
+  cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE);
+  cudaMemcpy(cudaPopulation, population_vector.data() + SAMPLE_LENGTH,
+             sizeof(key_type) * TEST_SIZE, cudaMemcpyHostToDevice);
+
+  // memcpy(sample_heap, sample_vector.cbegin(),sizeof(key_type) *(SAMPLE_LENGTH
+  // + 2));
+  puts("Copy successful");
+
+  initCuda<<<1, 1>>>(cudaSample, cudaPopulation);
+  cudaDeviceSynchronize();
+
+  dim3 grid_dim = 32, block_dim = 32;
+
+  unsigned long step = test_size / (grid_dim.x * block_dim.x);
+  printf("(old) step: %lu\n", step);
+  assert(!(test_size % (grid_dim.x * block_dim.x)));
+
+  for (int scale = 2; scale <= 8; scale += 2) {
+    const unsigned 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);
+    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/expt_0425.cu b/expt_0425.cu
new file mode 100644
index 0000000..68c52f3
--- /dev/null
+++ b/expt_0425.cu
@@ -0,0 +1,225 @@
+#include "sortlib.cuh"
+
+#include <algorithm>
+#include <cassert>
+#include <cstdio>
+#include <vector>
+
+std::vector<unsigned long long> population_vector, sample_vector;
+
+constexpr size_t SAMPLE_LENGTH = 1023;
+constexpr size_t TEST_LENGTH = 32'768;
+// constexpr size_t TEST_LENGTH = 4096;
+constexpr size_t RESERVED_BLOCK = 10'000;
+
+unsigned long long max_value = 0, min_value = 0xfffffffff;
+typedef std::pair<int, int> dim_pair_type;
+
+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);
+}
+// #define TEST_BOUNDS
+
+__device__ key_type *cudaSampleItem, *cudaPopulationItem;
+__device__ bool *cdf_result;
+#ifdef TEST_BOUNDS
+__device__ unsigned insert_value;
+__device__ unsigned int index_max, index_min;
+#endif
+
+__global__ void initSample(key_type *sample, key_type *population) {
+  cudaSampleItem = sample;
+  cudaPopulationItem = population;
+  // cudaMalloc(&cdf_result, sizeof(bool) * TEST_LENGTH);
+  cdf_result = nullptr;
+#ifdef TEST_BOUNDS
+  index_max = 0;
+  index_min = 0x7fffffff;
+#endif
+}
+
+__global__ void kernel(unsigned long step, const double slice_size,
+                       const unsigned long split_size) {
+  auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(key_type) * 8);
+  // printf("kernel1 step: %ld\n", step);
+  for (int i = 0; i < step; i++) {
+    auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i;
+
+    // printf("%d\n", tid);
+    auto index = (int)(custom_sort.sample_cdf_custom_version(
+                           cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH,
+                           cudaPopulationItem[tid]) /
+                       slice_size) -
+                 1;
+    // printf("%llu %d\n", cudaPopulationItem[tid], index);
+    /*if (cdf_result[index]) {
+      while (cdf_result[++index]) {
+        if (index >= split_size) {
+          //printf("escape: %llu %lu\n", cudaPopulationItem[tid], split_size);
+        }
+        assert(index < (split_size + RESERVED_BLOCK));
+      }
+    }*/
+    cdf_result[index] = true;
+  }
+}
+
+__global__ void kernel2_real_binary(unsigned long step, const double slice_size,
+                                    const unsigned long split_size) {
+  auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(key_type) * 8);
+  // printf("%d\n", custom_sort.MOVE_OFFSET);
+  //  printf("kernel2\n");
+
+  for (int i = 0; i < step; i++) {
+    auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i;
+    // printf("%lu ", tid);
+
+    // assert(tid < TEST_LENGTH);
+    auto index = (int)(custom_sort.sample_cdf(cudaSampleItem,
+                                              cudaSampleItem + SAMPLE_LENGTH,
+                                              cudaPopulationItem[tid]) /
+                       slice_size);
+    // printf("%llu %d\n", cudaPopulationItem[tid], index);
+    /*if (cdf_result[index]) {
+      while (cdf_result[++index]) {
+        assert(index < split_size);
+      }
+    }*/
+    cdf_result[index] = true;
+  }
+}
+
+__global__ void print2() {
+  /*for (int i = 0; i< SAMPLE_LENGTH; i++) {
+    printf("%llu ", cudaSampleItem[i]);
+  }
+  printf("\n");*/
+  printf("%lu\n", CustomSort::fast_log(SAMPLE_LENGTH));
+}
+
+__global__ void print_function() {
+#ifdef TEST_BOUNDS
+  printf("%u\n", insert_value);
+  printf("%u %u\n", index_min, index_max);
+#endif
+}
+
+__global__ void initStorage(unsigned scale, unsigned test_size) {
+  cudaFree(cdf_result);
+  // printf("test size: %u\n", test_size);
+  cudaMalloc(&cdf_result, (scale * test_size + RESERVED_BLOCK) * sizeof(bool));
+  memset(cdf_result, 0, (scale * test_size + RESERVED_BLOCK) * sizeof(bool));
+}
+
+__global__ void freeStorage() {
+  cudaFree(cudaSampleItem);
+  cudaFree(cudaPopulationItem);
+  cudaFree(cdf_result);
+}
+
+__global__ void initCustomSample() {
+  // printf("init\n");
+  key_type *tmp = nullptr;
+  auto custom_sort = CustomSort(SAMPLE_LENGTH, 0);
+  cudaMalloc(&tmp, sizeof(key_type) * SAMPLE_LENGTH);
+  for (size_t i = 0; i < SAMPLE_LENGTH; i++) {
+    tmp[i] = cudaSampleItem[custom_sort.calculate_index(i) - 1];
+  }
+  // printf("copy\n");
+  memcpy(cudaSampleItem, tmp, sizeof(key_type) * SAMPLE_LENGTH);
+  cudaFree(tmp);
+  tmp = nullptr;
+  // memset(cdf_result, 0, sizeof(key_type) * TEST_LENGTH);
+  cudaFree(cdf_result);
+  cdf_result = nullptr;
+  // printf("finalize\n");
+}
+
+const dim_pair_type DIM_PAIR[] = {dim_pair_type(32, 32)};
+
+void run_kernel(size_t test_size, bool custom = false) {
+  for (auto &pair : DIM_PAIR) {
+    dim3 grid_dim = pair.first, block_dim = pair.second;
+    unsigned long step = test_size / (grid_dim.x * block_dim.x);
+    printf("%scurrent dim: %d %d %lu\n", custom ? "custom " : "", pair.first,
+           pair.second, step);
+    for (int scale = 2; scale <= 16; scale += 2) {
+      const unsigned 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);
+      if (custom)
+        kernel<<<grid_dim, block_dim>>>(step, slice_size, split_size);
+      else
+        kernel2_real_binary<<<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);
+      if (time == 0) {
+        // printf("last error: %u\n", cudaGetLastError());
+      }
+      printf("time: %lf\n", time);
+
+      print_function<<<1, 1>>>();
+      cudaDeviceSynchronize();
+    }
+  }
+}
+
+int main() {
+
+  auto test_size = TEST_LENGTH;
+  FILE *file = fopen("normal_distribution.txt", "r");
+  assert(file);
+  for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i))
+    ;
+  fclose(file);
+
+  printf("test size: %d\n", test_size);
+  sample_vector = std::vector<key_type>(
+      population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH - 2);
+  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;
+  cudaMalloc(&cudaSample, sizeof(key_type) * SAMPLE_LENGTH);
+  assert(SAMPLE_LENGTH == sample_vector.size());
+  cudaMemcpy(cudaSample, sample_vector.data(), sizeof(key_type) * SAMPLE_LENGTH,
+             cudaMemcpyHostToDevice);
+  cudaMalloc(&cudaPopulation, sizeof(key_type) * test_size);
+  cudaMemcpy(cudaPopulation, population_vector.data() + SAMPLE_LENGTH,
+             sizeof(key_type) * test_size, cudaMemcpyHostToDevice);
+
+  // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8);
+  // custom_sort.testCalculation();
+  testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH);
+  cudaDeviceSynchronize();
+
+  initSample<<<1, 1>>>(cudaSample, cudaPopulation);
+  cudaDeviceSynchronize();
+  run_kernel(test_size);
+
+  initCustomSample<<<1, 1>>>();
+  cudaDeviceSynchronize();
+
+  run_kernel(test_size, true);
+  freeStorage<<<1, 1>>>();
+  cudaDeviceSynchronize();
+}