21 #define MIN_DEBUG_LEVEL 4
27 int minPartition,
int maxPartition,
int numPartitions)
29 minPart(minPartition), maxPart(maxPartition), numParts(numPartitions)
54 knl_fast_grad_table, knl_fast_ener_table,
55 knl_scor_grad_table, knl_scor_ener_table,
56 avx_tiles_eps4_sigma, avx_tiles_eps4_sigma_14,
59 knl_slow_ener_table, knl_excl_grad_table,
60 knl_excl_ener_table, avxTilesMode);
68 for (
int i=0; i<2; i++) {
80 #if defined(NAMD_CUDA) || defined(NAMD_HIP)
90 for (
int i=0; i<2; i++) {
120 if (
patch[0]->flags.doGBIS) {
124 #if !(defined(NAMD_CUDA) || defined(NAMD_HIP))
132 if (
patch[0]->flags.doGBIS) {
134 for (
int i=0; i<2; i++) {
138 if (
patch[0]->flags.doNonbonded)
return 1;
142 for (
int i=0; i<2; i++) {
146 if (
patch[0]->flags.doNonbonded)
return 1;
150 for (
int i=0; i<2; i++) {
157 for (
int i=0; i<2; i++) {
171 #if !(defined(NAMD_CUDA) || defined(NAMD_HIP))
184 doForceTiles(p, pExt, r);
194 int a = 0;
int b = 1;
205 #ifdef TRACE_COMPUTE_OBJECTS
206 double traceObjStartTime = CmiWallTimer();
209 DebugM(2,
"doForce() called.\n");
211 numAtoms[1] <<
" patch #2 atoms\n");
237 if (
patch[0]->flags.savePairlists ) {
242 (
patch[0]->flags.maxAtomMovement +
274 #if NAMD_ComputeNonbonded_SortAtoms != 0
291 DebugM(4,
"opening velocity boxes\n");
298 #if !(defined(NAMD_CUDA) || defined(NAMD_HIP))
299 params.
ff[0] = r[a]->
f[Results::nbond_virial];
300 params.
ff[1] = r[b]->
f[Results::nbond_virial];
307 #if NAMD_SeparateWaters != 0
308 params.numWaterAtoms[0] = numWaterAtoms[a];
309 params.numWaterAtoms[1] = numWaterAtoms[b];
319 if (
patch[0]->flags.doFullElectrostatics )
321 #if !(defined(NAMD_CUDA) || defined(NAMD_HIP))
325 if (
patch[0]->flags.doMolly ) {
354 DebugM(4,
"closing velocity boxes\n");
364 if (
patch[0]->flags.doGBIS) {
444 #ifdef TRACE_COMPUTE_OBJECTS
463 const int pid0 =
tileLists.patchOrder0();
464 const int pid1 =
tileLists.patchOrder1();
466 #ifdef TRACE_COMPUTE_OBJECTS
467 double traceObjStartTime = CmiWallTimer();
470 DebugM(2,
"doForceTiles() called.\n");
472 numAtoms[1] <<
" patch #2 atoms\n");
479 if (
patch[0]->flags.savePairlists ) {
497 if (
patch[0]->flags.savePairlists)
508 const int t1 =
trans[pid0];
509 const int t2 =
trans[pid1];
510 offset.
x += (t1%3-1) - (t2%3-1);
511 offset.
y += ((t1/3)%3-1) - ((t2/3)%3-1);
512 offset.
z += (t1/9-1) - (t2/9-1);
515 tileLists.nbForceAVX512(doEnergy, doVirial, doFull, doList);
550 #ifdef TRACE_COMPUTE_OBJECTS
Box< Patch, GBReal > * registerDEdaSumDeposit(Compute *cid)
static int pressureProfileSlabs
Box< Patch, CompAtom > * registerAvgPositionPickup(Compute *cid)
virtual void initialize()
void unregisterAvgPositionPickup(Compute *cid, Box< Patch, CompAtom > **const box)
BigReal solvent_dielectric
Box< Patch, Real > * dHdrPrefixBox[2]
static void submitReductionData(BigReal *, SubmitReduction *)
#define TRACE_COMPOBJ_IDOFFSET
void unregisterPsiSumDeposit(Compute *cid, Box< Patch, GBReal > **const box)
Box< Patch, CompAtom > * positionBox[2]
static PatchMap * Object()
void calcGBIS(nonbonded *params, GBISParamStruct *gbisParams)
SimParameters * simParameters
ComputeNonbondedWorkArrays *const workArrays
Box< Patch, Real > * intRadBox[2]
virtual void doForce(CompAtom *p[2], CompAtomExt *pExt[2], Results *r[2])
void unregisterDHdrPrefixPickup(Compute *cid, Box< Patch, Real > **const box)
static void submitPressureProfileData(BigReal *, SubmitReduction *)
static BigReal pressureProfileThickness
SubmitReduction * willSubmit(int setID, int size=-1)
Box< Patch, Real > * registerBornRadPickup(Compute *cid)
int get_table_dim() const
if(ComputeNonbondedUtil::goMethod==2)
static ReductionMgr * Object(void)
static void(* calcMergePair)(nonbonded *)
BigReal length(void) const
Pairlists gbisStepPairlists[numGBISPairlists]
static int pressureProfileAtomTypes
static void(* calcMergePairEnergy)(nonbonded *)
SimParameters * simParameters
BigReal * pressureProfileData
BigReal coulomb_radius_offset
static void(* calcSlowPairEnergy)(nonbonded *)
SubmitReduction * pressureProfileReduction
Box< Patch, CompAtom > * avgPositionBox[2]
__global__ void const int const TileList *__restrict__ tileLists
static void(* calcPair)(nonbonded *)
Box< Patch, GBReal > * registerPsiSumDeposit(Compute *cid)
static void(* calcSlowPair)(nonbonded *)
Box< Patch, Results > * forceBox[2]
const TableEntry * table_val(unsigned int i, unsigned int j) const
virtual void initialize()
Pairlists * gbisStepPairlists[4]
Box< Patch, CompAtom > * velocityBox[2]
static BigReal pressureProfileMin
Box< Patch, GBReal > * psiSumBox[2]
BigReal pairlistTolerance
Box< Patch, Real > * registerIntRadPickup(Compute *cid)
void skipWork(const LDObjHandle &handle)
static Bool pressureProfileOn
virtual void atomUpdate()
Box< Patch, Real > * bornRadBox[2]
static LdbCoordinator * Object()
Box< Patch, GBReal > * dEdaSumBox[2]
void unregisterBornRadPickup(Compute *cid, Box< Patch, Real > **const box)
static void(* calcPairEnergy)(nonbonded *)
ScaledPosition center(int pid) const
SubmitReduction * reduction
Vector offset(int i) const
Position unscale(ScaledPosition s) const
GBISParamStruct gbisParams
static void(* calcFullPair)(nonbonded *)
void unregisterIntRadPickup(Compute *cid, Box< Patch, Real > **const box)
Box< Patch, Real > * registerDHdrPrefixPickup(Compute *cid)
ComputeNonbondedPair(ComputeID c, PatchID pid[], int trans[], ComputeNonbondedWorkArrays *_workArrays, int minPartition=0, int maxPartition=1, int numPartitions=1)
static const LJTable * ljTable
static const int numGBISPairlists
static void(* calcFullPairEnergy)(nonbonded *)
BigReal pairlistTolerance
void unregisterVelocityPickup(Compute *cid, Box< Patch, CompAtom > **const box)
BigReal reductionData[reductionDataSize]
BigReal * pressureProfileReduction
void unregisterDEdaSumDeposit(Compute *cid, Box< Patch, GBReal > **const box)
Box< Patch, CompAtom > * registerVelocityPickup(Compute *cid)
void close(Data **const t)
ComputeNonbondedWorkArrays * workArrays
void register_cuda_compute_pair(ComputeID c, PatchID pid[], int t[])