diff options
| -rw-r--r-- | CMakeLists.txt | 10 | ||||
| -rw-r--r-- | expt_0821.cpp | 14 | ||||
| -rw-r--r-- | expt_0830.cu | 235 | ||||
| -rw-r--r-- | sortlib.cuh | 139 | ||||
| -rw-r--r-- | sortlib.h | 26 |
5 files changed, 406 insertions, 18 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index e05ef87..bdca176 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -131,4 +131,12 @@ add_executable(valid_sort valid_sort.cpp) target_link_libraries(valid_sort m stdc++) add_executable(expt_0821 expt_0821.cpp sortlib.h) -target_link_libraries(expt_0821 m stdc++)
\ No newline at end of file +target_link_libraries(expt_0821 m stdc++) + +add_executable(expt_0830 expt_0830.cu sortlib.cuh) +target_link_libraries(expt_0830 m stdc++) + +set_target_properties(expt_0830 PROPERTIES + CUDA_SEPARABLE_COMPILATION ON) +set_target_properties(expt_0830 PROPERTIES CUDA_ARCHITECTURES "75") +set_target_properties(expt_0830 PROPERTIES LINKER_LANGUAGE CUDA)
\ No newline at end of file diff --git a/expt_0821.cpp b/expt_0821.cpp index 330cf40..cd0fd07 100644 --- a/expt_0821.cpp +++ b/expt_0821.cpp @@ -24,7 +24,7 @@ inline void store_into_vector(unsigned long long value) { } void valid_sort(std::vector<key_type> &tmp) { - auto iter_end = population_vector.begin() + SAMPLE_LENGTH + 5; + auto iter_end = population_vector.begin() + SAMPLE_LENGTH + 1024; for (auto it = population_vector.begin() + SAMPLE_LENGTH; it != iter_end; it++) { // long std_search = @@ -33,15 +33,19 @@ void valid_sort(std::vector<key_type> &tmp) { tmp.data(), tmp.data() + tmp.size(), *it) - tmp.data(); - auto cuda_binary_search = + auto cuda_binary_search1 = custom_sort.binary_search(sample_vector.data(), *it) - sample_vector.data(); - cuda_binary_search = custom_sort.calculate_location(cuda_binary_search) - 1; + auto cuda_binary_search = + custom_sort.calculate_location(cuda_binary_search1) - 1; /*printf("%ld:%ld %llu:%llu\n", std_search, cuda_binary_search, *(tmp.begin() + std_search), *(cuda_binary_search + sample_vector.begin()));*/ - printf("%ld:%ld \n", std_search, cuda_binary_search); - // assert(std_search == cuda_binary_search); + printf("%ld:%ld %llu %llu %llu\n", std_search, cuda_binary_search, *it, + *(tmp.data() + std_search), + *(cuda_binary_search1 + sample_vector.data())); + // assert(*it > *(cuda_binary_search1 + sample_vector.data())); + // assert(std_search == cuda_binary_search); } } diff --git a/expt_0830.cu b/expt_0830.cu new file mode 100644 index 0000000..bbca9cf --- /dev/null +++ b/expt_0830.cu @@ -0,0 +1,235 @@ +#include "sortlib.cuh" + +#include <algorithm> +#include <cassert> +#include <cstdio> +#include <vector> + +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; + +unsigned long long max_value = 0, min_value = 0xfffffffff; +typedef std::pair<int, int> dim_pair_type; + +inline void store_into_vector(unsigned long long value) { + if (max_value < value) { + max_value = value; + } + if (min_value > value) { + min_value = value; + } + population_vector.push_back(value); +} +// #define TEST_BOUNDS + +__device__ key_type *cudaSampleItem, *cudaPopulationItem; +__device__ bool *cdf_result; +#ifdef TEST_BOUNDS +__device__ unsigned insert_value; +__device__ unsigned int index_max, index_min; +#endif + +__global__ void initSample(key_type *sample, key_type *population) { + cudaSampleItem = sample; + cudaPopulationItem = population; + // cudaMalloc(&cdf_result, sizeof(bool) * TEST_LENGTH); + cdf_result = nullptr; +#ifdef TEST_BOUNDS + index_max = 0; + index_min = 0x7fffffff; +#endif +} + +__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"); + 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); + if (cdf_result[index]) { + while (cdf_result[++index]) { + assert(index < split_size); + } + } + cdf_result[index] = true; + } +} + +__global__ void kernel2_real_binary(unsigned long step, const double slice_size, + const unsigned long split_size) { + auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(key_type) * 8); + // printf("%d\n", custom_sort.MOVE_OFFSET); + // printf("kernel2\n"); + + for (int i = 0; i < step; i++) { + auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i; + // printf("%lu ", tid); + + // assert(tid < TEST_LENGTH); + auto index = (int)(custom_sort.sample_cdf(cudaSampleItem, + cudaSampleItem + SAMPLE_LENGTH, + cudaPopulationItem[tid]) / + slice_size); + // printf("%llu %d\n", cudaPopulationItem[tid], index); + if (cdf_result[index]) { + while (cdf_result[++index]) { + assert(index < split_size); + } + } + cdf_result[index] = true; +#ifdef TEST_BOUNDS + atomicAdd(&insert_value, 1); + + while (true) { + unsigned tmp = index_min; + if (tmp < tid) { + break; + } + if (atomicCAS(&index_min, tmp, tid) == tmp) { + break; + } + } + while (true) { + unsigned tmp = index_max; + if (tmp > tid) { + break; + } + if (atomicCAS(&index_max, tmp, tid) == tmp) { + break; + } + } +#endif + } +} + +__global__ void print2() { + /*for (int i = 0; i< SAMPLE_LENGTH; i++) { + printf("%llu ", cudaSampleItem[i]); + } + printf("\n");*/ + printf("%lu\n", CustomSort::fast_log(SAMPLE_LENGTH)); +} + +__global__ void print_function() { +#ifdef TEST_BOUNDS + printf("%u\n", insert_value); + printf("%u %u\n", index_min, index_max); +#endif +} + +__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)); +} + +__global__ void freeStorage() { + cudaFree(cudaSampleItem); + cudaFree(cudaPopulationItem); + cudaFree(cdf_result); +} + +__global__ void initCustomSample() { + 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]; + } + 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; +} + +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) { + dim3 grid_dim = pair.first, block_dim = pair.second; + unsigned long step = test_size / (grid_dim.x * block_dim.x); + printf("%scurrent dim: %d %d %lu\n", custom ? "custom " : "", pair.first, + pair.second, step); + for (int scale = 2; scale <= 16; scale += 2) { + const unsigned long split_size = test_size * scale; + const auto slice_size = 1.0 / (double)split_size; + printf("scale: %i ", scale); + initStorage<<<1, 1>>>(scale, test_size); + cudaDeviceSynchronize(); + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + cudaEventRecord(start, nullptr); + if (custom) + kernel<<<grid_dim, block_dim>>>(step, slice_size, split_size); + else + kernel2_real_binary<<<grid_dim, block_dim>>>(step, slice_size, + split_size); + cudaDeviceSynchronize(); + cudaEventRecord(stop, nullptr); + cudaEventSynchronize(stop); + float time; + cudaEventElapsedTime(&time, start, stop); + cudaEventDestroy(start); + cudaEventDestroy(stop); + printf("time: %lf\n", time); + + print_function<<<1, 1>>>(); + cudaDeviceSynchronize(); + } + } +} + +int main() { + + auto test_size = TEST_LENGTH; + FILE *file = fopen("normal_distribution.txt", "r"); + assert(file); + for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i)) + ; + fclose(file); + + sample_vector = std::vector<key_type>( + population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH - 2); + sample_vector.push_back(min_value); + sample_vector.push_back(max_value); + std::sort(sample_vector.begin(), sample_vector.end()); + + key_type *cudaSample = nullptr, *cudaPopulation; + cudaMalloc(&cudaSample, sizeof(key_type) * SAMPLE_LENGTH); + assert(SAMPLE_LENGTH == sample_vector.size()); + cudaMemcpy(cudaSample, sample_vector.data(), sizeof(key_type) * SAMPLE_LENGTH, + cudaMemcpyHostToDevice); + cudaMalloc(&cudaPopulation, sizeof(key_type) * test_size); + cudaMemcpy(cudaPopulation, population_vector.data() + SAMPLE_LENGTH, + sizeof(key_type) * test_size, cudaMemcpyHostToDevice); + + // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8); + // custom_sort.testCalculation(); + testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH); + cudaDeviceSynchronize(); + + initSample<<<1, 1>>>(cudaSample, cudaPopulation); + cudaDeviceSynchronize(); + run_kernel(test_size); + + initCustomSample<<<1, 1>>>(); + cudaDeviceSynchronize(); + + run_kernel(test_size, true); + freeStorage<<<1, 1>>>(); + cudaDeviceSynchronize(); +}
\ No newline at end of file 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 @@ -3,7 +3,7 @@ #include <cstddef> #include <cstdio> -template <typename C> class CustomSort { +template <typename> class CustomSort { public: CustomSort(size_t length, int move_offset) : LENGTH(length), MOVE_OFFSET(move_offset) {} @@ -45,11 +45,12 @@ public: } // son = get_son_from_step(son, (*last_known_point > val)); - auto tmp = (int)((*last_known_point - val) >> MOVE_OFFSET); - printf("%llu %llu ", val, *last_known_point); + auto tmp = ((*last_known_point - val) >> MOVE_OFFSET); + // printf("%llu %llu ", val, *last_known_point); son = son * 2 + tmp; - puts(tmp == 0 ? "1:left" : "1:right"); - // if (son < 0) son = 0; + // 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; @@ -65,16 +66,17 @@ public: if (val == mid_val) { return begin + mid; } - printf("%llu %llu ", val, mid_val); - puts(val > mid_val ? "right" : "left"); - /*begin += (int)((mid_val - val) >> MOVE_OFFSET) & (mid + 1); - end -= (int)(~((mid_val - val) >> MOVE_OFFSET)) & (mid + 1);*/ - if (val > mid_val) { + 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); + /*if (val > mid_val) { begin = begin + mid + 1; } else { end = end - mid - 1; - } - last_known_point = begin; + }*/ } return last_known_point; } |