From e7de21daf6c674633ef1159c2fd042523c1ec32e Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Mon, 12 Jun 2023 00:51:30 +0800 Subject: refector: Rename last year experimental Signed-off-by: KunoiSayami --- CMakeLists.txt | 18 ++-- expt_0528.cu | 2 +- expt_0604.cpp | 1 + expt_0702.cpp | 102 -------------------- expt_0702.py | 23 ----- expt_0722.cpp | 129 ------------------------- expt_0726.cpp | 126 ------------------------ expt_0729.cu | 221 ------------------------------------------ expt_0802.cpp | 76 --------------- expt_0804.cu | 222 ------------------------------------------- expt_0809.cu | 219 ------------------------------------------ expt_0821.cpp | 81 ---------------- expt_0830.cu | 290 -------------------------------------------------------- expt_220702.cpp | 102 ++++++++++++++++++++ expt_220702.py | 23 +++++ expt_220722.cpp | 129 +++++++++++++++++++++++++ expt_220726.cpp | 126 ++++++++++++++++++++++++ expt_220729.cu | 221 ++++++++++++++++++++++++++++++++++++++++++ expt_220802.cpp | 76 +++++++++++++++ expt_220804.cu | 222 +++++++++++++++++++++++++++++++++++++++++++ expt_220809.cu | 219 ++++++++++++++++++++++++++++++++++++++++++ expt_220821.cpp | 81 ++++++++++++++++ expt_220830.cu | 290 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++ 23 files changed, 1500 insertions(+), 1499 deletions(-) delete mode 100644 expt_0702.cpp delete mode 100755 expt_0702.py delete mode 100644 expt_0722.cpp delete mode 100644 expt_0726.cpp delete mode 100644 expt_0729.cu delete mode 100644 expt_0802.cpp delete mode 100644 expt_0804.cu delete mode 100644 expt_0809.cu delete mode 100644 expt_0821.cpp delete mode 100644 expt_0830.cu create mode 100644 expt_220702.cpp create mode 100755 expt_220702.py create mode 100644 expt_220722.cpp create mode 100644 expt_220726.cpp create mode 100644 expt_220729.cu create mode 100644 expt_220802.cpp create mode 100644 expt_220804.cu create mode 100644 expt_220809.cu create mode 100644 expt_220821.cpp create mode 100644 expt_220830.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index 56323ee..d0292ff 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -84,30 +84,30 @@ set_target_properties(normal_distribution PROPERTIES LINKER_LANGUAGE CXX) add_executable(normal_distribution_expt3 expt_3.cpp) set_target_properties(normal_distribution_expt3 PROPERTIES LINKER_LANGUAGE CXX) -add_executable(normal_distribution_expt0702 expt_0702.cpp) +add_executable(normal_distribution_expt0702 expt_220702.cpp) set_target_properties(normal_distribution_expt0702 PROPERTIES LINKER_LANGUAGE CXX) add_executable(work_0719 work_0719.cu) set_cuda_target_base(work_0719) -add_executable(expt_0722 expt_0722.cpp) +add_executable(expt_0722 expt_220722.cpp) set_target_properties(expt_0722 PROPERTIES LINKER_LANGUAGE CXX) -add_executable(expt_0726 expt_0726.cpp) +add_executable(expt_0726 expt_220726.cpp) set_target_properties(expt_0726 PROPERTIES LINKER_LANGUAGE CXX) -add_executable(expt_0729 expt_0729.cu) +add_executable(expt_0729 expt_220729.cu) set_cuda_target_base(expt_0729) -add_executable(expt_0802 expt_0802.cpp) +add_executable(expt_0802 expt_220802.cpp) target_link_libraries(expt_0802 m stdc++) -add_executable(expt_0804 expt_0804.cu) +add_executable(expt_0804 expt_220804.cu) set_cuda_target_base(expt_0804) -add_executable(expt_0809 expt_0809.cu) +add_executable(expt_0809 expt_220809.cu) set_cuda_target_base(expt_0809) @@ -117,10 +117,10 @@ target_link_libraries(valid_sort_new m stdc++) add_executable(valid_sort valid_sort.cpp) target_link_libraries(valid_sort m stdc++) -add_executable(expt_0821 expt_0821.cpp sortlib.h) +add_executable(expt_0821 expt_220821.cpp sortlib.h) target_link_libraries(expt_0821 m stdc++) -add_executable(expt_0830 expt_0830.cu sortlib.cuh) +add_executable(expt_0830 expt_220830.cu sortlib.cuh) set_cuda_target_base(expt_0830) add_executable(normal_skiplist normal_skiplist.cpp) diff --git a/expt_0528.cu b/expt_0528.cu index d74c93a..fadf3b8 100644 --- a/expt_0528.cu +++ b/expt_0528.cu @@ -729,12 +729,12 @@ int main(int argc, char **argv) { // Launch main kernel + blocks = calcBlocks(search_length); cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, nullptr); - blocks = calcBlocks(search_length); kernel<<>>(cudaOperatorItems, search_length, cudaResult); CudaCheckError(); cudaDeviceSynchronize(); diff --git a/expt_0604.cpp b/expt_0604.cpp index a1dbf26..1109bcb 100644 --- a/expt_0604.cpp +++ b/expt_0604.cpp @@ -52,6 +52,7 @@ int main(int argc, char const *argv[]) { if (limit <= 0) { limit = max_limit; } + auto s = measure("stdio", stdio); auto f = measure("fstream", fstream); std::cout << (s < f ? "stdio" : "fstream") << " win\n"; diff --git a/expt_0702.cpp b/expt_0702.cpp deleted file mode 100644 index f7023f7..0000000 --- a/expt_0702.cpp +++ /dev/null @@ -1,102 +0,0 @@ -#include -#include -#include -#include -#include -#include -#include -#include - -constexpr size_t length = 1048576; -std::vector population_vector, sample_vector, - original_vector; -std::vector> result_storage; - -long double sample_cdf(long double x) { - auto it = lower_bound(sample_vector.begin(), sample_vector.end(), x); - if (it == sample_vector.end()) { - return 1; - } - if (it == sample_vector.begin()) { - return 0; - } - auto it_prev = it - 1; - return ((long double)(it_prev - sample_vector.begin()) + - (long double)(x - *it_prev) / (*it - *it_prev)) / - (long double)(sample_vector.size() - 1); -} - -inline long double safe_ceil(long double value) { - auto c = std::ceil(value); - return c == 0 ? 1 : c; -} - -long double calc_stddev(const long sample_length) { - std::vector elements; - std::vector tmp_vector; - for (auto &element : result_storage) { - elements.push_back(element.size()); - } - tmp_vector.resize(elements.size()); - // auto sum = std::accumulate(elements.begin(), - // elements.end(),decltype(elements)::value_type(0)); size_t sum = length; - // printf("\nsum: %ld\n", sum); - auto mu = length / (long double)elements.size(); - // long double mu = sample_length; - // printf("mu: %Lf\n", mu); - - std::transform(elements.begin(), elements.end(), tmp_vector.begin(), - [&mu](size_t element) { return std::pow(element - mu, 2); }); - - return std::sqrt(std::accumulate(tmp_vector.begin(), tmp_vector.end(), - decltype(tmp_vector)::value_type(0)) / - elements.size()); -} - -void mian(const long sample_length, const long split_size) { - sample_vector = std::vector( - original_vector.begin(), original_vector.begin() + sample_length); - result_storage.resize(sample_length, std::vector()); - - sample_vector.push_back(population_vector[0]); - sample_vector.push_back(*(population_vector.cend() - 1)); - std::sort(sample_vector.begin(), sample_vector.end()); - const auto slice_size = 1.0 / (long double)split_size; - - for (auto &element : original_vector) { - result_storage[(int)safe_ceil(sample_cdf(element) / slice_size) - 1] - .push_back(element); - /*auto tmp = sample_cdf(element) / slice_size; - if (tmp < 1) { - printf("%lld %Lf ", element, tmp); - }*/ - } - for (const auto &element : result_storage) { - printf("%ld ", element.size()); - - } - printf("\n%Lf\n", calc_stddev(sample_length)); - result_storage.clear(); -} - -int main(int argc, char const *argv[]) { - assert((length & 1) == 0); - assert(length > 1024); - - FILE *file = fopen("normal_distribution.txt", "r"); - assert(file); - for (long long i; fscanf(file, "%lld ", &i) != EOF; - population_vector.push_back(i)) - ; - fclose(file); - - assert(population_vector.size() == length); - - original_vector = population_vector; - - std::sort(population_vector.begin(), population_vector.end()); - auto sample_length = 0, split_size = 0; - while (std::cin >> sample_length >> split_size) { - mian(sample_length, split_size); - } -} \ No newline at end of file diff --git a/expt_0702.py b/expt_0702.py deleted file mode 100755 index 5ac42b0..0000000 --- a/expt_0702.py +++ /dev/null @@ -1,23 +0,0 @@ -#!/usr/bin/env python3 -import argparse -import os -import subprocess - -import matplotlib.pyplot as plt - - -def function_pdf(proc: str): - p = subprocess.Popen([proc], stdin=subprocess.PIPE, stdout=subprocess.PIPE) - output, _ = p.communicate(b"1024 1024\n32768 1024\n32768 32768\n") - p.stdin.close() - for index in range(0, len(lines := output.splitlines()), 2): - plt.hist(list(map(int, lines[index].split())), bins=100) - plt.title(f"stddev: {lines[index + 1].decode()}") - plt.show() - - -if __name__ == '__main__': - parser = argparse.ArgumentParser() - parser.add_argument('proc') - matches = parser.parse_args() - function_pdf(matches.proc) diff --git a/expt_0722.cpp b/expt_0722.cpp deleted file mode 100644 index db4b833..0000000 --- a/expt_0722.cpp +++ /dev/null @@ -1,129 +0,0 @@ -// Experimental content: test scale for element cell -/* -Sample output: -scale: 2, max: 7 2: 339 3: 86 4: 27 5: 13 6: 2 7: 2 -scale: 3, max: 6 2: 287 3: 70 4: 18 5: 2 6: 1 -scale: 4, max: 5 2: 257 3: 43 4: 14 5: 3 -scale: 5, max: 5 2: 241 3: 46 4: 6 5: 1 -scale: 6, max: 4 2: 226 3: 26 4: 3 -scale: 7, max: 4 2: 188 3: 26 4: 1 -scale: 8, max: 4 2: 182 3: 20 4: 3 -scale: 9, max: 4 2: 165 3: 17 4: 3 -scale: 10, max: 4 2: 154 3: 16 4: 1 -scale: 11, max: 4 2: 153 3: 9 4: 1 -scale: 12, max: 4 2: 127 3: 11 4: 1 -scale: 13, max: 3 2: 135 3: 5 -scale: 14, max: 3 2: 118 3: 7 -scale: 15, max: 3 2: 129 3: 4 -scale: 16, max: 4 2: 108 3: 5 4: 1 - */ -#include -#include -#include -#include -#include -#include - -constexpr size_t length = 2097152; -std::vector population_vector, sample_vector, - original_vector; -// std::vector> result_storage; -std::vector result_storage; -std::vector counter; - -long double sample_cdf(long double x) { - auto it = lower_bound(sample_vector.begin(), sample_vector.end(), x); - if (it == sample_vector.end()) { - return 1; - } - if (it == sample_vector.begin()) { - return 0; - } - auto it_prev = it - 1; - return ((long double)(it_prev - sample_vector.begin()) + - (long double)(x - *it_prev) / (*it - *it_prev)) / - (long double)(sample_vector.size() - 1); -} - -inline long double safe_ceil(long double value) { - auto c = std::ceil(value); - return c == 0 ? 1 : c; -} - -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); -} - -constexpr long sample_length = 1024; -constexpr long test_size = 2048; - -void mian(long scale_size) { - const long split_size = test_size * scale_size; - counter.resize(8); - - result_storage.resize(split_size, 0); - - const auto slice_size = 1.0 / (long double)(split_size); - - for (long i = 0; i < test_size; i++) { - auto index = - safe_ceil(sample_cdf(original_vector[i + sample_length]) / slice_size) - - 1; - // printf("%Lf\n", index); - result_storage[(int)index]++; - } - auto element_size_max = 0; - for (auto element : result_storage) { - if (element > 1) { - counter[element]++; - } - if (element_size_max < element) { - element_size_max = element; - } - } - printf("scale: %2ld, max: %d ", scale_size, element_size_max); - for (int i = 2; i < 8; i++) { - if (counter[i]) { - printf(" %d: %3d", i, counter[i]); - } - } - puts(""); - result_storage.clear(); - counter.clear(); -} - -int main(int _argc, char const *_argv[]) { - assert((length & 1) == 0); - assert(length > 1024); - - 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("%zu %zu\n", population_vector.size(), length); - - assert(population_vector.size() == length); - - original_vector = population_vector; - - sample_vector = std::vector( - original_vector.begin(), original_vector.begin() + sample_length); - sample_vector.push_back(min_value); - sample_vector.push_back(max_value); - - std::sort(sample_vector.begin(), sample_vector.end()); - - for (int i = 2; i <= 16; i++) { - mian(i); - } -} \ No newline at end of file diff --git a/expt_0726.cpp b/expt_0726.cpp deleted file mode 100644 index 397bc90..0000000 --- a/expt_0726.cpp +++ /dev/null @@ -1,126 +0,0 @@ -#include -#include -#include -#include -#include -#include -#include - -constexpr size_t length = 2097152; -std::vector population_vector, sample_vector, - original_vector; -// std::vector> result_storage; -std::vector result_storage; - -long double sample_cdf(long double x) { - auto it = lower_bound(sample_vector.begin(), sample_vector.end(), x); - if (it == sample_vector.end()) { - return 1; - } - if (it == sample_vector.begin()) { - return 0; - } - auto it_prev = it - 1; - return ((long double)(it_prev - sample_vector.begin()) + - (long double)(x - *it_prev) / (*it - *it_prev)) / - (long double)(sample_vector.size() - 1); -} - -inline long double safe_ceil(long double value) { - auto c = std::ceil(value); - return c == 0 ? 1 : c; -} - -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); -} - -constexpr long sample_length = 1024; -constexpr long test_size = 1000000; - -void mian(long scale_size) { - const long split_size = test_size * scale_size; - result_storage.clear(); - result_storage.resize(split_size, false); - - const auto slice_size = 1.0 / (double)(split_size); - - for (long i = 0; i < test_size; i++) { - auto index = (int)safe_ceil(sample_cdf(original_vector[i + sample_length]) / - slice_size) - - 1; - // printf("%Lf\n", index); - if (result_storage[index]) { - while (result_storage[++index]) { - assert(index < split_size); - } - } - result_storage[index] = true; - } -} - -int main(int _argc, char const *_argv[]) { - assert((length & 1) == 0); - assert(length > 1024); - - 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); - - original_vector = population_vector; - - sample_vector = std::vector( - original_vector.begin(), original_vector.begin() + sample_length); - sample_vector.push_back(min_value); - sample_vector.push_back(max_value); - - std::sort(sample_vector.begin(), sample_vector.end()); - - for (int i = 2; i <= 16; i++) { - - auto start = std::chrono::high_resolution_clock::now(); - mian(i); - auto end = std::chrono::high_resolution_clock::now(); - - auto element_size_max = 0; - for (auto element : result_storage) { - if (element_size_max < element) { - element_size_max = element; - } - } - assert(element_size_max == 1); - printf("scale: %d, time spend: %ldms\n", i, - std::chrono::duration_cast(end - start) - .count()); - } -} -/* - * Sample output: -scale: 2, time spend: 6749258ms -scale: 3, time spend: 1721671ms -scale: 4, time spend: 611112ms -scale: 5, time spend: 353036ms -scale: 6, time spend: 266654ms -scale: 7, time spend: 267346ms -scale: 8, time spend: 267377ms -scale: 9, time spend: 260396ms -scale: 10, time spend: 267724ms -scale: 11, time spend: 263023ms -scale: 12, time spend: 270249ms -scale: 13, time spend: 265440ms -scale: 14, time spend: 266876ms -scale: 15, time spend: 270683ms -scale: 16, time spend: 273216ms - */ \ No newline at end of file diff --git a/expt_0729.cu b/expt_0729.cu deleted file mode 100644 index 5660c6c..0000000 --- a/expt_0729.cu +++ /dev/null @@ -1,221 +0,0 @@ -// Experimental content: test default data layout and scale performance -#include -#include -#include -#include -#include -#include - -constexpr size_t length = 2097152; -std::vector population_vector, sample_vector; - -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, key_type *end, - const key_type val) { - auto begin = start; - key_type *last_known_point = nullptr; - while (begin < end) { - auto mid = (end - begin) / 2; - auto mid_val = *(start + mid); - if (val == mid_val) { - return start + mid; - } else if (val > mid_val) { - begin = begin + mid + 1; - } else { - end = end - mid - 1; - } - last_known_point = begin; - } - return last_known_point; -} - -__device__ double sample_cdf(double x) { - auto it = - cudaBinarySearch(cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH + 2, 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("init storage\n"); -#ifdef TEST_BOUNDS - index_max = 0; - index_min = 0x7fffffff; -#endif -} - -__global__ void init(key_type *sample_item, key_type *population_item) { - cudaSampleItem = sample_item; - cudaPopulationItem = population_item; - // 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( - population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH); - 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 = nullptr; - cudaMalloc(&cudaSample, sizeof(key_type) * (SAMPLE_LENGTH + 2)); - cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE); - cudaMemcpy(cudaSample, &sample_vector[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)); - - init<<<1, 1>>>(cudaSample, cudaPopulation); - cudaDeviceSynchronize(); - - dim3 grid_dim = 16, block_dim = 64; - - unsigned long step = test_size / (grid_dim.x * block_dim.x); - printf("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<<>>(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; -} -/* - * Sample output: -step: 1024 -scale: 2 time: 95.163521 -scale: 3 time: 98.489342 -scale: 4 time: 71.021828 -scale: 5 time: 69.743874 -scale: 6 time: 68.885506 -scale: 7 time: 68.294655 -scale: 8 time: 68.837410 - */ \ No newline at end of file diff --git a/expt_0802.cpp b/expt_0802.cpp deleted file mode 100644 index 094da65..0000000 --- a/expt_0802.cpp +++ /dev/null @@ -1,76 +0,0 @@ -// Experimental content: Test sort algorithm -#include -#include -#include -#include -#include - -std::vector> a; -std::vector result; -constexpr size_t LENGTH = 1024; - -void calculate(int l, int r, int dep) { - if (l > r) - return; - int mid = (l + r) >> 1; - a[dep].push_back(mid + 1); - calculate(l, mid - 1, dep + 1); - calculate(mid + 1, r, dep + 1); -} - -size_t calculate_location(size_t index) { - int step = (int)std::log2(++index) + 1; - auto location = - LENGTH / (1 << step) * ((index & ((1 << (step - 1)) - 1)) * 2 + 1); - return location ? location : LENGTH; -} -typedef int key_type; - -const key_type *cudaBinarySearch(key_type *const start, const key_type *end, - const key_type val) { - int step_limit = (int)std::log2(LENGTH) + 1; - 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 = get_son_from_step(son, (*last_known_point > val)); - son = son * 2 + -((*last_known_point - val) >> 31); - last_known_point = next_level_start + son; - printf("%d\n", *last_known_point); - } - return last_known_point; -} - -int main() { - int n = LENGTH; - - int c = 0; - for (int k = n; k; k >>= 1) - c++; - a.resize(c); - calculate(0, n - 1, 0); - for (auto &i : a) { - for (int &j : i) { - // cout << j << ' '; - result.push_back(j); - } - } - - for (int i = 0; i < 32; i++) { - printf("%3d ", i); - } - - puts(""); - - for (int i = 0; i < 32; i++) { - printf("%3d ", result[i]); - assert(calculate_location(i) == result[i]); - } - puts(""); - cudaBinarySearch(&result[0], &*result.end(), 928); - return 0; -} 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 -#include -#include -#include -#include -#include - -constexpr size_t length = 2097152; -std::vector 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( - 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<<>>(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 diff --git a/expt_0809.cu b/expt_0809.cu deleted file mode 100644 index 44d9992..0000000 --- a/expt_0809.cu +++ /dev/null @@ -1,219 +0,0 @@ -// Experimental content: test scale performance -/* - * Sample output: -scale: 2 time: 97.101822 -scale: 3 time: 80.391197 -scale: 4 time: 66.238525 -scale: 5 time: 70.442848 -scale: 6 time: 69.601219 -scale: 7 time: 69.934914 -scale: 8 time: 68.154366 - */ -#include -#include -#include -#include -#include -#include -#include - -constexpr size_t length = 2097152; -std::vector population_vector, sample_vector; - -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, key_type *end, - const key_type val) { - auto begin = start; - key_type *last_known_point = nullptr; - while (begin < end) { - auto mid = (end - begin) / 2; - auto mid_val = *(start + mid); - if (val == mid_val) { - return start + mid; - } - begin += ((mid_val - val) >> 63) & (mid + 1); - end -= (~((mid_val - val) >> 63)) & (mid + 1); - last_known_point = begin; - } - return last_known_point; -} - -__device__ double sample_cdf(double x) { - auto it = - cudaBinarySearch(cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH + 2, 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("init storage\n"); -#ifdef TEST_BOUNDS - index_max = 0; - index_min = 0x7fffffff; -#endif -} - -__global__ void init(key_type *sample_item, key_type *population_item) { - cudaSampleItem = sample_item; - cudaPopulationItem = population_item; - // 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( - population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH); - 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 = nullptr; - cudaMalloc(&cudaSample, sizeof(key_type) * (SAMPLE_LENGTH + 2)); - cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE); - cudaMemcpy(cudaSample, &sample_vector[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)); - - init<<<1, 1>>>(cudaSample, cudaPopulation); - cudaDeviceSynchronize(); - - dim3 grid_dim = 16, block_dim = 64; - - unsigned long step = test_size / (grid_dim.x * block_dim.x); - // printf("(0809)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<<>>(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 diff --git a/expt_0821.cpp b/expt_0821.cpp deleted file mode 100644 index 6bbcef0..0000000 --- a/expt_0821.cpp +++ /dev/null @@ -1,81 +0,0 @@ -// Experimental content: reliability test -#include "sortlib.h" -#include -#include -#include -#include -#include - -typedef unsigned long long key_type; -std::vector population_vector, sample_vector; -constexpr size_t SAMPLE_LENGTH = 1023, TEST_LENGTH = 1024; -constexpr int MOVE_OFFSET = sizeof(key_type) * 8 - 1; - -auto custom_sort = CustomSort(SAMPLE_LENGTH, MOVE_OFFSET); - -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); -} - -void valid_sort(std::vector &tmp) { - auto iter_end = population_vector.begin() + SAMPLE_LENGTH + 1024; - for (auto it = population_vector.begin() + SAMPLE_LENGTH; it != iter_end; - it++) { - // long std_search = - // std::lower_bound(tmp.begin(), tmp.end(), *it) - tmp.begin(); - long std_search = custom_sort.original_binary_search( - tmp.data(), tmp.data() + tmp.size(), *it) - - tmp.data(); - - auto cuda_binary_search1 = - custom_sort.binary_search(sample_vector.data(), *it) - - sample_vector.data(); - auto cuda_binary_search = - custom_sort.calculate_location(cuda_binary_search1) - 1; - /*printf("%ld:%ld %llu:%llu\n", std_search, cuda_binary_search, - *(tmp.begin() + std_search), - *(cuda_binary_search + sample_vector.begin()));*/ - printf("%ld:%ld %llu %llu %llu\n", std_search, cuda_binary_search, *it, - *(tmp.data() + std_search), - *(cuda_binary_search1 + sample_vector.data())); - // assert(*it > *(cuda_binary_search1 + sample_vector.data())); - // assert(std_search == cuda_binary_search); - } -} - -int main() { - FILE *file = fopen("normal_distribution.txt", "r"); - assert(file); - for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i)) - ; - fclose(file); - - std::vector tmp(population_vector.begin(), - population_vector.begin() + SAMPLE_LENGTH - 2); - tmp.push_back(max_value); - tmp.push_back(min_value); - - std::sort(tmp.begin(), tmp.end()); - - sample_vector.resize(SAMPLE_LENGTH); - for (size_t i = 0; i < SAMPLE_LENGTH; i++) { - sample_vector[custom_sort.calculate_location_inverse(i) - 1] = tmp[i]; - auto location = custom_sort.calculate_location(i); - auto location2 = custom_sort.calculate_location_inverse(location - 1); - // printf("%zu %zu\n", location, location2); - assert(location2 - 1 == i); - } - // valid_sort(tmp); - /*for (auto element : sample_vector) { - printf("%d ", element); - }*/ - valid_sort(tmp); - puts(""); -} \ No newline at end of file diff --git a/expt_0830.cu b/expt_0830.cu deleted file mode 100644 index fc05092..0000000 --- a/expt_0830.cu +++ /dev/null @@ -1,290 +0,0 @@ -#define ENABLE_SORT_TEST -#define SORT_FIRST_VERSION -#include "sortlib.cuh" - -#include -#include -#include -#include - -std::vector population_vector, 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 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, *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; -#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, cudaPopulationItem[tid]) / - slice_size) - - 1; - // 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_real_binary(unsigned long step, const double slice_size, - const unsigned long split_size) { - auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(key_type) * 8); - // printf("%d\n", custom_sort.MOVE_OFFSET); - // printf("kernel2\n"); - - for (int i = 0; i < step; i++) { - auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i; - // printf("%lu ", tid); - - // assert(tid < TEST_LENGTH); - auto index = (int)(custom_sort.sample_cdf(cudaSampleItem, - cudaSampleItem + SAMPLE_LENGTH, - cudaPopulationItem[tid]) / - slice_size); - // printf("%llu %d\n", cudaPopulationItem[tid], index); - /*if (cdf_result[index]) { - while (cdf_result[++index]) { - assert(index < split_size); - } - }*/ - cdf_result[index] = true; -#ifdef TEST_BOUNDS - atomicAdd(&insert_value, 1); - - 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 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 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)); -} - -__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(16, 64), dim_pair_type(32, 32)}; - -void run_kernel(size_t test_size, bool custom = false) { - 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", custom ? "custom " : "", 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); - if (custom) - kernel<<>>(step, slice_size, split_size); - else - kernel2_real_binary<<>>(step, slice_size, - split_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(); - - initSample<<<1, 1>>>(cudaSample, cudaPopulation); - cudaDeviceSynchronize(); - run_kernel(test_size); - - initCustomSample<<<1, 1>>>(); - cudaDeviceSynchronize(); - - run_kernel(test_size, true); - freeStorage<<<1, 1>>>(); - cudaDeviceSynchronize(); -} - -/* - * Sample output: -test size: 32768 -current dim: 16 64 32 -scale: 2 time: 0.100352 -scale: 4 time: 0.102176 -scale: 6 time: 0.104448 -scale: 8 time: 0.100352 -scale: 10 time: 0.100512 -scale: 12 time: 0.104384 -scale: 14 time: 0.101152 -scale: 16 time: 0.099232 -current dim: 32 32 32 -scale: 2 time: 0.099168 -scale: 4 time: 0.101376 -scale: 6 time: 0.106496 -scale: 8 time: 0.099584 -scale: 10 time: 0.111616 -scale: 12 time: 0.106464 -scale: 14 time: 0.100352 -scale: 16 time: 0.108544 -custom current dim: 16 64 32 -scale: 2 time: 0.113664 -scale: 4 time: 0.112512 -scale: 6 time: 0.115712 -scale: 8 time: 0.114720 -scale: 10 time: 0.107232 -scale: 12 time: 0.109568 -scale: 14 time: 0.107328 -scale: 16 time: 0.107520 -custom current dim: 32 32 32 -scale: 2 time: 0.106496 -scale: 4 time: 0.106368 -scale: 6 time: 0.108384 -scale: 8 time: 0.108384 -scale: 10 time: 0.109792 -scale: 12 time: 0.112704 -scale: 14 time: 0.108672 -scale: 16 time: 0.122848 - - */ \ No newline at end of file diff --git a/expt_220702.cpp b/expt_220702.cpp new file mode 100644 index 0000000..f7023f7 --- /dev/null +++ b/expt_220702.cpp @@ -0,0 +1,102 @@ +#include +#include +#include +#include +#include +#include +#include +#include + +constexpr size_t length = 1048576; +std::vector population_vector, sample_vector, + original_vector; +std::vector> result_storage; + +long double sample_cdf(long double x) { + auto it = lower_bound(sample_vector.begin(), sample_vector.end(), x); + if (it == sample_vector.end()) { + return 1; + } + if (it == sample_vector.begin()) { + return 0; + } + auto it_prev = it - 1; + return ((long double)(it_prev - sample_vector.begin()) + + (long double)(x - *it_prev) / (*it - *it_prev)) / + (long double)(sample_vector.size() - 1); +} + +inline long double safe_ceil(long double value) { + auto c = std::ceil(value); + return c == 0 ? 1 : c; +} + +long double calc_stddev(const long sample_length) { + std::vector elements; + std::vector tmp_vector; + for (auto &element : result_storage) { + elements.push_back(element.size()); + } + tmp_vector.resize(elements.size()); + // auto sum = std::accumulate(elements.begin(), + // elements.end(),decltype(elements)::value_type(0)); size_t sum = length; + // printf("\nsum: %ld\n", sum); + auto mu = length / (long double)elements.size(); + // long double mu = sample_length; + // printf("mu: %Lf\n", mu); + + std::transform(elements.begin(), elements.end(), tmp_vector.begin(), + [&mu](size_t element) { return std::pow(element - mu, 2); }); + + return std::sqrt(std::accumulate(tmp_vector.begin(), tmp_vector.end(), + decltype(tmp_vector)::value_type(0)) / + elements.size()); +} + +void mian(const long sample_length, const long split_size) { + sample_vector = std::vector( + original_vector.begin(), original_vector.begin() + sample_length); + result_storage.resize(sample_length, std::vector()); + + sample_vector.push_back(population_vector[0]); + sample_vector.push_back(*(population_vector.cend() - 1)); + std::sort(sample_vector.begin(), sample_vector.end()); + const auto slice_size = 1.0 / (long double)split_size; + + for (auto &element : original_vector) { + result_storage[(int)safe_ceil(sample_cdf(element) / slice_size) - 1] + .push_back(element); + /*auto tmp = sample_cdf(element) / slice_size; + if (tmp < 1) { + printf("%lld %Lf ", element, tmp); + }*/ + } + for (const auto &element : result_storage) { + printf("%ld ", element.size()); + + } + printf("\n%Lf\n", calc_stddev(sample_length)); + result_storage.clear(); +} + +int main(int argc, char const *argv[]) { + assert((length & 1) == 0); + assert(length > 1024); + + FILE *file = fopen("normal_distribution.txt", "r"); + assert(file); + for (long long i; fscanf(file, "%lld ", &i) != EOF; + population_vector.push_back(i)) + ; + fclose(file); + + assert(population_vector.size() == length); + + original_vector = population_vector; + + std::sort(population_vector.begin(), population_vector.end()); + auto sample_length = 0, split_size = 0; + while (std::cin >> sample_length >> split_size) { + mian(sample_length, split_size); + } +} \ No newline at end of file diff --git a/expt_220702.py b/expt_220702.py new file mode 100755 index 0000000..5ac42b0 --- /dev/null +++ b/expt_220702.py @@ -0,0 +1,23 @@ +#!/usr/bin/env python3 +import argparse +import os +import subprocess + +import matplotlib.pyplot as plt + + +def function_pdf(proc: str): + p = subprocess.Popen([proc], stdin=subprocess.PIPE, stdout=subprocess.PIPE) + output, _ = p.communicate(b"1024 1024\n32768 1024\n32768 32768\n") + p.stdin.close() + for index in range(0, len(lines := output.splitlines()), 2): + plt.hist(list(map(int, lines[index].split())), bins=100) + plt.title(f"stddev: {lines[index + 1].decode()}") + plt.show() + + +if __name__ == '__main__': + parser = argparse.ArgumentParser() + parser.add_argument('proc') + matches = parser.parse_args() + function_pdf(matches.proc) diff --git a/expt_220722.cpp b/expt_220722.cpp new file mode 100644 index 0000000..db4b833 --- /dev/null +++ b/expt_220722.cpp @@ -0,0 +1,129 @@ +// Experimental content: test scale for element cell +/* +Sample output: +scale: 2, max: 7 2: 339 3: 86 4: 27 5: 13 6: 2 7: 2 +scale: 3, max: 6 2: 287 3: 70 4: 18 5: 2 6: 1 +scale: 4, max: 5 2: 257 3: 43 4: 14 5: 3 +scale: 5, max: 5 2: 241 3: 46 4: 6 5: 1 +scale: 6, max: 4 2: 226 3: 26 4: 3 +scale: 7, max: 4 2: 188 3: 26 4: 1 +scale: 8, max: 4 2: 182 3: 20 4: 3 +scale: 9, max: 4 2: 165 3: 17 4: 3 +scale: 10, max: 4 2: 154 3: 16 4: 1 +scale: 11, max: 4 2: 153 3: 9 4: 1 +scale: 12, max: 4 2: 127 3: 11 4: 1 +scale: 13, max: 3 2: 135 3: 5 +scale: 14, max: 3 2: 118 3: 7 +scale: 15, max: 3 2: 129 3: 4 +scale: 16, max: 4 2: 108 3: 5 4: 1 + */ +#include +#include +#include +#include +#include +#include + +constexpr size_t length = 2097152; +std::vector population_vector, sample_vector, + original_vector; +// std::vector> result_storage; +std::vector result_storage; +std::vector counter; + +long double sample_cdf(long double x) { + auto it = lower_bound(sample_vector.begin(), sample_vector.end(), x); + if (it == sample_vector.end()) { + return 1; + } + if (it == sample_vector.begin()) { + return 0; + } + auto it_prev = it - 1; + return ((long double)(it_prev - sample_vector.begin()) + + (long double)(x - *it_prev) / (*it - *it_prev)) / + (long double)(sample_vector.size() - 1); +} + +inline long double safe_ceil(long double value) { + auto c = std::ceil(value); + return c == 0 ? 1 : c; +} + +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); +} + +constexpr long sample_length = 1024; +constexpr long test_size = 2048; + +void mian(long scale_size) { + const long split_size = test_size * scale_size; + counter.resize(8); + + result_storage.resize(split_size, 0); + + const auto slice_size = 1.0 / (long double)(split_size); + + for (long i = 0; i < test_size; i++) { + auto index = + safe_ceil(sample_cdf(original_vector[i + sample_length]) / slice_size) - + 1; + // printf("%Lf\n", index); + result_storage[(int)index]++; + } + auto element_size_max = 0; + for (auto element : result_storage) { + if (element > 1) { + counter[element]++; + } + if (element_size_max < element) { + element_size_max = element; + } + } + printf("scale: %2ld, max: %d ", scale_size, element_size_max); + for (int i = 2; i < 8; i++) { + if (counter[i]) { + printf(" %d: %3d", i, counter[i]); + } + } + puts(""); + result_storage.clear(); + counter.clear(); +} + +int main(int _argc, char const *_argv[]) { + assert((length & 1) == 0); + assert(length > 1024); + + 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("%zu %zu\n", population_vector.size(), length); + + assert(population_vector.size() == length); + + original_vector = population_vector; + + sample_vector = std::vector( + original_vector.begin(), original_vector.begin() + sample_length); + sample_vector.push_back(min_value); + sample_vector.push_back(max_value); + + std::sort(sample_vector.begin(), sample_vector.end()); + + for (int i = 2; i <= 16; i++) { + mian(i); + } +} \ No newline at end of file diff --git a/expt_220726.cpp b/expt_220726.cpp new file mode 100644 index 0000000..397bc90 --- /dev/null +++ b/expt_220726.cpp @@ -0,0 +1,126 @@ +#include +#include +#include +#include +#include +#include +#include + +constexpr size_t length = 2097152; +std::vector population_vector, sample_vector, + original_vector; +// std::vector> result_storage; +std::vector result_storage; + +long double sample_cdf(long double x) { + auto it = lower_bound(sample_vector.begin(), sample_vector.end(), x); + if (it == sample_vector.end()) { + return 1; + } + if (it == sample_vector.begin()) { + return 0; + } + auto it_prev = it - 1; + return ((long double)(it_prev - sample_vector.begin()) + + (long double)(x - *it_prev) / (*it - *it_prev)) / + (long double)(sample_vector.size() - 1); +} + +inline long double safe_ceil(long double value) { + auto c = std::ceil(value); + return c == 0 ? 1 : c; +} + +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); +} + +constexpr long sample_length = 1024; +constexpr long test_size = 1000000; + +void mian(long scale_size) { + const long split_size = test_size * scale_size; + result_storage.clear(); + result_storage.resize(split_size, false); + + const auto slice_size = 1.0 / (double)(split_size); + + for (long i = 0; i < test_size; i++) { + auto index = (int)safe_ceil(sample_cdf(original_vector[i + sample_length]) / + slice_size) - + 1; + // printf("%Lf\n", index); + if (result_storage[index]) { + while (result_storage[++index]) { + assert(index < split_size); + } + } + result_storage[index] = true; + } +} + +int main(int _argc, char const *_argv[]) { + assert((length & 1) == 0); + assert(length > 1024); + + 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); + + original_vector = population_vector; + + sample_vector = std::vector( + original_vector.begin(), original_vector.begin() + sample_length); + sample_vector.push_back(min_value); + sample_vector.push_back(max_value); + + std::sort(sample_vector.begin(), sample_vector.end()); + + for (int i = 2; i <= 16; i++) { + + auto start = std::chrono::high_resolution_clock::now(); + mian(i); + auto end = std::chrono::high_resolution_clock::now(); + + auto element_size_max = 0; + for (auto element : result_storage) { + if (element_size_max < element) { + element_size_max = element; + } + } + assert(element_size_max == 1); + printf("scale: %d, time spend: %ldms\n", i, + std::chrono::duration_cast(end - start) + .count()); + } +} +/* + * Sample output: +scale: 2, time spend: 6749258ms +scale: 3, time spend: 1721671ms +scale: 4, time spend: 611112ms +scale: 5, time spend: 353036ms +scale: 6, time spend: 266654ms +scale: 7, time spend: 267346ms +scale: 8, time spend: 267377ms +scale: 9, time spend: 260396ms +scale: 10, time spend: 267724ms +scale: 11, time spend: 263023ms +scale: 12, time spend: 270249ms +scale: 13, time spend: 265440ms +scale: 14, time spend: 266876ms +scale: 15, time spend: 270683ms +scale: 16, time spend: 273216ms + */ \ No newline at end of file diff --git a/expt_220729.cu b/expt_220729.cu new file mode 100644 index 0000000..5660c6c --- /dev/null +++ b/expt_220729.cu @@ -0,0 +1,221 @@ +// Experimental content: test default data layout and scale performance +#include +#include +#include +#include +#include +#include + +constexpr size_t length = 2097152; +std::vector population_vector, sample_vector; + +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, key_type *end, + const key_type val) { + auto begin = start; + key_type *last_known_point = nullptr; + while (begin < end) { + auto mid = (end - begin) / 2; + auto mid_val = *(start + mid); + if (val == mid_val) { + return start + mid; + } else if (val > mid_val) { + begin = begin + mid + 1; + } else { + end = end - mid - 1; + } + last_known_point = begin; + } + return last_known_point; +} + +__device__ double sample_cdf(double x) { + auto it = + cudaBinarySearch(cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH + 2, 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("init storage\n"); +#ifdef TEST_BOUNDS + index_max = 0; + index_min = 0x7fffffff; +#endif +} + +__global__ void init(key_type *sample_item, key_type *population_item) { + cudaSampleItem = sample_item; + cudaPopulationItem = population_item; + // 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( + population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH); + 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 = nullptr; + cudaMalloc(&cudaSample, sizeof(key_type) * (SAMPLE_LENGTH + 2)); + cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE); + cudaMemcpy(cudaSample, &sample_vector[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)); + + init<<<1, 1>>>(cudaSample, cudaPopulation); + cudaDeviceSynchronize(); + + dim3 grid_dim = 16, block_dim = 64; + + unsigned long step = test_size / (grid_dim.x * block_dim.x); + printf("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<<>>(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; +} +/* + * Sample output: +step: 1024 +scale: 2 time: 95.163521 +scale: 3 time: 98.489342 +scale: 4 time: 71.021828 +scale: 5 time: 69.743874 +scale: 6 time: 68.885506 +scale: 7 time: 68.294655 +scale: 8 time: 68.837410 + */ \ No newline at end of file diff --git a/expt_220802.cpp b/expt_220802.cpp new file mode 100644 index 0000000..094da65 --- /dev/null +++ b/expt_220802.cpp @@ -0,0 +1,76 @@ +// Experimental content: Test sort algorithm +#include +#include +#include +#include +#include + +std::vector> a; +std::vector result; +constexpr size_t LENGTH = 1024; + +void calculate(int l, int r, int dep) { + if (l > r) + return; + int mid = (l + r) >> 1; + a[dep].push_back(mid + 1); + calculate(l, mid - 1, dep + 1); + calculate(mid + 1, r, dep + 1); +} + +size_t calculate_location(size_t index) { + int step = (int)std::log2(++index) + 1; + auto location = + LENGTH / (1 << step) * ((index & ((1 << (step - 1)) - 1)) * 2 + 1); + return location ? location : LENGTH; +} +typedef int key_type; + +const key_type *cudaBinarySearch(key_type *const start, const key_type *end, + const key_type val) { + int step_limit = (int)std::log2(LENGTH) + 1; + 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 = get_son_from_step(son, (*last_known_point > val)); + son = son * 2 + -((*last_known_point - val) >> 31); + last_known_point = next_level_start + son; + printf("%d\n", *last_known_point); + } + return last_known_point; +} + +int main() { + int n = LENGTH; + + int c = 0; + for (int k = n; k; k >>= 1) + c++; + a.resize(c); + calculate(0, n - 1, 0); + for (auto &i : a) { + for (int &j : i) { + // cout << j << ' '; + result.push_back(j); + } + } + + for (int i = 0; i < 32; i++) { + printf("%3d ", i); + } + + puts(""); + + for (int i = 0; i < 32; i++) { + printf("%3d ", result[i]); + assert(calculate_location(i) == result[i]); + } + puts(""); + cudaBinarySearch(&result[0], &*result.end(), 928); + return 0; +} diff --git a/expt_220804.cu b/expt_220804.cu new file mode 100644 index 0000000..66f66ec --- /dev/null +++ b/expt_220804.cu @@ -0,0 +1,222 @@ +// Experimental content: Test cdf performance (Deprecated) +#include +#include +#include +#include +#include +#include + +constexpr size_t length = 2097152; +std::vector 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( + 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<<>>(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 diff --git a/expt_220809.cu b/expt_220809.cu new file mode 100644 index 0000000..44d9992 --- /dev/null +++ b/expt_220809.cu @@ -0,0 +1,219 @@ +// Experimental content: test scale performance +/* + * Sample output: +scale: 2 time: 97.101822 +scale: 3 time: 80.391197 +scale: 4 time: 66.238525 +scale: 5 time: 70.442848 +scale: 6 time: 69.601219 +scale: 7 time: 69.934914 +scale: 8 time: 68.154366 + */ +#include +#include +#include +#include +#include +#include +#include + +constexpr size_t length = 2097152; +std::vector population_vector, sample_vector; + +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, key_type *end, + const key_type val) { + auto begin = start; + key_type *last_known_point = nullptr; + while (begin < end) { + auto mid = (end - begin) / 2; + auto mid_val = *(start + mid); + if (val == mid_val) { + return start + mid; + } + begin += ((mid_val - val) >> 63) & (mid + 1); + end -= (~((mid_val - val) >> 63)) & (mid + 1); + last_known_point = begin; + } + return last_known_point; +} + +__device__ double sample_cdf(double x) { + auto it = + cudaBinarySearch(cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH + 2, 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("init storage\n"); +#ifdef TEST_BOUNDS + index_max = 0; + index_min = 0x7fffffff; +#endif +} + +__global__ void init(key_type *sample_item, key_type *population_item) { + cudaSampleItem = sample_item; + cudaPopulationItem = population_item; + // 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( + population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH); + 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 = nullptr; + cudaMalloc(&cudaSample, sizeof(key_type) * (SAMPLE_LENGTH + 2)); + cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE); + cudaMemcpy(cudaSample, &sample_vector[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)); + + init<<<1, 1>>>(cudaSample, cudaPopulation); + cudaDeviceSynchronize(); + + dim3 grid_dim = 16, block_dim = 64; + + unsigned long step = test_size / (grid_dim.x * block_dim.x); + // printf("(0809)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<<>>(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 diff --git a/expt_220821.cpp b/expt_220821.cpp new file mode 100644 index 0000000..6bbcef0 --- /dev/null +++ b/expt_220821.cpp @@ -0,0 +1,81 @@ +// Experimental content: reliability test +#include "sortlib.h" +#include +#include +#include +#include +#include + +typedef unsigned long long key_type; +std::vector population_vector, sample_vector; +constexpr size_t SAMPLE_LENGTH = 1023, TEST_LENGTH = 1024; +constexpr int MOVE_OFFSET = sizeof(key_type) * 8 - 1; + +auto custom_sort = CustomSort(SAMPLE_LENGTH, MOVE_OFFSET); + +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); +} + +void valid_sort(std::vector &tmp) { + auto iter_end = population_vector.begin() + SAMPLE_LENGTH + 1024; + for (auto it = population_vector.begin() + SAMPLE_LENGTH; it != iter_end; + it++) { + // long std_search = + // std::lower_bound(tmp.begin(), tmp.end(), *it) - tmp.begin(); + long std_search = custom_sort.original_binary_search( + tmp.data(), tmp.data() + tmp.size(), *it) - + tmp.data(); + + auto cuda_binary_search1 = + custom_sort.binary_search(sample_vector.data(), *it) - + sample_vector.data(); + auto cuda_binary_search = + custom_sort.calculate_location(cuda_binary_search1) - 1; + /*printf("%ld:%ld %llu:%llu\n", std_search, cuda_binary_search, + *(tmp.begin() + std_search), + *(cuda_binary_search + sample_vector.begin()));*/ + printf("%ld:%ld %llu %llu %llu\n", std_search, cuda_binary_search, *it, + *(tmp.data() + std_search), + *(cuda_binary_search1 + sample_vector.data())); + // assert(*it > *(cuda_binary_search1 + sample_vector.data())); + // assert(std_search == cuda_binary_search); + } +} + +int main() { + FILE *file = fopen("normal_distribution.txt", "r"); + assert(file); + for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i)) + ; + fclose(file); + + std::vector tmp(population_vector.begin(), + population_vector.begin() + SAMPLE_LENGTH - 2); + tmp.push_back(max_value); + tmp.push_back(min_value); + + std::sort(tmp.begin(), tmp.end()); + + sample_vector.resize(SAMPLE_LENGTH); + for (size_t i = 0; i < SAMPLE_LENGTH; i++) { + sample_vector[custom_sort.calculate_location_inverse(i) - 1] = tmp[i]; + auto location = custom_sort.calculate_location(i); + auto location2 = custom_sort.calculate_location_inverse(location - 1); + // printf("%zu %zu\n", location, location2); + assert(location2 - 1 == i); + } + // valid_sort(tmp); + /*for (auto element : sample_vector) { + printf("%d ", element); + }*/ + valid_sort(tmp); + puts(""); +} \ No newline at end of file diff --git a/expt_220830.cu b/expt_220830.cu new file mode 100644 index 0000000..fc05092 --- /dev/null +++ b/expt_220830.cu @@ -0,0 +1,290 @@ +#define ENABLE_SORT_TEST +#define SORT_FIRST_VERSION +#include "sortlib.cuh" + +#include +#include +#include +#include + +std::vector population_vector, 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 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, *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; +#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, cudaPopulationItem[tid]) / + slice_size) - + 1; + // 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_real_binary(unsigned long step, const double slice_size, + const unsigned long split_size) { + auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(key_type) * 8); + // printf("%d\n", custom_sort.MOVE_OFFSET); + // printf("kernel2\n"); + + for (int i = 0; i < step; i++) { + auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i; + // printf("%lu ", tid); + + // assert(tid < TEST_LENGTH); + auto index = (int)(custom_sort.sample_cdf(cudaSampleItem, + cudaSampleItem + SAMPLE_LENGTH, + cudaPopulationItem[tid]) / + slice_size); + // printf("%llu %d\n", cudaPopulationItem[tid], index); + /*if (cdf_result[index]) { + while (cdf_result[++index]) { + assert(index < split_size); + } + }*/ + cdf_result[index] = true; +#ifdef TEST_BOUNDS + atomicAdd(&insert_value, 1); + + 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 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 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)); +} + +__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(16, 64), dim_pair_type(32, 32)}; + +void run_kernel(size_t test_size, bool custom = false) { + 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", custom ? "custom " : "", 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); + if (custom) + kernel<<>>(step, slice_size, split_size); + else + kernel2_real_binary<<>>(step, slice_size, + split_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(); + + initSample<<<1, 1>>>(cudaSample, cudaPopulation); + cudaDeviceSynchronize(); + run_kernel(test_size); + + initCustomSample<<<1, 1>>>(); + cudaDeviceSynchronize(); + + run_kernel(test_size, true); + freeStorage<<<1, 1>>>(); + cudaDeviceSynchronize(); +} + +/* + * Sample output: +test size: 32768 +current dim: 16 64 32 +scale: 2 time: 0.100352 +scale: 4 time: 0.102176 +scale: 6 time: 0.104448 +scale: 8 time: 0.100352 +scale: 10 time: 0.100512 +scale: 12 time: 0.104384 +scale: 14 time: 0.101152 +scale: 16 time: 0.099232 +current dim: 32 32 32 +scale: 2 time: 0.099168 +scale: 4 time: 0.101376 +scale: 6 time: 0.106496 +scale: 8 time: 0.099584 +scale: 10 time: 0.111616 +scale: 12 time: 0.106464 +scale: 14 time: 0.100352 +scale: 16 time: 0.108544 +custom current dim: 16 64 32 +scale: 2 time: 0.113664 +scale: 4 time: 0.112512 +scale: 6 time: 0.115712 +scale: 8 time: 0.114720 +scale: 10 time: 0.107232 +scale: 12 time: 0.109568 +scale: 14 time: 0.107328 +scale: 16 time: 0.107520 +custom current dim: 32 32 32 +scale: 2 time: 0.106496 +scale: 4 time: 0.106368 +scale: 6 time: 0.108384 +scale: 8 time: 0.108384 +scale: 10 time: 0.109792 +scale: 12 time: 0.112704 +scale: 14 time: 0.108672 +scale: 16 time: 0.122848 + + */ \ No newline at end of file -- cgit v1.3.1