diff options
| -rw-r--r-- | CMakeLists.txt | 20 | ||||
| -rw-r--r-- | expt_0502.cu | 278 | ||||
| -rw-r--r-- | expt_0503.cu | 235 | ||||
| -rw-r--r-- | sortlib.cuh | 43 |
4 files changed, 575 insertions, 1 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index 161d8ea..6389a30 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -161,10 +161,28 @@ set_target_properties(expt_0425 PROPERTIES set_target_properties(expt_0425 PROPERTIES CUDA_ARCHITECTURES "75") set_target_properties(expt_0425 PROPERTIES LINKER_LANGUAGE CUDA) + add_executable(expt_0501 expt_0501.cu) target_link_libraries(expt_0501 m stdc++) set_target_properties(expt_0501 PROPERTIES CUDA_SEPARABLE_COMPILATION ON) set_target_properties(expt_0501 PROPERTIES CUDA_ARCHITECTURES "75") -set_target_properties(expt_0501 PROPERTIES LINKER_LANGUAGE CUDA)
\ No newline at end of file +set_target_properties(expt_0501 PROPERTIES LINKER_LANGUAGE CUDA) + +add_executable(expt_0502 expt_0502.cu) +target_link_libraries(expt_0502 m stdc++) + +set_target_properties(expt_0502 PROPERTIES + CUDA_SEPARABLE_COMPILATION ON) +set_target_properties(expt_0502 PROPERTIES CUDA_ARCHITECTURES "75") +set_target_properties(expt_0502 PROPERTIES LINKER_LANGUAGE CUDA) + + +add_executable(expt_0503 expt_0503.cu) +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 diff --git a/expt_0502.cu b/expt_0502.cu new file mode 100644 index 0000000..eb4ff5b --- /dev/null +++ b/expt_0502.cu @@ -0,0 +1,278 @@ +// Experimental content: Test the correctness of binary search and +// special(single thread) search +#include "sortlib.cuh" + +#include <algorithm> +#include <cassert> +#include <cstdio> +#include <vector> + +std::vector<unsigned long long> population_vector, sample_vector, + le_sample_vector; + +constexpr size_t SAMPLE_LENGTH = 1023; +constexpr size_t TEST_LENGTH = 32'768; +// constexpr size_t TEST_LENGTH = 4096; +constexpr size_t RESERVED_BLOCK = 10'000; + +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__ key_type *cudaSampleItem, *cudaNormalSampleItem, *cudaPopulationItem; +__device__ bool *cdf_result, *cdf_normal_result; +#ifdef TEST_BOUNDS +__device__ unsigned insert_value; +__device__ unsigned int index_max, index_min; +#endif + +__global__ void initSample(key_type *sample, key_type *normal_sample, + key_type *population) { + cudaSampleItem = sample; + cudaPopulationItem = population; + cudaNormalSampleItem = normal_sample; + // 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 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.sample_cdf_custom_version( + cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH, + cudaPopulationItem[tid]) / + slice_size) - + 1; + + auto index2 = + (int)(FactorySort::sample_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); + for (int i = 0; i < step; i++) { + auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i; + + // printf("%d\n", tid); + auto index = custom_sort.calculate_index( + custom_sort.binary_search(cudaSampleItem, + cudaSampleItem + SAMPLE_LENGTH, + cudaPopulationItem[tid]) - + cudaSampleItem); + + auto index2 = + FactorySort::cudaBinarySearch(cudaNormalSampleItem, + cudaNormalSampleItem + SAMPLE_LENGTH, + cudaPopulationItem[tid]) - + cudaNormalSampleItem; + printf("x => %llu %ld %ld\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 print2() { + /*for (int i = 0; i< SAMPLE_LENGTH; i++) { + printf("%llu ", cudaSampleItem[i]); + } + printf("\n");*/ + printf("%lu\n", CustomSort::fast_log(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 < SAMPLE_LENGTH; i++) { + printf("%lld ", cudaSampleItem[i]); + } + printf("\n"); +} + +__global__ void initStorage(unsigned scale, unsigned test_size) { + cudaFree(cdf_result); + // printf("test size: %u\n", test_size); + cudaMalloc(&cdf_result, (scale * test_size + RESERVED_BLOCK) * sizeof(bool)); + memset(cdf_result, 0, (scale * test_size + RESERVED_BLOCK) * sizeof(bool)); + cudaFree(cdf_normal_result); + cudaMalloc(&cdf_normal_result, + (scale * test_size + RESERVED_BLOCK) * sizeof(bool)); + memset(cdf_normal_result, 0, + (scale * test_size + RESERVED_BLOCK) * sizeof(bool)); +} + +__global__ void compareCdf(unsigned scale, unsigned test_size) { + auto compare_size = scale * test_size * sizeof(bool); + for (int i = 0; i < compare_size; i++) { + if (cdf_result[i] != cdf_normal_result[i]) { + printf("Compare error at %d\n", i); + return; + } + } + printf("Compare successful"); +} + +__global__ void freeStorage() { + cudaFree(cudaSampleItem); + cudaFree(cudaPopulationItem); + 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) { + for (auto &pair : DIM_PAIR) { + dim3 grid_dim = pair.first, block_dim = pair.second; + unsigned long step = test_size / (grid_dim.x * block_dim.x); + printf("%scurrent dim: %d %d %lu\n", "", pair.first, pair.second, step); + for (int scale = 2; scale <= 16; scale += 2) { + const unsigned long split_size = test_size * scale; + const auto slice_size = 1.0 / (double)split_size; + printf("scale: %i ", scale); + initStorage<<<1, 1>>>(scale, test_size); + cudaDeviceSynchronize(); + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + cudaEventRecord(start, nullptr); + 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("time: %lf\n", time); + compareCdf<<<1, 1>>>(scale, test_size); + cudaDeviceSynchronize(); + + print_function<<<1, 1>>>(); + cudaDeviceSynchronize(); + } + } +} + +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); + + printf("test size: %d\n", test_size); + 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; + + 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); + + key_type *cudaNormalSample = nullptr; + cudaMalloc(&cudaNormalSample, sizeof(key_type) * SAMPLE_LENGTH); + cudaMemcpy(cudaNormalSample, le_sample_vector.data(), + sizeof(key_type) * SAMPLE_LENGTH, cudaMemcpyHostToDevice); + + // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8); + // custom_sort.testCalculation(); + 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); + + // check_items<<<1, 1>>>(); + // cudaDeviceSynchronize(); + run_kernel(test_size); + // check_items<<<1, 1>>>(); + // cudaDeviceSynchronize(); + + freeStorage<<<1, 1>>>(); + cudaDeviceSynchronize(); +} diff --git a/expt_0503.cu b/expt_0503.cu new file mode 100644 index 0000000..96cf44e --- /dev/null +++ b/expt_0503.cu @@ -0,0 +1,235 @@ +// Experimental content: Test branch performance +#include "sortlib.cuh" + +#include <algorithm> +#include <cassert> +#include <cstdio> +#include <vector> + +std::vector<unsigned long long> population_vector, sample_vector, + le_sample_vector; + +constexpr size_t SAMPLE_LENGTH = 1023; +constexpr size_t TEST_LENGTH = 32'768; +// constexpr size_t TEST_LENGTH = 4096; +constexpr size_t RESERVED_BLOCK = 10'000; + +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__ key_type *cudaSampleItem, *cudaNormalSampleItem, *cudaPopulationItem; +__device__ bool *cdf_result, *cdf_normal_result; +#ifdef TEST_BOUNDS +__device__ unsigned insert_value; +__device__ unsigned int index_max, index_min; +#endif + +__global__ void initSample(key_type *sample, key_type *normal_sample, + key_type *population) { + cudaSampleItem = sample; + cudaPopulationItem = population; + cudaNormalSampleItem = normal_sample; + // 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 kernel(unsigned long step, const double slice_size, + const unsigned long split_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.sample_cdf_custom_version( + cudaSampleItem, cudaSampleItem + 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, SAMPLE_LENGTH, + cudaPopulationItem[tid]) / + slice_size) - + 1; + cdf_normal_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(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 < SAMPLE_LENGTH; i++) { + printf("%lld ", cudaSampleItem[i]); + } + printf("\n"); +} + +__global__ void initStorage(unsigned scale, unsigned test_size) { + cudaFree(cdf_result); + // printf("test size: %u\n", test_size); + cudaMalloc(&cdf_result, (scale * test_size + RESERVED_BLOCK) * sizeof(bool)); + memset(cdf_result, 0, (scale * test_size + RESERVED_BLOCK) * sizeof(bool)); + cudaFree(cdf_normal_result); + cudaMalloc(&cdf_normal_result, + (scale * test_size + RESERVED_BLOCK) * sizeof(bool)); + memset(cdf_normal_result, 0, + (scale * test_size + RESERVED_BLOCK) * sizeof(bool)); +} + +__global__ void freeStorage() { + cudaFree(cudaSampleItem); + cudaFree(cudaPopulationItem); + 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"); +} + +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; + initStorage<<<1, 1>>>(scale, test_size); + cudaDeviceSynchronize(); + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + cudaEventRecord(start, nullptr); + if (normal) { + kernel<<<grid_dim, block_dim>>>(step, slice_size, split_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("time: %lf\n", time); + cudaDeviceSynchronize(); + + print_function<<<1, 1>>>(); + cudaDeviceSynchronize(); +} + +int main(int argc, char const *argv[]) { + + auto test_size = TEST_LENGTH; + + if (argc >= 2) { + test_size = strtol(argv[1], 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); + + printf("population: %zu, test size: %zu\n", population_vector.size(), + test_size); + + 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; + + 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); + + key_type *cudaNormalSample = nullptr; + cudaMalloc(&cudaNormalSample, sizeof(key_type) * SAMPLE_LENGTH); + cudaMemcpy(cudaNormalSample, le_sample_vector.data(), + sizeof(key_type) * SAMPLE_LENGTH, cudaMemcpyHostToDevice); + + // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8); + // custom_sort.testCalculation(); + 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); + + // check_items<<<1, 1>>>(); + // cudaDeviceSynchronize(); + run_kernel(test_size); + run_kernel(test_size, false); + // check_items<<<1, 1>>>(); + // cudaDeviceSynchronize(); + + freeStorage<<<1, 1>>>(); + cudaDeviceSynchronize(); +} diff --git a/sortlib.cuh b/sortlib.cuh index ce40e0f..ca5525b 100644 --- a/sortlib.cuh +++ b/sortlib.cuh @@ -85,6 +85,49 @@ public: 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) { + 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; + } + /// Should be correct version __device__ __host__ double sample_cdf_custom_version(key_type *start, const key_type *end, key_type x) { |