1 files changed, 56 insertions, 290 deletions

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;
 }