1 files changed, 37 insertions, 27 deletions

diff --git a/expt_0425.cu b/expt_0425.cu
index 68c52f3..1a8c4e0 100644
--- a/expt_0425.cu
+++ b/expt_0425.cu
@@ -5,7 +5,8 @@
 #include <cstdio>
 #include <vector>
 
-std::vector<unsigned long long> population_vector, sample_vector;
+std::vector<unsigned long long> population_vector, sample_vector,
+    le_sample_vector;
 
 constexpr size_t SAMPLE_LENGTH = 1023;
 constexpr size_t TEST_LENGTH = 32'768;
@@ -44,6 +45,8 @@ __global__ void initSample(key_type *sample, key_type *population) {
 #endif
 }
 
+__global__ void initNormalStorage(key_type *sample) { cudaSampleItem = sample; }
+
 __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);
@@ -70,27 +73,14 @@ __global__ void kernel(unsigned long step, const double slice_size,
   }
 }
 
-__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");
+__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;
-    // 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);
-      }
-    }*/
+    auto index = (int)(FactorySort::sample_cdf(cudaSampleItem, SAMPLE_LENGTH,
+                                               cudaPopulationItem[tid]) /
+                       slice_size) -
+                 1;
     cdf_result[index] = true;
   }
 }
@@ -110,6 +100,13 @@ __global__ void print_function() {
 #endif
 }
 
+__global__ void check_items() {
+  for (int i = 0; i < SAMPLE_LENGTH; i++) {
+    printf("%lld ", cudaSampleItem[i]);
+  }
+  printf("\n");
+}
+
 __global__ void initStorage(unsigned scale, unsigned test_size) {
   cudaFree(cdf_result);
   // printf("test size: %u\n", test_size);
@@ -143,12 +140,11 @@ __global__ void initCustomSample() {
 
 const dim_pair_type DIM_PAIR[] = {dim_pair_type(32, 32)};
 
-void run_kernel(size_t test_size, bool custom = false) {
+void run_kernel(size_t test_size, int stage = 1) {
   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);
+    printf("%scurrent dim: %d %d %lu\n", "", 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;
@@ -159,11 +155,13 @@ void run_kernel(size_t test_size, bool custom = false) {
       cudaEventCreate(&start);
       cudaEventCreate(&stop);
       cudaEventRecord(start, nullptr);
-      if (custom)
+      switch (stage) {
+      case 1:
         kernel<<<grid_dim, block_dim>>>(step, slice_size, split_size);
-      else
-        kernel2_real_binary<<<grid_dim, block_dim>>>(step, slice_size,
-                                                     split_size);
+      default:
+        kernel2<<<grid_dim, block_dim>>>(step, slice_size);
+      }
+
       cudaDeviceSynchronize();
       cudaEventRecord(stop, nullptr);
       cudaEventSynchronize(stop);
@@ -197,6 +195,7 @@ int main() {
   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;
 
   key_type *cudaSample = nullptr, *cudaPopulation;
   cudaMalloc(&cudaSample, sizeof(key_type) * SAMPLE_LENGTH);
@@ -219,7 +218,18 @@ int main() {
   initCustomSample<<<1, 1>>>();
   cudaDeviceSynchronize();
 
-  run_kernel(test_size, true);
+  key_type *cudaNormalSample = nullptr;
+  cudaMalloc(&cudaNormalSample, sizeof(key_type) * SAMPLE_LENGTH);
+  cudaMemcpy(cudaNormalSample, le_sample_vector.data(),
+             sizeof(key_type) * SAMPLE_LENGTH, cudaMemcpyHostToDevice);
+
+  initNormalStorage<<<1, 1>>>(cudaNormalSample);
+  cudaDeviceSynchronize();
+  check_items<<<1, 1>>>();
+  cudaDeviceSynchronize();
+
+  run_kernel(test_size, 0);
+
   freeStorage<<<1, 1>>>();
   cudaDeviceSynchronize();
 }