diff options
| -rw-r--r-- | CMakeLists.txt | 8 | ||||
| -rwxr-xr-x | analysis0617.py | 3 | ||||
| -rw-r--r-- | expt_0525.cu | 48 | ||||
| -rw-r--r-- | expt_0528.cu | 12 | ||||
| -rw-r--r-- | expt_0618_2.cpp | 33 | ||||
| -rw-r--r-- | expt_0618_3.cu | 117 |
6 files changed, 164 insertions, 57 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index d0292ff..bef2d9d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -177,3 +177,11 @@ add_executable(expt_0604 expt_0604.cpp) set_target_properties(expt_0604 PROPERTIES LINKER_LANGUAGE CXX) +add_executable(expt_0618 expt_0618.cu) +set_cuda_target_base(expt_0618) + +add_executable(expt_0618_2 expt_0618_2.cpp) +set_target_properties(expt_0618_2 PROPERTIES LINKER_LANGUAGE CXX) + +add_executable(expt_0618_3 expt_0618_3.cu) +set_cuda_target_base(expt_0618_3) diff --git a/analysis0617.py b/analysis0617.py index 58356f8..46a36b9 100755 --- a/analysis0617.py +++ b/analysis0617.py @@ -76,6 +76,7 @@ async def main(matches: argparse.Namespace) -> None: if search not in result[sample]: result[sample].update({search: {'normal': [], 'custom': []}}) for _ in range(3): + await asyncio.sleep(1) retries = 3 while (normal := await run_exec(matches.exec2, search, insertion)) is None: retries -= 1 @@ -83,6 +84,7 @@ async def main(matches: argparse.Namespace) -> None: if not retries: break total_run += 4 - retries + await asyncio.sleep(1) retries = 3 while (custom := await run_custom_exec(matches.exec1, sample, search, insertion)) is None: retries -= 1 @@ -106,6 +108,7 @@ if __name__ == '__main__': run_.add_argument("exec2") run_.add_argument("sample_limit") run_.add_argument("limit") + run_.add_argument("--test-times", default=3) parser_ = arg_.parse_args() if parser_.sub == 'run': asyncio.run(main(parser_)) diff --git a/expt_0525.cu b/expt_0525.cu index 5af1fc0..8e7cd10 100644 --- a/expt_0525.cu +++ b/expt_0525.cu @@ -204,52 +204,6 @@ public: } }; -/*struct MapNode { - LL key; - // Node *point[MAX_LEVEL + 1]; - Node *point; -}; - -class MemMap { -public: - size_t size; - size_t real_size; - MapNode *store; - - MemMap() : size(0), store(nullptr), real_size(0) {} - - __device__ bool insert(MapNode node) { - bool need_extend = this->size + 1 > this->real_size; - if (need_extend) { - bool need_copy = this->real_size == 0; - if (!need_copy) { - this->real_size += 1; - } - this->real_size *= 2; - MapNode *old = this->store; - this->store = new MapNode[this->real_size]; - if (need_copy) { - memcpy(this->store, old, this->real_size * sizeof(MapNode *)); - } - delete[] old; - } - // need sort after insert - this->store[size] = node; - this->size += 1; - } - - __device__ Node *search(LL key) { - for (int offset = 0; offset < this->size; offset++) { - if (this->store[offset].key >= key) { - return this->store[offset].point; - } - } - return nullptr; - } - - __device__ ~MemMap() { delete[] store; } -};*/ - // Definition of lock-free skip list class LockFreeSkipList { @@ -736,13 +690,13 @@ int main(int argc, char **argv) { sizeof(key_type) * search_length, cudaMemcpyHostToDevice); // Launch main kernel + blocks = calcBlocks(search_length); cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, nullptr); - blocks = calcBlocks(search_length); kernel<<<blocks, NUM_THREADS>>>(cudaOperatorItems, search_length, cudaResult); CudaCheckError(); cudaDeviceSynchronize(); diff --git a/expt_0528.cu b/expt_0528.cu index cbc6c91..27bc29f 100644 --- a/expt_0528.cu +++ b/expt_0528.cu @@ -263,13 +263,10 @@ class LockFreeSkipList { int *randoms = nullptr; - double scaleSize; - public: Node *head = nullptr; Node *tail = nullptr; - LockFreeSkipList(int *randoms, size_t random_length, double scale_size) - : scaleSize(scale_size) { + LockFreeSkipList(int *randoms, size_t random_length) { Node *h = new Node(0); // size_ = 0; Node *t = new Node(std::numeric_limits<key_type>::max() - 1); @@ -622,10 +619,6 @@ __global__ void print_function() { memcpy(spend_time, lockFreeSkipList->spend_time, sizeof(int) * NUM_ITEMS); }*/ -inline double calcSliceSize(size_t insertion_size) { - return 1.0 / (double)insertion_size; -} - inline size_t calcBlocks(size_t input) { return (input % (NUM_THREADS * FACTOR) == 0) ? input / (NUM_THREADS * FACTOR) @@ -700,8 +693,7 @@ int main(int argc, char **argv) { // Allocate the skip list LockFreeSkipList *Clist; - auto *list = new LockFreeSkipList(random_level, insertion_length, - calcSliceSize(insertion_length)); + auto *list = new LockFreeSkipList(random_level, insertion_length); cudaMalloc(&Clist, sizeof(LockFreeSkipList)); cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice); diff --git a/expt_0618_2.cpp b/expt_0618_2.cpp new file mode 100644 index 0000000..4c104ad --- /dev/null +++ b/expt_0618_2.cpp @@ -0,0 +1,33 @@ +#include <iostream> +#include <random> + +constexpr size_t MAX_LEVEL = 32; + +void generateRandomLevel(std::vector<unsigned int> &v, long run_times) { + // auto levels = new int[size]; + std::random_device randomDevice; + std::mt19937 randomEngine(randomDevice()); + std::geometric_distribution<> distribution(0.5); + for (int i = 0; i < run_times; i++) { + v[(int)std::min(MAX_LEVEL, (size_t)distribution(randomEngine))]++; + } +} + +int main(int argc, char const *argv[]) { + + auto run_times = 65536L; + + if (argc >= 2) { + run_times = strtol(argv[1], nullptr, 10); + } + + std::vector<unsigned int> result(MAX_LEVEL + 3); + generateRandomLevel(result, run_times); + + for (int i = 0; i < 32; i++) { + if (!result[i]) { + continue; + } + printf("%d: %d\n", i, result[i]); + } +}
\ No newline at end of file diff --git a/expt_0618_3.cu b/expt_0618_3.cu new file mode 100644 index 0000000..4558f54 --- /dev/null +++ b/expt_0618_3.cu @@ -0,0 +1,117 @@ +// Experimental content: Print level by cdf + +#include "sortlib.cuh" +#define READ_NO_OUTPUT +#include "read_helper_p.h" +#include <random> + +long pow_for_sample(long n) { + auto x = 2; + for (int i = 1; i < n; i++) { + x *= 2; + } + return x - 1; +} +constexpr size_t MAX_LEVEL = 32; +constexpr auto RESULT_LENGTH = MAX_LEVEL + 3; + +static unsigned trailing_zeroes(size_t index) { + constexpr auto block_size = 2; + unsigned bits = 0; + auto x = index / block_size; + + if (x) { + while (x % block_size == 0) { + ++bits; + x /= block_size; + } + } + return bits; +} + +inline double calcSliceSize(size_t insertion_size) { + return 1.0 / (double)insertion_size; +} + +void generateRandomLevel(std::vector<unsigned int> &v, long run_times) { + // auto levels = new int[size]; + std::random_device randomDevice; + std::mt19937 randomEngine(randomDevice()); + std::geometric_distribution<> distribution(0.5); + for (int i = 0; i < run_times; i++) { + v[(int)std::min(MAX_LEVEL, (size_t)distribution(randomEngine))]++; + } +} + +void multiple_generate(std::vector<unsigned int> &v, long run_times, + const int outer_times = 10) { + std::vector<std::vector<unsigned int>> results(outer_times); + std::random_device randomDevice; + std::mt19937 randomEngine(randomDevice()); + std::geometric_distribution<> distribution(0.5); + for (int x = 0; x < outer_times; x++) { + results[x].resize(RESULT_LENGTH, 0); + for (int i = 0; i < run_times; i++) { + results[x] + [(int)std::min(MAX_LEVEL, (size_t)distribution(randomEngine))]++; + } + } + for (int level = 0; level < MAX_LEVEL; level++) { + auto sum = 0L; + for (int i = 0; i < outer_times; i++) { + sum += results[i][level]; + } + v[level] = sum / outer_times; + } +} + +int main(int argc, char const *argv[]) { + + auto multiple_times = 0L; + if (argc < 3) { + printf("Usage %s [sample(pow)] [population]\n", argv[0]); + return 1; + } + + if (argc == 4) { + multiple_times = strtol(argv[3], nullptr, 10); + printf("multiple_times defined: %ld\n", multiple_times); + } + + auto sample_length = pow_for_sample(strtol(argv[1], nullptr, 10)); + auto population_length = strtol(argv[2], nullptr, 10); + + printf("sample length: %ld, population length: %ld\n", sample_length, + population_length); + + ReadHelper<key_type> readHelper("normal_distribution.txt", sample_length, 0); + readHelper.readFile(); + + std::vector<key_type> sample, population; + std::vector<unsigned int> result(RESULT_LENGTH), result2(RESULT_LENGTH); + readHelper.split_into(sample, population); + rebuildSort(sample); + + auto scale_size = calcSliceSize(population_length); + + auto sort = CustomSort(sample_length, sizeof(key_type) * 8); + + for (auto element : readHelper.population_vector) { + auto ret = sort.sample_cdf_custom_version(sample.data(), element); + auto cdf_index = ret / scale_size; + auto index = trailing_zeroes((size_t)cdf_index); + result[index]++; + } + + if (!multiple_times) + generateRandomLevel(result2, population_length); + else + multiple_generate(result2, population_length); + + for (int i = 0; i < 32; i++) { + if (!result[i] && !result2[i]) { + continue; + } + printf("%d: %d %d\n", i, result[i], result2[i]); + } +} |
