#ifndef LOCKFREE_SORTLIB_CUH #define LOCKFREE_SORTLIB_CUH #include #include #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), STEP_LIMIT(fast_log(LENGTH)) {} const size_t LENGTH; /// MOVE_OFFSET means bit to select branch const int MOVE_OFFSET; const unsigned int STEP_LIMIT; __device__ __host__ static size_t fast_log(size_t a) { #ifdef __CUDA_ARCH__ return (size_t)log2((double)a); #else return (size_t)std::log2(a); #endif } __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; } #ifdef ENABLE_SORT_TEST __device__ __host__ __attribute__((unused)) void testCalculation() const { for (size_t i = 0; i < LENGTH; i++) { auto left = calculate_index(i); auto right = calculate_rank(left - 1); /*if (i + 1 == LENGTH || i + 1 != right) { printf("%lu %lu %lu\n", i, left, right); printf("%lu %lu %lu\n", LENGTH, fast_log(i + 1), (LENGTH + 1) >> fast_log(i + 1) >> 1); }*/ assert(i + 1 == right); } } #endif __device__ __host__ const key_type *binary_search(key_type *const start, // const key_type *end, const key_type val) const { // int step_limit = (int)fast_log(LENGTH); // printf("Binary search\n"); key_type *last_known_point = start; auto son = 0UL; 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 branch_selector = ((*last_known_point - val) >> MOVE_OFFSET); // printf("tmp: %llu\n", tmp); // printf("%llu %llu ", val, *last_known_point); son = son * 2 + branch_selector; // 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) >> brenchSelector));*/ 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; } /// Should be correct version __device__ __host__ double cdf(key_type *start, // const key_type *end, key_type x) const { // printf("custom version:\n"); auto it = this->binary_search(start, x); // printf("search result: %ld\n", it - start); // assert(it <= start + this->LENGTH); /*if (it < start) { printf("it: %p, start: %p\n", it, start); }*/ // assert(it >= start); auto prev_real_location = calculate_rank(it - start) - 1; if (prev_real_location == this->LENGTH) { return 1; } if (prev_real_location == 0) { return 0; } // printf("cal rank: %lu\n", prev_real_location); auto it_prev = start + calculate_index(prev_real_location - 1) - 1; return ((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); } #ifdef SORT_FIRST_VERSION /// First version __device__ __host__ key_type *original_binary_search(key_type *start, const key_type *end, key_type &val) const { 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; } /// CDF original version (should only work on default data layout) __device__ __host__ double sample_cdf(key_type *start, key_type *end, key_type x) const { 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); } #endif #ifdef ENABLE_SORT_TEST static void testSelf(size_t length) { CustomSort(length, sizeof(long) * 8).testCalculation(); #ifndef DISABLE_TEST_WARNING puts("You can remove this function if passed already"); #endif } #endif __host__ __device__ size_t length() const { return this->LENGTH; } }; template void rebuild(std::vector &original) { auto sample_length = original.size(); auto sorter = CustomSort(sample_length, sizeof(T) * 8); auto tmp = new T[sample_length]; for (size_t i = 0; i < sample_length; i++) { tmp[i] = original[sorter.calculate_index(i) - 1]; } memcpy(original.data(), tmp, sizeof(T) * sample_length); delete[] tmp; } template void rebuildSort(std::vector &original) { std::sort(original.begin(), original.end()); rebuild(original); } #ifdef ENABLE_SORT_TEST __global__ void testCustomCalculation(size_t length) { CustomSort(length, sizeof(long) * 8).testCalculation(); } __global__ void testCustomCalculationWithPrint(size_t length) { CustomSort(length, sizeof(long) * 8).testCalculation(); printf("%lu pass\n", length); } #endif class FactorySort { public: __device__ __host__ static const key_type * binarySearch(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 cdf(key_type *begin, unsigned long length, key_type x) { // printf("%f\n", x); auto end = begin + length; auto it = binarySearch(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