#ifndef CUDAUTILS_H
#define CUDAUTILS_H

#include <cstdio>
#include <cuda_runtime.h>
#include <iostream>

#ifdef WIN32
#include <intrin.h>
#endif

#define safeCall(err) __safeCall(err, __FILE__, __LINE__)
#define checkMsg(msg) __checkMsg(msg, __FILE__, __LINE__)

namespace cudasift {

inline void __safeCall(cudaError err, const char *file, const int line) {
  if (cudaSuccess != err) {
    fprintf(stderr,
            "safeCall() Runtime API error in file <%s>, line %i : %s.\n", file,
            line, cudaGetErrorString(err));
    exit(-1);
  }
}

inline void __checkMsg(const char *errorMessage, const char *file,
                       const int line) {
  cudaError_t err = cudaGetLastError();
  if (cudaSuccess != err) {
    fprintf(stderr, "checkMsg() CUDA error: %s in file <%s>, line %i : %s.\n",
            errorMessage, file, line, cudaGetErrorString(err));
    exit(-1);
  }
}

inline bool deviceInit(int dev) {
  int deviceCount;
  safeCall(cudaGetDeviceCount(&deviceCount));
  if (deviceCount == 0) {
    fprintf(stderr, "CUDA error: no devices supporting CUDA.\n");
    return false;
  }
  if (dev < 0)
    dev = 0;
  if (dev > deviceCount - 1)
    dev = deviceCount - 1;
  cudaDeviceProp deviceProp;
  safeCall(cudaGetDeviceProperties(&deviceProp, dev));
  if (deviceProp.major < 1) {
    fprintf(stderr, "error: device does not support CUDA.\n");
    return false;
  }
  safeCall(cudaSetDevice(dev));
  return true;
}

class TimerGPU {
public:
  cudaEvent_t start, stop;
  cudaStream_t stream;
  TimerGPU(cudaStream_t stream_) : stream(stream_) {
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, stream);
  }
  ~TimerGPU() {
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
  }
  float read() {
    cudaEventRecord(stop, stream);
    cudaEventSynchronize(stop);
    float time;
    cudaEventElapsedTime(&time, start, stop);
    return time;
  }
};

class TimerCPU {
  static const int bits = 10;

public:
  long long beg_clock;
  float freq;
  TimerCPU(float freq_) : freq(freq_) { // freq = clock frequency in MHz
    beg_clock = getTSC(bits);
  }
  long long getTSC(int bits) {
#ifdef WIN32
    return __rdtsc() / (1LL << bits);
#else
    unsigned int low, high;
    __asm__(".byte 0x0f, 0x31" : "=a"(low), "=d"(high));
    return ((long long)high << (32 - bits)) | ((long long)low >> bits);
#endif
  }
  float read() {
    long long end_clock = getTSC(bits);
    long long Kcycles = end_clock - beg_clock;
    float time = (float)(1 << bits) * Kcycles / freq / 1e3f;
    return time;
  }
};

template <class T>
__device__ __inline__ T ShiftDown(T var, unsigned int delta, int width = 32) {
#if (CUDART_VERSION >= 9000)
  return __shfl_down_sync(0xffffffff, var, delta, width);
#else
  return __shfl_down(var, delta, width);
#endif
}

template <class T>
__device__ __inline__ T ShiftUp(T var, unsigned int delta, int width = 32) {
#if (CUDART_VERSION >= 9000)
  return __shfl_up_sync(0xffffffff, var, delta, width);
#else
  return __shfl_up(var, delta, width);
#endif
}

template <class T>
__device__ __inline__ T Shuffle(T var, unsigned int lane, int width = 32) {
#if (CUDART_VERSION >= 9000)
  return __shfl_sync(0xffffffff, var, lane, width);
#else
  return __shfl(var, lane, width);
#endif
}

} // namespace cudasift

#endif
