From 1fbb943005ad78335f0b921be516b3df8140406d Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Sun, 4 Jun 2023 18:34:21 +0800 Subject: --- CMakeLists.txt | 11 +- analysis0510.py | 10 +- expt_0501.cu | 1 + expt_0503.cu | 4 +- expt_0516.cu | 2 +- expt_0516_2.cu | 1 + expt_0520.cu | 4 +- expt_0604.cu | 73 +++++ read_helper.h | 5 +- read_helper_p.h | 111 +++++++ work.cu | 893 -------------------------------------------------------- work_0719.cu | 893 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++ 12 files changed, 1098 insertions(+), 910 deletions(-) create mode 100644 expt_0604.cu create mode 100644 read_helper_p.h delete mode 100644 work.cu create mode 100644 work_0719.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index 6fb4891..02efadf 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -87,12 +87,8 @@ set_target_properties(normal_distribution_expt3 PROPERTIES LINKER_LANGUAGE CXX) add_executable(normal_distribution_expt0702 expt_0702.cpp) set_target_properties(normal_distribution_expt0702 PROPERTIES LINKER_LANGUAGE CXX) -add_executable(work work.cu) -target_link_libraries(work m stdc++) - -set_target_properties(work PROPERTIES - CUDA_SEPARABLE_COMPILATION ON) -set_target_properties(work PROPERTIES CUDA_ARCHITECTURES "75") +add_executable(work_0719 work_0719.cu) +set_cuda_target_base(work_0719) add_executable(expt_0722 expt_0722.cpp) set_target_properties(expt_0722 PROPERTIES LINKER_LANGUAGE CXX) @@ -177,5 +173,6 @@ set_cuda_target_base(expt_0528) add_executable(expt_0529 expt_0529.cu read_helper.h) set_cuda_target_base(expt_0529) - +add_executable(expt_0604 expt_0604.cu read_helper.h) +set_cuda_target_base(expt_0604) diff --git a/analysis0510.py b/analysis0510.py index 59cc65c..78290ab 100755 --- a/analysis0510.py +++ b/analysis0510.py @@ -6,7 +6,7 @@ def read_from_line(line: str) -> dict[str, int | float]: ret = {} if 'skip:' in line: ret.update({'skip': int(line.split('skip:', 1)[1].split(',')[0])}) - ret.update({'sample': int(line.split('length:', 1)[1].split('|', 1)[0])}) + ret.update({'sample': int(line.split('length:', 1)[1].split('custom time:', 1)[0])}) ret.update({'custom': float(line.split('custom time: ', 1)[1].split('time')[0])}) ret.update({'normal': float(line.rsplit('time:', 1)[1])}) return ret @@ -30,6 +30,11 @@ def get_array_summary(sz: list[float]) -> tuple[str, str]: return f'{middle:.6f}', f'{avg:.6f}' +def calc_improve(element: list[int]) -> float: + a, b = list(map(float, element[:2])) + return a / b if a > b else b / a + + def read(file: str, need_calc: bool = False, output_latex: bool = False) -> None: items = {} with open(file) as fin: @@ -50,8 +55,9 @@ def read(file: str, need_calc: bool = False, output_latex: bool = False) -> None new_result[key][x].append(element[x]) print(*keys) for key, value in new_result.items(): - column = list(get_array_summary(value[x]) for x in keys) + column = list(get_array_summary(value[x])[1] for x in keys) if output_latex: + column.append(f'{calc_improve(column):.6f}') column = list(map(str, column)) column.insert(0, str(key)) print('\\hline\n', ' & '.join(column), '\\\\') diff --git a/expt_0501.cu b/expt_0501.cu index 2d740e4..dcd25f5 100644 --- a/expt_0501.cu +++ b/expt_0501.cu @@ -1,5 +1,6 @@ // Experimental content: Test the correctness of binary search and special // search +#define ENABLE_SORT_TEST #include "sortlib.cuh" #include diff --git a/expt_0503.cu b/expt_0503.cu index 58f35aa..0ada96c 100644 --- a/expt_0503.cu +++ b/expt_0503.cu @@ -1,4 +1,5 @@ // Experimental content: Test branch performance +#define ENABLE_SORT_TEST #include "sortlib.cuh" #include @@ -59,8 +60,7 @@ __global__ void kernel(unsigned long step, const double slice_size, // printf("%d\n", tid); auto index = (int)(custom_sort.sample_cdf_custom_version( - cudaSampleItem, cudaSampleItem + cuda_sample_length, - cudaPopulationItem[tid]) / + cudaSampleItem, cudaPopulationItem[tid]) / slice_size) - 1; cdf_result[index] = true; diff --git a/expt_0516.cu b/expt_0516.cu index 351a8c4..48b3811 100644 --- a/expt_0516.cu +++ b/expt_0516.cu @@ -1,6 +1,6 @@ // Experimental content: Test branch performance (optimized memory and output) // #define PRINT_READ_PROCESS -#include "read_helper.h" +#include "read_helper_p.h" #include "sortlib.cuh" #include diff --git a/expt_0516_2.cu b/expt_0516_2.cu index 860e146..9171d01 100644 --- a/expt_0516_2.cu +++ b/expt_0516_2.cu @@ -1,5 +1,6 @@ // Experimental content: Test CustomSort Calculation #define DISABLE_TEST_WARNING +#define ENABLE_SORT_TEST #include "sortlib.cuh" #include diff --git a/expt_0520.cu b/expt_0520.cu index bd1439f..42cdab3 100644 --- a/expt_0520.cu +++ b/expt_0520.cu @@ -202,14 +202,12 @@ int main(int argc, char **argv) { } 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); - ReadHelper reader("normal_distribution.txt", sample_length, insertion_length, - total_row); + ReadHelper reader("normal_distribution.txt", sample_length, insertion_length); reader.readFile(); key_type *cudaPopulation; diff --git a/expt_0604.cu b/expt_0604.cu new file mode 100644 index 0000000..bd64ad0 --- /dev/null +++ b/expt_0604.cu @@ -0,0 +1,73 @@ +// Experimental content: Print level by cdf + +#include "sortlib.cuh" +#define READ_NO_OUTPUT +#include "read_helper.h" + +long pow_for_sample(long n) { + auto x = 2; + for (int i = 1; i < n; i++) { + x *= 2; + } + return x - 1; +} + +static unsigned trailing_zeroes(size_t index) { + constexpr auto block_size = 2; + unsigned bits = 0; + auto x = index / block_size; + + if (x) { + while (x % block_size == 0) { + ++bits; + x /= block_size; + } + } + return bits; +} + +inline double calcSliceSize(size_t insertion_size) { + return 1.0 / (double)insertion_size; +} + +constexpr auto RESULT_LENGTH = 35; + +int main(int argc, char const *argv[]) { + + if (argc != 3) { + printf("Usage %s [sample(pow)] [population]\n", argv[0]); + return 1; + } + + auto sample_length = pow_for_sample(strtol(argv[1], nullptr, 10)); + auto population_length = strtol(argv[2], nullptr, 10); + + printf("sample length: %ld, population length: %ld\n", sample_length, + population_length); + + ReadHelper readHelper("normal_distribution.txt", sample_length, 0); + readHelper.readFile(); + + std::vector sample, population; + std::vector result(RESULT_LENGTH); + readHelper.split_into(sample, population); + rebuildSort(sample); + + auto scale_size = calcSliceSize(population_length); + + auto sort = CustomSort(sample_length, sizeof(key_type) * 8); + + for (auto element : readHelper.population_vector) { + auto ret = sort.sample_cdf_custom_version(sample.data(), element); + auto cdf_index = ret / scale_size; + auto index = trailing_zeroes((size_t)cdf_index); + result[index]++; + } + + for (int i = 0; i < 32; i++) { + if (!result[i]) { + continue; + } + printf("%d: %d\n", i, result[i]); + } +} diff --git a/read_helper.h b/read_helper.h index f8396f2..2256ef8 100644 --- a/read_helper.h +++ b/read_helper.h @@ -65,14 +65,14 @@ public: return dst(mt19937); } - void readFile() { + bool readFile() { this->population_vector.clear(); auto read_number = 0UL; FILE *file = fopen(filename, "r"); if (file == nullptr) { fprintf(stderr, "Unable to open file %s\n", filename); - exit(1); + return false; } P_ERR("Reading sample"); @@ -98,6 +98,7 @@ public: read_number++; fclose(file); P_ERR("\r"); + return true; } void split_into(std::vector &sample, std::vector &p) { diff --git a/read_helper_p.h b/read_helper_p.h new file mode 100644 index 0000000..1b1d0eb --- /dev/null +++ b/read_helper_p.h @@ -0,0 +1,111 @@ + +#ifndef LOCKFREE_READ_HELPER_H +#define LOCKFREE_READ_HELPER_H + +#include +#include +#include +#include + +#ifndef P_ERR +#ifdef PRINT_READ_PROCESS +#define P_ERR(...) fprintf(stderr, __VA_ARGS__) +#else +#define P_ERR(...) +#endif +#endif + +template class ReadHelper { + // typedef unsigned long long key_type; + + key_type max_value = std::numeric_limits::min(), + min_value = std::numeric_limits::max(); + + inline void store_into_vector(key_type value) { + if (max_value < value) { + max_value = value; + } + if (min_value > value) { + min_value = value; + } + population_vector.push_back(value); + } + + char const *filename; + static constexpr size_t REVERSED_BLOCK = 256; + + auto genRandomRow(size_t total_row) const { + if (total_row != 0) { + assert(total_row > (population_length + sample_length + REVERSED_BLOCK)); + return randomRow(total_row - population_length - sample_length - + REVERSED_BLOCK); + } + return 0UL; + } + +public: + const size_t sample_length, population_length, random_number; + ReadHelper(char const *filename, size_t sample_length, + size_t population_length, size_t total_row = 0) + : filename(filename), sample_length(sample_length), + population_length(population_length), + random_number(genRandomRow(total_row)) {} + key_type maxValue() const { return max_value; } + key_type minValue() const { return min_value; } + + std::vector population_vector; + + static unsigned long randomRow(unsigned long max_value_) { + std::random_device randomDevice; + std::mt19937 mt19937(randomDevice()); + std::uniform_int_distribution dst(0, max_value_); + return dst(mt19937); + } + + bool readFile() { + this->population_vector.clear(); + auto read_number = 0UL; + std::ifstream fin(filename); + + if (!fin.is_open()) { + return false; + } + + for (key_type i; read_number < sample_length && !fin.eof(); + store_into_vector(i)) { + fin >> i; + read_number++; + } + + if (random_number > 0) { + read_number = random_number; + for (key_type i; read_number > 0 && !fin.eof(); fin >> i) + read_number--; + } + + read_number = 0; + + auto remain = population_length + REVERSED_BLOCK; + for (key_type i; read_number < remain && !fin.eof(); store_into_vector(i)) { + fin >> i; + read_number++; + } + fin.close(); + return true; + } + + void split_into(std::vector &sample, std::vector &p) { + sample.resize(sample_length - 2); + p.resize(population_length); + memcpy(sample.data(), population_vector.data(), + sizeof(key_type) * (sample_length - 2)); + sample.push_back(this->max_value); + sample.push_back(this->min_value); + memcpy(p.data(), population_vector.data() + sample_length, + sizeof(key_type) * population_length); + } + + size_t size() const { return this->population_vector.size(); } +}; + +#endif \ No newline at end of file diff --git a/work.cu b/work.cu deleted file mode 100644 index 46bdfb4..0000000 --- a/work.cu +++ /dev/null @@ -1,893 +0,0 @@ -/* - -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 -#include -#include -#include -#include -#include - -#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 -} - -class Node; - -// Definition of generic node class - -class __attribute__((aligned(16))) Node { -public: - int topLevel; // Level of the node - LL key; // Key value - LL next[MAX_LEVEL + 1]; // Array of next links - - // Create a next field from a reference and mark bit - __device__ __host__ LL CreateRef(Node *ref, bool mark) { - LL val = (LL)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) { - LL 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) { - LL oldVal = (LL)expectedRef | oldMark; - LL newVal = (LL)newRef | newMark; - LL *ref = &(next[index]); - LL oldValOut = atomicCAS(ref, oldVal, newVal); - if (oldValOut == oldVal) - return true; - return false; - } - - // Constructor for sentinel nodes - Node(LL k) { - key = k; - topLevel = MAX_LEVEL; - int i; - for (i = 0; i < MAX_LEVEL + 1; i++) { - next[i] = CreateRef((Node *)nullptr, false); - } - } -}; - -// Definition of lock-free skip list - -class LockFreeSkipList { -public: - Node *head; - Node *tail; - LockFreeSkipList() { - Node *h = new Node(0); - // size_ = 0; -#if __WORDSIZE == 64 - Node *t = new Node((LL)NUM_ITEMS + 10); -#else - Node *t = new Node((LL)0xffffffff); -#endif - cudaMalloc((void **)&head, sizeof(Node)); - - cudaMalloc((void **)&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); - } - __device__ bool find(LL, Node **, Node **); // Helping method - __device__ bool Add(LL); - __device__ bool Delete(LL); - __device__ bool Search(LL); - -#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__ LL - *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(LL key) { - LL 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; -} - -__device__ LockFreeSkipList *l; // The lock-free skip list - -__device__ LL KeyIndex[KEY_INDEX_SIZE]; - -__device__ key_type SampleStorage[SAMPLE_SIZE]; - -// Kernel for initializing device memory - -__global__ void init(LockFreeSkipList *l1, Node **n, LL *rands) { - randoms = rands; - nodes = n; - l = l1; -} - -// Find the window holding key -// On the way clean up logically deleted nodes (those with set marked bit) - -__device__ bool -LockFreeSkipList::find(LL 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(LL 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(LL 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(LL key) { - Node *newNode = GetNewNode(key); - 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); - } - } - return true; - } - } -} - -__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 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 -} - -__global__ void copy_function(int *spend_time) { - memcpy(spend_time, l->spend_time, sizeof(int) * NUM_ITEMS); -} - -__device__ key_type *cudaBinarySearch(key_type *start, key_type *end, - key_type val) { - auto begin = start; - key_type *last_known_point = nullptr; - while (begin < end) { - auto mid = (end - begin) / 2; - auto mid_val = *(start + mid); - if (val == mid_val) { - return start + mid; - } else if (val > mid_val) { - begin = begin + mid + 1; - } else { - end = end - mid - 1; - } - last_known_point = begin; - } - return last_known_point; -} - -__device__ long double sample_cdf(long double x) { - auto it = cudaBinarySearch(SampleStorage, SampleStorage + SAMPLE_SIZE, x); - if (it == SampleStorage + SAMPLE_SIZE) { - return 1; - } - if (it == SampleStorage) { - return 0; - } - auto it_prev = it - 1; - return (double(it_prev - SampleStorage) + - (x - (long double)*it_prev) / (long double)(*it - *it_prev)) / - double(SAMPLE_SIZE - 1); -} - -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 uniformIntDistributionArray(0, - NUM_ITEMS - 1); - // std::vector storage; - - // std::normal_distribution 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(); - - 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, Clevels); - cudaDeviceSynchronize(); - - // Launch main kernel - - cudaEvent_t start, stop; - cudaEventCreate(&start); - cudaEventCreate(&stop); - cudaEventRecord(start, nullptr); - - kernel<<>>(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<<>>(Citems, Cop2, Cresult); -#else - kernel<<>>(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/work_0719.cu b/work_0719.cu new file mode 100644 index 0000000..46bdfb4 --- /dev/null +++ b/work_0719.cu @@ -0,0 +1,893 @@ +/* + +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 +#include +#include +#include +#include +#include + +#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 +} + +class Node; + +// Definition of generic node class + +class __attribute__((aligned(16))) Node { +public: + int topLevel; // Level of the node + LL key; // Key value + LL next[MAX_LEVEL + 1]; // Array of next links + + // Create a next field from a reference and mark bit + __device__ __host__ LL CreateRef(Node *ref, bool mark) { + LL val = (LL)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) { + LL 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) { + LL oldVal = (LL)expectedRef | oldMark; + LL newVal = (LL)newRef | newMark; + LL *ref = &(next[index]); + LL oldValOut = atomicCAS(ref, oldVal, newVal); + if (oldValOut == oldVal) + return true; + return false; + } + + // Constructor for sentinel nodes + Node(LL k) { + key = k; + topLevel = MAX_LEVEL; + int i; + for (i = 0; i < MAX_LEVEL + 1; i++) { + next[i] = CreateRef((Node *)nullptr, false); + } + } +}; + +// Definition of lock-free skip list + +class LockFreeSkipList { +public: + Node *head; + Node *tail; + LockFreeSkipList() { + Node *h = new Node(0); + // size_ = 0; +#if __WORDSIZE == 64 + Node *t = new Node((LL)NUM_ITEMS + 10); +#else + Node *t = new Node((LL)0xffffffff); +#endif + cudaMalloc((void **)&head, sizeof(Node)); + + cudaMalloc((void **)&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); + } + __device__ bool find(LL, Node **, Node **); // Helping method + __device__ bool Add(LL); + __device__ bool Delete(LL); + __device__ bool Search(LL); + +#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__ LL + *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(LL key) { + LL 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; +} + +__device__ LockFreeSkipList *l; // The lock-free skip list + +__device__ LL KeyIndex[KEY_INDEX_SIZE]; + +__device__ key_type SampleStorage[SAMPLE_SIZE]; + +// Kernel for initializing device memory + +__global__ void init(LockFreeSkipList *l1, Node **n, LL *rands) { + randoms = rands; + nodes = n; + l = l1; +} + +// Find the window holding key +// On the way clean up logically deleted nodes (those with set marked bit) + +__device__ bool +LockFreeSkipList::find(LL 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(LL 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(LL 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(LL key) { + Node *newNode = GetNewNode(key); + 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); + } + } + return true; + } + } +} + +__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 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 +} + +__global__ void copy_function(int *spend_time) { + memcpy(spend_time, l->spend_time, sizeof(int) * NUM_ITEMS); +} + +__device__ key_type *cudaBinarySearch(key_type *start, key_type *end, + key_type val) { + auto begin = start; + key_type *last_known_point = nullptr; + while (begin < end) { + auto mid = (end - begin) / 2; + auto mid_val = *(start + mid); + if (val == mid_val) { + return start + mid; + } else if (val > mid_val) { + begin = begin + mid + 1; + } else { + end = end - mid - 1; + } + last_known_point = begin; + } + return last_known_point; +} + +__device__ long double sample_cdf(long double x) { + auto it = cudaBinarySearch(SampleStorage, SampleStorage + SAMPLE_SIZE, x); + if (it == SampleStorage + SAMPLE_SIZE) { + return 1; + } + if (it == SampleStorage) { + return 0; + } + auto it_prev = it - 1; + return (double(it_prev - SampleStorage) + + (x - (long double)*it_prev) / (long double)(*it - *it_prev)) / + double(SAMPLE_SIZE - 1); +} + +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 uniformIntDistributionArray(0, + NUM_ITEMS - 1); + // std::vector storage; + + // std::normal_distribution 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(); + + 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, Clevels); + cudaDeviceSynchronize(); + + // Launch main kernel + + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + cudaEventRecord(start, nullptr); + + kernel<<>>(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<<>>(Citems, Cop2, Cresult); +#else + kernel<<>>(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; +} -- cgit v1.3.1