From 0e537b32ae8f80985d3fd6b898ad6151cd8ac86c Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Tue, 2 Aug 2022 18:45:00 +0800 Subject: feat: Implement expt_0729 Signed-off-by: KunoiSayami --- CMakeLists.txt | 9 +++ expt_0726.cpp | 7 +- expt_0729.cu | 198 ++++++++++++++++++++++++++++++++++++++++++++++++ normal_distribution.cpp | 16 +++- 4 files changed, 224 insertions(+), 6 deletions(-) create mode 100644 expt_0729.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index 858cea0..3bd9244 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -84,3 +84,12 @@ set_target_properties(expt_0722 PROPERTIES LINKER_LANGUAGE CXX) add_executable(expt_0726 expt_0726.cpp) set_target_properties(expt_0726 PROPERTIES LINKER_LANGUAGE CXX) +add_executable(expt_0729 expt_0729.cu) +#set_target_properties(expt_0729 PROPERTIES LINKER_LANGUAGE CXX) +target_link_libraries(expt_0729 m stdc++) + +set_target_properties(expt_0729 PROPERTIES + CUDA_SEPARABLE_COMPILATION ON) +set_target_properties(expt_0729 PROPERTIES CUDA_ARCHITECTURES "75") +set_target_properties(expt_0726 PROPERTIES LINKER_LANGUAGE CUDA) + diff --git a/expt_0726.cpp b/expt_0726.cpp index 5b37df3..1043044 100644 --- a/expt_0726.cpp +++ b/expt_0726.cpp @@ -44,14 +44,14 @@ inline void store_into_vector(unsigned long long value) { } constexpr long sample_length = 1024; -constexpr long test_size = 2048; +constexpr long test_size = 1000000; void mian(long scale_size) { const long split_size = test_size * scale_size; result_storage.clear(); result_storage.resize(split_size, false); - const auto slice_size = 1.0 / (long double)(split_size); + const auto slice_size = 1.0 / (double)(split_size); for (long i = 0; i < test_size; i++) { auto index = (int)safe_ceil(sample_cdf(original_vector[i + sample_length]) / @@ -102,6 +102,7 @@ int main(int _argc, char const *_argv[]) { } assert(element_size_max == 1); printf("scale: %d, time spend: %ldms\n", i, - duration_cast((end - start)).count()); + std::chrono::duration_cast(end - start) + .count()); } } \ No newline at end of file diff --git a/expt_0729.cu b/expt_0729.cu new file mode 100644 index 0000000..4e28281 --- /dev/null +++ b/expt_0729.cu @@ -0,0 +1,198 @@ +#include +#include +#include +#include +#include +#include + +constexpr size_t length = 2097152; +std::vector population_vector, sample_vector; + +unsigned long long max_value = 0, min_value = 0xfffffffff; + +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); +} + +typedef unsigned long long key_type; + +typedef unsigned long long *key_type_ptr; + +constexpr long SAMPLE_LENGTH = 1024; + +constexpr long TEST_SIZE = 1024; + +__device__ key_type *SampleItem, *PopulationItem; +__device__ bool *cdf_result; +__device__ unsigned insert_value; +#ifdef TEST_BOUNDS +__device__ unsigned int index_max, index_min; +#endif + +__device__ const key_type *cudaBinarySearch(key_type *start, key_type *end, + 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; + } + last_known_point = begin; + } + return last_known_point; +} + +__device__ double sample_cdf(double x) { + auto it = cudaBinarySearch(SampleItem, SampleItem + SAMPLE_LENGTH, x); + if (it == SampleItem + SAMPLE_LENGTH) { + return 1; + } + if (it == SampleItem) { + return 0; + } + auto it_prev = it - 1; + return (double(it_prev - SampleItem) + + (x - (double)*it_prev) / (double)(*it - *it_prev)) / + double(SAMPLE_LENGTH - 1); +} + +__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)); + insert_value = 0; + printf("init storage\n"); +#ifdef TEST_BOUNDS + index_max = 0; + index_min = 0x7fffffff; +#endif +} + +__global__ void init(key_type *sample_item, key_type *population_item) { + SampleItem = sample_item; + PopulationItem = population_item; + // cudaMalloc(&SampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long long)); + // cudaMalloc(&Storage, TEST_SIZE * sizeof(key_type)); + cdf_result = nullptr; +} + +__global__ void kernel(unsigned long step, const double slice_size, + const long split_size) { + for (int i = 1; i <= step; i++) { + auto tid = i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + + threadIdx.x - 1024; + + // printf("%d\n", tid); + auto index = (int)(sample_cdf(PopulationItem[tid - 1024]) / slice_size); + if (cdf_result[index]) { + while (cdf_result[++index]) { + assert(index < split_size); + } + } + cdf_result[index] = true; + atomicAdd(&insert_value, 1); + +#ifdef TEST_BOUNDS + 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 print_function() { + printf("%u\n", insert_value); +#ifdef TEST_BOUNDS + printf("%u %u\n", index_min, index_max); +#endif +} + +__device__ inline long double safe_ceil(double value) { + auto c = ceil(value); + return c == 0 ? 1 : c; +} + +int main(int argc, char const *argv[]) { + size_t test_size; + if (argc == 1) { + test_size = 16384; + } else { + try { + test_size = std::stol(argv[1]); + } catch (...) { + return 1; + } + } + + FILE *file = fopen("normal_distribution.txt", "r"); + assert(file); + for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i)) + ; + fclose(file); + + assert(population_vector.size() == length); + + sample_vector = std::vector( + population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH); + 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 = nullptr; + cudaMalloc(&cudaSample, sizeof(key_type) * (SAMPLE_LENGTH + 2)); + cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE); + cudaMemcpy(cudaSample, &sample_vector[0], + sizeof(key_type) * sample_vector.size(), cudaMemcpyHostToDevice); + cudaMemcpy(cudaPopulation, &population_vector[SAMPLE_LENGTH], + sizeof(key_type) * TEST_SIZE, cudaMemcpyHostToDevice); + // memcpy(sample_heap, sample_vector.cbegin(),sizeof(key_type) *(SAMPLE_LENGTH + // + 2)); + + init<<<1, 1>>>(cudaSample, cudaPopulation); + cudaDeviceSynchronize(); + + unsigned long step = test_size / 1024; + printf("step: %lu\n", step); + assert(!(test_size % 1024)); + + for (int scale = 2; scale <= 16; scale++) { + const long split_size = test_size * scale; + const auto slice_size = 1.0 / (double)split_size; + printf("scale: %i\n", scale); + initStorage<<<1, 1>>>(scale, test_size); + cudaDeviceSynchronize(); + kernel<<<2, 512>>>(step, slice_size, split_size); + cudaDeviceSynchronize(); + print_function<<<1, 1>>>(); + cudaDeviceSynchronize(); + } + return 0; +} \ No newline at end of file diff --git a/normal_distribution.cpp b/normal_distribution.cpp index 002beba..72f914c 100644 --- a/normal_distribution.cpp +++ b/normal_distribution.cpp @@ -5,15 +5,25 @@ #include #include -constexpr size_t length = 1048576; - -int main() { +int main(int argc, char const *argv[]) { std::random_device randomDevice; std::mt19937 randomEngine(randomDevice()); std::normal_distribution normalDistribution(2147483648, 2147483648); + size_t length; + if (argc == 1) { + length = 1048576; + } else { + try { + length = std::stol(argv[1]); + } catch (...) { + fprintf(stderr, "Wrong arguments\n"); + return 1; + } + } + std::set set; std::vector vector; -- cgit v1.3.1