From 12a9246ace908d8b9e3020c7137922eb1391ef22 Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Sat, 28 May 2022 01:32:24 +0800 Subject: feat: Use std::random Signed-off-by: KunoiSayami --- CMakeLists.txt | 10 ++++++++++ main.cu | 62 ++++++++++++++++++++++++++++++++++++++++++---------------- 2 files changed, 55 insertions(+), 17 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index d7e3860..f4b7297 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -3,8 +3,18 @@ project(lockfree CUDA) set(CMAKE_CUDA_STANDARD 14) +#set(GXX_INCLUDE /usr/lib/gcc/x86_64-pc-linux-gnu/11.3.0/include) + add_executable(lockfree main.cu) #add_definitions(-DNUM_ITEMS=1048576 -DKEYS=1048576 -DFACTOR=1) +target_compile_options(lockfree PRIVATE $<$: + -v + >) + +# from: Peter +target_link_libraries(lockfree m) +target_link_libraries(lockfree stdc++) + set_target_properties(lockfree PROPERTIES CUDA_SEPARABLE_COMPILATION ON) diff --git a/main.cu b/main.cu index 7c0aae3..80d75a8 100644 --- a/main.cu +++ b/main.cu @@ -64,10 +64,12 @@ Conference on Parallel and Distributed Systems, December 2012. //#include"cutil.h" // Comment this if cutil.h is not available #include "cuda_runtime.h" +#include #include #include #include #include +#include #if __WORDSIZE == 64 typedef unsigned long long LL; @@ -472,12 +474,17 @@ __global__ void kernel(LL *items, LL *op, LL *result) { } // Generate the level of a newly created node -LL Randomlevel() { +/*LL Randomlevel() { LL v = 1; double p = 0.5; while (((rand() / (double)(RAND_MAX)) < p) && (v < MAX_LEVEL)) v++; return v; +}*/ + +LL Randomlevel(std::mt19937 &randomEngine) { + std::geometric_distribution<> distribution(0.5); + return std::min(MAX_LEVEL, (size_t)distribution(randomEngine) + 1); } int main(int argc, char **argv) { @@ -492,37 +499,45 @@ int main(int argc, char **argv) { int deletes = atoi(argv[2]); if (adds + deletes > 100) { - printf("Sum of add and delete precentages exceeds 100.\nAborting...\n"); + printf("Sum of add and delete percentages exceeds 100.\nAborting...\n"); exit(1); } // Allocate necessary arrays - LL *op = (LL *)malloc(sizeof(LL) * NUM_ITEMS); - LL *levels = (LL *)malloc(sizeof(LL) * NUM_ITEMS); - LL *items = (LL *)malloc(sizeof(LL) * NUM_ITEMS); - LL *result = (LL *)malloc(sizeof(LL) * NUM_ITEMS); + LL *op = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); + LL *levels = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); + LL *items = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); + LL *result = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); int i; // NUM_ITEMS is the total number of operations to execute - srand(0); + // srand(0); + + std::random_device randomDevice; + std::mt19937 randomEngine(randomDevice()); + std::uniform_int_distribution uniformIntDistributionArray(0, NUM_ITEMS); + for (i = 0; i < NUM_ITEMS; i++) { items[i] = i + 3; // 10+rand()%KEYS; // Keys associated with // operations } for (i = 0; i < NUM_ITEMS; i++) { - int first = rand() % NUM_ITEMS; - int second = rand() % NUM_ITEMS; - LL temp; + /*int first = rand() % NUM_ITEMS; + int second = rand() % NUM_ITEMS;*/ + + std::swap(items[uniformIntDistributionArray(randomEngine)], + items[uniformIntDistributionArray(randomEngine)]); + /*LL temp; temp = items[first]; items[first] = items[second]; - items[second] = temp; + items[second] = temp;*/ } // Pre-generated levels of skip list nodes (relevant only if op[i] is add) - srand(0); + // srand(0); for (i = 0; i < NUM_ITEMS; i++) { - levels[i] = Randomlevel() - 1; + levels[i] = Randomlevel(randomEngine) - 1; } // Populate the sequence of operations @@ -578,6 +593,7 @@ int main(int argc, char **argv) { cudaMalloc((void **)&pointers[i], sizeof(Node)); #endif } + #ifdef _CUTIL_H_ CUDA_SAFE_CALL(cudaMalloc((void **)&Cpointers, sizeof(Node *) * adds)); CUDA_SAFE_CALL(cudaMemcpy(Cpointers, pointers, sizeof(Node *) * adds, @@ -592,6 +608,7 @@ int main(int argc, char **argv) { LockFreeSkipList *Clist; auto *list = new LockFreeSkipList(); + #ifdef _CUTIL_H_ CUDA_SAFE_CALL(cudaMalloc((void **)&Clist, sizeof(LockFreeSkipList))); CUDA_SAFE_CALL(cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), @@ -626,7 +643,7 @@ int main(int argc, char **argv) { cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); - cudaEventRecord(start, 0); + cudaEventRecord(start, nullptr); kernel<<>>(Citems, Cop, Cresult); CudaCheckError(); @@ -636,7 +653,7 @@ int main(int argc, char **argv) { // exit(-1); } cudaDeviceSynchronize(); - cudaEventRecord(stop, 0); + cudaEventRecord(stop, nullptr); cudaEventSynchronize(stop); float time; cudaEventElapsedTime(&time, start, stop); @@ -660,7 +677,7 @@ int main(int argc, char **argv) { cudaEventCreate(&start); cudaEventCreate(&stop); - cudaEventRecord(start, 0); + cudaEventRecord(start, nullptr); kernel<<>>(Citems, Cop2, Cresult); CudaCheckError(); @@ -670,7 +687,7 @@ int main(int argc, char **argv) { // exit(-1); } cudaDeviceSynchronize(); - cudaEventRecord(stop, 0); + cudaEventRecord(stop, nullptr); cudaEventSynchronize(stop); cudaEventElapsedTime(&time, start, stop); cudaEventDestroy(start); @@ -701,5 +718,16 @@ int main(int argc, char **argv) { // print<<<1,32>>>(); cudaDeviceSynchronize(); + /*cudaFree(Clist); + cudaFree(Cop2); + cudaFree(Clevels); + cudaFree(Cop); + cudaFree(Citems); + cudaFree(Cresult); + free(pointers); + delete [] op; + delete [] levels; + delete [] items; + delete [] result;*/ return 0; } -- cgit v1.3.1