#include <cstdlib>
#include <cstdio>
#include <iostream>
#include <iomanip>
#include <cuda.h>
#include <gauge_field.h>
#include "quda_matrix.h"
#include "svd_quda.h"

Classes
class	hisq::fermion_force::DerivativeCoefficients< Real >
Namespaces
namespace	hisq
namespace	hisq::fermion_force
Defines
#define	HISQ_UNITARIZE_PI 3.14159265358979323846
#define	HISQ_UNITARIZE_PI23 HISQ_UNITARIZE_PI*2.0/3.0
#define	REUNIT_SVD_ONLY HOST_REUNIT_SVD_ONLY
#define	HISQ_UNITARIZE_EPS HOST_HISQ_UNITARIZE_EPS
#define	REUNIT_ALLOW_SVD HOST_REUNIT_ALLOW_SVD
#define	REUNIT_SVD_REL_ERROR HOST_REUNIT_SVD_REL_ERROR
#define	REUNIT_SVD_ABS_ERROR HOST_REUNIT_SVD_ABS_ERROR
#define	MAX_DET_ERROR HOST_MAX_DET_ERROR
#define	HISQ_FORCE_FILTER HOST_HISQ_FORCE_FILTER
Functions
void	hisq::fermion_force::setUnitarizeForceConstants (double unitarize_eps, double hisq_force_filter, double max_det_error, bool allow_svd, bool svd_only, double svd_rel_error, double svd_abs_error)
template<class Cmplx >
__device__ __host__ void	hisq::fermion_force::accumBothDerivatives (Matrix< Cmplx, 3 > *result, const Matrix< Cmplx, 3 > &left, const Matrix< Cmplx, 3 > &right, const Matrix< Cmplx, 3 > &outer_prod)
template<class Cmplx >
__device__ __host__ void	hisq::fermion_force::accumDerivatives (Matrix< Cmplx, 3 > *result, const Matrix< Cmplx, 3 > &left, const Matrix< Cmplx, 3 > &right, const Matrix< Cmplx, 3 > &outer_prod)
template<class T >
__device__ __host__ T	hisq::fermion_force::getAbsMin (const T *const array, int size)
template<class Real >
__device__ __host__ bool	hisq::fermion_force::checkAbsoluteError (Real a, Real b, Real epsilon)
template<class Real >
__device__ __host__ bool	hisq::fermion_force::checkRelativeError (Real a, Real b, Real epsilon)
template<class Cmplx >
__device__ __host__ void	hisq::fermion_force::reciprocalRoot (Matrix< Cmplx, 3 > res, DerivativeCoefficients< typename RealTypeId< Cmplx >::Type > deriv_coeffs, typename RealTypeId< Cmplx >::Type f[3], Matrix< Cmplx, 3 > &q, int *unitarization_failed)
template<class Cmplx >
__device__ __host__ void	hisq::fermion_force::getUnitarizeForceSite (const Matrix< Cmplx, 3 > &v, const Matrix< Cmplx, 3 > &outer_prod, Matrix< Cmplx, 3 > result, int unitarization_failed)
template<class Cmplx >
__global__ void	hisq::fermion_force::getUnitarizeForceField (Cmplx link_even, Cmplx link_odd, Cmplx old_force_even, Cmplx old_force_odd, Cmplx force_even, Cmplx force_odd, int *unitarization_failed)
void	hisq::fermion_force::unitarizeForceCPU (const QudaGaugeParam &param, cpuGaugeField &cpuOldForce, cpuGaugeField &cpuGauge, cpuGaugeField *cpuNewForce)
void	hisq::fermion_force::unitarizeForceCuda (const QudaGaugeParam &param, cudaGaugeField &cudaOldForce, cudaGaugeField &cudaGauge, cudaGaugeField cudaNewForce, int unitarization_failed)
Variables
__constant__ double	DEV_HISQ_UNITARIZE_EPS
__constant__ double	DEV_HISQ_FORCE_FILTER
__constant__ double	DEV_MAX_DET_ERROR
__constant__ bool	DEV_REUNIT_ALLOW_SVD
__constant__ bool	DEV_REUNIT_SVD_ONLY
__constant__ double	DEV_REUNIT_SVD_REL_ERROR
__constant__ double	DEV_REUNIT_SVD_ABS_ERROR