#ifndef LOCKFREE_SORTLIB_CUH #define LOCKFREE_SORTLIB_CUH #include #include 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