diff options
| -rw-r--r-- | CMakeLists.txt | 7 | ||||
| -rw-r--r-- | main.cu | 57 | ||||
| -rw-r--r-- | normal_distribution.cpp | 17 |
3 files changed, 76 insertions, 5 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index f816aed..e249fba 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,6 +1,5 @@ cmake_minimum_required(VERSION 3.22) -project(lockfree CUDA) - +project(lockfree CUDA CXX) set(CMAKE_CUDA_STANDARD 14) #set(GXX_INCLUDE /usr/lib/gcc/x86_64-pc-linux-gnu/11.3.0/include) @@ -19,6 +18,7 @@ target_link_libraries(lockfree m stdc++) set_target_properties(lockfree PROPERTIES CUDA_SEPARABLE_COMPILATION ON) +set_target_properties(lockfree PROPERTIES CUDA_ARCHITECTURES "75") function(build_different_target is_random build_size step_size) @@ -62,3 +62,6 @@ build_different_target(true 1048576 2) build_different_target(true 1048576 8) +add_executable(normal_distribution normal_distribution.cpp) +set_target_properties(normal_distribution PROPERTIES LINKER_LANGUAGE CXX) + @@ -84,13 +84,19 @@ typedef unsigned int LL; #define STEP_SIZE 2 #endif -//#define MEASURE_TIME +#define MEASURE_TIME //#define MEASURE_ACCESS #if (defined(MEASURE_ACCESS) && defined(MEASURE_TIME)) #error "Shouldn't define MEASURE_TIME and MEASURE_ACCESS at the same time" #endif +#ifdef MEASURE_TIME +#include <cuda/atomic> +#undef BUILD_SIZE +#define BUILD_SIZE 1024 +#endif + // Maximum level of a node in the skip list //#define MAX_LEVEL 32 constexpr size_t MAX_LEVEL = 16; @@ -258,6 +264,43 @@ public: __device__ ~MemMap() { delete[] store; } }; +class CudaSpinLock { + static constexpr int UNLOCKED = 0; + static constexpr int LOCKED = 1; + + cuda::atomic<int> m_value; + bool isFake; + +public: + __device__ __host__ explicit CudaSpinLock() + : m_value(UNLOCKED), isFake(false) {} + + __device__ __host__ explicit CudaSpinLock(bool fake) + : m_value(UNLOCKED), isFake(fake) {} + + __device__ void lock() { + if (!isFake) { + while (true) { + int expected = UNLOCKED; + // this->m_value.wait(LOCKED); + if (this->m_value.compare_exchange_weak(expected, LOCKED)) + break; + } + } + } + + __device__ void unlock() { + if (!isFake) { + m_value.store(UNLOCKED); + } + } + + __device__ bool isLock() { + // printf("%d\n", this->m_value.load()); + return !isFake && this->m_value.load() == LOCKED; + } +}; + // Definition of lock-free skip list class LockFreeSkipList { @@ -302,9 +345,11 @@ public: #ifdef MEASURE_TIME unsigned round = 0; + CudaSpinLock spinLock; __device__ void increaseRoundCount(unsigned count = 1) { atomicAdd(&this->round, count); } + unsigned long long total_time = 0; __device__ unsigned getRoundCount() const { return this->round; } #endif }; @@ -399,6 +444,8 @@ __device__ bool LockFreeSkipList::Search(LL key) { Node *succ; int level; #ifdef MEASURE_TIME + this->spinLock.lock(); + auto start_time = clock64(); #endif for (level = MAX_LEVEL; level >= bottomLevel; level--) { curr = pred->GetReference(level); @@ -426,7 +473,11 @@ __device__ bool LockFreeSkipList::Search(LL key) { } } #ifdef MEASURE_TIME + this->spinLock.unlock(); + unsigned long long end_time = clock64() - start_time; + // printf("%lld\n", end_time); atomicAdd(&this->round, 1); + atomicAdd(&this->total_time, end_time); #endif return (curr != nullptr && curr->key == key); } @@ -626,7 +677,8 @@ __global__ void print_function() { printf("count: %u\n", l->getAccessCount()); #endif #ifdef MEASURE_TIME - printf("round: %u\n", l->getRoundCount()); + printf("round: %u time: %llu avg: %.2lf\n", l->getRoundCount(), l->total_time, + l->total_time / (double)l->getRoundCount()); #endif } @@ -811,7 +863,6 @@ int main(int argc, char **argv) { cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, nullptr); - kernel<<<blocks, NUM_THREADS>>>(Citems, Cop2, Cresult); CudaCheckError(); error = cudaGetLastError(); diff --git a/normal_distribution.cpp b/normal_distribution.cpp new file mode 100644 index 0000000..c9cf701 --- /dev/null +++ b/normal_distribution.cpp @@ -0,0 +1,17 @@ +#include <algorithm> +#include <cstdio> +#include <random> +int main() { + + std::random_device randomDevice; + std::mt19937 randomEngine(randomDevice()); + std::normal_distribution<long double> normalDistribution(2147483647, + 2147483647); + + for (int i = 0; i < 1048576; i++) { + printf("%llu\n", + (unsigned long long)std::round(normalDistribution(randomEngine))); + } + + return 0; +}
\ No newline at end of file |