// Experimental content: Test branch performance (optimized memory and output)
// #define PRINT_READ_PROCESS
#include "read_helper.h"
#include "sortlib.cuh"

#include <algorithm>
#include <cassert>
#include <cstdio>
#include <vector>

constexpr auto NUM_THREADS = 512;
constexpr auto FACTOR = 1;

std::vector<unsigned long long> population_vector, sample_vector,
    le_sample_vector;

constexpr long DEFAULT_SAMPLE_LENGTH = 65535;
constexpr size_t TEST_LENGTH = 32'768;
// constexpr size_t TEST_LENGTH = 4096;
// constexpr size_t RESERVED_BLOCK = 100;

typedef std::pair<int, int> dim_pair_type;

__device__ unsigned long cuda_sample_length;
__device__ key_type *cudaSampleItem, *cudaNormalSampleItem, *cudaPopulationItem;

__global__ void initSample(key_type *sample, key_type *population) {
  cudaSampleItem = sample;
  cudaPopulationItem = population;
  // cudaMalloc(&cdf_result, sizeof(bool) * TEST_LENGTH);
  // cdf_result = nullptr;
  // cdf_normal_result = nullptr;
}

__global__ void initNormalSample(key_type *normal_sample) {
  cudaNormalSampleItem = normal_sample;
}

/*
#define CudaCheckError() cudaCheckError(__FILE__, __LINE__)
inline void cudaCheckError(const char *file, const int line) {
  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);
  }
}
*/

__global__ void kernel(const size_t test_length) {
  auto custom_sort = CustomSort(cuda_sample_length, sizeof(key_type) * 8);
  // printf("kernel1 step: %ld\n", step);
  for (int i = 0; i < FACTOR; i++) {
    auto tid =
        i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;

    if (tid >= test_length)
      return;
    // printf("%d\n", tid);
    custom_sort.sample_cdf_custom_version(cudaSampleItem,
                                          cudaPopulationItem[tid]);
    // cdf_result[index] = true;
  }
}

__global__ void kernel2(const size_t test_length) {

  for (int i = 0; i < FACTOR; i++) {
    auto tid =
        i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;

    if (tid >= test_length)
      return;

    FactorySort::sample_cdf(cudaNormalSampleItem, cuda_sample_length,
                            cudaPopulationItem[tid]);
    // cdf_result[index] = true;
  }
}

inline size_t calcBlocks(size_t input) {
  return (input % (NUM_THREADS * FACTOR) == 0)
             ? input / (NUM_THREADS * FACTOR)
             : (input / (NUM_THREADS * FACTOR)) + 1;
}

void run_kernel(size_t test_size, bool custom = true) {
  // need description

  // Only scale mode will use
  auto blocks = calcBlocks(test_size);

  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start, nullptr);
  if (custom) {
    kernel<<<blocks, NUM_THREADS>>>(test_size);
  } else {
    kernel2<<<blocks, NUM_THREADS>>>(test_size);
  }

  cudaDeviceSynchronize();
  cudaEventRecord(stop, nullptr);
  cudaEventSynchronize(stop);
  float time;
  cudaEventElapsedTime(&time, start, stop);
  cudaEventDestroy(start);
  cudaEventDestroy(stop);
  if (time == 0) {
    // printf("last error: %u\n", cudaGetLastError());
  }
  printf("%stime: %lf ", custom ? "custom " : "", time);
  cudaDeviceSynchronize();
}

__global__ void freeStorageStage1() { cudaFree(cudaSampleItem); }

__global__ void freeStorageStage2() {
  cudaFree(cudaNormalSampleItem);
  cudaFree(cudaPopulationItem);
  // cudaFree(cdf_result);
}

long pow_for_sample(long n) {
  auto x = 2;
  for (int i = 1; i < n; i++) {
    x *= 2;
  }
  return x - 1;
}

__global__ void applyCudaSampleLength(unsigned long length) {
  cuda_sample_length = length;
}

int main(int argc, char const *argv[]) {

  auto test_size = TEST_LENGTH;
  auto sample_length = DEFAULT_SAMPLE_LENGTH;
  auto total_row = 0UL;

  if (argc >= 2) {
    test_size = strtol(argv[1], nullptr, 10);
  }
  if (argc >= 3) {
    sample_length = pow_for_sample(strtol(argv[2], nullptr, 10));
  }
  if (argc >= 4) {
    total_row = strtol(argv[3], nullptr, 10);
  }

  // printf("%ld %ld %ld\n", test_size, sample_length, total_row);

  ReadHelper readHelper("normal_distribution.txt", sample_length, test_size,
                        total_row);

  printf("population: %zu, skip: %lu, sample length: %zu ", test_size,
         readHelper.random_number, sample_length);
  readHelper.readFile();

  readHelper.split_into(sample_vector, population_vector);
  assert(sample_vector.size() == sample_length);
  assert(population_vector.size() == test_size);

  std::sort(sample_vector.begin(), sample_vector.end());
  le_sample_vector = sample_vector;
  rebuild(sample_vector);

  /*printf("population: %zu sample length: %zu\n", population_vector.size(),
         sample_vector.size());*/

  key_type *cudaSample = nullptr, *cudaPopulation;

  P_ERR("Coping sample");
  cudaMalloc(&cudaSample, sizeof(key_type) * sample_vector.size());
  assert(sample_length == sample_vector.size());
  cudaMemcpy(cudaSample, sample_vector.data(),
             sizeof(key_type) * sample_vector.size(), cudaMemcpyHostToDevice);
  cudaMalloc(&cudaPopulation, sizeof(key_type) * population_vector.size());

  assert(population_vector.size() == test_size);

  P_ERR("\rCoping population");
  cudaMemcpy(cudaPopulation, population_vector.data(),
             sizeof(key_type) * population_vector.size(),
             cudaMemcpyHostToDevice);

  // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8);
  // custom_sort.testCalculation();
  P_ERR("\rApply custom length and test calculation");
  applyCudaSampleLength<<<1, 1>>>(sample_vector.size());
  // testCustomCalculation<<<1, 1>>>(sample_length);
  cudaDeviceSynchronize();
  P_ERR("\rApply custom length and test calculation completed");

  // Can remove this function if pass
  // CustomSort::testSelf(sample_length);

  initSample<<<1, 1>>>(cudaSample, cudaPopulation);
  cudaDeviceSynchronize();

  // check_items<<<1, 1>>>();
  // cudaDeviceSynchronize();
  P_ERR("\r                                                   \rRunning "
        "kernel\n");
  run_kernel(test_size);

  cudaFree(cudaSample);

  key_type *cudaNormalSample = nullptr;

  P_ERR("\nCoping new sample");
  cudaMalloc(&cudaNormalSample, sizeof(key_type) * le_sample_vector.size());
  cudaMemcpy(cudaNormalSample, le_sample_vector.data(),
             sizeof(key_type) * le_sample_vector.size(),
             cudaMemcpyHostToDevice);

  /*cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) *
     sample_length, cudaMemcpyDeviceToHost);*/

  initNormalSample<<<1, 1>>>(cudaNormalSample);
  cudaDeviceSynchronize();

  P_ERR("\r                          ");
  P_ERR("\rRunning kernel2\n");
  run_kernel(test_size, false);
  // check_items<<<1, 1>>>();
  // cudaDeviceSynchronize();
  puts("");

  freeStorageStage2<<<1, 1>>>();
  cudaDeviceSynchronize();
}