diff options
| -rw-r--r-- | CMakeLists.txt | 8 | ||||
| -rw-r--r-- | main.cu | 76 |
2 files changed, 16 insertions, 68 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index f4b7297..faf0f29 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -8,9 +8,11 @@ set(CMAKE_CUDA_STANDARD 14) add_executable(lockfree main.cu) #add_definitions(-DNUM_ITEMS=1048576 -DKEYS=1048576 -DFACTOR=1) -target_compile_options(lockfree PRIVATE $<$<COMPILE_LANGUAGE:CUDA>: - -v - >) +if (false) + target_compile_options(lockfree PRIVATE $<$<COMPILE_LANGUAGE:CUDA>: + -v + >) +endif () # from: Peter target_link_libraries(lockfree m) @@ -68,7 +68,6 @@ Conference on Parallel and Distributed Systems, December 2012. #include <cassert> #include <cstdio> #include <cstdlib> -#include <ctime> #include <random> #if __WORDSIZE == 64 @@ -86,13 +85,13 @@ constexpr size_t MAX_LEVEL = 32; constexpr size_t NUM_THREADS = 512; constexpr size_t NUM_ITEMS = 1048576; -constexpr size_t KEYS = 1048576; +// constexpr size_t KEYS = 1048576; constexpr size_t FACTOR = 1; // Supported operations -#define ADD (0) -#define DELETE (1) -#define SEARCH (2) +constexpr int ADD = 0; +constexpr int DELETE = 1; +constexpr int SEARCH = 2; #define CUDA_ERROR_CHECK @@ -198,32 +197,16 @@ public: #else Node *t = new Node((LL)0xffffffff); #endif -#ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMalloc((void **)&head, sizeof(Node))); -#else cudaMalloc((void **)&head, sizeof(Node)); -#endif -#ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMalloc((void **)&tail, sizeof(Node))); -#else cudaMalloc((void **)&tail, sizeof(Node)); -#endif int i; for (i = 0; i < h->topLevel + 1; i++) { h->SetRef(i, tail, false); } -#ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice)); -#else cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice); -#endif -#ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice)); -#else cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice); -#endif } __device__ bool find(LL, Node **, Node **); // Helping method __device__ bool Add(LL); @@ -291,7 +274,7 @@ LockFreeSkipList::find(LL key, Node **preds, beenThereDoneThat = false; // printf("find key is %d \n",(int)key); } - if (beenThereDoneThat && !snip) + if (beenThereDoneThat) break; if (curr->key <= key) { pred = curr; @@ -300,12 +283,12 @@ LockFreeSkipList::find(LL key, Node **preds, break; } } - if (beenThereDoneThat && !snip) + if (beenThereDoneThat) break; preds[level] = pred; succs[level] = curr; } - if (beenThereDoneThat && !snip) + if (beenThereDoneThat) continue; return ((curr->key == key)); } @@ -334,7 +317,7 @@ __device__ bool LockFreeSkipList::Search(LL key) { } } } - return (curr->key == key); + return (curr != nullptr && curr->key == key); } __device__ bool LockFreeSkipList::Delete(LL key) { @@ -473,7 +456,6 @@ __global__ void kernel(LL *items, LL *op, LL *result) { } } -// Generate the level of a newly created node /*LL Randomlevel() { LL v = 1; double p = 0.5; @@ -482,6 +464,7 @@ __global__ void kernel(LL *items, LL *op, LL *result) { return v; }*/ +// Generate the level of a newly created node LL Randomlevel(std::mt19937 &randomEngine) { std::geometric_distribution<> distribution(0.5); return std::min(MAX_LEVEL, (size_t)distribution(randomEngine) + 1); @@ -495,8 +478,8 @@ int main(int argc, char **argv) { } // Extract operations ratio - int adds = atoi(argv[1]); - int deletes = atoi(argv[2]); + long adds = strtol(argv[1], nullptr, 10); + long deletes = strtol(argv[2], nullptr, 10); if (adds + deletes > 100) { printf("Sum of add and delete percentages exceeds 100.\nAborting...\n"); @@ -560,18 +543,6 @@ int main(int argc, char **argv) { LL *Cresult; LL *Clevels; -#ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMalloc((void **)&Cresult, sizeof(LL) * NUM_ITEMS)); - CUDA_SAFE_CALL(cudaMalloc((void **)&Citems, sizeof(LL) * NUM_ITEMS)); - CUDA_SAFE_CALL(cudaMalloc((void **)&Cop, sizeof(LL) * NUM_ITEMS)); - CUDA_SAFE_CALL(cudaMalloc((void **)&Clevels, sizeof(LL) * NUM_ITEMS)); - CUDA_SAFE_CALL(cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS, - cudaMemcpyHostToDevice)); - CUDA_SAFE_CALL(cudaMemcpy(Citems, items, sizeof(LL) * NUM_ITEMS, - cudaMemcpyHostToDevice)); - CUDA_SAFE_CALL( - cudaMemcpy(Cop, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice)); -#else cudaMalloc((void **)&Cresult, sizeof(LL) * NUM_ITEMS); cudaMalloc((void **)&Citems, sizeof(LL) * NUM_ITEMS); cudaMalloc((void **)&Cop, sizeof(LL) * NUM_ITEMS); @@ -579,45 +550,25 @@ int main(int argc, char **argv) { cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice); cudaMemcpy(Citems, items, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice); cudaMemcpy(Cop, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice); -#endif - Node **pointers = (Node **)malloc(sizeof(LL) * adds); Node **Cpointers; // Allocate the pool of free nodes for (i = 0; i < adds; i++) { -#ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMalloc((void **)&pointers[i], sizeof(Node))); -#else 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, - cudaMemcpyHostToDevice)); -#else cudaMalloc((void **)&Cpointers, sizeof(Node *) * adds); cudaMemcpy(Cpointers, pointers, sizeof(Node *) * adds, cudaMemcpyHostToDevice); -#endif // Allocate the skip list 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), - cudaMemcpyHostToDevice)); -#else cudaMalloc((void **)&Clist, sizeof(LockFreeSkipList)); cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice); -#endif - // Calculate the number of thread blocks // NUM_ITEMS = total number of operations to execute // NUM_THREADS = number of threads per block @@ -707,12 +658,7 @@ int main(int argc, char **argv) { // Move results back to host memory -#ifdef _CUTIL_H_ - CUDA_SAFE_CALL(cudaMemcpy(result, Cresult, sizeof(LL) * NUM_ITEMS, - cudaMemcpyDeviceToHost)); -#else cudaMemcpy(result, Cresult, sizeof(LL) * NUM_ITEMS, cudaMemcpyDeviceToHost); -#endif // Uncomment the following for debugging // print<<<1,32>>>(); |