diff options
| -rw-r--r-- | .gitignore | 1 | ||||
| -rw-r--r-- | CMakeLists.txt | 32 | ||||
| -rwxr-xr-x | image.py | 16 | ||||
| -rw-r--r-- | main.cu | 41 |
4 files changed, 79 insertions, 11 deletions
@@ -11,3 +11,4 @@ CTestTestfile.cmake _deps cmake-*/ .idea/ +test
\ No newline at end of file diff --git a/CMakeLists.txt b/CMakeLists.txt index faf0f29..7e8ff4d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -15,8 +15,36 @@ if (false) endif () # from: Peter -target_link_libraries(lockfree m) -target_link_libraries(lockfree stdc++) +target_link_libraries(lockfree m stdc++) set_target_properties(lockfree PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + +function(build_different_target is_random build_size step_size) + + string(CONCAT basic_name "lockfree_" "${build_size}" "_" "${step_size}") + + message(STATUS "different target: ${basic_name} ${is_random}") + add_executable("${basic_name}" "") + + target_sources("${basic_name}" + PRIVATE + main.cu + ) + + #add_definitions("${basic_name}" -DBUILD_SIZE=${build_size} -DSTEP_SIZE=${step_size}) + target_compile_definitions("${basic_name}" PRIVATE -DBUILD_SIZE=${build_size} -DSTEP_SIZE=${step_size}) + if ("${is_random}") + target_compile_definitions("${basic_name} PRIVATE -DRANDOM_TARGET") + endif () + + #target_link_libraries(basic_name m stdc++) + + set_target_properties("${basic_name}" PROPERTIES + CUDA_SEPARABLE_COMPILATION ON) + +endfunction() + +build_different_target(false 1024 2) +build_different_target(false 1024 4) +build_different_target(false 1024 8) diff --git a/image.py b/image.py new file mode 100755 index 0000000..199ba9f --- /dev/null +++ b/image.py @@ -0,0 +1,16 @@ +#!/usr/bin/env python3 +import matplotlib.pyplot as plt + +x = [] +y = [] + +with open('test') as fin: + for line in fin.read().splitlines(): + x_, y_ = list(map(int, line.strip().split(','))) + x.append(x_) + y.append(y_) + +plt.plot(x, y) +plt.xlabel('x') +plt.ylabel('y') +plt.show() @@ -76,6 +76,14 @@ typedef unsigned long long LL; typedef unsigned int LL; #endif +#ifndef BUILD_SIZE +#define BUILD_SIZE 1048576 +#endif + +#ifndef STEP_SIZE +#define STEP_SIZE 2 +#endif + // Maximum level of a node in the skip list //#define MAX_LEVEL 32 constexpr size_t MAX_LEVEL = 32; @@ -84,18 +92,26 @@ constexpr size_t MAX_LEVEL = 32; //#define NUM_THREADS 512 constexpr size_t NUM_THREADS = 512; -constexpr size_t NUM_ITEMS = 1048576; +constexpr size_t NUM_ITEMS = BUILD_SIZE; // constexpr size_t KEYS = 1048576; constexpr size_t FACTOR = 1; // should change this to dynamic next time constexpr size_t KEY_INDEX_SIZE = 32; +constexpr int block_size = STEP_SIZE; + // Supported operations constexpr int ADD = 0; constexpr int DELETE = 1; constexpr int SEARCH = 2; +#ifdef RANDOM_TARGET +constexpr const char *TARGET_STRING = "RANDOM"; +#else +constexpr const char *TARGET_STRING = "PERFECT"; +#endif + #define CUDA_ERROR_CHECK #define CudaSafeCall(err) __cudaSafeCall(err, __FILE__, __LINE__) @@ -544,8 +560,6 @@ unsigned trailing_zeroes(LL n) { return bits; } -constexpr int block_size = 8; - LL CustomLevel(LL value) { auto left = std::lower_bound(storage.begin(), storage.end(), value); auto right = std::upper_bound(storage.begin(), storage.end(), value); @@ -559,7 +573,7 @@ LL CustomLevel(LL value) { if (index % block_size == 0) { auto level = trailing_zeroes(index / block_size + 1) + 1; - // printf("%ld %u\n", index, level); + // printf("%ld,%u\n", index, level); return level; } @@ -626,8 +640,11 @@ int main(int argc, char **argv) { // Pre-generated levels of skip list nodes (relevant only if op[i] is add) // srand(0); for (i = 0; i < NUM_ITEMS; i++) { - // levels[i] = Randomlevel(randomEngine) - 1; //36/14 +#ifdef RANDOM_HEIGHT + // levels[i] = Randomlevel(randomEngine) - 1; // 36/14 +#else levels[i] = CustomLevel(items[i]) - 1; // 31/18 +#endif } // Populate the sequence of operations @@ -720,7 +737,13 @@ int main(int argc, char **argv) { // Print kernel execution time in milliseconds - printf("%lf\n", time); + printf("%s ", TARGET_STRING); + +#ifndef RANDOM_TARGET + printf("%d ", block_size); +#endif + + printf("%lu: %lf", NUM_ITEMS, time); // Launch main kernel for query @@ -753,7 +776,7 @@ int main(int argc, char **argv) { // Print kernel execution time in milliseconds - printf("%lf\n", time); + printf(" %lf\n", time); // Check for errors @@ -769,9 +792,9 @@ int main(int argc, char **argv) { // Uncomment the following for debugging // print<<<1,32>>>(); - cudaDeviceSynchronize(); + // cudaDeviceSynchronize(); - print_function<<<1, 1>>>(); + // print_function<<<1, 1>>>(); cudaDeviceSynchronize(); /*cudaFree(Clist); |