diff options
| -rw-r--r-- | CMakeLists.txt | 12 | ||||
| -rw-r--r-- | expt_0510.cu | 7 | ||||
| -rw-r--r-- | expt_0516.cu | 304 | ||||
| -rw-r--r-- | normal_distribution.cpp | 2 | ||||
| -rw-r--r-- | sortlib.cuh | 90 |
5 files changed, 348 insertions, 67 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index 4f66661..7ad8949 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -194,4 +194,14 @@ 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 +set_target_properties(expt_0510 PROPERTIES LINKER_LANGUAGE CUDA) + + +add_executable(expt_0516 expt_0516.cu) +target_link_libraries(expt_0516 m stdc++) + +set_target_properties(expt_0516 PROPERTIES + CUDA_SEPARABLE_COMPILATION ON) +set_target_properties(expt_0516 PROPERTIES CUDA_ARCHITECTURES "75") +set_target_properties(expt_0516 PROPERTIES LINKER_LANGUAGE CUDA) + diff --git a/expt_0510.cu b/expt_0510.cu index 7616c13..3efad47 100644 --- a/expt_0510.cu +++ b/expt_0510.cu @@ -183,8 +183,7 @@ void readFile(char const *filename, long sample_length, read_number = randomRow(total_row - sample_length - max_number - 256); total_row = read_number; // fprintf(stderr, "Skip %lu\n", read_number); - for (long long i; read_number > 0 && fscanf(file, "%lld ", &i) != EOF; - store_into_vector(i)) + for (long long i; read_number > 0 && fscanf(file, "%lld ", &i) != EOF;) read_number--; } @@ -225,7 +224,9 @@ int main(int argc, char const *argv[]) { total_row = strtol(argv[3], nullptr, 10); } - readFile("normal_distribution.txt", test_size, total_row); + CustomSort::testSelf(sample_length); + + readFile("normal_distribution.txt", sample_length, total_row, test_size); printf("population: %zu, skip: %lu, test size: %zu, sample length: %zu | ", population_vector.size(), total_row, test_size, sample_length); diff --git a/expt_0516.cu b/expt_0516.cu new file mode 100644 index 0000000..4d9be45 --- /dev/null +++ b/expt_0516.cu @@ -0,0 +1,304 @@ +// Experimental content: Test branch performance (optimized memory and output) +#include "sortlib.cuh" + +#include <algorithm> +#include <cassert> +#include <cstdio> +#include <random> +#include <vector> + +#ifndef NDEBUG +#define P_ERR(...) fprintf(stderr, __VA_ARGS__) +#else +#define P_ERR(...) +#endif + +std::vector<unsigned long long> 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<int, int> 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; +} + +__global__ void initNormalSample(key_type *normal_sample) { + cudaNormalSampleItem = normal_sample; +} + +__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++) { + auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i; + + // printf("%d\n", tid); + auto _index = (int)(custom_sort.sample_cdf_custom_version( + cudaSampleItem, 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 printFunction() {} + +__global__ void checkItems() { + 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; +} + +constexpr dim3 grid_dim = 32, block_dim = 32; + +bool checkTestSize() {} + +void run_kernel(size_t test_size, bool custom = true) { + // need description + unsigned long step = test_size / (grid_dim.x * block_dim.x); + + // Only scale mode will use + 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 (custom) { + kernel<<<grid_dim, block_dim>>>(step, slice_size); + } else { + kernel2<<<grid_dim, block_dim>>>(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("%stime: %lf ", custom ? "custom " : "", time); + cudaDeviceSynchronize(); + + printFunction<<<1, 1>>>(); + cudaDeviceSynchronize(); +} + +__global__ void freeStorageStage1() { cudaFree(cudaSampleItem); } + +__global__ void freeStorageStage2() { + cudaFree(cudaNormalSampleItem); + cudaFree(cudaPopulationItem); + // cudaFree(cdf_result); +} + +unsigned long randomRow(unsigned long max_value_) { + std::random_device randomDevice; + std::mt19937 mt19937(randomDevice()); + std::uniform_int_distribution<std::mt19937::result_type> dst(0, max_value_); + return dst(mt19937); +} + +void readFile(char const *filename, long sample_length, + unsigned long &total_row, + unsigned long max_number = TEST_LENGTH) { + auto read_number = 0UL; + FILE *file = fopen(filename, "r"); + assert(file); + P_ERR("Reading sample"); + for (long long i; + read_number < max_number && fscanf(file, "%lld ", &i) != EOF; + store_into_vector(i)) + read_number++; + + if (total_row > 0) { + P_ERR("\rReading skip"); + read_number = randomRow(total_row - sample_length - max_number - 256); + total_row = read_number; + // fprintf(stderr, "Skip %lu\n", read_number); + for (long long i; read_number > 0 && fscanf(file, "%lld ", &i) != EOF;) + read_number--; + } + + P_ERR("\rReading population"); + read_number = 0; + + auto remain = sample_length + 256; + for (long long i; read_number < remain && fscanf(file, "%lld ", &i) != EOF; + store_into_vector(i)) + read_number++; + fclose(file); + P_ERR("\r"); +} + +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; + auto total_row = 0UL; + + if (argc >= 2) { + test_size = strtol(argv[1], nullptr, 10); + } + if (argc >= 3) { + sample_length = pow_for_sample(strtol(argv[2], nullptr, 10)); + } + if (argc >= 4) { + total_row = strtol(argv[3], nullptr, 10); + } + + readFile("normal_distribution.txt", sample_length, total_row, test_size); + + printf("population: %zu, skip: %lu, test size: %zu, sample length: %zu\n", + population_vector.size(), total_row, test_size, sample_length); + + sample_vector = std::vector<key_type>( + 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; + + P_ERR("Coping sample"); + 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); + + P_ERR("\rCoping population"); + 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(); + P_ERR("\rApply custom length and test calculation"); + applyCudaSampleLength<<<1, 1>>>(sample_length); + testCustomCalculation<<<1, 1>>>(sample_length, sizeof(key_type) * 8); + cudaDeviceSynchronize(); + P_ERR("\rApply custom length and test calculation completed"); + + // Can remove this function if pass + CustomSort::testSelf(sample_length, sizeof(key_type) * 8); + + initSample<<<1, 1>>>(cudaSample, cudaPopulation); + cudaDeviceSynchronize(); + + // check_items<<<1, 1>>>(); + // cudaDeviceSynchronize(); + P_ERR("\r \rRunning " + "kernel\n"); + run_kernel(test_size); + + freeStorageStage1<<<1, 1>>>(); + cudaDeviceSynchronize(); + + key_type *cudaNormalSample = nullptr; + + P_ERR("\nCoping new sample"); + 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(); + + P_ERR("\r "); + P_ERR("\rRunning kernel2\n"); + run_kernel(test_size, false); + // check_items<<<1, 1>>>(); + // cudaDeviceSynchronize(); + puts(""); + + freeStorageStage2<<<1, 1>>>(); + cudaDeviceSynchronize(); +} diff --git a/normal_distribution.cpp b/normal_distribution.cpp index 211721a..55ffe33 100644 --- a/normal_distribution.cpp +++ b/normal_distribution.cpp @@ -69,7 +69,7 @@ int main(int argc, char const *argv[]) { length = 1048576; } else { try { - length = std::stol(argv[1]); + length = std::strtol(argv[1], nullptr, 10); } catch (...) { fprintf(stderr, "Wrong arguments\n"); return 1; diff --git a/sortlib.cuh b/sortlib.cuh index ca5525b..58aa7c1 100644 --- a/sortlib.cuh +++ b/sortlib.cuh @@ -15,8 +15,9 @@ public: /// MOVE_OFFSET means bit to select branch const int MOVE_OFFSET; - const int STEP_LIMIT; + const unsigned int STEP_LIMIT; + // IEEE 754 __device__ __host__ static size_t fast_log(size_t a) { float t = a; return (((*(int *)&t) >> 23) + 1) & 127; @@ -35,20 +36,20 @@ public: __device__ __host__ __attribute__((unused)) void testCalculation() const { for (size_t i = 0; i < LENGTH; i++) { - auto l = calculate_index(i); - auto l2 = calculate_rank(l - 1); + auto left = calculate_index(i); + auto right = calculate_rank(left - 1); // printf("%lu %lu\n", l, l2); - assert(i + 1 == l2); + assert(i + 1 == right); } } __device__ __host__ const key_type *binary_search(key_type *const start, - const key_type *end, - const key_type val) { + // const key_type *end, + const key_type val) const { // int step_limit = (int)fast_log(LENGTH); key_type *last_known_point = start; - auto son = 0; + auto son = 0UL; for (int i = 0; i < STEP_LIMIT; i++) { const auto next_level_start = start + (1 << (i + 1)) - 1; @@ -56,59 +57,16 @@ public: // printf("start: %ld, last: %ld\n", next_level_start - start, // last_known_point - start); - if (last_known_point > end) { + /*if (last_known_point > end) { printf("%ld\n", last_known_point - start); - } - if (*last_known_point == val) { - return last_known_point; - } - - if (next_level_start > end) { - printf("%ld\n", next_level_start - start); - } - - // son = get_son_from_step(son, (*last_known_point > val)); - auto branch_selector = ((*last_known_point - val) >> MOVE_OFFSET); - // printf("tmp: %llu\n", tmp); - // printf("%llu %llu ", val, *last_known_point); - son = son * 2 + branch_selector; - // printf("%d\n", son); - // puts(tmp == 0 ? "1:left" : "1:right"); - // if (son < 0) son = 0; - /*printf("%d %d %d\n", (1 << (i + 1)), son, - -(int)((*last_known_point - val) >> brenchSelector));*/ - last_known_point = next_level_start + son; - if (last_known_point > end) { - printf("%p %ld\n", end, next_level_start - start); - } - } - return last_known_point; - } - - __device__ __host__ const key_type *binary_search2(key_type *const start, - const key_type *end, - const key_type val) { - - // int step_limit = (int)fast_log(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; - - // printf("start: %ld, last: %ld\n", next_level_start - start, - // last_known_point - start); - - if (last_known_point > end) { - printf("%ld\n", last_known_point - start); - } + }*/ if (*last_known_point == val) { return last_known_point; } - if (next_level_start > end) { + /*if (next_level_start > end) { printf("%ld\n", next_level_start - start); - } + }*/ // son = get_son_from_step(son, (*last_known_point > val)); auto branch_selector = ((*last_known_point - val) >> MOVE_OFFSET); @@ -121,18 +79,19 @@ public: /*printf("%d %d %d\n", (1 << (i + 1)), son, -(int)((*last_known_point - val) >> brenchSelector));*/ last_known_point = next_level_start + son; - if (last_known_point > end) { + /*if (last_known_point > end) { printf("%p %ld\n", end, next_level_start - start); - } + }*/ } return last_known_point; } /// Should be correct version - __device__ __host__ double - sample_cdf_custom_version(key_type *start, const key_type *end, key_type x) { + __device__ __host__ double sample_cdf_custom_version(key_type *start, + // const key_type *end, + key_type x) const { // printf("custom version:\n"); - auto it = this->binary_search(start, end, x); + auto it = this->binary_search(start, x); // printf("search result: %ld\n", it - start); // assert(it <= start + this->LENGTH); /*if (it < start) { @@ -145,6 +104,7 @@ public: if (prev_real_location == this->LENGTH) { return 1; } + if (prev_real_location == 0) { return 0; } @@ -160,8 +120,9 @@ public: } /// First version - __device__ __host__ key_type * - original_binary_search(key_type *start, const key_type *end, key_type &val) { + __device__ __host__ key_type *original_binary_search(key_type *start, + const key_type *end, + key_type &val) const { auto begin = start; key_type *last_known_point = nullptr; while (begin < end) { @@ -188,7 +149,7 @@ public: /// CDF original version (should only work on default data layout) __device__ __host__ double sample_cdf(key_type *start, key_type *end, - key_type x) { + key_type x) const { auto it = this->original_binary_search(start, end, x); if (it == end) { return 1; @@ -202,6 +163,11 @@ public: (double)(x - *it_prev) / (double)(*it - *it_prev)) / (double)((end - start) - 1); } + + static void testSelf(size_t length) { + CustomSort(length, sizeof(long) * 8).testCalculation(); + puts("You can remove this function if passed already"); + } }; __global__ void testCustomCalculation(size_t length) { |