summaryrefslogtreecommitdiff
path: root/main.cu
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2022-06-01 14:57:24 +0800
committerKunoiSayami <[email protected]>2022-06-01 14:57:24 +0800
commit1f0a9046ff2570eec52855130556279a2e332116 (patch)
treea308394baf3c7a6025ca5e41881045c58180149d /main.cu
parent6a37727e6b1f9c94dac8ce8c53392bbe1c67674a (diff)
feat: Add size counter
Signed-off-by: KunoiSayami <[email protected]>
Diffstat (limited to 'main.cu')
-rw-r--r--main.cu24
1 files changed, 18 insertions, 6 deletions
diff --git a/main.cu b/main.cu
index f43e4bc..fe801d4 100644
--- a/main.cu
+++ b/main.cu
@@ -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);