#include "sortlib.cuh" #include #include #include #include constexpr size_t length = 2097152; constexpr long TEST_LENGTH = 32'768; std::vector population_vector, sample_vector, sample_vector_into_cuda; #define CHECK_CUDA_ERR(x) \ { \ auto ret_ = (x); \ if (ret_ != 0) { \ printf("%d: CUDA ERR: %d\n", __LINE__, ret_); \ } \ } unsigned long long max_value = 0, min_value = 0xfffffffff; inline void store_into_vector(unsigned long long value) { if (max_value < value) { max_value = value; } else if (min_value > value) { min_value = value; } population_vector.push_back(value); } typedef unsigned long long key_type; typedef unsigned long long *key_type_ptr; constexpr long SAMPLE_LENGTH = 1024; __device__ key_type *cudaSampleItem, *cudaPopulationItem; __device__ bool *cdf_result; __device__ unsigned insert_value; __global__ void initStorage(unsigned scale, unsigned test_size) { CHECK_CUDA_ERR(cudaFree(cdf_result)); CHECK_CUDA_ERR(cudaMalloc(&cdf_result, scale * test_size * sizeof(bool))); memset(cdf_result, 0, scale * test_size * sizeof(bool)); insert_value = 0; } __global__ void initCuda(key_type *sample_item, key_type *population_item) { cudaSampleItem = sample_item; cudaPopulationItem = population_item; cdf_result = nullptr; } __global__ void kernel(unsigned long step, const double slice_size) { for (int i = 0; i < step; i++) { auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i; // printf("%d\n", tid); auto index = (int)(FactorySort::cdf(cudaSampleItem, length, cudaPopulationItem[tid]) / slice_size); cdf_result[index] = true; } } __global__ void print_function() { // printf("%u\n", insert_value); #ifdef TEST_BOUNDS printf("%u %u\n", index_min, index_max); #endif } __global__ void freeStorage() { cudaFree(cudaSampleItem); cudaFree(cudaPopulationItem); cudaFree(cdf_result); } typedef std::pair dim_pair_type; 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("current 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); kernel<<>>(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); 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( 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()); 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(); // testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH); // cudaDeviceSynchronize(); run_kernel(test_size); freeStorage<<<1, 1>>>(); cudaDeviceSynchronize(); }