// Experimental content: Test log performance
#include "sortlib.cuh"
#include <algorithm>
#include <cstdio>

#ifndef LOCKFREE_SORTLIB_CUH
#include <cassert>
class CustomSort {
public:
  __device__ __host__ static size_t fast_log(size_t a) {
#ifdef __CUDA_ARCH__
    return (size_t)log2((double)a);
#else
    return (size_t)std::log2(a);
#endif
  }
};

#endif

class OriginalFastLog {
public:
  __device__ __host__ static size_t fast_log(size_t a) {
    float t = a;
    return (((*(int *)&t) >> 23) + 1) & 127;
  }
};

constexpr int TEST_SIZE = 262144;

__global__ void init(int *result1, int *result2) {
  memset(result1, 0, sizeof(int) * 32);
  memset(result2, 0, sizeof(int) * 32);
}

__global__ void testNew(int *result) {
  for (int i = 2; i < TEST_SIZE; i++) {
    result[CustomSort::fast_log(i)]++;
  }
}

__global__ void testOld(int *result) {
  for (int i = 2; i < TEST_SIZE; i++) {
    result[OriginalFastLog::fast_log(i)]++;
  }
}

__global__ void testResult(const int *result1, const int *result2) {
  for (int i = 0; i < 32; i++) {
    assert(result1[i] == result2[i]);
  }
}

void run_kernel(int *dst, bool custom = false) {

  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start, nullptr);
  if (custom) {
    testNew<<<1, 1>>>(dst);
  } else {
    testOld<<<1, 1>>>(dst);
  }

  cudaDeviceSynchronize();
  cudaEventRecord(stop, nullptr);
  cudaEventSynchronize(stop);
  float time;
  cudaEventElapsedTime(&time, start, stop);
  cudaEventDestroy(start);
  cudaEventDestroy(stop);
  printf("%stime: %lf ", custom ? "custom " : "", time);
  cudaDeviceSynchronize();
}

int main() {
  int *result1, *result2;
  cudaMalloc(&result1, sizeof(int) * 32);
  cudaMalloc(&result2, sizeof(int) * 32);

  init<<<1, 1>>>(result1, result2);
  cudaDeviceSynchronize();

  run_kernel(result1);
  run_kernel(result2);

  cudaFree(result1);
  cudaFree(result2);
  puts("");
}