summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--CMakeLists.txt11
-rw-r--r--expt_0520.cu567
-rw-r--r--main.cu10
-rw-r--r--skiplist.cuh397
-rw-r--r--skiplistcustom.cuh397
5 files changed, 1376 insertions, 6 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt
index baf37e7..883dbc4 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -214,7 +214,6 @@ set_target_properties(expt_0516_2 PROPERTIES
set_target_properties(expt_0516_2 PROPERTIES CUDA_ARCHITECTURES "75")
set_target_properties(expt_0516_2 PROPERTIES LINKER_LANGUAGE CUDA)
-
add_executable(expt_0517 expt_0517.cu)
target_link_libraries(expt_0517 m stdc++)
@@ -223,4 +222,14 @@ set_target_properties(expt_0517 PROPERTIES
set_target_properties(expt_0517 PROPERTIES CUDA_ARCHITECTURES "75")
set_target_properties(expt_0517 PROPERTIES LINKER_LANGUAGE CUDA)
+add_executable(expt_0520 expt_0520.cu skiplistcustom.cuh)
+target_link_libraries(expt_0520 m stdc++)
+
+set_target_properties(expt_0520 PROPERTIES
+ CUDA_SEPARABLE_COMPILATION ON)
+set_target_properties(expt_0520 PROPERTIES CUDA_ARCHITECTURES "75")
+set_target_properties(expt_0520 PROPERTIES LINKER_LANGUAGE CUDA)
+
+
+
diff --git a/expt_0520.cu b/expt_0520.cu
new file mode 100644
index 0000000..839fef8
--- /dev/null
+++ b/expt_0520.cu
@@ -0,0 +1,567 @@
+/*
+
+Copyright 2012-2013 Indian Institute of Technology Kanpur. All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions are met:
+
+1. Redistributions of source code must retain the above copyright notice, this
+list of conditions, and the following disclaimer.
+2. Redistributions in binary form must reproduce the above copyright notice,
+this list of conditions, and the following disclaimer in the documentation
+and/or other materials provided with the distribution.
+
+THIS SOFTWARE IS PROVIDED BY INDIAN INSTITUTE OF TECHNOLOGY KANPUR ``AS IS''
+AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+DISCLAIMED. IN NO EVENT SHALL INDIAN INSTITUTE OF TECHNOLOGY KANPUR OR
+THE CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT
+OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+POSSIBILITY OF SUCH DAMAGE.
+
+The views and conclusions contained in the software and documentation are
+those of the authors and should not be interpreted as representing official
+policies, either expressed or implied, of Indian Institute of Technology Kanpur.
+
+*/
+
+/**********************************************************************************
+
+ Lock-free skip list for CUDA; tested for CUDA 4.2 on 32-bit Ubuntu 10.10 and
+64-bit Ubuntu 12.04. Developed at IIT Kanpur.
+
+ Inputs: Percentage of add and delete operations (e.g., 30 50 for 30% add and
+50% delete) Output: Prints the total time (in milliseconds) to execute the the
+sequence of operations
+
+ Compilation flags: -O3 -arch sm_20 -I ~/NVIDIA_GPU_Computing_SDK/C/common/inc/
+-DNUM_ITEMS=num_ops -DFACTOR=num_ops_per_thread -DKEYS=num_keys
+
+ NUM_ITEMS is the total number of operations (mix of add, delete, search) to
+execute.
+
+ FACTOR is the number of operations per thread.
+
+ KEYS is the number of integer keys assumed in the range [10, 9+KEYS].
+ The paper cited below states that the key range is [0, KEYS-1]. However, we
+have shifted the range by +10 so that the head sentinel key (the minimum key)
+can be chosen as zero. Any positive shift other than +10 would also work.
+
+ The include path ~/NVIDIA_GPU_Computing_SDK/C/common/inc/ is needed for
+cutil.h.
+
+ Related work:
+
+ Prabhakar Misra and Mainak Chaudhuri. Performance Evaluation of Concurrent
+Lock-free Data Structures on GPUs. In Proceedings of the 18th IEEE International
+Conference on Parallel and Distributed Systems, December 2012.
+
+***************************************************************************************/
+
+// #include"cutil.h" // Comment this if cutil.h is not available
+// #include "cuda_runtime.h"
+#include "skiplist.cuh"
+#include <algorithm>
+#include <cassert>
+#include <cstdio>
+#include <cstdlib>
+#include <random>
+#include <set>
+
+#if __WORDSIZE == 64
+typedef unsigned long long LL;
+#else
+typedef unsigned int LL;
+#endif
+
+#ifndef BUILD_SIZE
+#define BUILD_SIZE 1048576
+#endif
+
+#ifndef STEP_SIZE
+#define STEP_SIZE 2
+#endif
+
+// #define MEASURE_TIME
+// #define MEASURE_ACCESS
+
+#if (defined(MEASURE_ACCESS) && defined(MEASURE_TIME))
+#error "Shouldn't define MEASURE_TIME and MEASURE_ACCESS at the same time"
+#endif
+
+#ifdef MEASURE_TIME
+#undef BUILD_SIZE
+#define BUILD_SIZE 1024
+#endif
+
+// Maximum level of a node in the skip list
+// #define MAX_LEVEL 32
+// constexpr size_t MAX_LEVEL = 16;
+
+// Number of threads per block
+// #define NUM_THREADS 512
+constexpr size_t NUM_THREADS = 512;
+
+// constexpr size_t NUM_ITEMS = BUILD_SIZE;
+// constexpr size_t KEYS = 1048576;
+constexpr size_t FACTOR = 1;
+
+// should change this to dynamic next time
+// constexpr size_t KEY_INDEX_SIZE = 32;
+constexpr size_t SAMPLE_SIZE = 1024;
+
+constexpr int block_size = STEP_SIZE;
+
+// Supported operations
+constexpr int ADD = 0;
+constexpr int DELETE = 1;
+constexpr int SEARCH = 2;
+
+typedef LL key_type;
+
+#ifdef RANDOM_TARGET
+constexpr const char *TARGET_STRING = "RANDOM";
+#else
+constexpr const char *TARGET_STRING = "PERFECT";
+#endif
+
+#define CUDA_ERROR_CHECK
+
+#define CudaSafeCall(err) __cudaSafeCall(err, __FILE__, __LINE__)
+#define CudaCheckError() __cudaCheckError(__FILE__, __LINE__)
+
+inline void __cudaSafeCall(cudaError err, const char *file, const int line) {
+#ifdef CUDA_ERROR_CHECK
+ if (cudaSuccess != err) {
+ fprintf(stderr, "cudaSafeCall() failed at %s:%i : %s\n", file, line,
+ cudaGetErrorString(err));
+ exit(-1);
+ }
+#endif
+}
+
+inline void __cudaCheckError(const char *file, const int line) {
+#ifdef CUDA_ERROR_CHECK
+ cudaError err = cudaGetLastError();
+ if (cudaSuccess != err) {
+ fprintf(stderr, "cudaCheckError() failed at %s:%i : %s\n", file, line,
+ cudaGetErrorString(err));
+ exit(-1);
+ }
+
+ // More careful checking. However, this will affect performance.
+ // Comment away if needed.
+ err = cudaDeviceSynchronize();
+ if (cudaSuccess != err) {
+ fprintf(stderr, "cudaCheckError() with sync failed at %s:%i : %s\n", file,
+ line, cudaGetErrorString(err));
+ exit(-1);
+ }
+#endif
+}
+
+__device__ LockFreeSkipList *l; // The lock-free skip list
+
+__device__ key_type SampleStorage[SAMPLE_SIZE];
+
+// Kernel for initializing device memory
+
+__global__ void init(LockFreeSkipList *lockFreeSkipList, Node **n) {
+ nodes = n;
+ l = lockFreeSkipList;
+}
+
+__global__ void print() {
+ // For debugging
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid == 0) {
+ Node *p = l->head;
+ bool marked = false;
+ while (p != nullptr) {
+#if __WORDSIZE == 64
+ printf("%#llx, %u, marked=%u, address is %p : ", p->key, p->topLevel,
+ marked, p);
+#else
+ printf("%#x, %u, marked=%u, address is %p\n", p->key, p->topLevel, marked,
+ p);
+#endif
+ for (int i = 0; i < p->topLevel + 1; i++) {
+ printf(" %d ", (int)(p->GetReference(i)->key));
+ }
+ printf("\n");
+ p = p->Get(0, &marked);
+ }
+ printf("\n");
+ }
+}
+
+// The main kernel
+
+__global__ void kernel(LL *items, LL *op, LL *result) {
+ // The array items holds the sequence of keys
+ // The array op holds the sequence of operations
+ // The array result, at the end, will hold the outcome of the operations
+
+ int tid, i;
+ for (i = 0; i < FACTOR;
+ i++) { // FACTOR is the number of operations per thread
+ tid = i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid >= NUM_ITEMS)
+ return;
+
+ // Grab the operation and the associated key and execute
+ LL item = items[tid];
+ if (op[tid] == ADD) {
+ result[tid] = l->Add(item);
+ }
+ if (op[tid] == DELETE) {
+ result[tid] = l->Delete(item);
+ }
+ if (op[tid] == SEARCH) {
+#ifdef MEASURE_TIME
+ unsigned long long start_time = clock64();
+#endif
+ result[tid] = l->Search(item);
+#ifdef MEASURE_TIME
+ unsigned long long end_time = clock64() - start_time;
+ if (l->spend_time[tid]) {
+ printf("conflict: %d\n", tid);
+ }
+ l->spend_time[tid] = (int)end_time;
+#endif
+ }
+ }
+}
+
+/*LL Randomlevel() {
+ LL v = 1;
+ double p = 0.5;
+ while (((rand() / (double)(RAND_MAX)) < p) && (v < MAX_LEVEL))
+ v++;
+ return v;
+}*/
+
+// Generate the level of a newly created node
+LL RandomLevel(std::mt19937 &randomEngine, double p) {
+ std::geometric_distribution<> distribution(p);
+ return std::min(MAX_LEVEL, (size_t)distribution(randomEngine));
+}
+
+std::vector<LL> storage;
+
+unsigned trailing_zeroes(size_t index) {
+ unsigned bits = 0;
+ LL x = index / block_size;
+
+ if (x) {
+ while (x % block_size == 0) {
+ ++bits;
+ x /= block_size;
+ }
+ }
+ return bits;
+}
+
+LL CustomLevel(LL value) {
+ auto left = std::lower_bound(storage.begin(), storage.end(), value);
+ auto right = std::upper_bound(storage.begin(), storage.end(), value);
+
+ if (right - left != 1) {
+ printf("%ld\n", right - left);
+ }
+ assert(right - left == 1);
+
+ auto index = left - storage.begin();
+
+ if (index % block_size == 0) {
+ auto level = trailing_zeroes(index) + 1;
+ // printf("%ld,%u\n", index, level);
+ return level;
+ }
+
+ return 1;
+}
+
+__global__ void print_function() {
+#ifdef MEASURE_ACCESS
+ printf("count: %u\n", l->getAccessCount());
+#endif
+}
+
+std::vector<double> population;
+std::vector<double> sample;
+
+void initialize(const std::vector<double> &input_population,
+ const std::vector<double> &input_sample) {
+ population = input_population;
+ std::sort(population.begin(), population.end());
+ sample = input_sample;
+ std::sort(sample.begin(), sample.end());
+}
+
+int main(int argc, char **argv) {
+ if (argc != 3) {
+ printf("Need two arguments: percent add ops and percent delete ops (e.g., "
+ "30 50 for 30%% add and 50%% delete).\nAborting...\n");
+ exit(1);
+ }
+
+ // Extract operations ratio
+ long adds = strtol(argv[1], nullptr, 10);
+ long deletes = strtol(argv[2], nullptr, 10);
+
+ storage.reserve(NUM_ITEMS);
+
+ if (adds + deletes > 100) {
+ printf("Sum of add and delete percentages exceeds 100.\nAborting...\n");
+ exit(1);
+ }
+
+ // Allocate necessary arrays
+ LL *op = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS);
+ LL *levels = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS);
+ LL *items = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS);
+ LL *result = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS);
+ int i;
+
+ // NUM_ITEMS is the total number of operations to execute
+ // srand(0);
+
+ std::random_device randomDevice;
+ std::mt19937 randomEngine(randomDevice());
+ std::uniform_int_distribution<int> uniformIntDistributionArray(0,
+ NUM_ITEMS - 1);
+ // std::vector<LL> storage;
+
+ // std::normal_distribution<long double> normalDistribution{2147483647,
+ // 2147483647};
+
+ for (i = 0; i < NUM_ITEMS; i++) {
+ items[i] = i + 3; // 10+rand()%KEYS;
+ // Keys associated with
+ // operations
+ storage.push_back(i + 3);
+ // auto key = (key_type)std::round(normalDistribution(randomEngine));
+ // items[i] = key;
+ // storage.push_back(key);
+ }
+
+ std::sort(storage.begin(), storage.end());
+
+#if 0
+ for (i = 0; i < NUM_ITEMS; i++) {
+ /*int first = rand() % NUM_ITEMS;
+ int second = rand() % NUM_ITEMS;*/
+
+ std::swap(items[uniformIntDistributionArray(randomEngine)],
+ items[uniformIntDistributionArray(randomEngine)]);
+ /*LL temp;
+ temp = items[first];
+ items[first] = items[second];
+ items[second] = temp;*/
+ }
+#endif
+
+ // Pre-generated levels of skip list nodes (relevant only if op[i] is add)
+ // srand(0);
+ for (i = 0; i < NUM_ITEMS; i++) {
+#ifdef RANDOM_HEIGHT
+ levels[i] = RandomLevel(1 / randomEngine) - 1; // 36/14
+#else
+ levels[i] = CustomLevel(items[i]) - 1; // 31/18
+#endif
+ }
+
+ // Populate the sequence of operations
+ for (i = 0; i < (NUM_ITEMS * adds) / 100; i++) {
+ op[i] = ADD;
+ }
+ for (; i < (NUM_ITEMS * (adds + deletes)) / 100; i++) {
+ op[i] = DELETE;
+ }
+ for (; i < NUM_ITEMS; i++) {
+ op[i] = SEARCH;
+ }
+
+ adds = (NUM_ITEMS * adds) / 100;
+
+ // Allocate device memory
+
+ LL *Citems;
+ LL *Cop;
+ LL *Cresult;
+ // LL *Clevels;
+
+ cudaMalloc((void **)&Cresult, sizeof(LL) * NUM_ITEMS);
+ cudaMalloc((void **)&Citems, sizeof(LL) * NUM_ITEMS);
+ cudaMalloc((void **)&Cop, sizeof(LL) * NUM_ITEMS);
+ // cudaMalloc((void **)&Clevels, sizeof(LL) * NUM_ITEMS);
+ // cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS,
+ // cudaMemcpyHostToDevice);
+ cudaMemcpy(Citems, items, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice);
+ cudaMemcpy(Cop, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice);
+ Node **pointers = (Node **)new LL[adds]; // malloc(sizeof(LL) * adds);
+ Node **Cpointers;
+
+ // Allocate the pool of free nodes
+
+ for (i = 0; i < adds; i++) {
+ cudaMalloc((void **)&pointers[i], sizeof(Node));
+ }
+ cudaMalloc((void **)&Cpointers, sizeof(Node *) * adds);
+ cudaMemcpy(Cpointers, pointers, sizeof(Node *) * adds,
+ cudaMemcpyHostToDevice);
+
+ // Allocate the skip list
+
+ LockFreeSkipList *Clist;
+ auto *list = new LockFreeSkipList(levels);
+
+ cudaMalloc((void **)&Clist, sizeof(LockFreeSkipList));
+ cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice);
+ // Calculate the number of thread blocks
+ // NUM_ITEMS = total number of operations to execute
+ // NUM_THREADS = number of threads per block
+ // FACTOR = number of operations per thread
+
+ int blocks = (NUM_ITEMS % (NUM_THREADS * FACTOR) == 0)
+ ? NUM_ITEMS / (NUM_THREADS * FACTOR)
+ : (NUM_ITEMS / (NUM_THREADS * FACTOR)) + 1;
+
+ // Error checking code
+ cudaError_t error = cudaGetLastError();
+ if (cudaSuccess != error) {
+ printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
+ exit(-1);
+ }
+
+ // Initialize the device memory
+ init<<<1, 32>>>(Clist, Cpointers);
+ cudaDeviceSynchronize();
+
+ // Launch main kernel
+
+ cudaEvent_t start, stop;
+ cudaEventCreate(&start);
+ cudaEventCreate(&stop);
+ cudaEventRecord(start, nullptr);
+
+ kernel<<<blocks, NUM_THREADS>>>(Citems, Cop, Cresult);
+ CudaCheckError();
+ error = cudaGetLastError();
+ if (cudaSuccess != error) {
+ printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
+ // exit(-1);
+ }
+ cudaDeviceSynchronize();
+ cudaEventRecord(stop, nullptr);
+ cudaEventSynchronize(stop);
+ float time;
+ cudaEventElapsedTime(&time, start, stop);
+ cudaEventDestroy(start);
+ cudaEventDestroy(stop);
+
+ // Print kernel execution time in milliseconds
+
+ printf("%s %d ", TARGET_STRING, block_size);
+
+ printf("%lu: %lf", NUM_ITEMS, time);
+
+ // Launch main kernel for query
+
+ // Populate the sequence of operations
+ for (i = 0; i < NUM_ITEMS; i++) {
+ op[i] = SEARCH;
+ }
+
+ LL *Cop2;
+ cudaMalloc((void **)&Cop2, sizeof(LL) * NUM_ITEMS);
+ cudaMemcpy(Cop2, op, sizeof(LL) * NUM_ITEMS, cudaMemcpyHostToDevice);
+
+ cudaEventCreate(&start);
+ cudaEventCreate(&stop);
+ cudaEventRecord(start, nullptr);
+#ifdef MEASURE_TIME
+ kernel<<<NUM_ITEMS, 1>>>(Citems, Cop2, Cresult);
+#else
+ kernel<<<blocks, NUM_THREADS>>>(Citems, Cop2, Cresult);
+#endif
+ CudaCheckError();
+ error = cudaGetLastError();
+ if (cudaSuccess != error) {
+ printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
+ // exit(-1);
+ }
+ cudaDeviceSynchronize();
+ cudaEventRecord(stop, nullptr);
+ cudaEventSynchronize(stop);
+ cudaEventElapsedTime(&time, start, stop);
+ cudaEventDestroy(start);
+ cudaEventDestroy(stop);
+
+ // Print kernel execution time in milliseconds
+
+ printf(" %lf\n", time);
+
+ // Check for errors
+
+ error = cudaGetLastError();
+ if (cudaSuccess != error) {
+ printf("error1:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
+ exit(-1);
+ }
+
+ // Move results back to host memory
+
+ cudaMemcpy(result, Cresult, sizeof(LL) * NUM_ITEMS, cudaMemcpyDeviceToHost);
+
+ // Uncomment the following for debugging
+ // print<<<1,32>>>();
+ cudaDeviceSynchronize();
+
+#if (defined(MEASURE_TIME) || defined(MEASURE_ACCESS))
+
+ print_function<<<1, 1>>>();
+
+ cudaDeviceSynchronize();
+#ifdef MEASURE_TIME
+ {
+ int *cuda_tmp = nullptr;
+ cudaMalloc(&cuda_tmp, sizeof(int) * NUM_ITEMS);
+ copy_function<<<1, 1>>>(cuda_tmp);
+ cudaDeviceSynchronize();
+ CudaCheckError();
+ FILE *file = fopen("spend_time.txt", "w");
+ int *tmp = new int[NUM_ITEMS];
+ cudaMemcpy(tmp, cuda_tmp, sizeof(int) * NUM_ITEMS, cudaMemcpyDeviceToHost);
+ // memcpy(tmp, SpendTime, sizeof(int) * NUM_ITEMS);
+ for (i = 0; i < NUM_ITEMS; i++) {
+ if (tmp[i] == 0)
+ break;
+ fprintf(file, "%d\n", tmp[i]);
+ }
+ // printf("%d\n", i);
+ delete[] tmp;
+ fclose(file);
+ }
+ // for (auto element : SpendTimeVec)
+ // printf("%d\n", element);
+#endif
+#endif
+ /*cudaFree(Clist);
+ cudaFree(Cop2);
+ cudaFree(Clevels);
+ cudaFree(Cop);
+ cudaFree(Citems);
+ cudaFree(Cresult);
+ free(pointers);
+ delete [] op;
+ delete [] levels;
+ delete [] items;
+ delete [] result;*/
+ return 0;
+}
diff --git a/main.cu b/main.cu
index 5539991..7f1b27e 100644
--- a/main.cu
+++ b/main.cu
@@ -62,8 +62,8 @@ Conference on Parallel and Distributed Systems, December 2012.
***************************************************************************************/
-//#include"cutil.h" // Comment this if cutil.h is not available
-#include "cuda_runtime.h"
+// #include"cutil.h" // Comment this if cutil.h is not available
+// #include "cuda_runtime.h"
#include <algorithm>
#include <cassert>
#include <cstdio>
@@ -86,7 +86,7 @@ typedef unsigned int LL;
#endif
#define MEASURE_TIME
-//#define MEASURE_ACCESS
+// #define MEASURE_ACCESS
#if (defined(MEASURE_ACCESS) && defined(MEASURE_TIME))
#error "Shouldn't define MEASURE_TIME and MEASURE_ACCESS at the same time"
@@ -98,11 +98,11 @@ typedef unsigned int LL;
#endif
// Maximum level of a node in the skip list
-//#define MAX_LEVEL 32
+// #define MAX_LEVEL 32
constexpr size_t MAX_LEVEL = 16;
// Number of threads per block
-//#define NUM_THREADS 512
+// #define NUM_THREADS 512
constexpr size_t NUM_THREADS = 512;
constexpr size_t NUM_ITEMS = BUILD_SIZE;
diff --git a/skiplist.cuh b/skiplist.cuh
new file mode 100644
index 0000000..54c4732
--- /dev/null
+++ b/skiplist.cuh
@@ -0,0 +1,397 @@
+#pragma once
+#ifndef LOCKFREE_SKIPLIST_CUH_
+#define LOCKFREE_SKIPLIST_CUH_
+
+#ifndef LOCKFREE_KEYTYPE_
+#define LOCKFREE_KEYTYPE_
+typedef unsigned long long key_type;
+#endif
+
+#ifndef BUILD_SIZE
+#define BUILD_SIZE 1048576
+#endif
+
+#ifdef MEASURE_TIME
+#undef BUILD_SIZE
+#define BUILD_SIZE 1024
+#endif
+
+#ifndef LOCKFREE_DEFINITIONS_
+#define LOCKFREE_DEFINITIONS_
+constexpr size_t MAX_LEVEL = 16;
+
+constexpr size_t NUM_ITEMS = BUILD_SIZE;
+#endif
+
+class Node;
+
+__device__ Node **nodes;
+
+// Definition of generic node class
+
+class __attribute__((aligned(16))) Node {
+public:
+ int topLevel; // Level of the node
+ key_type key; // Key value
+ key_type next[MAX_LEVEL + 1]; // Array of next links
+
+ // Create a next field from a reference and mark bit
+ __device__ __host__ key_type CreateRef(Node *ref, bool mark) {
+ auto val = (key_type)ref;
+ val = val | mark;
+ return val;
+ }
+
+ __device__ __host__ void SetRef(int index, Node *ref, bool mark) {
+ next[index] = CreateRef(ref, mark);
+ }
+
+ // Extract the reference from a next field
+ __device__ Node *GetReference(int index) {
+ key_type ref = next[index];
+ return (Node *)((ref >> 1) << 1);
+ }
+
+ // Extract the reference and mark bit from a next field
+ __device__ Node *Get(int index, bool *marked) {
+ marked[0] = next[index] % 2;
+ return (Node *)((next[index] >> 1) << 1);
+ }
+
+ // CompareAndSet wrapper
+ __device__ bool CompareAndSet(int index, Node *expectedRef, Node *newRef,
+ bool oldMark, bool newMark) {
+ key_type oldVal = (key_type)expectedRef | oldMark;
+ key_type newVal = (key_type)newRef | newMark;
+ key_type *ref = &(next[index]);
+ key_type oldValOut = atomicCAS(ref, oldVal, newVal);
+ if (oldValOut == oldVal)
+ return true;
+ return false;
+ }
+
+ // Constructor for sentinel nodes
+ Node(key_type k) {
+ key = k;
+ topLevel = MAX_LEVEL;
+ int i;
+ for (i = 0; i < MAX_LEVEL + 1; i++) {
+ next[i] = CreateRef((Node *)nullptr, false);
+ }
+ }
+};
+
+/*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 {
+ static constexpr size_t SAMPLE_LENGTH = 1024;
+
+public:
+ Node *head;
+ Node *tail;
+
+ unsigned int *pointerIndex;
+
+ key_type *randoms;
+
+ LockFreeSkipList(key_type *randoms) {
+ Node *h = new Node(0);
+ // size_ = 0;
+#if __WORDSIZE == 64
+ Node *t = new Node((key_type)NUM_ITEMS + 10);
+#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++) {
+ h->SetRef(i, tail, false);
+ }
+ cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice);
+
+ cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice);
+
+ cudaMalloc(&pointerIndex, sizeof(unsigned int));
+ initDeviceVariable();
+
+ cudaMalloc(&randoms, sizeof(key_type) * SAMPLE_LENGTH);
+ cudaMemcpy(this->randoms, randoms, sizeof(key_type) * SAMPLE_LENGTH,
+ cudaMemcpyHostToDevice);
+ }
+
+ void initDeviceVariable() const {
+ unsigned int x = 0;
+ cudaMemcpy(pointerIndex, &x, sizeof(unsigned int), cudaMemcpyHostToDevice);
+ }
+
+ __device__ bool find(key_type, Node **, Node **); // Helping method
+ __device__ bool Add(key_type);
+ __device__ bool Delete(key_type);
+ __device__ bool Search(key_type);
+
+#ifdef MEASURE_ACCESS
+ unsigned access_times = 0;
+
+ __device__ unsigned getAccessCount() const { return this->access_times; }
+ __device__ void increaseAccessCount(unsigned count = 1) {
+ atomicAdd(&this->access_times, count);
+ }
+#else
+ __device__ void increaseAccessCount(unsigned _count = 1) {}
+#endif
+
+#ifdef MEASURE_TIME
+ unsigned round = 0;
+ __device__ void increaseRoundCount(unsigned count = 1) {
+ atomicAdd(&this->round, count);
+ }
+ int spend_time[NUM_ITEMS]{0};
+ unsigned long long total_time = 0;
+ __device__ unsigned getRoundCount() const { return this->round; }
+#endif
+};
+
+//__device__ Node **nodes; // Pool of pre-allocated nodes
+//__device__ unsigned int pointerIndex = 0; // Index into pool of free nodes
+//__device__ key_type *randoms; // Array storing the levels of the nodes in the
+// free pool
+
+// 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) {
+ 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->SetRef(i, nullptr, false);
+ }
+ return n;
+}
+
+// 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
+ int bottomLevel = 0;
+ bool marked[] = {false};
+ bool snip;
+ Node *pred;
+ Node *curr = nullptr;
+ Node *succ;
+ bool beenThereDoneThat;
+ while (true) {
+ beenThereDoneThat = false;
+ pred = head;
+ int level;
+ for (level = MAX_LEVEL; level >= bottomLevel; level--) {
+ curr = pred->GetReference(level);
+ while (true) {
+ succ = curr->Get(level, marked);
+ while (marked[0]) {
+ snip = pred->CompareAndSet(level, curr, succ, false, false);
+ beenThereDoneThat = true;
+ if (!snip)
+ break;
+ curr = pred->GetReference(level);
+ succ = curr->Get(level, marked);
+ beenThereDoneThat = false;
+ // printf("find key is %d \n",(int)key);
+ }
+ if (beenThereDoneThat)
+ break;
+ if (curr->key <= key) {
+ pred = curr;
+ curr = succ;
+ } else {
+ break;
+ }
+ }
+ if (beenThereDoneThat)
+ break;
+ preds[level] = pred;
+ succs[level] = curr;
+ }
+ if (beenThereDoneThat)
+ continue;
+ return ((curr->key == key));
+ }
+}
+
+__device__ bool LockFreeSkipList::Search(key_type key) {
+ int bottomLevel = 0;
+ bool marked = false;
+ Node *pred = head;
+ Node *curr = nullptr;
+ Node *succ;
+ int level;
+ for (level = MAX_LEVEL; level >= bottomLevel; level--) {
+ curr = pred->GetReference(level);
+#ifdef MEASURE_ACCESS
+ this->increaseAccessCount();
+#endif
+ while (true) {
+ succ = curr->Get(level, &marked);
+#ifdef MEASURE_ACCESS
+ this->increaseAccessCount();
+#endif
+ while (marked) {
+ curr = curr->GetReference(level);
+ succ = curr->Get(level, &marked);
+#ifdef MEASURE_ACCESS
+ this->increaseAccessCount(2);
+#endif
+ }
+ if (curr->key < key) {
+ pred = curr;
+ curr = succ;
+ } else {
+ break;
+ }
+ }
+ }
+ return (curr != nullptr && curr->key == key);
+}
+
+__device__ bool LockFreeSkipList::Delete(key_type key) {
+ int bottomLevel = 0;
+ Node *preds[MAX_LEVEL + 1];
+ Node *succs[MAX_LEVEL + 1];
+ Node *succ;
+ bool marked[] = {false};
+ while (true) {
+ bool found = find(key, preds, succs);
+ if (!found) {
+ return false;
+ } else {
+ Node *nodeToDelete = succs[bottomLevel];
+ int level;
+ for (level = nodeToDelete->topLevel; level >= bottomLevel + 1; level--) {
+ succ = nodeToDelete->Get(level, marked);
+ while (!marked[0]) {
+ nodeToDelete->CompareAndSet(level, succ, succ, false, true);
+ succ = nodeToDelete->Get(level, marked);
+ }
+ }
+ succ = nodeToDelete->Get(bottomLevel, marked);
+ while (true) {
+ bool iMarkedIt =
+ nodeToDelete->CompareAndSet(bottomLevel, succ, succ, false, true);
+ 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;
+ }
+ }
+ }
+ }
+}
+
+__device__ bool LockFreeSkipList::Add(key_type key) {
+ Node *newNode = GetNewNode(nodes, key, this->pointerIndex, this->randoms);
+ int topLevel = newNode->topLevel;
+ int bottomLevel = 0;
+ Node *preds[MAX_LEVEL + 1];
+ Node *succs[MAX_LEVEL + 1];
+ int level;
+ while (true) {
+ bool found = find(key, preds, succs);
+ if (found) {
+ return false;
+ } else {
+ Node *pred;
+ Node *succ;
+ for (level = bottomLevel; level <= topLevel; level++) {
+ succ = succs[level];
+ newNode->SetRef(level, succ, false);
+ }
+ pred = preds[bottomLevel];
+ succ = succs[bottomLevel];
+ bool t;
+ // printf("--- key is %d pred is %d succ is %d level is %d
+ // \n",(int)key,(int)pred->key,(int)succ->key,0);
+ t = pred->CompareAndSet(bottomLevel, succ, newNode, false, false);
+ if (!t) {
+ continue;
+ }
+ for (level = bottomLevel + 1; level <= topLevel; level++) {
+
+ while (true) {
+ pred = preds[level];
+ succ = succs[level];
+ newNode->SetRef(level, succ, false);
+ // printf("-- key is %d pred is %d succ is %d level is %d
+ // \n",(int)key,(int)pred->key,(int)succ->key,(int)level);
+ if (pred->CompareAndSet(level, succ, newNode, false, false)) {
+ break;
+ }
+ // printf("key is %d pred is %d succ is %d level is %d
+ // \n",(int)key,(int)pred->key,(int)succ->key,(int)level);
+ find(key, preds, succs);
+ }
+ }
+ // size_ += 1;
+ // this->key_map.insert(MapNode(ll, newNode));
+ // atomicAdd(&size_, 1);
+ return true;
+ }
+ }
+}
+#endif \ No newline at end of file
diff --git a/skiplistcustom.cuh b/skiplistcustom.cuh
new file mode 100644
index 0000000..54c4732
--- /dev/null
+++ b/skiplistcustom.cuh
@@ -0,0 +1,397 @@
+#pragma once
+#ifndef LOCKFREE_SKIPLIST_CUH_
+#define LOCKFREE_SKIPLIST_CUH_
+
+#ifndef LOCKFREE_KEYTYPE_
+#define LOCKFREE_KEYTYPE_
+typedef unsigned long long key_type;
+#endif
+
+#ifndef BUILD_SIZE
+#define BUILD_SIZE 1048576
+#endif
+
+#ifdef MEASURE_TIME
+#undef BUILD_SIZE
+#define BUILD_SIZE 1024
+#endif
+
+#ifndef LOCKFREE_DEFINITIONS_
+#define LOCKFREE_DEFINITIONS_
+constexpr size_t MAX_LEVEL = 16;
+
+constexpr size_t NUM_ITEMS = BUILD_SIZE;
+#endif
+
+class Node;
+
+__device__ Node **nodes;
+
+// Definition of generic node class
+
+class __attribute__((aligned(16))) Node {
+public:
+ int topLevel; // Level of the node
+ key_type key; // Key value
+ key_type next[MAX_LEVEL + 1]; // Array of next links
+
+ // Create a next field from a reference and mark bit
+ __device__ __host__ key_type CreateRef(Node *ref, bool mark) {
+ auto val = (key_type)ref;
+ val = val | mark;
+ return val;
+ }
+
+ __device__ __host__ void SetRef(int index, Node *ref, bool mark) {
+ next[index] = CreateRef(ref, mark);
+ }
+
+ // Extract the reference from a next field
+ __device__ Node *GetReference(int index) {
+ key_type ref = next[index];
+ return (Node *)((ref >> 1) << 1);
+ }
+
+ // Extract the reference and mark bit from a next field
+ __device__ Node *Get(int index, bool *marked) {
+ marked[0] = next[index] % 2;
+ return (Node *)((next[index] >> 1) << 1);
+ }
+
+ // CompareAndSet wrapper
+ __device__ bool CompareAndSet(int index, Node *expectedRef, Node *newRef,
+ bool oldMark, bool newMark) {
+ key_type oldVal = (key_type)expectedRef | oldMark;
+ key_type newVal = (key_type)newRef | newMark;
+ key_type *ref = &(next[index]);
+ key_type oldValOut = atomicCAS(ref, oldVal, newVal);
+ if (oldValOut == oldVal)
+ return true;
+ return false;
+ }
+
+ // Constructor for sentinel nodes
+ Node(key_type k) {
+ key = k;
+ topLevel = MAX_LEVEL;
+ int i;
+ for (i = 0; i < MAX_LEVEL + 1; i++) {
+ next[i] = CreateRef((Node *)nullptr, false);
+ }
+ }
+};
+
+/*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 {
+ static constexpr size_t SAMPLE_LENGTH = 1024;
+
+public:
+ Node *head;
+ Node *tail;
+
+ unsigned int *pointerIndex;
+
+ key_type *randoms;
+
+ LockFreeSkipList(key_type *randoms) {
+ Node *h = new Node(0);
+ // size_ = 0;
+#if __WORDSIZE == 64
+ Node *t = new Node((key_type)NUM_ITEMS + 10);
+#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++) {
+ h->SetRef(i, tail, false);
+ }
+ cudaMemcpy(head, h, sizeof(Node), cudaMemcpyHostToDevice);
+
+ cudaMemcpy(tail, t, sizeof(Node), cudaMemcpyHostToDevice);
+
+ cudaMalloc(&pointerIndex, sizeof(unsigned int));
+ initDeviceVariable();
+
+ cudaMalloc(&randoms, sizeof(key_type) * SAMPLE_LENGTH);
+ cudaMemcpy(this->randoms, randoms, sizeof(key_type) * SAMPLE_LENGTH,
+ cudaMemcpyHostToDevice);
+ }
+
+ void initDeviceVariable() const {
+ unsigned int x = 0;
+ cudaMemcpy(pointerIndex, &x, sizeof(unsigned int), cudaMemcpyHostToDevice);
+ }
+
+ __device__ bool find(key_type, Node **, Node **); // Helping method
+ __device__ bool Add(key_type);
+ __device__ bool Delete(key_type);
+ __device__ bool Search(key_type);
+
+#ifdef MEASURE_ACCESS
+ unsigned access_times = 0;
+
+ __device__ unsigned getAccessCount() const { return this->access_times; }
+ __device__ void increaseAccessCount(unsigned count = 1) {
+ atomicAdd(&this->access_times, count);
+ }
+#else
+ __device__ void increaseAccessCount(unsigned _count = 1) {}
+#endif
+
+#ifdef MEASURE_TIME
+ unsigned round = 0;
+ __device__ void increaseRoundCount(unsigned count = 1) {
+ atomicAdd(&this->round, count);
+ }
+ int spend_time[NUM_ITEMS]{0};
+ unsigned long long total_time = 0;
+ __device__ unsigned getRoundCount() const { return this->round; }
+#endif
+};
+
+//__device__ Node **nodes; // Pool of pre-allocated nodes
+//__device__ unsigned int pointerIndex = 0; // Index into pool of free nodes
+//__device__ key_type *randoms; // Array storing the levels of the nodes in the
+// free pool
+
+// 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) {
+ 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->SetRef(i, nullptr, false);
+ }
+ return n;
+}
+
+// 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
+ int bottomLevel = 0;
+ bool marked[] = {false};
+ bool snip;
+ Node *pred;
+ Node *curr = nullptr;
+ Node *succ;
+ bool beenThereDoneThat;
+ while (true) {
+ beenThereDoneThat = false;
+ pred = head;
+ int level;
+ for (level = MAX_LEVEL; level >= bottomLevel; level--) {
+ curr = pred->GetReference(level);
+ while (true) {
+ succ = curr->Get(level, marked);
+ while (marked[0]) {
+ snip = pred->CompareAndSet(level, curr, succ, false, false);
+ beenThereDoneThat = true;
+ if (!snip)
+ break;
+ curr = pred->GetReference(level);
+ succ = curr->Get(level, marked);
+ beenThereDoneThat = false;
+ // printf("find key is %d \n",(int)key);
+ }
+ if (beenThereDoneThat)
+ break;
+ if (curr->key <= key) {
+ pred = curr;
+ curr = succ;
+ } else {
+ break;
+ }
+ }
+ if (beenThereDoneThat)
+ break;
+ preds[level] = pred;
+ succs[level] = curr;
+ }
+ if (beenThereDoneThat)
+ continue;
+ return ((curr->key == key));
+ }
+}
+
+__device__ bool LockFreeSkipList::Search(key_type key) {
+ int bottomLevel = 0;
+ bool marked = false;
+ Node *pred = head;
+ Node *curr = nullptr;
+ Node *succ;
+ int level;
+ for (level = MAX_LEVEL; level >= bottomLevel; level--) {
+ curr = pred->GetReference(level);
+#ifdef MEASURE_ACCESS
+ this->increaseAccessCount();
+#endif
+ while (true) {
+ succ = curr->Get(level, &marked);
+#ifdef MEASURE_ACCESS
+ this->increaseAccessCount();
+#endif
+ while (marked) {
+ curr = curr->GetReference(level);
+ succ = curr->Get(level, &marked);
+#ifdef MEASURE_ACCESS
+ this->increaseAccessCount(2);
+#endif
+ }
+ if (curr->key < key) {
+ pred = curr;
+ curr = succ;
+ } else {
+ break;
+ }
+ }
+ }
+ return (curr != nullptr && curr->key == key);
+}
+
+__device__ bool LockFreeSkipList::Delete(key_type key) {
+ int bottomLevel = 0;
+ Node *preds[MAX_LEVEL + 1];
+ Node *succs[MAX_LEVEL + 1];
+ Node *succ;
+ bool marked[] = {false};
+ while (true) {
+ bool found = find(key, preds, succs);
+ if (!found) {
+ return false;
+ } else {
+ Node *nodeToDelete = succs[bottomLevel];
+ int level;
+ for (level = nodeToDelete->topLevel; level >= bottomLevel + 1; level--) {
+ succ = nodeToDelete->Get(level, marked);
+ while (!marked[0]) {
+ nodeToDelete->CompareAndSet(level, succ, succ, false, true);
+ succ = nodeToDelete->Get(level, marked);
+ }
+ }
+ succ = nodeToDelete->Get(bottomLevel, marked);
+ while (true) {
+ bool iMarkedIt =
+ nodeToDelete->CompareAndSet(bottomLevel, succ, succ, false, true);
+ 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;
+ }
+ }
+ }
+ }
+}
+
+__device__ bool LockFreeSkipList::Add(key_type key) {
+ Node *newNode = GetNewNode(nodes, key, this->pointerIndex, this->randoms);
+ int topLevel = newNode->topLevel;
+ int bottomLevel = 0;
+ Node *preds[MAX_LEVEL + 1];
+ Node *succs[MAX_LEVEL + 1];
+ int level;
+ while (true) {
+ bool found = find(key, preds, succs);
+ if (found) {
+ return false;
+ } else {
+ Node *pred;
+ Node *succ;
+ for (level = bottomLevel; level <= topLevel; level++) {
+ succ = succs[level];
+ newNode->SetRef(level, succ, false);
+ }
+ pred = preds[bottomLevel];
+ succ = succs[bottomLevel];
+ bool t;
+ // printf("--- key is %d pred is %d succ is %d level is %d
+ // \n",(int)key,(int)pred->key,(int)succ->key,0);
+ t = pred->CompareAndSet(bottomLevel, succ, newNode, false, false);
+ if (!t) {
+ continue;
+ }
+ for (level = bottomLevel + 1; level <= topLevel; level++) {
+
+ while (true) {
+ pred = preds[level];
+ succ = succs[level];
+ newNode->SetRef(level, succ, false);
+ // printf("-- key is %d pred is %d succ is %d level is %d
+ // \n",(int)key,(int)pred->key,(int)succ->key,(int)level);
+ if (pred->CompareAndSet(level, succ, newNode, false, false)) {
+ break;
+ }
+ // printf("key is %d pred is %d succ is %d level is %d
+ // \n",(int)key,(int)pred->key,(int)succ->key,(int)level);
+ find(key, preds, succs);
+ }
+ }
+ // size_ += 1;
+ // this->key_map.insert(MapNode(ll, newNode));
+ // atomicAdd(&size_, 1);
+ return true;
+ }
+ }
+}
+#endif \ No newline at end of file
generated by cgit v1.3.1 (git 2.54.0) at 2026-09-04 07:01:16 +0000