fix: Fix expt_0830 SIGSEGV

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

2 files changed, 58 insertions, 24 deletions

diff --git a/expt_0830.cu b/expt_0830.cu
index bbca9cf..845c44e 100644
--- a/expt_0830.cu
+++ b/expt_0830.cu
@@ -8,8 +8,9 @@
 std::vector<unsigned long long> population_vector, sample_vector;
 
 constexpr size_t SAMPLE_LENGTH = 1023;
-// constexpr size_t TEST_LENGTH = 16'384;
-constexpr size_t TEST_LENGTH = 2048;
+constexpr size_t TEST_LENGTH = 16'384;
+// constexpr size_t TEST_LENGTH = 4096;
+constexpr size_t RESERVED_BLOCK = 10'000;
 
 unsigned long long max_value = 0, min_value = 0xfffffffff;
 typedef std::pair<int, int> dim_pair_type;
@@ -46,17 +47,23 @@ __global__ void initSample(key_type *sample, key_type *population) {
 __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);
-  // printf("kernel1\n");
+  // 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, cudaPopulationItem[tid]) /
-                       slice_size);
+                           cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH,
+                           cudaPopulationItem[tid]) /
+                       slice_size) -
+                 1;
+    // printf("%llu %d\n", cudaPopulationItem[tid], index);
     if (cdf_result[index]) {
       while (cdf_result[++index]) {
-        assert(index < split_size);
+        /*if (index >= split_size) {
+          printf("escape: %llu %lu\n", cudaPopulationItem[tid], split_size);
+        }*/
+        assert(index < (split_size + RESERVED_BLOCK));
       }
     }
     cdf_result[index] = true;
@@ -127,8 +134,9 @@ __global__ void print_function() {
 
 __global__ void initStorage(unsigned scale, unsigned test_size) {
   cudaFree(cdf_result);
-  cudaMalloc(&cdf_result, scale * test_size * sizeof(bool));
-  memset(cdf_result, 0, scale * test_size * sizeof(bool));
+  // printf("test size: %u\n", test_size);
+  cudaMalloc(&cdf_result, (scale * test_size + RESERVED_BLOCK) * sizeof(bool));
+  memset(cdf_result, 0, (scale * test_size + RESERVED_BLOCK) * sizeof(bool));
 }
 
 __global__ void freeStorage() {
@@ -138,24 +146,24 @@ __global__ void freeStorage() {
 }
 
 __global__ void initCustomSample() {
+  // printf("init\n");
   key_type *tmp = nullptr;
   auto custom_sort = CustomSort(SAMPLE_LENGTH, 0);
   cudaMalloc(&tmp, sizeof(key_type) * SAMPLE_LENGTH);
   for (size_t i = 0; i < SAMPLE_LENGTH; i++) {
     tmp[i] = cudaSampleItem[custom_sort.calculate_index(i) - 1];
   }
+  // printf("copy\n");
   memcpy(cudaSampleItem, tmp, sizeof(key_type) * SAMPLE_LENGTH);
   cudaFree(tmp);
   tmp = nullptr;
   // memset(cdf_result, 0, sizeof(key_type) * TEST_LENGTH);
   cudaFree(cdf_result);
   cdf_result = nullptr;
+  // printf("finalize\n");
 }
 
-const dim_pair_type DIM_PAIR[] = {
-    dim_pair_type(16, 64),
-    // dim_pair_type(32, 32)
-};
+const dim_pair_type DIM_PAIR[] = {dim_pair_type(16, 64), dim_pair_type(32, 32)};
 
 void run_kernel(size_t test_size, bool custom = false) {
   for (auto &pair : DIM_PAIR) {
@@ -185,6 +193,9 @@ void run_kernel(size_t test_size, bool custom = false) {
       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>>>();
@@ -202,6 +213,7 @@ int main() {
     ;
   fclose(file);
 
+  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);
diff --git a/sortlib.cuh b/sortlib.cuh
index c90e408..1fdac9c 100644
--- a/sortlib.cuh
+++ b/sortlib.cuh
@@ -39,6 +39,7 @@ public:
   }
 
   __device__ __host__ const key_type *binary_search(key_type *const start,
+                                                    const key_type *end,
                                                     const key_type val) {
 
     int step_limit = (int)fast_log(LENGTH);
@@ -47,10 +48,21 @@ public:
 
     for (int i = 0; i < step_limit; i++) {
       const auto next_level_start = start + (1 << (i + 1)) - 1;
+
+      // printf("start: %ld, last: %ld\n", next_level_start - start,
+      //        last_known_point - start);
+
+      if (last_known_point > end) {
+        printf("%ld\n", last_known_point - start);
+      }
       if (*last_known_point == val) {
         return last_known_point;
       }
 
+      if (next_level_start > end) {
+        printf("%ld\n", next_level_start - start);
+      }
+
       // son = get_son_from_step(son, (*last_known_point > val));
       auto tmp = ((*last_known_point - val) >> MOVE_OFFSET);
       // printf("tmp: %llu\n", tmp);
@@ -62,32 +74,42 @@ public:
       /*printf("%d %d %d\n", (1 << (i + 1)), son,
              -(int)((*last_known_point - val) >> MOVE_OFFSET));*/
       last_known_point = next_level_start + son;
+      if (last_known_point > end) {
+        printf("%p %ld\n", end, next_level_start - start);
+      }
     }
     return last_known_point;
   }
 
-  __device__ __host__ double sample_cdf_custom_version(key_type *start,
-                                                       key_type x) {
-    auto it = this->binary_search(start, x);
+  __device__ __host__ double
+  sample_cdf_custom_version(key_type *start, const key_type *end, key_type x) {
+    // printf("custom version:\n");
+    auto it = this->binary_search(start, end, x);
+    // printf("search result: %ld\n", it - start);
     assert(it <= start + this->LENGTH);
-    if (it < start) {
+    /*if (it < start) {
       printf("it: %p, start: %p\n", it, start);
-    }
+    }*/
     assert(it >= start);
-    if (it == start + this->LENGTH) {
+
+    auto prev_real_location = calculate_rank(it - start) - 1;
+
+    if (prev_real_location == this->LENGTH) {
       return 1;
     }
-    if (it == start) {
+    if (prev_real_location == 0) {
       return 0;
     }
 
-    auto prev_real_location = calculate_rank(it - start) - 2;
+    // printf("cal rank: %lu\n", prev_real_location);
 
-    auto it_prev = start + calculate_index(prev_real_location) - 1;
+    auto it_prev = start + calculate_index(prev_real_location - 1) - 1;
 
-    return ((double)prev_real_location +
-            (double)(x - *it_prev) / (*it - *it_prev)) /
-           (double)(this->LENGTH - 1);
+    auto tmp = ((double)prev_real_location +
+                (double)(x - *it_prev) / (double)(*it - *it_prev)) /
+               (double)(this->LENGTH - 1);
+    // printf("tmp: %llu %lf %lu\n", *it - *it_prev, tmp, prev_real_location);
+    return tmp;
   }
 
   __device__ __host__ key_type *