diff options
| -rw-r--r-- | main.cu | 24 |
1 files changed, 18 insertions, 6 deletions
@@ -190,8 +190,10 @@ class LockFreeSkipList { public: Node *head; Node *tail; + unsigned size_ = 0; LockFreeSkipList() { Node *h = new Node(0); + // size_ = 0; #if __WORDSIZE == 64 Node *t = new Node((LL)NUM_ITEMS + 10); #else @@ -251,11 +253,11 @@ LockFreeSkipList::find(LL key, Node **preds, Node **succs) { // preds and succs are arrays of pointers int bottomLevel = 0; bool marked[] = {false}; - bool snip = false; - Node *pred = nullptr; + bool snip; + Node *pred; Node *curr = nullptr; - Node *succ = nullptr; - bool beenThereDoneThat = false; + Node *succ; + bool beenThereDoneThat; while (true) { beenThereDoneThat = false; pred = head; @@ -299,7 +301,7 @@ __device__ bool LockFreeSkipList::Search(LL key) { bool marked = false; Node *pred = head; Node *curr = nullptr; - Node *succ = nullptr; + Node *succ; int level; for (level = MAX_LEVEL; level >= bottomLevel; level--) { curr = pred->GetReference(level); @@ -347,6 +349,8 @@ __device__ bool LockFreeSkipList::Delete(LL key) { succ = succs[bottomLevel]->Get(bottomLevel, marked); if (iMarkedIt) { find(key, preds, succs); + // size_ -= 1; + atomicDec(&size_, 1); return true; } else if (marked[0]) { return false; @@ -399,6 +403,8 @@ __device__ bool LockFreeSkipList::Add(LL key) { find(key, preds, succs); } } + // size_ += 1; + atomicAdd(&size_, 1); return true; } } @@ -470,6 +476,8 @@ LL Randomlevel(std::mt19937 &randomEngine) { return std::min(MAX_LEVEL, (size_t)distribution(randomEngine) + 1); } +__global__ void print_function() { printf("size: %u\n", l->size_); } + int main(int argc, char **argv) { if (argc != 3) { printf("Need two arguments: percent add ops and percent delete ops (e.g., " @@ -501,7 +509,8 @@ int main(int argc, char **argv) { std::uniform_int_distribution<int> uniformIntDistributionArray(0, NUM_ITEMS); for (i = 0; i < NUM_ITEMS; i++) { - items[i] = i + 3; // 10+rand()%KEYS; // Keys associated with + items[i] = i + 3; // 10+rand()%KEYS; + // Keys associated with // operations } @@ -662,6 +671,9 @@ int main(int argc, char **argv) { // Uncomment the following for debugging // print<<<1,32>>>(); + cudaDeviceSynchronize(); + + print_function<<<1, 1>>>(); cudaDeviceSynchronize(); /*cudaFree(Clist); |
