// Experimental content: Test branch performance
#include "sortlib.cuh"

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

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 = 10'000;

unsigned long long max_value = 0, min_value = 0xfffffffff;
typedef std::pair<int, int> dim_pair_type;

inline void store_into_vector(unsigned long long value) {
  if (max_value < value) {
    max_value = value;
  }
  if (min_value > value) {
    min_value = value;
  }
  population_vector.push_back(value);
}
// #define TEST_BOUNDS

__device__ long cuda_sample_length;
__device__ key_type *cudaSampleItem, *cudaNormalSampleItem, *cudaPopulationItem;
__device__ bool *cdf_result, *cdf_normal_result;
#ifdef TEST_BOUNDS
__device__ unsigned insert_value;
__device__ unsigned int index_max, index_min;
#endif

__global__ void initSample(key_type *sample, key_type *normal_sample,
                           key_type *population) {
  cudaSampleItem = sample;
  cudaPopulationItem = population;
  cudaNormalSampleItem = normal_sample;
  // cudaMalloc(&cdf_result, sizeof(bool) * TEST_LENGTH);
  cdf_result = nullptr;
  cdf_normal_result = nullptr;
#ifdef TEST_BOUNDS
  index_max = 0;
  index_min = 0x7fffffff;
#endif
}

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

    // printf("%d\n", tid);
    auto index = (int)(custom_sort.sample_cdf_custom_version(
                           cudaSampleItem, cudaSampleItem + cuda_sample_length,
                           cudaPopulationItem[tid]) /
                       slice_size) -
                 1;
    cdf_result[index] = true;
  }
}

__global__ void kernel2(unsigned long step, const double slice_size) {

  for (int i = 0; i < step; i++) {
    auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i;
    auto index =
        (int)(FactorySort::sample_cdf(cudaNormalSampleItem, cuda_sample_length,
                                      cudaPopulationItem[tid]) /
              slice_size) -
        1;
    cdf_normal_result[index] = true;
  }
}

__global__ void print2() {
  /*for (int i = 0; i< SAMPLE_LENGTH; i++) {
    printf("%llu ", cudaSampleItem[i]);
  }
  printf("\n");*/
  printf("%lu\n", CustomSort::fast_log(cuda_sample_length));
}

__global__ void print_function() {
#ifdef TEST_BOUNDS
  printf("%u\n", insert_value);
  printf("%u %u\n", index_min, index_max);
#endif
}

__global__ void check_items() {
  for (int i = 0; i < cuda_sample_length; i++) {
    printf("%lld ", cudaSampleItem[i]);
  }
  printf("\n");
}

__global__ void initStorage(unsigned scale, unsigned test_size) {
  cudaFree(cdf_result);
  // printf("test size: %u\n", test_size);
  cudaMalloc(&cdf_result, (scale * test_size + RESERVED_BLOCK) * sizeof(bool));
  memset(cdf_result, 0, (scale * test_size + RESERVED_BLOCK) * sizeof(bool));
  cudaFree(cdf_normal_result);
  cudaMalloc(&cdf_normal_result,
             (scale * test_size + RESERVED_BLOCK) * sizeof(bool));
  memset(cdf_normal_result, 0,
         (scale * test_size + RESERVED_BLOCK) * sizeof(bool));
}

__global__ void freeStorage() {
  cudaFree(cudaSampleItem);
  cudaFree(cudaPopulationItem);
  cudaFree(cdf_result);
}

__global__ void initCustomSample() {
  // printf("init\n");
  key_type *tmp = nullptr;
  auto custom_sort = CustomSort(cuda_sample_length, 0);
  cudaMalloc(&tmp, sizeof(key_type) * cuda_sample_length);
  for (size_t i = 0; i < cuda_sample_length; i++) {
    tmp[i] = cudaSampleItem[custom_sort.calculate_index(i) - 1];
  }
  // printf("copy\n");
  memcpy(cudaSampleItem, tmp, sizeof(key_type) * cuda_sample_length);
  cudaFree(tmp);
  tmp = nullptr;
  // memset(cdf_result, 0, sizeof(key_type) * TEST_LENGTH);
  cudaFree(cdf_result);
  cdf_result = nullptr;
  // printf("finalize\n");
}

void run_kernel(size_t test_size, bool normal = true) {
  dim3 grid_dim = 32, block_dim = 32;
  unsigned long step = test_size / (grid_dim.x * block_dim.x);
  auto scale = 2;
  const unsigned long split_size = test_size * scale;
  const auto slice_size = 1.0 / (double)split_size;
  initStorage<<<1, 1>>>(scale, test_size);
  cudaDeviceSynchronize();
  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start, nullptr);
  if (normal) {
    kernel<<<grid_dim, block_dim>>>(step, slice_size, split_size);
  } else {
    kernel2<<<grid_dim, block_dim>>>(step, slice_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("time: %lf\n", time);
  cudaDeviceSynchronize();

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

void read_file(char const *filename, long sample_length,
               unsigned long max_number = TEST_LENGTH) {
  max_number += sample_length + 256;
  auto read_number = 0;
  FILE *file = fopen(filename, "r");
  assert(file);
  for (long long i;
       read_number < max_number && fscanf(file, "%lld ", &i) != EOF;
       store_into_vector(i))
    read_number++;
  fclose(file);
}

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

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

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

  auto test_size = TEST_LENGTH;
  auto sample_length = DEFAULT_SAMPLE_LENGTH;

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

  read_file("normal_distribution.txt", test_size);

  printf("population: %zu, test size: %zu, sample length: %zu\n",
         population_vector.size(), test_size, sample_length);

  sample_vector = std::vector<key_type>(
      population_vector.begin(), population_vector.begin() + sample_length - 2);
  sample_vector.push_back(min_value);
  sample_vector.push_back(max_value);
  std::sort(sample_vector.begin(), sample_vector.end());
  le_sample_vector = sample_vector;

  key_type *cudaSample = nullptr, *cudaPopulation;
  cudaMalloc(&cudaSample, sizeof(key_type) * sample_length);
  assert(sample_length == sample_vector.size());
  cudaMemcpy(cudaSample, sample_vector.data(), sizeof(key_type) * sample_length,
             cudaMemcpyHostToDevice);
  cudaMalloc(&cudaPopulation, sizeof(key_type) * test_size);
  cudaMemcpy(cudaPopulation, population_vector.data() + sample_length,
             sizeof(key_type) * test_size, cudaMemcpyHostToDevice);

  key_type *cudaNormalSample = nullptr;
  cudaMalloc(&cudaNormalSample, sizeof(key_type) * sample_length);
  cudaMemcpy(cudaNormalSample, le_sample_vector.data(),
             sizeof(key_type) * sample_length, cudaMemcpyHostToDevice);

  // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8);
  // custom_sort.testCalculation();
  applyCudaSampleLength<<<1, 1>>>(sample_length);
  testCustomCalculation<<<1, 1>>>(sample_length);
  cudaDeviceSynchronize();

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

  initCustomSample<<<1, 1>>>();
  cudaDeviceSynchronize();
  cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) * sample_length,
             cudaMemcpyDeviceToHost);

  // check_items<<<1, 1>>>();
  // cudaDeviceSynchronize();
  run_kernel(test_size);
  run_kernel(test_size, false);
  // check_items<<<1, 1>>>();
  // cudaDeviceSynchronize();

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