mirror of
https://github.com/paboyle/Grid.git
synced 2025-06-17 07:17:06 +01:00
Merge branch 'feature/dirichlet' of https://github.com/paboyle/Grid into feature/dirichlet
This commit is contained in:
@ -42,6 +42,8 @@ public:
|
||||
bool is_smeared = false;
|
||||
RealD deriv_norm_sum;
|
||||
RealD deriv_max_sum;
|
||||
RealD Fdt_norm_sum;
|
||||
RealD Fdt_max_sum;
|
||||
int deriv_num;
|
||||
RealD deriv_us;
|
||||
RealD S_us;
|
||||
@ -51,12 +53,17 @@ public:
|
||||
deriv_num=0;
|
||||
deriv_norm_sum = deriv_max_sum=0.0;
|
||||
}
|
||||
void deriv_log(RealD nrm, RealD max) { deriv_max_sum+=max; deriv_norm_sum+=nrm; deriv_num++;}
|
||||
RealD deriv_max_average(void) { return deriv_max_sum/deriv_num; };
|
||||
RealD deriv_norm_average(void) { return deriv_norm_sum/deriv_num; };
|
||||
void deriv_log(RealD nrm, RealD max,RealD Fdt_nrm,RealD Fdt_max) {
|
||||
deriv_max_sum+=max; deriv_norm_sum+=nrm;
|
||||
Fdt_max_sum+=Fdt_max; Fdt_norm_sum+=Fdt_nrm; deriv_num++;
|
||||
}
|
||||
RealD deriv_max_average(void) { return deriv_max_sum/deriv_num; };
|
||||
RealD deriv_norm_average(void) { return deriv_norm_sum/deriv_num; };
|
||||
RealD Fdt_max_average(void) { return Fdt_max_sum/deriv_num; };
|
||||
RealD Fdt_norm_average(void) { return Fdt_norm_sum/deriv_num; };
|
||||
RealD deriv_timer(void) { return deriv_us; };
|
||||
RealD S_timer(void) { return deriv_us; };
|
||||
RealD refresh_timer(void) { return deriv_us; };
|
||||
RealD S_timer(void) { return S_us; };
|
||||
RealD refresh_timer(void) { return refresh_us; };
|
||||
void deriv_timer_start(void) { deriv_us-=usecond(); }
|
||||
void deriv_timer_stop(void) { deriv_us+=usecond(); }
|
||||
void refresh_timer_start(void) { refresh_us-=usecond(); }
|
||||
|
@ -39,7 +39,7 @@ struct GparityWilsonImplParams {
|
||||
Coordinate twists;
|
||||
//mu=Nd-1 is assumed to be the time direction and a twist value of 1 indicates antiperiodic BCs
|
||||
Coordinate dirichlet; // Blocksize of dirichlet BCs
|
||||
GparityWilsonImplParams() : twists(Nd, 0), dirichlet(Nd, 0) {};
|
||||
GparityWilsonImplParams() : twists(Nd, 0) { dirichlet.resize(0); };
|
||||
};
|
||||
|
||||
struct WilsonImplParams {
|
||||
@ -48,13 +48,13 @@ struct WilsonImplParams {
|
||||
AcceleratorVector<Real,Nd> twist_n_2pi_L;
|
||||
AcceleratorVector<Complex,Nd> boundary_phases;
|
||||
WilsonImplParams() {
|
||||
dirichlet.resize(Nd,0);
|
||||
dirichlet.resize(0);
|
||||
boundary_phases.resize(Nd, 1.0);
|
||||
twist_n_2pi_L.resize(Nd, 0.0);
|
||||
};
|
||||
WilsonImplParams(const AcceleratorVector<Complex,Nd> phi) : boundary_phases(phi), overlapCommsCompute(false) {
|
||||
twist_n_2pi_L.resize(Nd, 0.0);
|
||||
dirichlet.resize(Nd,0);
|
||||
dirichlet.resize(0);
|
||||
}
|
||||
};
|
||||
|
||||
@ -62,7 +62,7 @@ struct StaggeredImplParams {
|
||||
Coordinate dirichlet; // Blocksize of dirichlet BCs
|
||||
StaggeredImplParams()
|
||||
{
|
||||
dirichlet.resize(Nd,0);
|
||||
dirichlet.resize(0);
|
||||
};
|
||||
};
|
||||
|
||||
|
@ -67,6 +67,36 @@ NAMESPACE_BEGIN(Grid);
|
||||
virtual std::string action_name(){return "OneFlavourEvenOddRatioRationalPseudoFermionAction";}
|
||||
};
|
||||
|
||||
template<class Impl,class ImplF>
|
||||
class OneFlavourEvenOddRatioRationalMixedPrecPseudoFermionAction : public GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<Impl,ImplF> {
|
||||
public:
|
||||
typedef OneFlavourRationalParams Params;
|
||||
private:
|
||||
static RationalActionParams transcribe(const Params &in){
|
||||
RationalActionParams out;
|
||||
out.inv_pow = 2;
|
||||
out.lo = in.lo;
|
||||
out.hi = in.hi;
|
||||
out.MaxIter = in.MaxIter;
|
||||
out.action_tolerance = out.md_tolerance = in.tolerance;
|
||||
out.action_degree = out.md_degree = in.degree;
|
||||
out.precision = in.precision;
|
||||
out.BoundsCheckFreq = in.BoundsCheckFreq;
|
||||
return out;
|
||||
}
|
||||
|
||||
public:
|
||||
OneFlavourEvenOddRatioRationalMixedPrecPseudoFermionAction(FermionOperator<Impl> &_NumOp,
|
||||
FermionOperator<Impl> &_DenOp,
|
||||
FermionOperator<ImplF> &_NumOpF,
|
||||
FermionOperator<ImplF> &_DenOpF,
|
||||
const Params & p, Integer ReliableUpdateFreq
|
||||
) :
|
||||
GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<Impl,ImplF>(_NumOp, _DenOp,_NumOpF, _DenOpF, transcribe(p),ReliableUpdateFreq){}
|
||||
|
||||
virtual std::string action_name(){return "OneFlavourEvenOddRatioRationalPseudoFermionAction";}
|
||||
};
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
|
||||
#endif
|
||||
|
@ -153,7 +153,7 @@ protected:
|
||||
Real force_max = std::sqrt(maxLocalNorm2(force));
|
||||
Real impulse_max = force_max * ep * HMC_MOMENTUM_DENOMINATOR;
|
||||
|
||||
as[level].actions.at(a)->deriv_log(force_abs,force_max);
|
||||
as[level].actions.at(a)->deriv_log(force_abs,force_max,impulse_abs,impulse_max);
|
||||
|
||||
std::cout << GridLogIntegrator<< "["<<level<<"]["<<a<<"] Force average: " << force_abs <<" "<<name<<std::endl;
|
||||
std::cout << GridLogIntegrator<< "["<<level<<"]["<<a<<"] Force max : " << force_max <<" "<<name<<std::endl;
|
||||
@ -285,6 +285,8 @@ public:
|
||||
<<"["<<level<<"]["<< actionID<<"] : "
|
||||
<<" force max " << as[level].actions.at(actionID)->deriv_max_average()
|
||||
<<" norm " << as[level].actions.at(actionID)->deriv_norm_average()
|
||||
<<" Fdt max " << as[level].actions.at(actionID)->Fdt_max_average()
|
||||
<<" norm " << as[level].actions.at(actionID)->Fdt_norm_average()
|
||||
<<" calls " << as[level].actions.at(actionID)->deriv_num
|
||||
<< std::endl;
|
||||
}
|
||||
|
@ -711,7 +711,9 @@ public:
|
||||
this->_comms_recv.resize(npoints);
|
||||
this->same_node.resize(npoints);
|
||||
|
||||
if ( p.dirichlet.size() ) DirichletBlock(p.dirichlet); // comms send/recv set up
|
||||
if ( p.dirichlet.size() ==0 ) p.dirichlet.resize(grid->Nd(),0);
|
||||
|
||||
DirichletBlock(p.dirichlet); // comms send/recv set up
|
||||
|
||||
_unified_buffer_size=0;
|
||||
surface_list.resize(0);
|
||||
|
@ -190,7 +190,7 @@ void Lambda6Apply(uint64_t num1, uint64_t num2, uint64_t num3,
|
||||
|
||||
#define accelerator_barrier(dummy) \
|
||||
{ \
|
||||
cudaStreamSynchronize(cpuStream); \
|
||||
cudaDeviceSynchronize(); \
|
||||
cudaError err = cudaGetLastError(); \
|
||||
if ( cudaSuccess != err ) { \
|
||||
printf("accelerator_barrier(): Cuda error %s \n", \
|
||||
@ -362,11 +362,11 @@ accelerator_inline int acceleratorSIMTlane(int Nsimd) {
|
||||
dim3 hip_blocks ((num1+nt-1)/nt,num2,1); \
|
||||
if(hip_threads.x * hip_threads.y * hip_threads.z <= 64){ \
|
||||
hipLaunchKernelGGL(LambdaApply64,hip_blocks,hip_threads, \
|
||||
0,cpuStream, \
|
||||
0,0/*cpuStream*/, \
|
||||
num1,num2,nsimd, lambda); \
|
||||
} else { \
|
||||
hipLaunchKernelGGL(LambdaApply,hip_blocks,hip_threads, \
|
||||
0,cpuStream, \
|
||||
0,0/*cpuStream*/, \
|
||||
num1,num2,nsimd, lambda); \
|
||||
} \
|
||||
}
|
||||
|
Reference in New Issue
Block a user