1 files changed, 139 insertions, 0 deletions
diff --git a/sortlib.cuh b/sortlib.cuh
new file mode 100644
index 0000000..c90e408
--- /dev/null
+++ b/sortlib.cuh
@@ -0,0 +1,139 @@
+#ifndef LOCKFREE_SORTLIB_CUH
+#define LOCKFREE_SORTLIB_CUH
+
+#include <cassert>
+#include <cstdio>
+
+typedef unsigned long long key_type;
+
+class CustomSort {
+public:
+  explicit __device__ __host__ CustomSort(size_t length, int move_offset)
+      : LENGTH(length), MOVE_OFFSET(move_offset - 1) {}
+  const size_t LENGTH;
+  const int MOVE_OFFSET;
+
+  __device__ __host__ static size_t fast_log(size_t a) {
+    float t = a;
+    return (((*(int *)&t) >> 23) + 1) & 127;
+  }
+
+  __device__ __host__ size_t calculate_index(size_t rank) const {
+    size_t bit_low = (LENGTH + 1) >> fast_log(++rank) >> 1;
+    return (((rank << 1) | 1) * bit_low - LENGTH - 1);
+  }
+
+  __device__ __host__ size_t calculate_rank(size_t index) const {
+    index++;
+    size_t low_bit = index & (-index);
+    return ((LENGTH + index) / low_bit) >> 1;
+  }
+
+  __device__ __host__ __attribute__((unused)) void testCalculation() const {
+    for (size_t i = 0; i < LENGTH; i++) {
+      auto l = calculate_index(i);
+      auto l2 = calculate_rank(l - 1);
+      // printf("%lu %lu\n", l, l2);
+      assert(i + 1 == l2);
+    }
+  }
+
+  __device__ __host__ const key_type *binary_search(key_type *const start,
+                                                    const key_type val) {
+
+    int step_limit = (int)fast_log(LENGTH);
+    key_type *last_known_point = start;
+    auto son = 0;
+
+    for (int i = 0; i < step_limit; i++) {
+      const auto next_level_start = start + (1 << (i + 1)) - 1;
+      if (*last_known_point == val) {
+        return last_known_point;
+      }
+
+      // son = get_son_from_step(son, (*last_known_point > val));
+      auto tmp = ((*last_known_point - val) >> MOVE_OFFSET);
+      // printf("tmp: %llu\n", tmp);
+      //    printf("%llu %llu ", val, *last_known_point);
+      son = son * 2 + tmp;
+      // printf("%d\n", son);
+      // puts(tmp == 0 ? "1:left" : "1:right");
+      //   if (son < 0) son = 0;
+      /*printf("%d %d %d\n", (1 << (i + 1)), son,
+             -(int)((*last_known_point - val) >> MOVE_OFFSET));*/
+      last_known_point = next_level_start + son;
+    }
+    return last_known_point;
+  }
+
+  __device__ __host__ double sample_cdf_custom_version(key_type *start,
+                                                       key_type x) {
+    auto it = this->binary_search(start, x);
+    assert(it <= start + this->LENGTH);
+    if (it < start) {
+      printf("it: %p, start: %p\n", it, start);
+    }
+    assert(it >= start);
+    if (it == start + this->LENGTH) {
+      return 1;
+    }
+    if (it == start) {
+      return 0;
+    }
+
+    auto prev_real_location = calculate_rank(it - start) - 2;
+
+    auto it_prev = start + calculate_index(prev_real_location) - 1;
+
+    return ((double)prev_real_location +
+            (double)(x - *it_prev) / (*it - *it_prev)) /
+           (double)(this->LENGTH - 1);
+  }
+
+  __device__ __host__ key_type *
+  original_binary_search(key_type *start, const key_type *end, key_type &val) {
+    auto begin = start;
+    key_type *last_known_point = nullptr;
+    while (begin < end) {
+      auto mid = (end - begin) / 2;
+      auto mid_val = *(begin + mid);
+      if (val == mid_val) {
+        return begin + mid;
+      }
+      last_known_point = begin;
+      // printf("%llu %llu ", val, mid_val);
+      // puts(val > mid_val ? "right" : "left");
+      // printf("%llu\n", ((val - mid_val) >> MOVE_OFFSET) * (mid + 1));
+      begin += ((mid_val - val) >> MOVE_OFFSET) * (mid + 1);
+      end -= ((val - mid_val) >> MOVE_OFFSET) * (mid + 1);
+      // printf("%llu\t%p %p\n", val, begin, end);
+      /*if (val > mid_val) {
+        begin = begin + mid + 1;
+      } else {
+        end = end - mid - 1;
+      }*/
+    }
+    return last_known_point;
+  }
+
+  __device__ __host__ double sample_cdf(key_type *start, key_type *end,
+                                        key_type x) {
+    auto it = this->original_binary_search(start, end, x);
+    if (it == end) {
+      return 1;
+    }
+    if (it == start) {
+      return 0;
+    }
+    auto it_prev = it - 1;
+
+    return ((double)(it_prev - start) +
+            (double)(x - *it_prev) / (double)(*it - *it_prev)) /
+           (double)((end - start) - 1);
+  }
+};
+
+__global__ void testCustomCalculation(size_t length) {
+  CustomSort(length, sizeof(long) * 8).testCalculation();
+}
+#endif // LOCKFREE_SORTLIB_CUH