// Experimental content: Test CustomSort Calculation #define DISABLE_TEST_WARNING #include "sortlib.cuh" #include #include #ifndef LOCKFREE_SORTLIB_CUH #include class CustomSort { public: explicit CustomSort(size_t length) : LENGTH(length) {} const size_t LENGTH; __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; } __device__ __host__ 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); } } static void testSelf(size_t length) { CustomSort(length).testCalculation(); } }; __global__ void testCustomCalculationWithPrint(size_t length) { CustomSort(length).testCalculation(); printf("%lu pass\n", length); } #endif long pow_for_sample(long n) { auto x = 2; for (int i = 1; i < n; i++) { x *= 2; } return x - 1; } int main(int argc, char const *argv[]) { auto max = 32L; auto cuda_only = false; if (argc > 1) { max = strtol(argv[1], nullptr, 10); } if (argc > 2) { cuda_only = strtol(argv[2], nullptr, 10); } for (int i = 4; i <= max && !cuda_only; i++) { CustomSort::testSelf(pow_for_sample(i)); printf("%d pass\n", i); } for (int i = 4; i <= max; i++) { testCustomCalculationWithPrint<<<1, 1>>>(pow_for_sample(i)); cudaDeviceSynchronize(); } }