diff options
Diffstat (limited to 'skiplistcustom.cuh')
| -rw-r--r-- | skiplistcustom.cuh | 31 |
1 files changed, 23 insertions, 8 deletions
diff --git a/skiplistcustom.cuh b/skiplistcustom.cuh index a16e80a..80ecc19 100644 --- a/skiplistcustom.cuh +++ b/skiplistcustom.cuh @@ -1,5 +1,6 @@ #pragma once #include "sortlib.cuh" +#include <limits> #ifndef LOCKFREE_SKIPLIST_CUH_ #define LOCKFREE_SKIPLIST_CUH_ @@ -85,7 +86,7 @@ public: // Definition of lock-free skip list class LockFreeSkipList { - static constexpr size_t SAMPLE_LENGTH = 1024; + // static constexpr size_t SAMPLE_LENGTH = 1023; public: Node *head; @@ -99,20 +100,21 @@ public: double slice_size; + size_t sampleLength; + LockFreeSkipList(key_type *samples, size_t sample_length) - : searcher(sample_length, sizeof(key_type)) { + : searcher(sample_length, sizeof(key_type)), sampleLength(sample_length) { Node *h = new Node(0); // size_ = 0; #if __WORDSIZE == 64 - Node *t = new Node((key_type)NUM_ITEMS + 10); + Node *t = new Node(std::numeric_limits<key_type>::max()); #else Node *t = new Node(0xffffffffULL); #endif cudaMalloc(&head, sizeof(Node)); cudaMalloc(&tail, sizeof(Node)); - int i; - for (i = 0; i < h->topLevel + 1; i++) { + for (int i = 0; i < h->topLevel + 1; i++) { h->SetRef(i, tail, false); } cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice); @@ -123,8 +125,10 @@ public: initDeviceVariable(); this->slice_size = 1.0 / sample_length; + // printf("%lf\n", this->slice_size); - cudaMalloc(&samples, sizeof(key_type) * sample_length); + printf("%ld\n", sample_length); + cudaMalloc(&this->samples, sizeof(key_type) * sample_length); cudaMemcpy(this->samples, samples, sizeof(key_type) * sample_length, cudaMemcpyHostToDevice); } @@ -150,9 +154,11 @@ public: return bits; } - __device__ size_t calcLevel(key_type k) { + __device__ size_t calcLevel(key_type k) const { + // printf("Samples\n"); auto index = this->searcher.sample_cdf_custom_version(this->samples, k) / slice_size; + // printf("Samples finish\n"); auto level = trailing_zeroes(index); return level; } @@ -162,6 +168,12 @@ public: //__device__ bool Delete(key_type); __device__ bool Search(key_type) const; + __device__ void testSample() { + for (int i = 0; i < sampleLength; i++) { + printf("%lld\n", this->samples[i]); + } + } + #ifdef MEASURE_ACCESS unsigned access_times = 0; @@ -193,6 +205,7 @@ public: __device__ Node *GetNewNode(key_type key, unsigned int *pointerIndex, int topLevel) { + printf("%d\n", topLevel); key_type ind = atomicInc(pointerIndex, NUM_ITEMS); Node *n = nodes[ind]; n->key = key; @@ -328,7 +341,9 @@ __device__ bool LockFreeSkipList::Search(key_type key) const { }*/ __device__ bool LockFreeSkipList::Add(key_type key) { - Node *newNode = GetNewNode(key, this->pointerIndex, calcLevel(key)); + auto level1 = calcLevel(key) - 1; + printf("%lu\n", level1); + Node *newNode = GetNewNode(key, this->pointerIndex, level1); int topLevel = newNode->topLevel; int bottomLevel = 0; Node *preds[MAX_LEVEL + 1]; |
