1 files changed, 254 insertions, 0 deletions
diff --git a/expt_0510.cu b/expt_0510.cu
new file mode 100644
index 0000000..063c6ba
--- /dev/null
+++ b/expt_0510.cu
@@ -0,0 +1,254 @@
+// Experimental content: Test branch performance (optimized memory)
+#include "sortlib.cuh"
+
+#include <algorithm>
+#include <cassert>
+#include <cstdio>
+#include <vector>
+
+std::vector<unsigned long long> population_vector, sample_vector,
+    le_sample_vector;
+
+constexpr long DEFAULT_SAMPLE_LENGTH = 65535;
+constexpr size_t TEST_LENGTH = 32'768;
+// constexpr size_t TEST_LENGTH = 4096;
+// constexpr size_t RESERVED_BLOCK = 100;
+
+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__ long cuda_sample_length;
+__device__ key_type *cudaSampleItem, *cudaNormalSampleItem, *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;
+  // cdf_normal_result = nullptr;
+#ifdef TEST_BOUNDS
+  index_max = 0;
+  index_min = 0x7fffffff;
+#endif
+}
+
+__global__ void initNormalSample(key_type *normal_sample) {
+  cudaNormalSampleItem = normal_sample;
+}
+
+__global__ void kernel(unsigned long step, const double slice_size,
+                       const unsigned long split_size) {
+  auto custom_sort = CustomSort(cuda_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 + cuda_sample_length,
+                            cudaPopulationItem[tid]) /
+                        slice_size) -
+                  1;
+    // cdf_result[index] = true;
+  }
+}
+
+__global__ void kernel2(unsigned long step, const double slice_size) {
+
+  for (int i = 0; i < step; i++) {
+    auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i;
+    auto _index =
+        (int)(FactorySort::sample_cdf(cudaNormalSampleItem, cuda_sample_length,
+                                      cudaPopulationItem[tid]) /
+              slice_size) -
+        1;
+    // 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(cuda_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 check_items() {
+  for (int i = 0; i < cuda_sample_length; i++) {
+    printf("%lld ", cudaSampleItem[i]);
+  }
+  printf("\n");
+}
+
+void initCustomSampleCPU() {
+  auto sample_length = sample_vector.size();
+  auto tmp = new key_type[sample_length];
+
+  auto custom_sort = CustomSort(sample_length, 0);
+
+  for (size_t i = 0; i < sample_length; i++) {
+    tmp[i] = sample_vector[custom_sort.calculate_index(i) - 1];
+  }
+
+  memcpy(sample_vector.data(), tmp, sizeof(key_type) * sample_length);
+
+  delete[] tmp;
+}
+
+void run_kernel(size_t test_size, bool normal = true) {
+  dim3 grid_dim = 32, block_dim = 32;
+  unsigned long step = test_size / (grid_dim.x * block_dim.x);
+  auto scale = 2;
+  const unsigned long split_size = test_size * scale;
+  const auto slice_size = 1.0 / (double)split_size;
+  cudaEvent_t start, stop;
+  cudaEventCreate(&start);
+  cudaEventCreate(&stop);
+  cudaEventRecord(start, nullptr);
+  if (normal) {
+    kernel<<<grid_dim, block_dim>>>(step, slice_size, split_size);
+  } else {
+    kernel2<<<grid_dim, block_dim>>>(step, slice_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);
+  cudaDeviceSynchronize();
+
+  print_function<<<1, 1>>>();
+  cudaDeviceSynchronize();
+}
+
+__global__ void freeStorageStage1() { cudaFree(cudaSampleItem); }
+
+__global__ void freeStorageStage2() {
+  cudaFree(cudaNormalSampleItem);
+  cudaFree(cudaPopulationItem);
+  // cudaFree(cdf_result);
+}
+
+void read_file(char const *filename, long sample_length,
+               unsigned long max_number = TEST_LENGTH) {
+  max_number += sample_length + 256;
+  auto read_number = 0;
+  FILE *file = fopen(filename, "r");
+  assert(file);
+  for (long long i;
+       read_number < max_number && fscanf(file, "%lld ", &i) != EOF;
+       store_into_vector(i))
+    read_number++;
+  fclose(file);
+}
+
+long pow_for_sample(long n) {
+  auto x = 2;
+  for (int i = 0; i < n; i++) {
+    x *= 2;
+  }
+  return x - 1;
+}
+
+__global__ void applyCudaSampleLength(long length) {
+  cuda_sample_length = length;
+}
+
+int main(int argc, char const *argv[]) {
+
+  auto test_size = TEST_LENGTH;
+  auto sample_length = DEFAULT_SAMPLE_LENGTH;
+
+  if (argc >= 2) {
+    test_size = strtol(argv[1], nullptr, 10);
+  }
+  if (argc >= 3) {
+    sample_length = pow_for_sample(strtol(argv[2], nullptr, 10));
+  }
+
+  read_file("normal_distribution.txt", test_size);
+
+  printf("population: %zu, test size: %zu, sample length: %zu\n",
+         population_vector.size(), test_size, sample_length);
+
+  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());
+  le_sample_vector = sample_vector;
+  initCustomSampleCPU();
+
+  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();
+  applyCudaSampleLength<<<1, 1>>>(sample_length);
+  testCustomCalculation<<<1, 1>>>(sample_length);
+  cudaDeviceSynchronize();
+
+  initSample<<<1, 1>>>(cudaSample, cudaPopulation);
+  cudaDeviceSynchronize();
+
+  // check_items<<<1, 1>>>();
+  // cudaDeviceSynchronize();
+  run_kernel(test_size);
+
+  freeStorageStage1<<<1, 1>>>();
+  cudaDeviceSynchronize();
+
+  key_type *cudaNormalSample = nullptr;
+  cudaMalloc(&cudaNormalSample, sizeof(key_type) * sample_length);
+  cudaMemcpy(cudaNormalSample, le_sample_vector.data(),
+             sizeof(key_type) * sample_length, cudaMemcpyHostToDevice);
+
+  cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) * sample_length,
+             cudaMemcpyDeviceToHost);
+
+  initNormalSample<<<1, 1>>>(cudaNormalSample);
+  cudaDeviceSynchronize();
+  run_kernel(test_size, false);
+  // check_items<<<1, 1>>>();
+  // cudaDeviceSynchronize();
+
+  freeStorageStage2<<<1, 1>>>();
+  cudaDeviceSynchronize();
+}