
#ifndef COMMON_H
#define COMMON_H

#include "FloatImage.h"
#include <iostream>

// macro to check CUDA errors
#define CUDA_SAFE_CALL(CALL) {        \
  int cuda_safe_call_retval = (CALL); \
  if(cuda_safe_call_retval!=0) {      \
  std::cerr << "error in " << __FILE__ << ":" << __LINE__<< " -> " << cuda_safe_call_retval << " when calling " << #CALL << "\n"; } }

#include <sys/time.h>
#include <unistd.h>
#include <stdlib.h>

inline double Timer_getTime(void)
{
    struct timeval tv;
    struct timezone tz;
    gettimeofday(&tv, &tz);
    return (double)tv.tv_sec + 1.e-6 * (double)tv.tv_usec;
}

typedef double Timer;

inline void startTimer(Timer& timer) { timer = Timer_getTime(); }
inline void stopTimer(Timer& timer)  { timer = Timer_getTime() - timer; }

#ifdef __CUDACC__

inline int iogs_max_thread_per_block()
{
  cudaDeviceProp devProp;
  cudaGetDeviceProperties(&devProp, 0);
  return devProp.maxThreadsPerBlock;
}

/** Macro to run the CUDA kernel 'FUNC' using 'N' threads.
 *  In the kernel, the current thread ID can then get obtained using the iogs_thread_id_1D() function.
 */
#define iogs_launch_kernel_1D(FUNC,N) int iogs_mtpb_##__LINE__ = iogs_max_thread_per_block(); \
    FUNC<<< (N+(iogs_mtpb_##__LINE__)-1)/(iogs_mtpb_##__LINE__), iogs_mtpb_##__LINE__ >>>

/** \returns the current thread ID. */
__device__ inline int iogs_thread_id_1D() { return threadIdx.x + blockDim.x * blockIdx.x; }


/** Macro to run the CUDA kernel 'FUNC' using a lattice of 'NX' x 'NY' threads.
 *  The total number of threads will be NX*NY.
 *  In the kernel, the position (x,y) of the thread can then get obtained using the iogs_thread_id_2D() function.
 */
#define iogs_launch_kernel_2D(FUNC,NX,NY) int iogs_mtpb_##__LINE__ = iogs_max_thread_per_block()/32; \
    FUNC<<< dim3((NX+(iogs_mtpb_##__LINE__)-1)/(iogs_mtpb_##__LINE__),(NY+31)/32, 1) , dim3(iogs_mtpb_##__LINE__,32,1) >>>

/** \returns the current thread ID when threads are organized as a 2D lattice. */
__device__ inline int2 iogs_thread_id_2D()
{
  int2 ret;
  ret.x = threadIdx.x + blockDim.x * blockIdx.x;
  ret.y = threadIdx.y + blockDim.y * blockIdx.y;
  return ret;
}

#endif

#endif