feat: Add skiplist header

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

1 files changed, 567 insertions, 0 deletions

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