feat(exp): Remove useless arguments

Signed-off-by: KunoiSayami <[email protected]>

1 files changed, 27 insertions, 46 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 <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
+#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<key_type>::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<<<blocks, NUM_THREADS>>>(cudaOperatorItems, search_length, Cresult);
+  kernel<<<blocks, NUM_THREADS>>>(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