From cbe4695b8bd0a4f0e429af19694030edeaf047a5 Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Fri, 7 Jul 2023 02:33:31 +0800 Subject: feat: Use new read helper class Signed-off-by: KunoiSayami --- expt_0502.cu | 78 ++++++++++++------------------------------------------------ expt_0503.cu | 32 +++++++------------------ 2 files changed, 24 insertions(+), 86 deletions(-) diff --git a/expt_0502.cu b/expt_0502.cu index ef0414d..004813d 100644 --- a/expt_0502.cu +++ b/expt_0502.cu @@ -9,6 +9,8 @@ #include #include +#include "read_helper.h" + std::vector population_vector, sample_vector, le_sample_vector; @@ -52,36 +54,6 @@ __global__ void initSample(key_type *sample, key_type *normal_sample, #endif } -__global__ void kernel(unsigned long step, const double slice_size) { - auto custom_sort = CustomSort(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.cdf(cudaSampleItem, cudaPopulationItem[tid]) / - slice_size) - - 1; - - auto index2 = (int)(FactorySort::cdf(cudaNormalSampleItem, SAMPLE_LENGTH, - cudaPopulationItem[tid]) / - slice_size) - - 1; - printf("x => %llu %d %d\n", cudaPopulationItem[tid], index, index2); - // printf("%llu %d\n", cudaPopulationItem[tid], index); - /*if (cdf_result[index]) { - while (cdf_result[++index]) { - if (index >= split_size) { - //printf("escape: %llu %lu\n", cudaPopulationItem[tid], split_size); - } - assert(index < (split_size + RESERVED_BLOCK)); - } - }*/ - // cdf_result[index] = true; - } -} - __global__ void kernel2(unsigned long step, const double _slice_size) { auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(key_type) * 8); // printf("kernel1 step: %ld\n", step); @@ -165,24 +137,6 @@ __global__ void freeStorage() { cudaFree(cdf_result); } -__global__ void initCustomSample() { - // printf("init\n"); - 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]; - } - // printf("copy\n"); - 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; - // printf("finalize\n"); -} - const dim_pair_type DIM_PAIR[] = {dim_pair_type(32, 32)}; void run_kernel(size_t test_size) { @@ -226,19 +180,17 @@ 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); + ReadHelper readHelper("normal_distribution.txt", SAMPLE_LENGTH, TEST_LENGTH, + 0); + readHelper.readFile(); printf("test size: %d\n", test_size); - 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); + + readHelper.split_into(sample_vector, population_vector); + std::sort(sample_vector.begin(), sample_vector.end()); le_sample_vector = sample_vector; + rebuild(sample_vector); key_type *cudaSample = nullptr, *cudaPopulation; cudaMalloc(&cudaSample, sizeof(key_type) * SAMPLE_LENGTH); @@ -256,16 +208,16 @@ int main() { // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8); // custom_sort.testCalculation(); - testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH); - cudaDeviceSynchronize(); + // testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH); + // cudaDeviceSynchronize(); initSample<<<1, 1>>>(cudaSample, cudaNormalSample, cudaPopulation); cudaDeviceSynchronize(); - initCustomSample<<<1, 1>>>(); - cudaDeviceSynchronize(); - cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) * SAMPLE_LENGTH, - cudaMemcpyDeviceToHost); + // initCustomSample<<<1, 1>>>(); + // cudaDeviceSynchronize(); + // cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) * + // SAMPLE_LENGTH,cudaMemcpyDeviceToHost); // check_items<<<1, 1>>>(); // cudaDeviceSynchronize(); diff --git a/expt_0503.cu b/expt_0503.cu index c738e43..2785392 100644 --- a/expt_0503.cu +++ b/expt_0503.cu @@ -2,6 +2,7 @@ #define ENABLE_SORT_TEST #include "sortlib.cuh" +#include "read_helper.h" #include #include #include @@ -51,8 +52,7 @@ __global__ void initSample(key_type *sample, key_type *normal_sample, #endif } -__global__ void kernel(unsigned long step, const double slice_size, - const unsigned long split_size) { +__global__ void kernel(unsigned long step, const double slice_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++) { @@ -151,7 +151,7 @@ void run_kernel(size_t test_size, bool normal = true) { cudaEventCreate(&stop); cudaEventRecord(start, nullptr); if (normal) { - kernel<<>>(step, slice_size, split_size); + kernel<<>>(step, slice_size); } else { kernel2<<>>(step, slice_size); } @@ -173,19 +173,6 @@ void run_kernel(size_t test_size, bool normal = true) { cudaDeviceSynchronize(); } -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++) { @@ -210,17 +197,16 @@ int main(int argc, char const *argv[]) { sample_length = pow_for_sample(strtol(argv[2], nullptr, 10)); } - read_file("normal_distribution.txt", test_size); + ReadHelper readHelper("normal_distribution.txt", sample_length, test_size, 0); + + printf("test size: %zu, sample length: %zu\n", test_size, sample_length); - printf("population: %zu, test size: %zu, sample length: %zu\n", - population_vector.size(), test_size, sample_length); + readHelper.readFile(); + readHelper.split_into(sample_vector, population_vector); - 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; + rebuild(sample_vector); key_type *cudaSample = nullptr, *cudaPopulation; cudaMalloc(&cudaSample, sizeof(key_type) * sample_length); -- cgit v1.3.1