summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--CMakeLists.txt2
-rw-r--r--expt_0520.cu346
-rw-r--r--read_helper.h45
-rw-r--r--skiplistcustom.cuh117
-rw-r--r--sortlib.cuh16
5 files changed, 161 insertions, 365 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 883dbc4..118786c 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -222,7 +222,7 @@ 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)
+add_executable(expt_0520 expt_0520.cu skiplistcustom.cuh read_helper.h)
target_link_libraries(expt_0520 m stdc++)
set_target_properties(expt_0520 PROPERTIES
diff --git a/expt_0520.cu b/expt_0520.cu
index 839fef8..88815a7 100644
--- a/expt_0520.cu
+++ b/expt_0520.cu
@@ -64,12 +64,11 @@ 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 "read_helper.h"
+#include "skiplistcustom.cuh"
#include <cassert>
#include <cstdio>
#include <cstdlib>
-#include <random>
#include <set>
#if __WORDSIZE == 64
@@ -116,11 +115,6 @@ 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
@@ -164,126 +158,31 @@ inline void __cudaCheckError(const char *file, const int line) {
#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");
- }
-}
+__global__ void init(Node **n) { nodes = n; }
// The main kernel
-__global__ void kernel(LL *items, LL *op, LL *result) {
+__global__ void kernel(LockFreeSkipList *skipList, const key_type *population,
+ size_t insertion_length) {
// 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;
+ for (int 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)
+ auto tid =
+ i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid >= insertion_length)
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;
- }
+ key_type item = population[tid];
+ skipList->Add(item);
}
- 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() {
@@ -292,155 +191,83 @@ __global__ void print_function() {
#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., "
+ printf("Need two arguments: sample_length add insertion_length 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);
+ auto error = cudaGetLastError();
+ auto total_row = 0UL;
+ auto sample_length = strtol(argv[1], nullptr, 0);
+ auto insertion_length = strtol(argv[2], nullptr, 0);
+ printf("insert_length %ld, sample_length: %ld\n", insertion_length,
+ sample_length);
- storage.reserve(NUM_ITEMS);
+ ReadHelper reader("normal_distribution.txt", sample_length, insertion_length);
+ printf("%d\n", __LINE__);
+ reader.readFile(total_row);
- 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;*/
+ printf("%d\n", __LINE__);
+ key_type *cudaPopulation;
- 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;
+ printf("%d\n", __LINE__);
+ cudaMalloc(&cudaPopulation, sizeof(key_type) * insertion_length);
+ std::vector<key_type> _sample, _population;
+ reader.split_into(_sample, _population);
+ cudaMemcpy(cudaPopulation, _population.data(),
+ sizeof(key_type) * insertion_length, cudaMemcpyHostToDevice);
+ printf("%d\n", __LINE__);
// 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;
+ Node **pointers =
+ (Node **)new LL[insertion_length]; // malloc(sizeof(LL) * adds);
+ Node **cudaNodePointers;
+
+ printf("%d\n", __LINE__);
// Allocate the pool of free nodes
- for (i = 0; i < adds; i++) {
- cudaMalloc((void **)&pointers[i], sizeof(Node));
+ for (int i = 0; i < insertion_length; i++) {
+ cudaMalloc(&pointers[i], sizeof(Node));
}
- cudaMalloc((void **)&Cpointers, sizeof(Node *) * adds);
- cudaMemcpy(Cpointers, pointers, sizeof(Node *) * adds,
+ cudaMalloc((void **)&cudaNodePointers, sizeof(Node *) * insertion_length);
+ cudaMemcpy(cudaNodePointers, pointers, sizeof(Node *) * insertion_length,
cudaMemcpyHostToDevice);
+ CudaCheckError();
+ printf("%d\n", __LINE__);
// Allocate the skip list
LockFreeSkipList *Clist;
- auto *list = new LockFreeSkipList(levels);
+ auto *list = new LockFreeSkipList(_sample.data(), SAMPLE_SIZE);
cudaMalloc((void **)&Clist, sizeof(LockFreeSkipList));
cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice);
+ CudaCheckError();
// 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;
+ size_t blocks = (insertion_length % (NUM_THREADS * FACTOR) == 0)
+ ? insertion_length / (NUM_THREADS * FACTOR)
+ : (insertion_length / (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);
+ init<<<1, 32>>>(cudaNodePointers);
cudaDeviceSynchronize();
// Launch main kernel
@@ -450,13 +277,8 @@ int main(int argc, char **argv) {
cudaEventCreate(&stop);
cudaEventRecord(start, nullptr);
- kernel<<<blocks, NUM_THREADS>>>(Citems, Cop, Cresult);
+ kernel<<<blocks, NUM_THREADS>>>(Clist, cudaPopulation, insertion_length);
CudaCheckError();
- error = cudaGetLastError();
- if (cudaSuccess != error) {
- printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
- // exit(-1);
- }
cudaDeviceSynchronize();
cudaEventRecord(stop, nullptr);
cudaEventSynchronize(stop);
@@ -471,58 +293,6 @@ int main(int argc, char **argv) {
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>>>();
@@ -552,16 +322,12 @@ int main(int argc, char **argv) {
// 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;*/
+ cudaFree(Clist);
+ for (int i = 0; i < insertion_length; i++) {
+ cudaFree(pointers[i]);
+ }
+ cudaFree(cudaNodePointers);
+ delete list;
+ delete[] pointers;
return 0;
}
diff --git a/read_helper.h b/read_helper.h
index 311ebbb..292eaac 100644
--- a/read_helper.h
+++ b/read_helper.h
@@ -15,12 +15,18 @@
#endif
class ReadHelper {
- unsigned long long max_value = 0, min_value = 0x7fffffffff;
+ typedef unsigned long long key_type;
+ key_type max_value = 0, min_value = 0x7fffffffff;
char const *filename;
+ static constexpr size_t REVERSED_BLOCK = 256;
public:
- ReadHelper(char const *filename) : filename(filename) {}
+ size_t sample_length, needed_read_length;
+ ReadHelper(char const *filename, size_t sample_length,
+ size_t insertion_length)
+ : filename(filename), sample_length(sample_length),
+ needed_read_length(insertion_length + sample_length) {}
unsigned long long maxValue() const { return max_value; }
unsigned long long minValue() const { return min_value; }
@@ -35,6 +41,11 @@ public:
population_vector.push_back(value);
}
+ void finish_read() {
+ this->population_vector.push_back(this->max_value);
+ this->population_vector.push_back(this->min_value);
+ }
+
static unsigned long randomRow(unsigned long max_value_) {
std::random_device randomDevice;
std::mt19937 mt19937(randomDevice());
@@ -42,20 +53,21 @@ public:
return dst(mt19937);
}
- void readFile(long sample_length, unsigned long &total_row,
- unsigned long max_number) {
+ void readFile(unsigned long &total_row) {
+ this->population_vector.clear();
auto read_number = 0UL;
FILE *file = fopen(filename, "r");
assert(file);
P_ERR("Reading sample");
for (long long i;
- read_number < max_number && fscanf(file, "%lld ", &i) != EOF;
+ read_number < needed_read_length && fscanf(file, "%lld ", &i) != EOF;
store_into_vector(i))
read_number++;
if (total_row > 0) {
P_ERR("\rReading skip");
- read_number = randomRow(total_row - sample_length - max_number - 256);
+ read_number = randomRow(total_row - sample_length - needed_read_length -
+ REVERSED_BLOCK);
total_row = read_number;
// fprintf(stderr, "Skip %lu\n", read_number);
for (long long i; read_number > 0 && fscanf(file, "%lld ", &i) != EOF;)
@@ -65,12 +77,31 @@ public:
P_ERR("\rReading population");
read_number = 0;
- auto remain = sample_length + 256;
+ auto remain = sample_length + REVERSED_BLOCK;
for (long long i; read_number < remain && fscanf(file, "%lld ", &i) != EOF;
store_into_vector(i))
read_number++;
fclose(file);
P_ERR("\r");
+ finish_read();
+ }
+
+ void split_into(std::vector<key_type> &sample, std::vector<key_type> &p) {
+ sample.reserve(sample_length);
+ p.reserve(needed_read_length - sample_length);
+ memcpy(sample.data(), population_vector.data(),
+ sizeof(key_type) * sample_length);
+ memcpy(p.data(), population_vector.data() + sample_length,
+ sizeof(key_type) * (needed_read_length - sample_length));
+ }
+
+ void split_into(key_type *&sample, key_type *&p) {
+ sample = new key_type[sample_length];
+ p = new key_type[needed_read_length - sample_length];
+ memcpy(sample, population_vector.data(), sizeof(key_type) * sample_length);
+ memcpy(p, population_vector.data() + sample_length,
+ sizeof(key_type) * (needed_read_length - sample_length));
}
};
+
#endif \ No newline at end of file
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;
+
+ CustomSort searcher;
- LockFreeSkipList(key_type *randoms) {
+ 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];
diff --git a/sortlib.cuh b/sortlib.cuh
index 5bacc11..ae0ea67 100644
--- a/sortlib.cuh
+++ b/sortlib.cuh
@@ -3,6 +3,7 @@
#include <cassert>
#include <cstdio>
+#include <vector>
typedef unsigned long long key_type;
@@ -176,8 +177,23 @@ public:
puts("You can remove this function if passed already");
#endif
}
+
+ __host__ __device__ size_t length() const { return this->LENGTH; }
};
+template <typename T> void static rebuild(std::vector<T> original) {
+ auto sample_length = original.size();
+ auto sorter = CustomSort(sample_length);
+ auto tmp = new T[sample_length];
+
+ for (size_t i = 0; i < sample_length; i++) {
+ tmp[i] = original[sorter.calculate_index(i) - 1];
+ }
+ memcpy(original.data(), tmp, sizeof(T) * sample_length);
+
+ delete[] tmp;
+}
+
__global__ void testCustomCalculation(size_t length) {
CustomSort(length, sizeof(long) * 8).testCalculation();
}
generated by cgit v1.3.1 (git 2.54.0) at 2026-09-04 07:03:52 +0000