2023-05-01 19:50

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

6 files changed, 187 insertions, 201 deletions

diff --git a/expt_0406.cu b/expt_0406.cu
index fd7915d..96533c9 100644
--- a/expt_0406.cu
+++ b/expt_0406.cu
@@ -1,13 +1,21 @@
 #include <algorithm>
 #include <cassert>
 #include <cstdio>
-#include <iostream>
 #include <vector>
 
 constexpr size_t length = 2097152;
+constexpr long TEST_LENGTH = 32'768;
 std::vector<unsigned long long> population_vector, sample_vector,
     sample_vector_into_cuda;
 
+#define CHECK_CUDA_ERR(x)                                                      \
+  {                                                                            \
+    auto ret_ = (x);                                                           \
+    if (ret_ != 0) {                                                           \
+      printf("%d: CUDA ERR: %d\n", __LINE__, ret_);                            \
+    }                                                                          \
+  }
+
 unsigned long long max_value = 0, min_value = 0xfffffffff;
 
 inline void store_into_vector(unsigned long long value) {
@@ -25,68 +33,20 @@ typedef unsigned long long *key_type_ptr;
 
 constexpr long SAMPLE_LENGTH = 1024;
 
-constexpr long TEST_SIZE = 32'768;
-
-__device__ key_type *SampleItem, *PopulationItem;
+__device__ key_type *cudaSampleItem, *cudaPopulationItem;
 __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));
+  CHECK_CUDA_ERR(cudaFree(cdf_result));
+  CHECK_CUDA_ERR(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));
+  cudaSampleItem = sample_item;
+  cudaPopulationItem = population_item;
   cdf_result = nullptr;
 }
 
@@ -95,126 +55,95 @@ __global__ void kernel(unsigned long step, const double slice_size) {
     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);
+    auto index = (int)(sample_cdf(cudaPopulationItem[tid]) / slice_size);
 
-#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
+    cdf_result[index] = true;
   }
 }
 
 __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;
+__global__ void freeStorage() {
+  cudaFree(cudaSampleItem);
+  cudaFree(cudaPopulationItem);
+  cudaFree(cdf_result);
+}
+
+typedef std::pair<int, int> dim_pair_type;
+const dim_pair_type DIM_PAIR[] = {dim_pair_type(32, 32)};
+
+void run_kernel(size_t test_size) {
+  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("current 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;
+      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);
+      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);
 
-  assert(population_vector.size() == length);
-
-  puts("Read success");
-
-  sample_vector = std::vector<unsigned long long>(
+  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);
-
-  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));
+  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);
+  cudaMalloc(&cudaPopulation, sizeof(key_type) * test_size);
   cudaMemcpy(cudaPopulation, population_vector.data() + SAMPLE_LENGTH,
-             sizeof(key_type) * TEST_SIZE, cudaMemcpyHostToDevice);
+             sizeof(key_type) * test_size, cudaMemcpyHostToDevice);
 
-  // memcpy(sample_heap, sample_vector.cbegin(),sizeof(key_type) *(SAMPLE_LENGTH
-  // + 2));
-  puts("Copy successful");
+  // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8);
+  // custom_sort.testCalculation();
+  // testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH);
+  // cudaDeviceSynchronize();
+  run_kernel(test_size);
 
-  initCuda<<<1, 1>>>(cudaSample, cudaPopulation);
+  freeStorage<<<1, 1>>>();
   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
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();
 }
diff --git a/expt_0729.cu b/expt_0729.cu
index 2efa575..5660c6c 100644
--- a/expt_0729.cu
+++ b/expt_0729.cu
@@ -29,7 +29,7 @@ constexpr long SAMPLE_LENGTH = 1024;
 
 constexpr long TEST_SIZE = 1024;
 
-__device__ key_type *SampleItem, *PopulationItem;
+__device__ key_type *cudaSampleItem, *cudaPopulationItem;
 __device__ bool *cdf_result;
 __device__ unsigned insert_value;
 #ifdef TEST_BOUNDS
@@ -56,15 +56,16 @@ __device__ const key_type *cudaBinarySearch(key_type *start, key_type *end,
 }
 
 __device__ double sample_cdf(double x) {
-  auto it = cudaBinarySearch(SampleItem, SampleItem + SAMPLE_LENGTH + 2, x);
-  if (it == SampleItem + SAMPLE_LENGTH) {
+  auto it =
+      cudaBinarySearch(cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH + 2, x);
+  if (it == cudaSampleItem + SAMPLE_LENGTH) {
     return 1;
   }
-  if (it == SampleItem) {
+  if (it == cudaSampleItem) {
     return 0;
   }
   auto it_prev = it - 1;
-  return (double(it_prev - SampleItem) +
+  return (double(it_prev - cudaSampleItem) +
           (x - (double)*it_prev) / (double)(*it - *it_prev)) /
          double(SAMPLE_LENGTH - 1);
 }
@@ -82,9 +83,10 @@ __global__ void initStorage(unsigned scale, unsigned test_size) {
 }
 
 __global__ void init(key_type *sample_item, key_type *population_item) {
-  SampleItem = sample_item;
-  PopulationItem = population_item;
-  // cudaMalloc(&SampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long long));
+  cudaSampleItem = sample_item;
+  cudaPopulationItem = population_item;
+  // cudaMalloc(&cudaSampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long
+  // long));
   //  cudaMalloc(&Storage, TEST_SIZE * sizeof(key_type));
   cdf_result = nullptr;
 }
@@ -96,7 +98,7 @@ __global__ void kernel(unsigned long step, const double slice_size,
                threadIdx.x + i;
 
     // printf("%d\n", tid);
-    auto index = (int)(sample_cdf(PopulationItem[tid]) / slice_size);
+    auto index = (int)(sample_cdf(cudaPopulationItem[tid]) / slice_size);
     if (cdf_result[index]) {
       while (cdf_result[++index]) {
         assert(index < split_size);
diff --git a/expt_0804.cu b/expt_0804.cu
index a94dd2b..66f66ec 100644
--- a/expt_0804.cu
+++ b/expt_0804.cu
@@ -30,7 +30,7 @@ constexpr long SAMPLE_LENGTH = 1024;
 
 constexpr long TEST_SIZE = 1024;
 
-__device__ key_type *SampleItem, *PopulationItem;
+__device__ key_type *cudaSampleItem, *cudaPopulationItem;
 __device__ bool *cdf_result;
 __device__ unsigned insert_value;
 #ifdef TEST_BOUNDS
@@ -62,15 +62,15 @@ __device__ __host__ size_t calculate_location(size_t index) {
 }
 
 __device__ double sample_cdf(double x) {
-  auto it = cudaBinarySearch(SampleItem, x);
-  if (it == SampleItem + SAMPLE_LENGTH) {
+  auto it = cudaBinarySearch(cudaSampleItem, x);
+  if (it == cudaSampleItem + SAMPLE_LENGTH) {
     return 1;
   }
-  if (it == SampleItem) {
+  if (it == cudaSampleItem) {
     return 0;
   }
   auto it_prev = it - 1;
-  return (double(it_prev - SampleItem) +
+  return (double(it_prev - cudaSampleItem) +
           (x - (double)*it_prev) / (double)(*it - *it_prev)) /
          double(SAMPLE_LENGTH - 1);
 }
@@ -88,11 +88,12 @@ __global__ void initStorage(unsigned scale, unsigned test_size) {
 }
 
 __global__ void initCuda(key_type *sample_item, key_type *population_item) {
-  SampleItem = sample_item;
-  PopulationItem = population_item;
+  cudaSampleItem = sample_item;
+  cudaPopulationItem = population_item;
   /*for (int i = 0; i < SAMPLE_LENGTH; i++) {
   }*/
-  // cudaMalloc(&SampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long long));
+  // cudaMalloc(&cudaSampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long
+  // long));
   //  cudaMalloc(&Storage, TEST_SIZE * sizeof(key_type));
   cdf_result = nullptr;
 }
@@ -104,7 +105,7 @@ __global__ void kernel(unsigned long step, const double slice_size,
                threadIdx.x + i;
 
     // printf("%d\n", tid);
-    auto index = (int)(sample_cdf(PopulationItem[tid]) / slice_size);
+    auto index = (int)(sample_cdf(cudaPopulationItem[tid]) / slice_size);
     if (cdf_result[index]) {
       while (cdf_result[++index]) {
         assert(index < split_size);
diff --git a/expt_0809.cu b/expt_0809.cu
index 3f16a9a..44d9992 100644
--- a/expt_0809.cu
+++ b/expt_0809.cu
@@ -40,7 +40,7 @@ constexpr long SAMPLE_LENGTH = 1024;
 
 constexpr long TEST_SIZE = 1024;
 
-__device__ key_type *SampleItem, *PopulationItem;
+__device__ key_type *cudaSampleItem, *cudaPopulationItem;
 __device__ bool *cdf_result;
 __device__ unsigned insert_value;
 #ifdef TEST_BOUNDS
@@ -65,15 +65,16 @@ __device__ const key_type *cudaBinarySearch(key_type *start, key_type *end,
 }
 
 __device__ double sample_cdf(double x) {
-  auto it = cudaBinarySearch(SampleItem, SampleItem + SAMPLE_LENGTH + 2, x);
-  if (it == SampleItem + SAMPLE_LENGTH) {
+  auto it =
+      cudaBinarySearch(cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH + 2, x);
+  if (it == cudaSampleItem + SAMPLE_LENGTH) {
     return 1;
   }
-  if (it == SampleItem) {
+  if (it == cudaSampleItem) {
     return 0;
   }
   auto it_prev = it - 1;
-  return (double(it_prev - SampleItem) +
+  return (double(it_prev - cudaSampleItem) +
           (x - (double)*it_prev) / (double)(*it - *it_prev)) /
          double(SAMPLE_LENGTH - 1);
 }
@@ -91,9 +92,10 @@ __global__ void initStorage(unsigned scale, unsigned test_size) {
 }
 
 __global__ void init(key_type *sample_item, key_type *population_item) {
-  SampleItem = sample_item;
-  PopulationItem = population_item;
-  // cudaMalloc(&SampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long long));
+  cudaSampleItem = sample_item;
+  cudaPopulationItem = population_item;
+  // cudaMalloc(&cudaSampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long
+  // long));
   //  cudaMalloc(&Storage, TEST_SIZE * sizeof(key_type));
   cdf_result = nullptr;
 }
@@ -105,7 +107,7 @@ __global__ void kernel(unsigned long step, const double slice_size,
                threadIdx.x + i;
 
     // printf("%d\n", tid);
-    auto index = (int)(sample_cdf(PopulationItem[tid]) / slice_size);
+    auto index = (int)(sample_cdf(cudaPopulationItem[tid]) / slice_size);
     if (cdf_result[index]) {
       while (cdf_result[++index]) {
         assert(index < split_size);
diff --git a/sortlib.cuh b/sortlib.cuh
index f9bc7d7..ce40e0f 100644
--- a/sortlib.cuh
+++ b/sortlib.cuh
@@ -116,7 +116,7 @@ public:
     // printf("tmp: %llu %lf %lu\n", *it - *it_prev, tmp, prev_real_location);
   }
 
-  // First version
+  /// First version
   __device__ __host__ key_type *
   original_binary_search(key_type *start, const key_type *end, key_type &val) {
     auto begin = start;
@@ -164,4 +164,46 @@ public:
 __global__ void testCustomCalculation(size_t length) {
   CustomSort(length, sizeof(long) * 8).testCalculation();
 }
+
+class FactorySort {
+
+public:
+  __device__ __host__ static const key_type *
+  cudaBinarySearch(key_type *start, const key_type *end, const key_type val) {
+    auto begin = start;
+    key_type *last_known_point = begin;
+    assert(begin < end);
+    while (begin <= end) {
+      auto mid = begin + (end - begin) / 2;
+      auto mid_val = *mid;
+      if (val == mid_val) {
+        return mid;
+      } else if (val > mid_val) {
+        begin = mid + 1;
+      } else {
+        end = mid - 1;
+      }
+      last_known_point = begin;
+    }
+    return last_known_point;
+  }
+
+  __device__ __host__ double static sample_cdf(key_type *begin, long length,
+                                               key_type x) {
+    // printf("%f\n", x);
+    auto end = begin + length;
+    auto it = cudaBinarySearch(begin, end, x);
+    if (it == end) {
+      return 1;
+    }
+    if (it == begin) {
+      return 0;
+    }
+    auto it_prev = it - 1;
+    return (double(it_prev - begin) +
+            (x - (double)*it_prev) / (double)(*it - *it_prev)) /
+           double(length - 1);
+  }
+};
+
 #endif // LOCKFREE_SORTLIB_CUH