feat: Optimize read_helper library

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

8 files changed, 111 insertions, 162 deletions

diff --git a/expt_0510.cu b/expt_0510.cu
index 82f2624..eca1b92 100644
--- a/expt_0510.cu
+++ b/expt_0510.cu
@@ -1,6 +1,7 @@
 // Experimental content: Test branch performance (optimized memory)
 #include "sortlib.cuh"
 
+#include "read_helper.h"
 #include <algorithm>
 #include <cassert>
 #include <cstdio>
@@ -167,34 +168,6 @@ unsigned long randomRow(unsigned long 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);
-  for (long long i;
-       read_number < max_number && fscanf(file, "%lld ", &i) != EOF;
-       store_into_vector(i))
-    read_number++;
-
-  if (total_row > 0) {
-    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--;
-  }
-
-  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);
-}
-
 long pow_for_sample(long n) {
   auto x = 2;
   for (int i = 0; i < n; i++) {
@@ -225,10 +198,13 @@ int main(int argc, char const *argv[]) {
 
   // CustomSort::testSelf(sample_length);
 
-  readFile("normal_distribution.txt", sample_length, total_row, test_size);
+  ReadHelper readHelper("normal_distribution.txt", sample_length, test_size,
+                        total_row);
+  readHelper.readFile();
 
   printf("population: %zu, skip: %lu, test size: %zu, sample length: %zu | ",
-         population_vector.size(), total_row, test_size, sample_length);
+         population_vector.size(), readHelper.random_number, test_size,
+         sample_length);
 
   sample_vector = std::vector<key_type>(
       population_vector.begin(), population_vector.begin() + sample_length - 2);
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);*/
diff --git a/expt_0520.cu b/expt_0520.cu
index d31dc78..bd1439f 100644
--- a/expt_0520.cu
+++ b/expt_0520.cu
@@ -208,8 +208,9 @@ int main(int argc, char **argv) {
   printf("insert_length %ld, sample_length: %ld\n", insertion_length,
          sample_length);
 
-  ReadHelper reader("normal_distribution.txt", sample_length, insertion_length);
-  reader.readFile(total_row);
+  ReadHelper reader("normal_distribution.txt", sample_length, insertion_length,
+                    total_row);
+  reader.readFile();
 
   key_type *cudaPopulation;
 
diff --git a/expt_0525.cu b/expt_0525.cu
index c2ee041..8867cee 100644
--- a/expt_0525.cu
+++ b/expt_0525.cu
@@ -76,7 +76,6 @@ typedef unsigned long long LL;
 // #define MEASURE_TIME
 //  #define MEASURE_ACCESS
 
-#define READ_NO_OUTPUT
 #include "read_helper.h"
 
 #if (defined(MEASURE_ACCESS) && defined(MEASURE_TIME))
@@ -108,12 +107,6 @@ constexpr size_t FACTOR = 1;
 
 typedef LL key_type;
 
-#ifdef RANDOM_TARGET
-constexpr const char *TARGET_STRING = "RANDOM";
-#else
-// constexpr const char *TARGET_STRING = "PERFECT";
-#endif
-
 #define CUDA_ERROR_CHECK
 
 #define CudaSafeCall(err) __cudaSafeCall(err, __FILE__, __LINE__)
@@ -153,7 +146,13 @@ inline void __cudaCheckError(const char *file, const int line) {
 
 // Definition of generic node class
 
-class __attribute__((aligned(16))) Node {
+class
+#ifndef _MSC_VER
+    __attribute__((aligned(16)))
+#else
+    __declspec(align(16))
+#endif
+    Node {
 public:
   int topLevel;             // Level of the node
   LL key;                   // Key value
@@ -636,7 +635,7 @@ inline double calcSliceSize(size_t insertion_size) {
   return 1.0 / (double)insertion_size;
 }
 
-inline size_t calcBlocks(size_t input) {
+inline auto calcBlocks(size_t input) {
   return (input % (NUM_THREADS * FACTOR) == 0)
              ? input / (NUM_THREADS * FACTOR)
              : (input / (NUM_THREADS * FACTOR)) + 1;
@@ -660,8 +659,8 @@ int main(int argc, char **argv) {
          insertion_length, search_length);
 
   ReadHelper readHelper("normal_distribution.txt", sample_length,
-                        insertion_length);
-  readHelper.readFile(nullptr);
+                        insertion_length, 0);
+  readHelper.readFile();
   std::vector<key_type> _sample, _population;
   readHelper.split_into(_sample, _population);
   /*printf("%lu, Create search vector: %ld\n", _population.size(),
diff --git a/expt_0528.cu b/expt_0528.cu
index 9ac7059..d74c93a 100644
--- a/expt_0528.cu
+++ b/expt_0528.cu
@@ -153,7 +153,13 @@ inline void __cudaCheckError(const char *file, const int line) {
 
 // Definition of generic node class
 
-class __attribute__((aligned(16))) Node {
+class
+#ifndef _MSC_VER
+    __attribute__((aligned(16)))
+#else
+    __declspec(align(16))
+#endif
+    Node {
 public:
   int topLevel;             // Level of the node
   LL key;                   // Key value
@@ -652,8 +658,8 @@ int main(int argc, char **argv) {
 
   printf("Insertion: %ld, Search: %ld\n", insertion_length, search_length);
 
-  ReadHelper readHelper("normal_distribution.txt", 0, insertion_length);
-  readHelper.readFile(nullptr);
+  ReadHelper readHelper("normal_distribution.txt", 0, insertion_length, 0);
+  readHelper.readFile();
   std::vector<key_type> _search(readHelper.population_vector.begin(),
                                 readHelper.population_vector.begin() +
                                     search_length);
diff --git a/expt_0529.cu b/expt_0529.cu
index 985648f..bd64ad0 100644
--- a/expt_0529.cu
+++ b/expt_0529.cu
@@ -46,7 +46,7 @@ int main(int argc, char const *argv[]) {
          population_length);
 
   ReadHelper readHelper("normal_distribution.txt", sample_length, 0);
-  readHelper.readFile(nullptr);
+  readHelper.readFile();
 
   std::vector<key_type> sample, population;
   std::vector<unsigned int> result(RESULT_LENGTH);
diff --git a/read_helper.h b/read_helper.h
index d19834a..f8396f2 100644
--- a/read_helper.h
+++ b/read_helper.h
@@ -2,6 +2,7 @@
 #ifndef LOCKFREE_READ_HELPER_H
 #define LOCKFREE_READ_HELPER_H
 
+#include <cassert>
 #include <cstdio>
 #include <random>
 #include <vector>
@@ -34,12 +35,22 @@ class ReadHelper {
   char const *filename;
   static constexpr size_t REVERSED_BLOCK = 256;
 
+  auto genRandomRow(size_t total_row) const {
+    if (total_row != 0) {
+      assert(total_row > (population_length + sample_length + REVERSED_BLOCK));
+      return randomRow(total_row - population_length - sample_length -
+                       REVERSED_BLOCK);
+    }
+    return 0UL;
+  }
+
 public:
-  const size_t sample_length, needed_read_length;
+  const size_t sample_length, population_length, random_number;
   ReadHelper(char const *filename, size_t sample_length,
-             size_t insertion_length)
+             size_t population_length, size_t total_row = 0)
       : filename(filename), sample_length(sample_length),
-        needed_read_length(insertion_length + sample_length) {
+        population_length(population_length),
+        random_number(genRandomRow(total_row)) {
     // printf("%zu %zu\n", sample_length, needed_read_length);
   }
   key_type maxValue() const { return max_value; }
@@ -54,10 +65,11 @@ public:
     return dst(mt19937);
   }
 
-  void readFile(unsigned long *total_row) {
+  void readFile() {
     this->population_vector.clear();
     auto read_number = 0UL;
     FILE *file = fopen(filename, "r");
+
     if (file == nullptr) {
       fprintf(stderr, "Unable to open file %s\n", filename);
       exit(1);
@@ -69,12 +81,10 @@ public:
          store_into_vector(i))
       read_number++;
 
-    if (total_row != nullptr) {
+    if (random_number > 0) {
       P_ERR("\rReading skip");
-      read_number = randomRow(*total_row - sample_length - needed_read_length -
-                              REVERSED_BLOCK);
-      *total_row = read_number;
       // fprintf(stderr, "Skip %lu\n", read_number);
+      read_number = random_number;
       for (key_type i; read_number > 0 && fscanf(file, "%lld ", &i) != EOF;)
         read_number--;
     }
@@ -82,7 +92,7 @@ public:
     P_ERR("\rReading population");
     read_number = 0;
 
-    auto remain = needed_read_length - sample_length + REVERSED_BLOCK;
+    auto remain = population_length + REVERSED_BLOCK;
     for (key_type i; read_number < remain && fscanf(file, "%lld ", &i) != EOF;
          store_into_vector(i))
       read_number++;
@@ -92,25 +102,25 @@ public:
 
   void split_into(std::vector<key_type> &sample, std::vector<key_type> &p) {
     sample.resize(sample_length - 2);
-    p.resize(needed_read_length - sample_length);
+    p.resize(population_length);
     // printf("%zu\n", needed_read_length - sample_length);
     memcpy(sample.data(), population_vector.data(),
            sizeof(key_type) * (sample_length - 2));
     sample.push_back(this->max_value);
     sample.push_back(this->min_value);
     memcpy(p.data(), population_vector.data() + sample_length,
-           sizeof(key_type) * (needed_read_length - sample_length));
+           sizeof(key_type) * population_length);
   }
 
   void split_into(key_type *&sample, key_type *&p) {
     sample = new key_type[sample_length];
-    p = new key_type[needed_read_length - sample_length];
+    p = new key_type[population_length];
     memcpy(sample, population_vector.data(),
            sizeof(key_type) * (sample_length - 2));
     sample[sample_length - 2] = this->max_value;
     sample[sample_length - 1] = this->min_value;
     memcpy(p, population_vector.data() + sample_length,
-           sizeof(key_type) * (needed_read_length - sample_length));
+           sizeof(key_type) * (population_length));
   }
 
   size_t size() const { return this->population_vector.size(); }
diff --git a/sortlib.cuh b/sortlib.cuh
index 64f355b..8d6788e 100644
--- a/sortlib.cuh
+++ b/sortlib.cuh
@@ -239,7 +239,8 @@ public:
     return last_known_point;
   }
 
-  __device__ __host__ double static sample_cdf(key_type *begin, long length,
+  __device__ __host__ double static sample_cdf(key_type *begin,
+                                               unsigned long length,
                                                key_type x) {
     // printf("%f\n", x);
     auto end = begin + length;