diff options
| author | KunoiSayami <[email protected]> | 2023-06-12 00:51:30 +0800 |
|---|---|---|
| committer | KunoiSayami <[email protected]> | 2023-06-12 00:51:30 +0800 |
| commit | e7de21daf6c674633ef1159c2fd042523c1ec32e (patch) | |
| tree | e9fea9b2611fb832898900e7fb698e8a1a4c28f3 /expt_0809.cu | |
| parent | 5f38ca067108eea5a8c334bd8cbe5866300d585c (diff) | |
refector: Rename last year experimental
Signed-off-by: KunoiSayami <[email protected]>
Diffstat (limited to 'expt_0809.cu')
| -rw-r--r-- | expt_0809.cu | 219 |
1 files changed, 0 insertions, 219 deletions
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 <algorithm> -#include <cassert> -#include <cstdio> -#include <cstring> -#include <iostream> -#include <type_traits> -#include <vector> - -constexpr size_t length = 2097152; -std::vector<unsigned long long> 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<unsigned long long>( - 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<<<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 |
