feat: Optimize read_helper library

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

1 files changed, 57 insertions, 101 deletions

diff --git a/expt_0516.cu b/expt_0516.cu
index 82755ac..351a8c4 100644
--- a/expt_0516.cu
+++ b/expt_0516.cu
@@ -1,18 +1,13 @@
 // Experimental content: Test branch performance (optimized memory and output)
+// #define PRINT_READ_PROCESS
+#include "read_helper.h"
 #include "sortlib.cuh"
 
 #include <algorithm>
 #include <cassert>
 #include <cstdio>
-#include <random>
 #include <vector>
 
-#ifndef NDEBUG
-#define P_ERR(...) fprintf(stderr, __VA_ARGS__)
-#else
-#define P_ERR(...)
-#endif
-
 std::vector<unsigned long long> population_vector, sample_vector,
     le_sample_vector;
 
@@ -21,27 +16,10 @@ 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__ unsigned 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;
@@ -55,6 +33,27 @@ __global__ void initNormalSample(key_type *normal_sample) {
   cudaNormalSampleItem = normal_sample;
 }
 
+/*
+#define CudaCheckError() cudaCheckError(__FILE__, __LINE__)
+inline void cudaCheckError(const char *file, const int line) {
+  cudaError err = cudaGetLastError();
+  if (cudaSuccess != err) {
+    fprintf(stderr, "cudaCheckError() failed at %s:%i : %s\n", file, line,
+            cudaGetErrorString(err));
+    exit(-1);
+  }
+
+  // More careful checking. However, this will affect performance.
+  // Comment away if needed.
+  err = cudaDeviceSynchronize();
+  if (cudaSuccess != err) {
+    fprintf(stderr, "cudaCheckError() with sync failed at %s:%i : %s\n", file,
+            line, cudaGetErrorString(err));
+    exit(-1);
+  }
+}
+*/
+
 __global__ void kernel(unsigned long step, const double slice_size) {
   auto custom_sort = CustomSort(cuda_sample_length, sizeof(key_type) * 8);
   // printf("kernel1 step: %ld\n", step);
@@ -100,25 +99,8 @@ __global__ void checkItems() {
   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;
-}
-
 constexpr dim3 grid_dim = 32, block_dim = 32;
 
-bool checkTestSize() {}
-
 void run_kernel(size_t test_size, bool custom = true) {
   // need description
   unsigned long step = test_size / (grid_dim.x * block_dim.x);
@@ -163,45 +145,6 @@ __global__ void freeStorageStage2() {
   // cudaFree(cdf_result);
 }
 
-unsigned long randomRow(unsigned long max_value_) {
-  std::random_device randomDevice;
-  std::mt19937 mt19937(randomDevice());
-  std::uniform_int_distribution<std::mt19937::result_type> dst(0, max_value_);
-  return dst(mt19937);
-}
-
-void readFile(char const *filename, long sample_length,
-              unsigned long &total_row,
-              unsigned long max_number = TEST_LENGTH) {
-  auto read_number = 0UL;
-  FILE *file = fopen(filename, "r");
-  assert(file);
-  P_ERR("Reading sample");
-  for (long long i;
-       read_number < max_number && fscanf(file, "%lld ", &i) != EOF;
-       store_into_vector(i))
-    read_number++;
-
-  if (total_row > 0) {
-    P_ERR("\rReading skip");
-    read_number = randomRow(total_row - sample_length - max_number - 256);
-    total_row = read_number;
-    // fprintf(stderr, "Skip %lu\n", read_number);
-    for (long long i; read_number > 0 && fscanf(file, "%lld ", &i) != EOF;)
-      read_number--;
-  }
-
-  P_ERR("\rReading population");
-  read_number = 0;
-
-  auto remain = sample_length + 256;
-  for (long long i; read_number < remain && fscanf(file, "%lld ", &i) != EOF;
-       store_into_vector(i))
-    read_number++;
-  fclose(file);
-  P_ERR("\r");
-}
-
 long pow_for_sample(long n) {
   auto x = 2;
   for (int i = 1; i < n; i++) {
@@ -210,7 +153,7 @@ long pow_for_sample(long n) {
   return x - 1;
 }
 
-__global__ void applyCudaSampleLength(long length) {
+__global__ void applyCudaSampleLength(unsigned long length) {
   cuda_sample_length = length;
 }
 
@@ -230,36 +173,46 @@ int main(int argc, char const *argv[]) {
     total_row = strtol(argv[3], nullptr, 10);
   }
 
-  readFile("normal_distribution.txt", sample_length, total_row, test_size);
+  // printf("%ld %ld %ld\n", test_size, sample_length, total_row);
+
+  ReadHelper readHelper("normal_distribution.txt", sample_length, test_size,
+                        total_row);
 
-  printf("population: %zu, skip: %lu, test size: %zu, sample length: %zu\n",
-         population_vector.size(), total_row, test_size, sample_length);
+  printf("population: %zu, skip: %lu, sample length: %zu ", test_size,
+         readHelper.random_number, sample_length);
+  readHelper.readFile();
+
+  readHelper.split_into(sample_vector, population_vector);
+  assert(sample_vector.size() == sample_length);
+  assert(population_vector.size() == 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());
   le_sample_vector = sample_vector;
-  initCustomSampleCPU();
+  rebuild(sample_vector);
+
+  /*printf("population: %zu sample length: %zu\n", population_vector.size(),
+         sample_vector.size());*/
 
   key_type *cudaSample = nullptr, *cudaPopulation;
 
   P_ERR("Coping sample");
-  cudaMalloc(&cudaSample, sizeof(key_type) * sample_length);
+  cudaMalloc(&cudaSample, sizeof(key_type) * sample_vector.size());
   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(cudaSample, sample_vector.data(),
+             sizeof(key_type) * sample_vector.size(), cudaMemcpyHostToDevice);
+  cudaMalloc(&cudaPopulation, sizeof(key_type) * population_vector.size());
+
+  assert(population_vector.size() == test_size);
 
   P_ERR("\rCoping population");
-  cudaMemcpy(cudaPopulation, population_vector.data() + sample_length,
-             sizeof(key_type) * test_size, cudaMemcpyHostToDevice);
+  cudaMemcpy(cudaPopulation, population_vector.data(),
+             sizeof(key_type) * population_vector.size(),
+             cudaMemcpyHostToDevice);
 
   // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8);
   // custom_sort.testCalculation();
   P_ERR("\rApply custom length and test calculation");
-  applyCudaSampleLength<<<1, 1>>>(sample_length);
+  applyCudaSampleLength<<<1, 1>>>(sample_vector.size());
   // testCustomCalculation<<<1, 1>>>(sample_length);
   cudaDeviceSynchronize();
   P_ERR("\rApply custom length and test calculation completed");
@@ -276,15 +229,18 @@ int main(int argc, char const *argv[]) {
         "kernel\n");
   run_kernel(test_size);
 
-  freeStorageStage1<<<1, 1>>>();
-  cudaDeviceSynchronize();
+  // freeStorageStage1<<<1, 1>>>();
+  // cudaDeviceSynchronize();
+
+  cudaFree(cudaSample);
 
   key_type *cudaNormalSample = nullptr;
 
   P_ERR("\nCoping new sample");
-  cudaMalloc(&cudaNormalSample, sizeof(key_type) * sample_length);
+  cudaMalloc(&cudaNormalSample, sizeof(key_type) * le_sample_vector.size());
   cudaMemcpy(cudaNormalSample, le_sample_vector.data(),
-             sizeof(key_type) * sample_length, cudaMemcpyHostToDevice);
+             sizeof(key_type) * le_sample_vector.size(),
+             cudaMemcpyHostToDevice);
 
   /*cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) *
      sample_length, cudaMemcpyDeviceToHost);*/