diff options
| -rw-r--r-- | expt_0503.cu | 79 |
1 files changed, 53 insertions, 26 deletions
diff --git a/expt_0503.cu b/expt_0503.cu index 96cf44e..58f35aa 100644 --- a/expt_0503.cu +++ b/expt_0503.cu @@ -9,7 +9,7 @@ std::vector<unsigned long long> population_vector, sample_vector, le_sample_vector; -constexpr size_t SAMPLE_LENGTH = 1023; +constexpr long DEFAULT_SAMPLE_LENGTH = 65535; constexpr size_t TEST_LENGTH = 32'768; // constexpr size_t TEST_LENGTH = 4096; constexpr size_t RESERVED_BLOCK = 10'000; @@ -28,6 +28,7 @@ inline void store_into_vector(unsigned long long value) { } // #define TEST_BOUNDS +__device__ long cuda_sample_length; __device__ key_type *cudaSampleItem, *cudaNormalSampleItem, *cudaPopulationItem; __device__ bool *cdf_result, *cdf_normal_result; #ifdef TEST_BOUNDS @@ -51,14 +52,14 @@ __global__ void initSample(key_type *sample, key_type *normal_sample, __global__ void kernel(unsigned long step, const double slice_size, const unsigned long split_size) { - auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(key_type) * 8); + 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 + SAMPLE_LENGTH, + cudaSampleItem, cudaSampleItem + cuda_sample_length, cudaPopulationItem[tid]) / slice_size) - 1; @@ -71,7 +72,7 @@ __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, SAMPLE_LENGTH, + (int)(FactorySort::sample_cdf(cudaNormalSampleItem, cuda_sample_length, cudaPopulationItem[tid]) / slice_size) - 1; @@ -84,7 +85,7 @@ __global__ void print2() { printf("%llu ", cudaSampleItem[i]); } printf("\n");*/ - printf("%lu\n", CustomSort::fast_log(SAMPLE_LENGTH)); + printf("%lu\n", CustomSort::fast_log(cuda_sample_length)); } __global__ void print_function() { @@ -95,7 +96,7 @@ __global__ void print_function() { } __global__ void check_items() { - for (int i = 0; i < SAMPLE_LENGTH; i++) { + for (int i = 0; i < cuda_sample_length; i++) { printf("%lld ", cudaSampleItem[i]); } printf("\n"); @@ -122,13 +123,13 @@ __global__ void freeStorage() { __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++) { + auto custom_sort = CustomSort(cuda_sample_length, 0); + cudaMalloc(&tmp, sizeof(key_type) * cuda_sample_length); + for (size_t i = 0; i < cuda_sample_length; i++) { tmp[i] = cudaSampleItem[custom_sort.calculate_index(i) - 1]; } // printf("copy\n"); - memcpy(cudaSampleItem, tmp, sizeof(key_type) * SAMPLE_LENGTH); + memcpy(cudaSampleItem, tmp, sizeof(key_type) * cuda_sample_length); cudaFree(tmp); tmp = nullptr; // memset(cdf_result, 0, sizeof(key_type) * TEST_LENGTH); @@ -172,47 +173,73 @@ 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++) { + 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)); + } - FILE *file = fopen("normal_distribution.txt", "r"); - assert(file); - for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i)) - ; - fclose(file); + read_file("normal_distribution.txt", test_size); - printf("population: %zu, test size: %zu\n", population_vector.size(), - test_size); + printf("population: %zu, test size: %zu, sample length: %zu\n", + population_vector.size(), test_size, sample_length); sample_vector = std::vector<key_type>( - population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH - 2); + 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; 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, + 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, + cudaMemcpy(cudaPopulation, population_vector.data() + sample_length, sizeof(key_type) * test_size, cudaMemcpyHostToDevice); key_type *cudaNormalSample = nullptr; - cudaMalloc(&cudaNormalSample, sizeof(key_type) * SAMPLE_LENGTH); + cudaMalloc(&cudaNormalSample, sizeof(key_type) * sample_length); cudaMemcpy(cudaNormalSample, le_sample_vector.data(), - sizeof(key_type) * SAMPLE_LENGTH, cudaMemcpyHostToDevice); + sizeof(key_type) * sample_length, cudaMemcpyHostToDevice); // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8); // custom_sort.testCalculation(); - testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH); + applyCudaSampleLength<<<1, 1>>>(sample_length); + testCustomCalculation<<<1, 1>>>(sample_length); cudaDeviceSynchronize(); initSample<<<1, 1>>>(cudaSample, cudaNormalSample, cudaPopulation); @@ -220,7 +247,7 @@ int main(int argc, char const *argv[]) { initCustomSample<<<1, 1>>>(); cudaDeviceSynchronize(); - cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) * SAMPLE_LENGTH, + cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) * sample_length, cudaMemcpyDeviceToHost); // check_items<<<1, 1>>>(); |