diff options
| -rw-r--r-- | expt_0802.cpp | 15 | ||||
| -rw-r--r-- | expt_0804.cu | 39 |
2 files changed, 24 insertions, 30 deletions
diff --git a/expt_0802.cpp b/expt_0802.cpp index fbd6c23..b1a906a 100644 --- a/expt_0802.cpp +++ b/expt_0802.cpp @@ -25,8 +25,6 @@ size_t calculate_location(size_t index) { } typedef int key_type; -inline double safeStep(double step) { return step < 0 ? 0 : step; } - int get_son_from_step(int point, bool negative) { return negative ? point * 2 : point * 2 + 1; } @@ -34,31 +32,18 @@ int get_son_from_step(int point, bool negative) { const key_type *cudaBinarySearch(key_type *const start, const key_type *end, const key_type val) { int step_limit = (int)std::log2(LENGTH) + 1; - const auto length = (end - start); 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; - /*const auto current_time = (1.0 / (1 << (i + 1))); - printf("%d ", *level_start); - printf("%d %d\n", *last_known_point, (1 << i)); if (*last_known_point == val) { - puts("find"); - return last_known_point; - } - sum = safeStep(sum + - ((*last_known_point > val) ? -current_time : current_time)); - last_known_point = level_start + (int)(sum * (double)(1 << i));*/ - if (*last_known_point == val) { - puts("found"); return last_known_point; } son = get_son_from_step(son, (*last_known_point > val)); last_known_point = next_level_start + son; printf("%d\n", *last_known_point); } - puts(""); return last_known_point; } diff --git a/expt_0804.cu b/expt_0804.cu index 47c269e..3a599ef 100644 --- a/expt_0804.cu +++ b/expt_0804.cu @@ -36,21 +36,23 @@ __device__ unsigned insert_value; __device__ unsigned int index_max, index_min; #endif -__device__ const key_type *cudaBinarySearch(key_type *start, key_type *end, +inline __device__ int get_son_from_step(int point, bool negative) { + return negative ? point * 2 : point * 2 + 1; +} + +__device__ const key_type *cudaBinarySearch(key_type *start, const key_type val) { - auto begin = start; - key_type *last_known_point = nullptr; - while (begin < end) { - auto mid = (end - begin) / 2; - auto mid_val = *(start + mid); - if (val == mid_val) { - return start + mid; - } else if (val > mid_val) { - begin = begin + mid + 1; - } else { - end = end - mid - 1; + int step_limit = (int)log2f(SAMPLE_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; } - last_known_point = begin; + son = get_son_from_step(son, (*last_known_point > val)); + last_known_point = next_level_start + son; } return last_known_point; } @@ -63,7 +65,7 @@ __device__ __host__ size_t calculate_location(size_t index) { } __device__ double sample_cdf(double x) { - auto it = cudaBinarySearch(SampleItem, SampleItem + SAMPLE_LENGTH + 2, x); + auto it = cudaBinarySearch(SampleItem, x); if (it == SampleItem + SAMPLE_LENGTH) { return 1; } @@ -91,6 +93,8 @@ __global__ void initStorage(unsigned scale, unsigned test_size) { __global__ void initCuda(key_type *sample_item, key_type *population_item) { SampleItem = sample_item; PopulationItem = population_item; + /*for (int i = 0; i < SAMPLE_LENGTH; i++) { + }*/ // cudaMalloc(&SampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long long)); // cudaMalloc(&Storage, TEST_SIZE * sizeof(key_type)); cdf_result = nullptr; @@ -169,10 +173,15 @@ int main(int argc, char const *argv[]) { std::sort(sample_vector.begin(), sample_vector.end()); + sample_vector_into_cuda.resize(SAMPLE_LENGTH); + for (int i = 0; i < SAMPLE_LENGTH; i++) { + sample_vector_into_cuda[i] = sample_vector[calculate_location(i) - 1]; + } + key_type *cudaSample = nullptr, *cudaPopulation = nullptr; cudaMalloc(&cudaSample, sizeof(key_type) * (SAMPLE_LENGTH)); cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE); - cudaMemcpy(cudaSample, &sample_vector[0], + cudaMemcpy(cudaSample, &sample_vector_into_cuda[0], sizeof(key_type) * sample_vector.size(), cudaMemcpyHostToDevice); cudaMemcpy(cudaPopulation, &population_vector[SAMPLE_LENGTH], sizeof(key_type) * TEST_SIZE, cudaMemcpyHostToDevice); |