diff options
| -rw-r--r-- | expt_0830.cu | 16 | ||||
| -rw-r--r-- | sortlib.cuh | 20 |
2 files changed, 18 insertions, 18 deletions
diff --git a/expt_0830.cu b/expt_0830.cu index 845c44e..e48d543 100644 --- a/expt_0830.cu +++ b/expt_0830.cu @@ -8,7 +8,7 @@ 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 = 32'768; // constexpr size_t TEST_LENGTH = 4096; constexpr size_t RESERVED_BLOCK = 10'000; @@ -58,14 +58,14 @@ __global__ void kernel(unsigned long step, const double slice_size, slice_size) - 1; // printf("%llu %d\n", cudaPopulationItem[tid], index); - if (cdf_result[index]) { + /*if (cdf_result[index]) { while (cdf_result[++index]) { - /*if (index >= split_size) { - printf("escape: %llu %lu\n", cudaPopulationItem[tid], split_size); - }*/ + if (index >= split_size) { + //printf("escape: %llu %lu\n", cudaPopulationItem[tid], split_size); + } assert(index < (split_size + RESERVED_BLOCK)); } - } + }*/ cdf_result[index] = true; } } @@ -86,11 +86,11 @@ __global__ void kernel2_real_binary(unsigned long step, const double slice_size, cudaPopulationItem[tid]) / slice_size); // printf("%llu %d\n", cudaPopulationItem[tid], index); - if (cdf_result[index]) { + /*if (cdf_result[index]) { while (cdf_result[++index]) { assert(index < split_size); } - } + }*/ cdf_result[index] = true; #ifdef TEST_BOUNDS atomicAdd(&insert_value, 1); diff --git a/sortlib.cuh b/sortlib.cuh index 1fdac9c..0aba73b 100644 --- a/sortlib.cuh +++ b/sortlib.cuh @@ -9,9 +9,10 @@ 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) {} + : LENGTH(length), MOVE_OFFSET(move_offset - 1), + STEP_LIMIT(fast_log(LENGTH)) {} const size_t LENGTH; - const int MOVE_OFFSET; + const int MOVE_OFFSET, STEP_LIMIT; __device__ __host__ static size_t fast_log(size_t a) { float t = a; @@ -42,11 +43,11 @@ public: const key_type *end, const key_type val) { - int step_limit = (int)fast_log(LENGTH); + // int step_limit = (int)fast_log(LENGTH); key_type *last_known_point = start; auto son = 0; - for (int i = 0; i < step_limit; i++) { + 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, @@ -86,11 +87,11 @@ public: // printf("custom version:\n"); auto it = this->binary_search(start, end, x); // printf("search result: %ld\n", it - start); - assert(it <= start + this->LENGTH); + // assert(it <= start + this->LENGTH); /*if (it < start) { printf("it: %p, start: %p\n", it, start); }*/ - assert(it >= start); + // assert(it >= start); auto prev_real_location = calculate_rank(it - start) - 1; @@ -105,11 +106,10 @@ public: auto it_prev = start + calculate_index(prev_real_location - 1) - 1; - auto tmp = ((double)prev_real_location + - (double)(x - *it_prev) / (double)(*it - *it_prev)) / - (double)(this->LENGTH - 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); - return tmp; } __device__ __host__ key_type * |