From a0476481fdd1e0e97f9ca888d494257a83ba56cc Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Thu, 11 May 2023 01:47:05 +0800 Subject: feat(exp): Add new experimental for save memory and test Signed-off-by: KunoiSayami --- CMakeLists.txt | 11 ++- expt_0510.cu | 254 ++++++++++++++++++++++++++++++++++++++++++++++++ normal_distribution.cpp | 3 +- 3 files changed, 266 insertions(+), 2 deletions(-) create mode 100644 expt_0510.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index 6389a30..4f66661 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -185,4 +185,13 @@ target_link_libraries(expt_0503 m stdc++) set_target_properties(expt_0503 PROPERTIES CUDA_SEPARABLE_COMPILATION ON) set_target_properties(expt_0503 PROPERTIES CUDA_ARCHITECTURES "75") -set_target_properties(expt_0503 PROPERTIES LINKER_LANGUAGE CUDA) \ No newline at end of file +set_target_properties(expt_0503 PROPERTIES LINKER_LANGUAGE CUDA) + + +add_executable(expt_0510 expt_0510.cu) +target_link_libraries(expt_0510 m stdc++) + +set_target_properties(expt_0510 PROPERTIES + CUDA_SEPARABLE_COMPILATION ON) +set_target_properties(expt_0510 PROPERTIES CUDA_ARCHITECTURES "75") +set_target_properties(expt_0510 PROPERTIES LINKER_LANGUAGE CUDA) \ No newline at end of file diff --git a/expt_0510.cu b/expt_0510.cu new file mode 100644 index 0000000..063c6ba --- /dev/null +++ b/expt_0510.cu @@ -0,0 +1,254 @@ +// Experimental content: Test branch performance (optimized memory) +#include "sortlib.cuh" + +#include +#include +#include +#include + +std::vector population_vector, sample_vector, + le_sample_vector; + +constexpr long DEFAULT_SAMPLE_LENGTH = 65535; +constexpr size_t TEST_LENGTH = 32'768; +// constexpr size_t TEST_LENGTH = 4096; +// constexpr size_t RESERVED_BLOCK = 100; + +unsigned long long max_value = 0, min_value = 0xfffffffff; +typedef std::pair 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__ long cuda_sample_length; +__device__ key_type *cudaSampleItem, *cudaNormalSampleItem, *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; + // cdf_normal_result = nullptr; +#ifdef TEST_BOUNDS + index_max = 0; + index_min = 0x7fffffff; +#endif +} + +__global__ void initNormalSample(key_type *normal_sample) { + cudaNormalSampleItem = normal_sample; +} + +__global__ void kernel(unsigned long step, const double slice_size, + const unsigned long split_size) { + auto custom_sort = CustomSort(cuda_sample_length, sizeof(key_type) * 8); + // printf("kernel1 step: %ld\n", step); + 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, cudaSampleItem + cuda_sample_length, + cudaPopulationItem[tid]) / + slice_size) - + 1; + // cdf_result[index] = true; + } +} + +__global__ void kernel2(unsigned long step, const double slice_size) { + + for (int i = 0; i < step; i++) { + auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i; + auto _index = + (int)(FactorySort::sample_cdf(cudaNormalSampleItem, cuda_sample_length, + cudaPopulationItem[tid]) / + slice_size) - + 1; + // cdf_result[index] = true; + } +} + +__global__ void print2() { + /*for (int i = 0; i< SAMPLE_LENGTH; i++) { + printf("%llu ", cudaSampleItem[i]); + } + printf("\n");*/ + printf("%lu\n", CustomSort::fast_log(cuda_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 check_items() { + for (int i = 0; i < cuda_sample_length; i++) { + printf("%lld ", cudaSampleItem[i]); + } + printf("\n"); +} + +void initCustomSampleCPU() { + auto sample_length = sample_vector.size(); + auto tmp = new key_type[sample_length]; + + auto custom_sort = CustomSort(sample_length, 0); + + for (size_t i = 0; i < sample_length; i++) { + tmp[i] = sample_vector[custom_sort.calculate_index(i) - 1]; + } + + memcpy(sample_vector.data(), tmp, sizeof(key_type) * sample_length); + + delete[] tmp; +} + +void run_kernel(size_t test_size, bool normal = true) { + dim3 grid_dim = 32, block_dim = 32; + unsigned long step = test_size / (grid_dim.x * block_dim.x); + auto scale = 2; + const unsigned long split_size = test_size * scale; + const auto slice_size = 1.0 / (double)split_size; + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + cudaEventRecord(start, nullptr); + if (normal) { + kernel<<>>(step, slice_size, split_size); + } else { + kernel2<<>>(step, slice_size); + } + + cudaDeviceSynchronize(); + cudaEventRecord(stop, nullptr); + cudaEventSynchronize(stop); + float time; + cudaEventElapsedTime(&time, start, stop); + cudaEventDestroy(start); + cudaEventDestroy(stop); + if (time == 0) { + // printf("last error: %u\n", cudaGetLastError()); + } + printf("time: %lf\n", time); + cudaDeviceSynchronize(); + + print_function<<<1, 1>>>(); + cudaDeviceSynchronize(); +} + +__global__ void freeStorageStage1() { cudaFree(cudaSampleItem); } + +__global__ void freeStorageStage2() { + cudaFree(cudaNormalSampleItem); + cudaFree(cudaPopulationItem); + // cudaFree(cdf_result); +} + +void read_file(char const *filename, long sample_length, + unsigned long max_number = TEST_LENGTH) { + max_number += sample_length + 256; + auto read_number = 0; + FILE *file = fopen(filename, "r"); + assert(file); + for (long long i; + read_number < max_number && fscanf(file, "%lld ", &i) != EOF; + store_into_vector(i)) + read_number++; + fclose(file); +} + +long pow_for_sample(long n) { + auto x = 2; + for (int i = 0; i < n; i++) { + x *= 2; + } + return x - 1; +} + +__global__ void applyCudaSampleLength(long length) { + cuda_sample_length = length; +} + +int main(int argc, char const *argv[]) { + + auto test_size = TEST_LENGTH; + auto sample_length = DEFAULT_SAMPLE_LENGTH; + + if (argc >= 2) { + test_size = strtol(argv[1], nullptr, 10); + } + if (argc >= 3) { + sample_length = pow_for_sample(strtol(argv[2], nullptr, 10)); + } + + read_file("normal_distribution.txt", test_size); + + printf("population: %zu, test size: %zu, sample length: %zu\n", + population_vector.size(), test_size, sample_length); + + sample_vector = std::vector( + 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()); + le_sample_vector = sample_vector; + initCustomSampleCPU(); + + 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(); + applyCudaSampleLength<<<1, 1>>>(sample_length); + testCustomCalculation<<<1, 1>>>(sample_length); + cudaDeviceSynchronize(); + + initSample<<<1, 1>>>(cudaSample, cudaPopulation); + cudaDeviceSynchronize(); + + // check_items<<<1, 1>>>(); + // cudaDeviceSynchronize(); + run_kernel(test_size); + + freeStorageStage1<<<1, 1>>>(); + cudaDeviceSynchronize(); + + key_type *cudaNormalSample = nullptr; + cudaMalloc(&cudaNormalSample, sizeof(key_type) * sample_length); + cudaMemcpy(cudaNormalSample, le_sample_vector.data(), + sizeof(key_type) * sample_length, cudaMemcpyHostToDevice); + + cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) * sample_length, + cudaMemcpyDeviceToHost); + + initNormalSample<<<1, 1>>>(cudaNormalSample); + cudaDeviceSynchronize(); + run_kernel(test_size, false); + // check_items<<<1, 1>>>(); + // cudaDeviceSynchronize(); + + freeStorageStage2<<<1, 1>>>(); + cudaDeviceSynchronize(); +} diff --git a/normal_distribution.cpp b/normal_distribution.cpp index 20ef61c..211721a 100644 --- a/normal_distribution.cpp +++ b/normal_distribution.cpp @@ -40,13 +40,14 @@ void checkFileValid(const std::string &name, size_t length) { } void progress_bar(const std::set &v, size_t length) { + constexpr auto sleep_time = std::chrono::seconds(1); size_t previous_rate = 0; while (v.size() < length) { fprintf(stderr, "\rProgress %.02f%%(%lu) rate: %lu/s", (double)v.size() * 100 / (double)length, v.size(), v.size() - previous_rate); previous_rate = v.size(); - std::this_thread::sleep_for(std::chrono::seconds(1)); + std::this_thread::sleep_for(sleep_time); } fputs("", stderr); } -- cgit v1.3.1