diff options
Diffstat (limited to 'expt_0804.cu')
| -rw-r--r-- | expt_0804.cu | 222 |
1 files changed, 0 insertions, 222 deletions
diff --git a/expt_0804.cu b/expt_0804.cu deleted file mode 100644 index 66f66ec..0000000 --- a/expt_0804.cu +++ /dev/null @@ -1,222 +0,0 @@ -// Experimental content: Test cdf performance (Deprecated) -#include <algorithm> -#include <cassert> -#include <cstdio> -#include <cstring> -#include <iostream> -#include <vector> - -constexpr size_t length = 2097152; -std::vector<unsigned long long> population_vector, sample_vector, - sample_vector_into_cuda; - -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; - } - 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; - -constexpr long TEST_SIZE = 1024; - -__device__ key_type *cudaSampleItem, *cudaPopulationItem; -__device__ bool *cdf_result; -__device__ unsigned insert_value; -#ifdef TEST_BOUNDS -__device__ unsigned int index_max, index_min; -#endif - -__device__ const key_type *cudaBinarySearch(key_type *start, - const key_type val) { - int step_limit = (int)log2f(SAMPLE_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; - if (*last_known_point == val) { - return last_known_point; - } - son = son * 2 + -((*last_known_point - val) >> 63); - last_known_point = next_level_start + son; - } - return last_known_point; -} - -__device__ __host__ size_t calculate_location(size_t index) { - int step = (int)std::log2(++index) + 1; - auto location = - SAMPLE_LENGTH / (1 << step) * ((index & ((1 << (step - 1)) - 1)) * 2 + 1); - return location ? location : SAMPLE_LENGTH; -} - -__device__ double sample_cdf(double x) { - auto it = cudaBinarySearch(cudaSampleItem, x); - if (it == cudaSampleItem + SAMPLE_LENGTH) { - return 1; - } - if (it == cudaSampleItem) { - return 0; - } - auto it_prev = it - 1; - return (double(it_prev - cudaSampleItem) + - (x - (double)*it_prev) / (double)(*it - *it_prev)) / - double(SAMPLE_LENGTH - 1); -} - -__global__ void initStorage(unsigned scale, unsigned test_size) { - cudaFree(cdf_result); - cudaMalloc(&cdf_result, scale * test_size * sizeof(bool)); - memset(cdf_result, 0, scale * test_size * sizeof(bool)); - insert_value = 0; - // printf("initCuda storage\n"); -#ifdef TEST_BOUNDS - index_max = 0; - index_min = 0x7fffffff; -#endif -} - -__global__ void initCuda(key_type *sample_item, key_type *population_item) { - cudaSampleItem = sample_item; - cudaPopulationItem = population_item; - /*for (int i = 0; i < SAMPLE_LENGTH; i++) { - }*/ - // cudaMalloc(&cudaSampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long - // long)); - // cudaMalloc(&Storage, TEST_SIZE * sizeof(key_type)); - cdf_result = nullptr; -} - -__global__ void kernel(unsigned long step, const double slice_size, - const long split_size) { - for (int i = 0; i < step; i++) { - auto tid = step * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + - threadIdx.x + i; - - // printf("%d\n", tid); - auto index = (int)(sample_cdf(cudaPopulationItem[tid]) / slice_size); - if (cdf_result[index]) { - while (cdf_result[++index]) { - assert(index < split_size); - } - } - cdf_result[index] = true; - // atomicAdd(&insert_value, 1); - -#ifdef TEST_BOUNDS - while (true) { - unsigned tmp = index_min; - if (tmp < tid) { - break; - } - if (atomicCAS(&index_min, tmp, tid) == tmp) { - break; - } - } - while (true) { - unsigned tmp = index_max; - if (tmp > tid) { - break; - } - if (atomicCAS(&index_max, tmp, tid) == tmp) { - break; - } - } -#endif - } -} - -__global__ void print_function() { - // printf("%u\n", insert_value); -#ifdef TEST_BOUNDS - printf("%u %u\n", index_min, index_max); -#endif -} - -int main(int argc, char const *argv[]) { - size_t test_size; - if (argc == 1) { - test_size = 1048576; - } else { - try { - test_size = std::stol(argv[1]); - } catch (...) { - return 1; - } - } - - FILE *file = fopen("normal_distribution.txt", "r"); - assert(file); - for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i)) - ; - fclose(file); - - assert(population_vector.size() == length); - - sample_vector = std::vector<unsigned long long>( - 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()); - - sample_vector_into_cuda.resize(SAMPLE_LENGTH); - for (int i = 0; i < SAMPLE_LENGTH; i++) { - sample_vector_into_cuda[i] = sample_vector[calculate_location(i) - 1]; - } - - key_type *cudaSample = nullptr, *cudaPopulation = nullptr; - cudaMalloc(&cudaSample, sizeof(key_type) * (SAMPLE_LENGTH)); - cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE); - cudaMemcpy(cudaSample, &sample_vector_into_cuda[0], - sizeof(key_type) * sample_vector.size(), cudaMemcpyHostToDevice); - cudaMemcpy(cudaPopulation, &population_vector[SAMPLE_LENGTH], - sizeof(key_type) * TEST_SIZE, cudaMemcpyHostToDevice); - // memcpy(sample_heap, sample_vector.cbegin(),sizeof(key_type) *(SAMPLE_LENGTH - // + 2)); - - initCuda<<<1, 1>>>(cudaSample, cudaPopulation); - cudaDeviceSynchronize(); - - dim3 grid_dim = 64, block_dim = 16; - - unsigned long step = test_size / (grid_dim.x * block_dim.x); - printf("(new) step: %lu\n", step); - assert(!(test_size % (grid_dim.x * block_dim.x))); - - for (int scale = 2; scale <= 8; scale++) { - const 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<<<grid_dim, block_dim>>>(step, slice_size, split_size); - cudaDeviceSynchronize(); - cudaEventRecord(stop, nullptr); - cudaEventSynchronize(stop); - float time; - cudaEventElapsedTime(&time, start, stop); - cudaEventDestroy(start); - cudaEventDestroy(stop); - - printf("time: %lf\n", time); - print_function<<<1, 1>>>(); - cudaDeviceSynchronize(); - } - return 0; -}
\ No newline at end of file |