13 #define MAX(a,b) ((a)>(b)?(a):(b))
15 #ifdef BUILD_MILC_INTERFACE
17 static bool initialized =
false;
18 static int gridDim[4];
19 static int localDim[4];
20 static int clover_alloc = 0;
21 static bool invalidate_quda_gauge =
true;
22 static bool create_quda_gauge =
false;
26 using namespace quda::fermion_force;
30 #define QUDAMILC_VERBOSE 1
32 void inline qudamilc_called(
const char* func,
QudaVerbosity verb){
33 #ifdef QUDAMILC_VERBOSE
36 printf(
"QUDA_MILC_INTERFACE: %s (called) \n",func);
38 printf(
"QUDA_MILC_INTERFACE: %s (return) \n",func);
44 void inline qudamilc_called(
const char * func){
51 qudamilc_called<true>(__func__);
52 if(initialized)
return;
56 qudamilc_called<false>(__func__);
62 qudamilc_called<true>(__func__);
64 qudamilc_called<false>(__func__);
71 static int rankFromCoords(
const int *coords,
void *fdata)
73 int *
dims =
static_cast<int *
>(fdata);
76 for (
int i = 2; i >= 0; i--) {
77 rank = dims[i] * rank + coords[i];
86 for(
int dir=0; dir<4; ++dir){ local_dim[dir] = input.
latsize[dir]; }
88 for(
int dir=0; dir<4; ++dir){ local_dim[dir] /= input.
machsize[dir]; }
90 for(
int dir=0; dir<4; ++dir){
91 if(local_dim[dir]%2 != 0){
92 printf(
"Error: Odd lattice dimensions are not supported\n");
97 for(
int dir=0; dir<4; ++dir) localDim[dir] = local_dim[dir];
104 for(
int dir=0; dir<4; ++dir) gridDim[dir] = input.
machsize[dir];
108 for(
int dir=0; dir<4; ++dir) gridDim[dir] = 1;
109 static int device = input.
device;
119 static bool initialized =
false;
121 if(initialized)
return;
122 qudamilc_called<true>(__func__);
128 const double unitarize_eps = 1e-14;
129 const double max_error = 1e-10;
131 #ifdef GPU_HISQ_FORCE
148 #endif // UNITARIZE_GPU
151 qudamilc_called<false>(__func__);
160 for(
int dir=0; dir<4; ++dir) gParam.
X[dir] = dim[dir];
162 gParam.
type = link_type;
180 static void invalidateGaugeQuda() {
182 invalidate_quda_gauge =
true;
190 const double act_path_coeff[6],
void* inlink,
void*
fatlink,
void*
longlink)
192 qudamilc_called<true>(__func__);
205 computeKSLinkQuda(fatlink, longlink, NULL, inlink, const_cast<double*>(act_path_coeff), ¶m, method);
206 qudamilc_called<false>(__func__);
214 const double act_path_coeff[6],
void* inlink,
void* fatlink,
void* ulink)
216 qudamilc_called<true>(__func__);
227 computeKSLinkQuda(fatlink, NULL, ulink, inlink, const_cast<double*>(act_path_coeff), ¶m, method);
228 qudamilc_called<false>(__func__);
235 #ifdef GPU_HISQ_FORCE
237 void qudaHisqForce(
int prec,
const double level2_coeff[6],
const double fat7_coeff[6],
238 const void*
const staple_src[4],
const void*
const one_link_src[4],
const void*
const naik_src[4],
239 const void*
const w_link,
const void*
const v_link,
const void*
const u_link,
240 void*
const milc_momentum)
242 qudamilc_called<true>(__func__);
250 staple_src, one_link_src, naik_src,
251 w_link, v_link, u_link, &gParam);
252 qudamilc_called<false>(__func__);
258 const void*
const one_link_src[4],
const void*
const naik_src[4],
259 const void*
const link,
void*
const milc_momentum)
261 qudamilc_called<true>(__func__);
271 qudamilc_called<false>(__func__);
278 void** quark_field,
void* oprod[2])
280 qudamilc_called<true>(__func__);
286 qudamilc_called<false>(__func__);
290 #endif // GPU_HISQ_FORCE
292 #ifdef GPU_GAUGE_FORCE
293 void qudaUpdateU(
int prec,
double eps,
void* momentum,
void* link)
295 qudamilc_called<true>(__func__);
302 qudamilc_called<false>(__func__);
308 static int getVolume(
const int dim[4])
310 return dim[0]*dim[1]*dim[2]*dim[3];
315 static void setLoopCoeffs(
const double milc_loop_coeff[3],
319 for(
int i=0; i<6; ++i) loop_coeff[i] = milc_loop_coeff[0];
322 for(
int i=0; i<18; ++i) loop_coeff[i+6] = milc_loop_coeff[1];
324 for(
int i=0; i<24; ++i) loop_coeff[i+24] = milc_loop_coeff[2];
331 static inline int opp(
int dir){
336 static void createGaugeForcePaths(
int **paths,
int dir){
340 for(
int i=0; i<4; ++i){
342 paths[
index][0] = i; paths[
index][1] = opp(dir); paths[index++][2] = opp(i);
343 paths[
index][0] = opp(i); paths[
index][1] = opp(dir); paths[index++][2] = i;
347 for(
int i=0; i<4; ++i){
365 for(
int i=0; i<4; ++i){
366 for(
int j=0; j<4; ++j){
367 if(i==dir || j==dir || i==j)
continue;
369 paths[
index][0] = i; paths[
index][1] = j; paths[
index][2] = opp(dir); paths[
index][3] = opp(i), paths[
index][4] = opp(j);
371 paths[
index][0] = i; paths[
index][1] = opp(j); paths[
index][2] = opp(dir); paths[
index][3] = opp(i), paths[
index][4] = j;
373 paths[
index][0] = opp(i); paths[
index][1] = j; paths[
index][2] = opp(dir); paths[
index][3] = i, paths[
index][4] = opp(j);
375 paths[
index][0] = opp(i); paths[
index][1] = opp(j); paths[
index][2] = opp(dir); paths[
index][3] = i, paths[
index][4] = j;
386 double milc_loop_coeff[3],
392 qudamilc_called<true>(__func__);
393 const int numPaths = 48;
398 static double loop_coeff[numPaths];
399 static int length[numPaths];
401 if((milc_loop_coeff[0] != loop_coeff[0]) ||
402 (milc_loop_coeff[1] != loop_coeff[6]) ||
403 (milc_loop_coeff[2] != loop_coeff[24]) ){
405 setLoopCoeffs(milc_loop_coeff, loop_coeff);
406 for(
int i=0; i<6; ++i) length[i] = 3;
407 for(
int i=6; i<48; ++i) length[i] = 5;
411 int** input_path_buf[4];
412 for(
int dir=0; dir<4; ++dir){
413 input_path_buf[dir] =
new int*[numPaths];
414 for(
int i=0; i<numPaths; ++i){
415 input_path_buf[dir][i] =
new int[length[i]];
417 createGaugeForcePaths(input_path_buf[dir], dir);
423 memset(milc_momentum, 0, 4*getVolume(localDim)*10*qudaGaugeParam.
cpu_prec);
425 loop_coeff, numPaths, max_length, eb3,
426 &qudaGaugeParam, timeinfo);
428 for(
int dir=0; dir<4; ++dir){
429 for(
int i=0; i<numPaths; ++i){
430 delete[] input_path_buf[dir][i];
432 delete[] input_path_buf[dir];
434 qudamilc_called<false>(__func__);
438 #endif // GPU_GAUGE_FORCE
440 static int getFatLinkPadding(
const int dim[4])
442 int padding =
MAX(dim[1]*dim[2]*dim[3]/2, dim[0]*dim[2]*dim[3]/2);
443 padding =
MAX(padding, dim[0]*dim[1]*dim[3]/2);
444 padding =
MAX(padding, dim[0]*dim[1]*dim[2]/2);
449 #ifdef GPU_STAGGERED_DIRAC
451 static void setInvertParams(
const int dim[4],
457 double target_residual,
458 double target_residual_hq,
460 double reliable_delta,
469 invertParam->
tol = target_residual;
470 invertParam->
tol_hq =target_residual_hq;
474 invertParam->
maxiter = maxiter;
490 invertParam->
gflops = 0.0;
506 invertParam->
sp_pad = dim[0]*dim[1]*dim[2]/2;
524 static void setInvertParams(
const int dim[4],
530 const double offset[],
531 const double target_residual_offset[],
532 const double target_residual_hq_offset[],
534 double reliable_delta,
541 const double null_mass = -1;
542 const double null_residual = -1;
545 setInvertParams(dim, cpu_prec, cuda_prec, cuda_prec_sloppy, cuda_prec_precondition,
546 null_mass, null_residual, null_residual, maxiter, reliable_delta, parity, verbosity, inverter, invertParam);
549 for(
int i=0; i<num_offset; ++i){
550 invertParam->
offset[i] = offset[i];
551 invertParam->
tol_offset[i] = target_residual_offset[i];
553 invertParam->
tol_hq_offset[i] = target_residual_hq_offset[i];
560 static void setGaugeParams(
const int dim[4],
565 const double tadpole,
569 for(
int dir=0; dir<4; ++dir){
570 gaugeParam->
X[dir] = dim[dir];
584 gaugeParam->
ga_pad = getFatLinkPadding(dim);
597 static void setColorSpinorParams(
const int dim[4],
606 for(
int dir=0; dir<4; ++dir) param->
x[dir] = dim[dir];
620 static size_t getColorVectorOffset(
QudaParity local_parity,
bool even_odd_exchange,
const int dim[4])
623 int volume = dim[0]*dim[1]*dim[2]*dim[3];
626 offset = even_odd_exchange ? volume*6/2 : 0;
628 offset = even_odd_exchange ? 0 : volume*6/2;
637 double*
const offset,
639 const double target_residual[],
640 const double target_fermilab_residual[],
641 const void*
const fatlink,
642 const void*
const longlink,
643 const double tadpole,
645 void** solutionArray,
646 double*
const final_residual,
647 double*
const final_fermilab_residual,
659 qudamilc_called<true>(__func__, verbosity);
661 if(target_residual[0] == 0){
662 errorQuda(
"qudaMultishiftInvert: target residual cannot be zero\n");
669 const bool use_mixed_precision = (((quda_precision==2) && inv_args.
mixed_precision) ||
670 ((quda_precision==1) && (inv_args.
mixed_precision==2)) ) ?
true :
false;
677 device_precision_sloppy = device_precision;
680 QudaPrecision device_precision_precondition = device_precision_sloppy;
700 setGaugeParams(localDim, host_precision, device_precision, device_precision_sloppy, device_precision_precondition, tadpole, &gaugeParam);
708 printfQuda(
"Using QUDA_L2_RELATIVE_RESIDUAL only");
711 printfQuda(
"Using QUDA_L2_RELATIVE_RESIDUAL");
713 printfQuda(
"Using QUDA_HEAVY_QUARK_RESIDUAL");
719 const double ignore_mass = 1.0;
724 const double reliable_delta = (use_mixed_precision ? 1e-1 :0.0);
725 setInvertParams(localDim, host_precision, device_precision, device_precision_sloppy, device_precision_precondition,
726 num_offsets, offset, target_residual, target_fermilab_residual,
731 setColorSpinorParams(localDim, host_precision, &csParam);
736 const int fat_pad = getFatLinkPadding(localDim);
738 gaugeParam.
ga_pad = fat_pad;
742 const int long_pad = 3*fat_pad;
744 gaugeParam.
ga_pad = long_pad;
750 void** sln_pointer = (
void**)malloc(num_offsets*
sizeof(
void*));
751 int quark_offset = getColorVectorOffset(local_parity,
false, gaugeParam.
X)*host_precision;
752 void* src_pointer = (
char*)source + quark_offset;
754 for(
int i=0; i<num_offsets; ++i) sln_pointer[i] = (
char*)solutionArray[i] + quark_offset;
760 *num_iters = invertParam.
iter;
761 for(
int i=0; i<num_offsets; ++i){
768 qudamilc_called<false>(__func__, verbosity);
779 double target_residual,
780 double target_fermilab_residual,
781 const void*
const fatlink,
782 const void*
const longlink,
783 const double tadpole,
786 double*
const final_residual,
787 double*
const final_fermilab_residual,
792 qudamilc_called<true>(__func__, verbosity);
794 if(target_fermilab_residual == 0 && target_residual == 0){
795 errorQuda(
"qudaInvert: requesting zero residual\n");
799 const bool use_mixed_precision = (((quda_precision==2) && inv_args.
mixed_precision) ||
800 ((quda_precision==1) && (inv_args.
mixed_precision==2) ) ) ?
true :
false;
812 device_precision_sloppy = device_precision;
817 QudaPrecision device_precision_precondition = device_precision_sloppy;
820 setGaugeParams(localDim, host_precision, device_precision, device_precision_sloppy, device_precision_precondition, tadpole, &gaugeParam);
830 double& target_res = target_residual;
831 double& target_res_hq = target_fermilab_residual;
832 const double reliable_delta = 1e-1;
834 setInvertParams(localDim, host_precision, device_precision, device_precision_sloppy, device_precision_precondition,
835 mass, target_res, target_res_hq, inv_args.
max_iter, reliable_delta, local_parity, verbosity,
QUDA_CG_INVERTER, &invertParam);
839 setColorSpinorParams(localDim, host_precision, &csParam);
843 const int fat_pad = getFatLinkPadding(localDim);
844 const int long_pad = 3*fat_pad;
848 gaugeParam.
ga_pad = fat_pad;
853 gaugeParam.
ga_pad = long_pad;
857 invalidate_quda_gauge =
false;
860 int quark_offset = getColorVectorOffset(local_parity,
false, gaugeParam.
X);
862 invertQuda(((
char*)solution + quark_offset*host_precision),
863 ((
char*)source + quark_offset*host_precision),
867 *num_iters = invertParam.
iter;
868 *final_residual = invertParam.
true_res;
869 *final_fermilab_residual = invertParam.
true_res_hq;
872 qudamilc_called<false>(__func__, verbosity);
879 #ifdef GPU_CLOVER_DIRAC
883 qudamilc_called<true>(__func__);
885 QudaGaugeParam gaugeParam = newMILCGaugeParam(localDim, qudaPrecision,
894 qudamilc_called<true>(__func__);
896 QudaGaugeParam gaugeParam = newMILCGaugeParam(localDim, qudaPrecision,
905 qudamilc_called<true>(__func__);
909 qudamilc_called<false>(__func__);
916 qudamilc_called<true>(__func__);
918 qudamilc_called<false>(__func__);
931 void qudaCloverDerivative(
void* out,
void* gauge,
void* oprod,
int mu,
int nu,
double coeff,
int precision,
int parity,
int conjugate)
947 int external_precision,
int quda_precision) {
958 device_precision_sloppy = device_precision;
961 for(
int dir=0; dir<4; ++dir) gaugeParam.
X[dir] = dim[dir];
970 bool trivial_phase =
true;
971 for(
int dir=0; dir<3; ++dir){
986 gaugeParam.
cpu_prec = host_precision;
991 gaugeParam.
ga_pad = getFatLinkPadding(dim);
997 int external_precision,
int quda_precision,
double kappa) {
1007 device_precision_sloppy = device_precision;
1023 invertParam.
cpu_prec = host_precision;
1024 invertParam.
cuda_prec = device_precision;
1043 const void* milc_link) {
1044 qudamilc_called<true>(__func__);
1046 setGaugeParams(gaugeParam, localDim, inv_args, external_precision, quda_precision);
1049 qudamilc_called<false>(__func__);
1054 qudamilc_called<true>(__func__);
1056 qudamilc_called<false>(__func__);
1064 void* milc_clover_inv,
1067 double clover_coeff,
1072 setInvertParam(invertParam, inv_args, external_precision, quda_precision, 0.0);
1080 if (clover_alloc == 0) {
1084 errorQuda(
"Clover term already allocated");
1088 if (compute_trlog) {
1089 trlog[0] = invertParam.
trlogA[0];
1090 trlog[1] = invertParam.
trlogA[1];
1097 if (clover_alloc==1) {
1101 errorQuda(
"Trying to free non-allocated clover term");
1110 double clover_coeff,
1112 double target_residual,
1113 double target_fermilab_residual,
1116 void* cloverInverse,
1119 double*
const final_residual,
1120 double*
const final_fermilab_residual,
1124 if(target_fermilab_residual !=0 && target_residual != 0){
1125 errorQuda(
"qudaCloverInvert: conflicting residuals requested\n");
1127 }
else if(target_fermilab_residual == 0 && target_residual == 0){
1128 errorQuda(
"qudaCloverInvert: requesting zero residual\n");
1139 setInvertParam(invertParam, inv_args, external_precision, quda_precision, kappa);
1141 invertParam.
tol = (target_residual != 0) ? target_residual : target_fermilab_residual;
1150 *num_iters = invertParam.
iter;
1151 *final_residual = invertParam.
true_res;
1152 *final_fermilab_residual = invertParam.
true_res_hq;
1164 double*
const offset,
1166 double clover_coeff,
1168 const double* target_residual_offset,
1169 const void* milc_link,
1171 void* milc_clover_inv,
1173 void** solutionArray,
1174 double*
const final_residual,
1179 qudamilc_called<true>(__func__, verbosity);
1181 for(
int i=0; i<num_offsets; ++i){
1182 if(target_residual_offset[i] == 0){
1183 errorQuda(
"qudaMultishiftInvert: target residual cannot be zero\n");
1189 setInvertParam(invertParam, inv_args, external_precision, quda_precision, kappa);
1192 for(
int i=0; i<num_offsets; ++i){
1193 invertParam.
offset[i] = offset[i];
1194 invertParam.
tol_offset[i] = target_residual_offset[i];
1196 invertParam.
tol = target_residual_offset[0];
1211 *num_iters = invertParam.
iter;
1212 for(
int i=0; i<num_offsets; ++i) final_residual[i] = invertParam.
true_res_offset[i];
1224 double*
const offset,
1226 double clover_coeff,
1228 const double* target_residual_offset,
1229 const void* milc_link,
1231 void* milc_clover_inv,
1237 double*
const final_residual,
1241 qudamilc_called<true>(__func__, verbosity);
1243 for(
int i=0; i<num_offsets; ++i){
1244 if(target_residual_offset[i] == 0){
1245 errorQuda(
"qudaMultishiftInvert: target residual cannot be zero\n");
1251 setInvertParam(invertParam, inv_args, external_precision, quda_precision, kappa);
1254 for(
int i=0; i<num_offsets; ++i){
1255 invertParam.
offset[i] = offset[i];
1256 invertParam.
tol_offset[i] = target_residual_offset[i];
1258 invertParam.
tol = target_residual_offset[0];
1273 *num_iters = invertParam.
iter;
1274 for(
int i=0; i<num_offsets; ++i) final_residual[i] = invertParam.
true_res_offset[i];
1279 #endif // GPU_CLOVER_DIRAC
1282 #endif // BUILD_MILC_INTERFACE
QudaDiracFieldOrder dirac_order
void qudaGaugeForce(int precision, int num_loop_types, double milc_loop_coeff[3], double eb3, void *milc_sitelink, void *milc_momentum)
QudaMassNormalization mass_normalization
double tol_hq_offset[QUDA_MAX_MULTI_SHIFT]
QudaReconstructType reconstruct_sloppy
void freeCloverQuda(void)
void destroyQudaGaugeField(void *gauge)
void computeCloverTraceQuda(void *out, void *clover, int mu, int nu, int dim[4])
void invertMultiShiftQuda(void **_hp_x, void *_hp_b, QudaInvertParam *param)
void setVerbosityQuda(QudaVerbosity verbosity, const char prefix[], FILE *outfile)
void qudaHisqParamsInit(QudaHisqParams_t hisq_params)
QudaVerbosity verbosity_precondition
void computeHISQForceQuda(void *momentum, long long *flops, const double level2_coeff[6], const double fat7_coeff[6], const void *const staple_src[4], const void *const one_link_src[4], const void *const naik_src[4], const void *const w_link, const void *const v_link, const void *const u_link, const QudaGaugeParam *param)
enum QudaPrecision_s QudaPrecision
void * qudaCreateExtendedGaugeField(void *gauge, int geometry, int precision)
void 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)
QudaInverterType inv_type_precondition
QudaVerbosity getVerbosity()
void invertQuda(void *h_x, void *h_b, QudaInvertParam *param)
QudaDslashType dslash_type
void setUnitarizeLinksConstants(double unitarize_eps, double max_error, bool allow_svd, bool svd_only, double svd_rel_error, double svd_abs_error, bool check_unitarization=true)
QudaReconstructType reconstruct_precondition
QudaInverterType inv_type
void qudaComputeOprod(int precision, int num_terms, double **coeff, void **quark_field, void *oprod[2])
QudaGaugeParam gaugeParam
void computeKSLinkQuda(void *fatlink, void *longlink, void *ulink, void *inlink, double *path_coeff, QudaGaugeParam *param, QudaComputeFatMethod method)
enum QudaSolveType_s QudaSolveType
void loadGaugeQuda(void *h_gauge, QudaGaugeParam *param)
int computeGaugeForceQuda(void *mom, void *sitelink, int ***input_path_buf, int *path_length, double *loop_coeff, int num_paths, int max_length, double dt, QudaGaugeParam *qudaGaugeParam, double *timeinfo)
__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int mu
void qudaInit(QudaInitArgs_t input)
double reunit_svd_rel_error
void * createGaugeField(void *gauge, int geometry, QudaGaugeParam *param)
void qudaLoadGaugeField(int external_precision, int quda_precision, QudaInvertArgs_t inv_args, const void *milc_link)
void qudaCloverMultishiftInvert(int external_precision, int quda_precision, int num_offsets, double *const offset, double kappa, double clover_coeff, QudaInvertArgs_t inv_args, const double *target_residual, const void *milc_link, void *milc_clover, void *milc_clover_inv, void *source, void **solutionArray, double *const final_residual, int *num_iters)
QudaGaugeFieldOrder gauge_order
void computeAsqtadForceQuda(void *const momentum, long long *flops, const double act_path_coeff[6], const void *const one_link_src[4], const void *const naik_src[4], const void *const link, const QudaGaugeParam *param)
void qudaSaveGaugeField(void *gauge, void *inGauge)
void qudaMultishiftInvert(int external_precision, int precision, int num_offsets, double *const offset, QudaInvertArgs_t inv_args, const double *target_residual, const double *target_relative_residual, const void *const milc_fatlink, const void *const milc_longlink, const double tadpole, void *source, void **solutionArray, double *const final_residual, double *const final_relative_residual, int *num_iters)
void loadCloverQuda(void *h_clover, void *h_clovinv, QudaInvertParam *inv_param)
QudaSiteSubset siteSubset
cudaGaugeField * cudaGauge
void qudaSetLayout(QudaLayout_t layout)
QudaPrecision clover_cuda_prec_sloppy
void invertMultiShiftMDQuda(void **_hp_xe, void **_hp_xo, void **_hp_ye, void **_hp_yo, void *_hp_b, QudaInvertParam *param)
QudaFieldLocation input_location
void initCommsGridQuda(int nDim, const int *dims, QudaCommsMap func, void *fdata)
double true_res_hq_offset[QUDA_MAX_MULTI_SHIFT]
QudaUseInitGuess use_init_guess
QudaSolutionType solution_type
QudaPrecision clover_cuda_prec
void initQuda(int device)
QudaFieldLocation output_location
QudaPrecision clover_cuda_prec_precondition
void qudaCloverTrace(void *out, void *clover, int mu, int nu)
VOLATILE spinorFloat kappa
__device__ __host__ int index(int i, int j)
void qudaFreeCloverField()
QudaFieldOrder fieldOrder
QudaPrecision cuda_prec_sloppy
double tol_offset[QUDA_MAX_MULTI_SHIFT]
double true_res_offset[QUDA_MAX_MULTI_SHIFT]
QudaInvertParam newQudaInvertParam(void)
void * qudaCreateGaugeField(void *gauge, int geometry, int precision)
void computeCloverDerivativeQuda(void *out, void *gauge, void *oprod, int mu, int nu, double coeff, QudaParity parity, QudaGaugeParam *param, int conjugate)
QudaPrecision cuda_prec_precondition
QudaCloverFieldOrder clover_order
QudaGammaBasis gammaBasis
void qudaUpdateU(int precision, double eps, void *momentum, void *link)
enum QudaSolutionType_s QudaSolutionType
QudaGammaBasis gamma_basis
__constant__ double coeff
QudaPrecision cuda_prec_sloppy
double reunit_svd_abs_error
double offset[QUDA_MAX_MULTI_SHIFT]
void qudaFreeGaugeField()
int use_sloppy_partial_accumulator
enum QudaParity_s QudaParity
QudaReconstructType reconstruct
enum QudaLinkType_s QudaLinkType
void qudaLoadCloverField(int external_precision, int quda_precision, QudaInvertArgs_t inv_args, void *milc_clover, void *milc_clover_inv, QudaSolutionType solution_type, QudaSolveType solve_type, double clover_coeff, int compute_trlog, double *trlog)
void qudaHisqForce(int precision, const double level2_coeff[6], const double fat7_coeff[6], const void *const staple_src[4], const void *const one_link_src[4], const void *const naik_src[4], const void *const w_link, const void *const v_link, const void *const u_link, void *const milc_momentum)
void qudaCloverInvert(int external_precision, int quda_precision, double kappa, double clover_coeff, QudaInvertArgs_t inv_args, double target_residual, double target_fermilab_residual, const void *milc_link, void *milc_clover, void *milc_clover_inv, void *source, void *solution, double *const final_residual, double *const final_fermilab_residual, int *num_iters)
void qudaCloverDerivative(void *out, void *gauge, void *oprod, int mu, int nu, double coeff, int precision, int parity, int conjugate)
void qudaLoadUnitarizedLink(int precision, QudaFatLinkArgs_t fatlink_args, const double path_coeff[6], void *inlink, void *fatlink, void *ulink)
void computeStaggeredOprodQuda(void **oprod, void **quark, int num, double **coeff, QudaGaugeParam *param)
cpuColorSpinorField * out
QudaPrecision cuda_prec_precondition
void * memset(void *s, int c, size_t n)
void * createExtendedGaugeField(void *gauge, int geometry, QudaGaugeParam *param)
void updateGaugeFieldQuda(void *gauge, void *momentum, double dt, int conj_mom, int exact, QudaGaugeParam *param)
void qudaCloverMultishiftMDInvert(int external_precision, int quda_precision, int num_offsets, double *const offset, double kappa, double clover_coeff, QudaInvertArgs_t inv_args, const double *target_residual, const void *milc_link, void *milc_clover, void *milc_clover_inv, void *source, void **psiEven, void **psiOdd, void **pEven, void **pOdd, double *const final_residual, int *num_iters)
Main header file for the QUDA library.
void qudaInvert(int external_precision, int quda_precision, double mass, QudaInvertArgs_t inv_args, double target_resid, double target_relresid, const void *const milc_fatlink, const void *const milc_longlink, const double tadpole, void *source, void *solution, double *const final_resid, double *const final_rel_resid, int *num_iters)
void saveGaugeField(void *outGauge, void *inGauge, QudaGaugeParam *param)
void setKernelPackT(bool pack)
QudaResidualType residual_type
enum QudaVerbosity_s QudaVerbosity
void qudaLoadKSLink(int precision, QudaFatLinkArgs_t fatlink_args, const double act_path_coeff[6], void *inlink, void *fatlink, void *longlink)
QudaPrecision clover_cpu_prec
enum QudaComputeFatMethod_s QudaComputeFatMethod
enum QudaInverterType_s QudaInverterType
void qudaAsqtadForce(int precision, const double act_path_coeff[6], const void *const one_link_src[4], const void *const naik_src[4], const void *const link, void *const milc_momentum)
void qudaDestroyGaugeField(void *gauge)
QudaGaugeParam newQudaGaugeParam(void)
QudaPreserveSource preserve_source