From be0483ef4b6935324f85f124da8f4c1b93e918f6 Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Wed, 24 May 2023 14:12:56 +0800 Subject: fix(exp): Fix expt_0520 malloc failure Signed-off-by: KunoiSayami --- skiplistcustom.cuh | 117 +++++++++++++++++++++++------------------------------ 1 file changed, 50 insertions(+), 67 deletions(-) (limited to 'skiplistcustom.cuh') diff --git a/skiplistcustom.cuh b/skiplistcustom.cuh index 54c4732..a16e80a 100644 --- a/skiplistcustom.cuh +++ b/skiplistcustom.cuh @@ -1,4 +1,5 @@ #pragma once +#include "sortlib.cuh" #ifndef LOCKFREE_SKIPLIST_CUH_ #define LOCKFREE_SKIPLIST_CUH_ @@ -81,52 +82,6 @@ public: } }; -/*struct MapNode { - LL key; - // Node *point[MAX_LEVEL + 1]; - Node *point; -}; - -class MemMap { -public: - size_t size; - size_t real_size; - MapNode *store; - - MemMap() : size(0), store(nullptr), real_size(0) {} - - __device__ bool insert(MapNode node) { - bool need_extend = this->size + 1 > this->real_size; - if (need_extend) { - bool need_copy = this->real_size == 0; - if (!need_copy) { - this->real_size += 1; - } - this->real_size *= 2; - MapNode *old = this->store; - this->store = new MapNode[this->real_size]; - if (need_copy) { - memcpy(this->store, old, this->real_size * sizeof(MapNode *)); - } - delete[] old; - } - // need sort after insert - this->store[size] = node; - this->size += 1; - } - - __device__ Node *search(LL key) { - for (int offset = 0; offset < this->size; offset++) { - if (this->store[offset].key >= key) { - return this->store[offset].point; - } - } - return nullptr; - } - - __device__ ~MemMap() { delete[] store; } -};*/ - // Definition of lock-free skip list class LockFreeSkipList { @@ -138,9 +93,14 @@ public: unsigned int *pointerIndex; - key_type *randoms; + key_type *samples; - LockFreeSkipList(key_type *randoms) { + CustomSort searcher; + + double slice_size; + + LockFreeSkipList(key_type *samples, size_t sample_length) + : searcher(sample_length, sizeof(key_type)) { Node *h = new Node(0); // size_ = 0; #if __WORDSIZE == 64 @@ -162,8 +122,10 @@ public: cudaMalloc(&pointerIndex, sizeof(unsigned int)); initDeviceVariable(); - cudaMalloc(&randoms, sizeof(key_type) * SAMPLE_LENGTH); - cudaMemcpy(this->randoms, randoms, sizeof(key_type) * SAMPLE_LENGTH, + this->slice_size = 1.0 / sample_length; + + cudaMalloc(&samples, sizeof(key_type) * sample_length); + cudaMemcpy(this->samples, samples, sizeof(key_type) * sample_length, cudaMemcpyHostToDevice); } @@ -172,10 +134,33 @@ public: cudaMemcpy(pointerIndex, &x, sizeof(unsigned int), cudaMemcpyHostToDevice); } - __device__ bool find(key_type, Node **, Node **); // Helping method + ~LockFreeSkipList() { cudaFree(samples); } + + __device__ unsigned static trailing_zeroes(size_t index) { + constexpr auto block_size = 2; + unsigned bits = 0; + unsigned x = index / block_size; + + if (x) { + while (x % block_size == 0) { + ++bits; + x /= block_size; + } + } + return bits; + } + + __device__ size_t calcLevel(key_type k) { + auto index = + this->searcher.sample_cdf_custom_version(this->samples, k) / slice_size; + auto level = trailing_zeroes(index); + return level; + } + + __device__ bool find(key_type, Node **, Node **) const; // Helping method __device__ bool Add(key_type); - __device__ bool Delete(key_type); - __device__ bool Search(key_type); + //__device__ bool Delete(key_type); + __device__ bool Search(key_type) const; #ifdef MEASURE_ACCESS unsigned access_times = 0; @@ -206,15 +191,13 @@ public: // Function for creating a new node when requested by an add operation -__device__ Node *GetNewNode(Node **nodes, key_type key, - unsigned int *pointerIndex, - const key_type *randoms) { +__device__ Node *GetNewNode(key_type key, unsigned int *pointerIndex, + int topLevel) { key_type ind = atomicInc(pointerIndex, NUM_ITEMS); Node *n = nodes[ind]; n->key = key; - n->topLevel = randoms[ind]; - int i; - for (i = 0; i < n->topLevel + 1; i++) { + n->topLevel = topLevel; + for (int i = 0; i < n->topLevel + 1; i++) { n->SetRef(i, nullptr, false); } return n; @@ -223,14 +206,14 @@ __device__ Node *GetNewNode(Node **nodes, key_type key, // Find the window holding key // On the way clean up logically deleted nodes (those with set marked bit) -__device__ bool -LockFreeSkipList::find(key_type key, Node **preds, - Node **succs) { // preds and succs are arrays of pointers +__device__ bool LockFreeSkipList::find( + key_type key, Node **preds, + Node **succs) const { // preds and succs are arrays of pointers int bottomLevel = 0; bool marked[] = {false}; bool snip; Node *pred; - Node *curr = nullptr; + Node *curr; Node *succ; bool beenThereDoneThat; while (true) { @@ -271,7 +254,7 @@ LockFreeSkipList::find(key_type key, Node **preds, } } -__device__ bool LockFreeSkipList::Search(key_type key) { +__device__ bool LockFreeSkipList::Search(key_type key) const { int bottomLevel = 0; bool marked = false; Node *pred = head; @@ -306,7 +289,7 @@ __device__ bool LockFreeSkipList::Search(key_type key) { return (curr != nullptr && curr->key == key); } -__device__ bool LockFreeSkipList::Delete(key_type key) { +/*__device__ bool LockFreeSkipList::Delete(key_type key) { int bottomLevel = 0; Node *preds[MAX_LEVEL + 1]; Node *succs[MAX_LEVEL + 1]; @@ -342,10 +325,10 @@ __device__ bool LockFreeSkipList::Delete(key_type key) { } } } -} +}*/ __device__ bool LockFreeSkipList::Add(key_type key) { - Node *newNode = GetNewNode(nodes, key, this->pointerIndex, this->randoms); + Node *newNode = GetNewNode(key, this->pointerIndex, calcLevel(key)); int topLevel = newNode->topLevel; int bottomLevel = 0; Node *preds[MAX_LEVEL + 1]; -- cgit v1.3.1