From 44a035ec3ebb094adb162ba0c0d374a3fe153094 Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Sun, 28 May 2023 17:31:57 +0800 Subject: feat(exp): Remove useless arguments Signed-off-by: KunoiSayami --- expt_0525.cu | 73 ++++++++++++++++++++++------------------------------------- read_helper.h | 11 +++++---- 2 files changed, 34 insertions(+), 50 deletions(-) diff --git a/expt_0525.cu b/expt_0525.cu index 3c5e5c2..b4e4da5 100644 --- a/expt_0525.cu +++ b/expt_0525.cu @@ -64,28 +64,17 @@ Conference on Parallel and Distributed Systems, December 2012. // #include"cutil.h" // Comment this if cutil.h is not available // #include "cuda_runtime.h" -#include "read_helper.h" #include "sortlib.cuh" -#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 +#define READ_NO_OUTPUT +#include "read_helper.h" -#ifndef STEP_SIZE -#define STEP_SIZE 2 -#endif +typedef unsigned long long LL; // #define MEASURE_TIME // #define MEASURE_ACCESS @@ -107,8 +96,8 @@ constexpr size_t MAX_LEVEL = 16; // #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 NUM_ITEMS = BUILD_SIZE; +// constexpr size_t KEYS = 1048576; constexpr size_t FACTOR = 1; // should change this to dynamic next time @@ -117,11 +106,6 @@ constexpr size_t FACTOR = 1; // 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 @@ -271,24 +255,20 @@ public: class LockFreeSkipList { - key_type *sample; + key_type *sample = nullptr; size_t sampleLength; CustomSort customSort; public: - Node *head; - Node *tail; + Node *head = nullptr; + Node *tail = nullptr; LockFreeSkipList(key_type *_sample, size_t sample_length) : sampleLength(sample_length), customSort(sample_length, sizeof(key_type) * 8) { Node *h = new Node(0); // size_ = 0; -#if __WORDSIZE == 64 Node *t = new Node(std::numeric_limits::max() - 1); -#else - Node *t = new Node((LL)0xffffffff); -#endif cudaMalloc(&head, sizeof(Node)); cudaMalloc(&tail, sizeof(Node)); @@ -337,7 +317,7 @@ public: atomicAdd(&this->access_times, count); } #else - __device__ void increaseAccessCount(unsigned _count = 1) {} + __device__ void increaseAccessCount(unsigned = 1) {} #endif #ifdef MEASURE_TIME @@ -353,17 +333,18 @@ public: __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 +//__device__ LL *randoms; // Array storing the levels of the nodes in the free +// pool +__device__ unsigned int NODE_LIMIT; // Function for creating a new node when requested by an add operation __device__ Node *GetNewNode(LL key, size_t topLevel) { - LL ind = atomicInc(&pointerIndex, NUM_ITEMS); + LL ind = atomicInc(&pointerIndex, NODE_LIMIT); Node *n = nodes[ind]; n->key = key; // n->topLevel = randoms[ind]; - n->topLevel = topLevel; + n->topLevel = (int)topLevel; int i; for (i = 0; i < n->topLevel + 1; i++) { n->SetRef(i, nullptr, false); @@ -379,10 +360,12 @@ __device__ LockFreeSkipList *lockFreeSkipList; // The lock-free skip list // Kernel for initializing device memory -__global__ void init(LockFreeSkipList *l1, Node **n, LL *rands) { - randoms = rands; +__global__ void init(LockFreeSkipList *l1, Node **n, + unsigned int insertion_limit) { + // randoms = rands; nodes = n; lockFreeSkipList = l1; + NODE_LIMIT = insertion_limit; } // Find the window holding key @@ -395,7 +378,7 @@ LockFreeSkipList::find(LL key, Node **preds, bool marked[] = {false}; bool snip; Node *pred; - Node *curr = nullptr; + Node *curr; Node *succ; bool beenThereDoneThat; while (true) { @@ -510,8 +493,8 @@ __device__ bool LockFreeSkipList::Delete(LL key) { } __device__ bool LockFreeSkipList::Add(LL key) { - Node *newNode = - GetNewNode(key, this->trailing_zeroes(this->searchIndex(key))); + Node *newNode = GetNewNode( + key, LockFreeSkipList::trailing_zeroes(this->searchIndex(key))); int topLevel = newNode->topLevel; int bottomLevel = 0; Node *preds[MAX_LEVEL + 1]; @@ -682,19 +665,17 @@ int main(int argc, char **argv) { // 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); + LL *result = new LL[search_length]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS); // Allocate device memory LL *cudaOperatorItems; // LL *Cop; - LL *Cresult; - LL *Clevels; + LL *cudaResult; - cudaMalloc(&Cresult, sizeof(LL) * NUM_ITEMS); + cudaMalloc(&cudaResult, sizeof(LL) * search_length); cudaMalloc(&cudaOperatorItems, sizeof(LL) * insertion_length); // cudaMalloc(&Cop, sizeof(LL) * NUM_ITEMS); - cudaMalloc(&Clevels, sizeof(LL) * NUM_ITEMS); // cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS, // cudaMemcpyHostToDevice); cudaMemcpy(cudaOperatorItems, _population.data(), @@ -730,7 +711,7 @@ int main(int argc, char **argv) { CudaCheckError(); // Initialize the device memory - init<<<1, 32>>>(Clist, Cpointers, Clevels); + init<<<1, 32>>>(Clist, Cpointers, insertion_length); cudaDeviceSynchronize(); // Insertion to skiplist @@ -751,7 +732,7 @@ int main(int argc, char **argv) { cudaEventRecord(start, nullptr); blocks = calcBlocks(search_length); - kernel<<>>(cudaOperatorItems, search_length, Cresult); + kernel<<>>(cudaOperatorItems, search_length, cudaResult); CudaCheckError(); cudaDeviceSynchronize(); cudaEventRecord(stop, nullptr); @@ -768,7 +749,7 @@ int main(int argc, char **argv) { // Move results back to host memory - cudaMemcpy(result, Cresult, sizeof(LL) * search_length, + cudaMemcpy(result, cudaResult, sizeof(LL) * search_length, cudaMemcpyDeviceToHost); // Uncomment the following for debugging diff --git a/read_helper.h b/read_helper.h index f12f4d9..eb7d2e3 100644 --- a/read_helper.h +++ b/read_helper.h @@ -2,12 +2,11 @@ #ifndef LOCKFREE_READ_HELPER_H #define LOCKFREE_READ_HELPER_H -#include #include #include #include -#ifndef P_EFF -#ifndef NDEBUG +#ifndef P_ERR +#if !(defined(NDEBUG) || defined(READ_NO_OUTPUT)) #define P_ERR(...) fprintf(stderr, __VA_ARGS__) #else #define P_ERR(...) @@ -62,7 +61,11 @@ public: this->population_vector.clear(); auto read_number = 0UL; FILE *file = fopen(filename, "r"); - assert(file); + if (file == nullptr) { + fprintf(stderr, "Unable to open file %s\n", filename); + exit(1); + } + P_ERR("Reading sample"); for (key_type i; read_number < sample_length && fscanf(file, "%lld ", &i) != EOF; -- cgit v1.3.1