/*

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 "read_helper.h"
#include "skiplistcustom.cuh"
#include <cassert>
#include <cstdio>
#include <cstdlib>
#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;

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__ key_type SampleStorage[SAMPLE_SIZE];

// Kernel for initializing device memory

__global__ void init(Node **n) { nodes = n; }

// The main kernel

__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

  for (int i = 0; i < FACTOR;
       i++) { // FACTOR is the number of operations per thread
    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
    key_type item = population[tid];
    skipList->Add(item);
  }
}

__global__ void print_function() {
#ifdef MEASURE_ACCESS
  printf("count: %u\n", l->getAccessCount());
#endif
}

int main(int argc, char **argv) {
  if (argc != 3) {
    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);
  }

  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);
  printf("%d\n", __LINE__);
  reader.readFile(total_row);

  printf("%d\n", __LINE__);
  key_type *cudaPopulation;

  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

  // cudaMalloc((void **)&Clevels, sizeof(LL) * NUM_ITEMS);
  // cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS,
  // cudaMemcpyHostToDevice);
  Node **pointers =
      (Node **)new LL[insertion_length]; // malloc(sizeof(LL) * adds);
  Node **cudaNodePointers;

  printf("%d\n", __LINE__);

  // Allocate the pool of free nodes

  for (int i = 0; i < insertion_length; i++) {
    cudaMalloc(&pointers[i], sizeof(Node));
  }
  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(_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

  size_t blocks = (insertion_length % (NUM_THREADS * FACTOR) == 0)
                      ? insertion_length / (NUM_THREADS * FACTOR)
                      : (insertion_length / (NUM_THREADS * FACTOR)) + 1;

  // Error checking code
  if (cudaSuccess != error) {
    printf("error0:CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
    exit(-1);
  }

  // Initialize the device memory
  init<<<1, 32>>>(cudaNodePointers);
  cudaDeviceSynchronize();

  // Launch main kernel

  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start, nullptr);

  kernel<<<blocks, NUM_THREADS>>>(Clist, cudaPopulation, insertion_length);
  CudaCheckError();
  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);

#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);
  for (int i = 0; i < insertion_length; i++) {
    cudaFree(pointers[i]);
  }
  cudaFree(cudaNodePointers);
  delete list;
  delete[] pointers;
  return 0;
}