NAMD
ComputeNonbondedPair.C
Go to the documentation of this file.
1 
7 #include "ComputeNonbondedPair.h"
8 #include "ReductionMgr.h"
9 #include "Patch.h"
10 #include "LdbCoordinator.h"
11 #include "PatchMap.h"
12 #include "ComputeMgr.h"
13 #include "Molecule.h"
14 
15 #include "Node.h"
16 #include "SimParameters.h"
17 
18 #include "Parameters.h"
19 #include "LJTable.h"
20 
21 #define MIN_DEBUG_LEVEL 4
22 // #define DEBUGM
23 #include "Debug.h"
24 
26  ComputeNonbondedWorkArrays* _workArrays,
27  int minPartition, int maxPartition, int numPartitions)
28  : ComputePatchPair(c,pid,trans), workArrays(_workArrays),
29  minPart(minPartition), maxPart(maxPartition), numParts(numPartitions)
30 {
32  if (pressureProfileOn) {
33  int n = pressureProfileAtomTypes;
37  } else {
39  pressureProfileData = NULL;
40  }
41  pairlistsValid = 0;
42  pairlistTolerance = 0.;
46 }
47 
48 #ifdef NAMD_AVXTILES
51 
52  if (avxTilesMode) {
53  tileLists.setSimParams(scaling, scale14, c1, c3, switchOn2,
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,
57  (float *)ljTable->table_val(0,0),
58  ljTable->get_table_dim(), knl_slow_grad_table,
59  knl_slow_ener_table, knl_excl_grad_table,
60  knl_excl_ener_table, avxTilesMode);
61  tileLists.atomUpdate(patch[0]->getTiles(), patch[1]->getTiles());
62  }
63 }
64 #endif
65 
68  for (int i=0; i<2; i++) {
70  // BEGIN LA
72  // END LA
73 
74  psiSumBox[i] = patch[i]->registerPsiSumDeposit(this);
75  intRadBox[i] = patch[i]->registerIntRadPickup(this);
76  bornRadBox[i] = patch[i]->registerBornRadPickup(this);
79  }
80 #if defined(NAMD_CUDA) || defined(NAMD_HIP)
82 #endif
83 }
84 
86 {
87  delete reduction;
89  delete [] pressureProfileData;
90  for (int i=0; i<2; i++) {
91  if (avgPositionBox[i] != NULL) {
93  }
94  // BEGIN LA
95  if (velocityBox[i] != NULL) {
97  }
98  // END LA
99 
100  if (psiSumBox[i] != NULL) {
102  }
103  if (intRadBox[i] != NULL) {
104  patch[i]->unregisterIntRadPickup(this,&intRadBox[i]);
105  }
106  if (bornRadBox[i] != NULL) {
108  }
109  if (dEdaSumBox[i] != NULL) {
111  }
112  if (dHdrPrefixBox[i] != NULL) {
114  }
115  }
116 }
117 
119 
120  if (patch[0]->flags.doGBIS) {
121  gbisPhase = 1 + (gbisPhase % 3);//1->2->3->1...
122  }
123 
124 #if !(defined(NAMD_CUDA) || defined(NAMD_HIP))
125  if ( patch[0]->flags.doNonbonded && (numAtoms[0] && numAtoms[1]) ) {
126  return 0; // work to do, enqueue as usual
127  } else {
128 #else
129  {
130 #endif
131 
132  if (patch[0]->flags.doGBIS) {
133  if (gbisPhase == 1) {
134  for (int i=0; i<2; i++) {
135  psiSumBox[i]->skip();
136  intRadBox[i]->skip();
137  }
138  if (patch[0]->flags.doNonbonded) return 1;
139  else gbisPhase = 2;
140  }
141  if (gbisPhase == 2) {
142  for (int i=0; i<2; i++) {
143  bornRadBox[i]->skip();
144  dEdaSumBox[i]->skip();
145  }
146  if (patch[0]->flags.doNonbonded) return 1;
147  else gbisPhase = 3;
148  }
149  if (gbisPhase == 3) {
150  for (int i=0; i<2; i++) {
151  dHdrPrefixBox[i]->skip();
152  }
153  }
154  }
155 
156  // skip all boxes
157  for (int i=0; i<2; i++) {
158  positionBox[i]->skip();
159  forceBox[i]->skip();
160  if ( patch[0]->flags.doMolly ) avgPositionBox[i]->skip();
161  // BEGIN LA
162  if (patch[0]->flags.doLoweAndersen) velocityBox[i]->skip();
163  // END LA
164  }
165 
167  reduction->submit();
168  if (pressureProfileOn)
170 
171 #if !(defined(NAMD_CUDA) || defined(NAMD_HIP))
172  // Inform load balancer
174 #endif
175 
176  return 1; // no work to do, do not enqueue
177  }
178 }
179 
181 {
182  #ifdef NAMD_AVXTILES
183  if (avxTilesMode) {
184  doForceTiles(p, pExt, r);
185  return;
186  }
187  #endif
188 
189  // Inform load balancer.
190  // I assume no threads will suspend until endWork is called
191 
192  //single phase declarations
193  int doEnergy = patch[0]->flags.doEnergy;
194  int a = 0; int b = 1;
195  // swap to place more atoms in inner loop (second patch)
196  if ( numAtoms[0] > numAtoms[1] ) { a = 1; b = 0; }
197  CompAtom* v[2];
198 
199 
200 /*******************************************************************************
201  * Prepare Parameters
202 *******************************************************************************/
203  if (!patch[0]->flags.doGBIS || gbisPhase == 1) {
204 
205 #ifdef TRACE_COMPUTE_OBJECTS
206  double traceObjStartTime = CmiWallTimer();
207 #endif
208 
209  DebugM(2,"doForce() called.\n");
210  DebugM(2, numAtoms[0] << " patch #1 atoms and " <<
211  numAtoms[1] << " patch #2 atoms\n");
212 
213 
214  for ( int i = 0; i < reductionDataSize; ++i )
215  reductionData[i] = 0;
216  if (pressureProfileOn) {
217  int n = pressureProfileAtomTypes;
218  memset(pressureProfileData, 0, 3*n*n*pressureProfileSlabs*sizeof(BigReal));
219  // adjust lattice dimensions to allow constant pressure
220  const Lattice &lattice = patch[0]->lattice;
222  pressureProfileMin = lattice.origin().z - 0.5*lattice.c().z;
223  }
224 
227 
231 
233 
235  params.savePairlists = 0;
236  params.usePairlists = 0;
237  if ( patch[0]->flags.savePairlists ) {
238  params.savePairlists = 1;
239  params.usePairlists = 1;
240  } else if ( patch[0]->flags.usePairlists && patch[1]->flags.usePairlists ) {
241  if ( ! pairlistsValid ||
242  ( patch[0]->flags.maxAtomMovement +
245  } else {
246  params.usePairlists = 1;
247  }
248  }
249  if ( ! params.usePairlists ) {
250  pairlistsValid = 0;
251  }
255  if ( params.savePairlists ) {
256  pairlistsValid = 1;
261  }
262 
263 
264  const Lattice &lattice = patch[0]->lattice;
265  params.offset = lattice.offset(trans[a]) - lattice.offset(trans[b]);
266 
267  PatchMap* patchMap = PatchMap::Object();
268  params.offset_f = params.offset + lattice.unscale(patchMap->center(patchID[a]))
269  - lattice.unscale(patchMap->center(patchID[b]));
270 
271  // Atom Sorting : If we are sorting the atoms along the line connecting
272  // the patch centers, then calculate a normalized vector pointing from
273  // patch a to patch b (i.e. outer loop patch to inner loop patch).
274  #if NAMD_ComputeNonbonded_SortAtoms != 0
275 
277 
278  #endif
279 
280  params.p[0] = p[a];
281  params.p[1] = p[b];
282  params.pExt[0] = pExt[a];
283  params.pExt[1] = pExt[b];
284 #ifdef NAMD_KNL
285  params.pFlt[0] = patch[a]->getCompAtomFlt();
286  params.pFlt[1] = patch[b]->getCompAtomFlt();
287 #endif
288  // BEGIN LA
290  if (params.doLoweAndersen) {
291  DebugM(4, "opening velocity boxes\n");
292  v[0] = velocityBox[0]->open();
293  v[1] = velocityBox[1]->open();
294  params.v[0] = v[a];
295  params.v[1] = v[b];
296  }
297  // END LA
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];
301 #endif
302  params.numAtoms[0] = numAtoms[a];
303  params.numAtoms[1] = numAtoms[b];
304  params.step = patch[0]->flags.step;
305 
306  // DMK - Atom Separation (water vs. non-water)
307  #if NAMD_SeparateWaters != 0
308  params.numWaterAtoms[0] = numWaterAtoms[a];
309  params.numWaterAtoms[1] = numWaterAtoms[b];
310  #endif
311 
312 
313 /*******************************************************************************
314  * Call Nonbonded Functions
315 *******************************************************************************/
316  if (numAtoms[0] && numAtoms[1]) {//only do if has atoms since gbis noWork doesn't account for no atoms
317 
318  //force calculation calls
319  if ( patch[0]->flags.doFullElectrostatics )
320  {
321 #if !(defined(NAMD_CUDA) || defined(NAMD_HIP))
322  params.fullf[0] = r[a]->f[Results::slow_virial];
323  params.fullf[1] = r[b]->f[Results::slow_virial];
324 #endif
325  if ( patch[0]->flags.doMolly ) {
326  if ( doEnergy )
328  else calcPair(&params);
329  CompAtom *p_avg[2];
330  p_avg[0] = avgPositionBox[0]->open();
331  p_avg[1] = avgPositionBox[1]->open();
332  params.p[0] = p_avg[a];
333  params.p[1] = p_avg[b];
334  if ( doEnergy ) calcSlowPairEnergy(&params);
335  else calcSlowPair(&params);
336  avgPositionBox[0]->close(&p_avg[0]);
337  avgPositionBox[1]->close(&p_avg[1]);
338  } else if ( patch[0]->flags.maxForceMerged == Results::slow ) {
339  if ( doEnergy ) calcMergePairEnergy(&params);
340  else calcMergePair(&params);
341  } else {
342  if ( doEnergy ) calcFullPairEnergy(&params);
343  else calcFullPair(&params);
344  }
345  }
346  else
347  if ( doEnergy ) calcPairEnergy(&params);
348  else calcPair(&params);
349 
350  }//end if has atoms
351 
352  // BEGIN LA
353  if (params.doLoweAndersen) {
354  DebugM(4, "closing velocity boxes\n");
355  velocityBox[0]->close(&v[0]);
356  velocityBox[1]->close(&v[1]);
357  }
358  // END LA
359  }// end not gbis
360 
361 /*******************************************************************************
362  * gbis Loop
363 *******************************************************************************/
364 if (patch[0]->flags.doGBIS) {
368  gbisParams.numPatches = 2;//pair
372  gbisParams.epsilon_p = simParams->dielectric;
374  gbisParams.kappa = simParams->kappa;
375  gbisParams.cutoff = simParams->cutoff;
377  gbisParams.a_cut = simParams->alpha_cutoff;
378  gbisParams.delta = simParams->gbis_delta;
379  gbisParams.beta = simParams->gbis_beta;
380  gbisParams.gamma = simParams->gbis_gamma;
381  gbisParams.alpha_max = simParams->alpha_max;
382  gbisParams.cid = cid;
383  gbisParams.patchID[0] = patch[a]->getPatchID();
384  gbisParams.patchID[1] = patch[b]->getPatchID();
386  if (patch[1]->flags.maxGroupRadius > gbisParams.maxGroupRadius)
388  gbisParams.doEnergy = doEnergy;
389  gbisParams.fsMax = simParams->fsMax;
390  for (int i = 0; i < numGBISPairlists; i++)
392 
393  //open boxes
394  if (gbisPhase == 1) {
395  gbisParams.intRad[0] = intRadBox[a]->open();
396  gbisParams.intRad[1] = intRadBox[b]->open();
397  gbisParams.psiSum[0] = psiSumBox[a]->open();
398  gbisParams.psiSum[1] = psiSumBox[b]->open();
401 
402  } else if (gbisPhase == 2) {
403  gbisParams.bornRad[0] = bornRadBox[a]->open();
404  gbisParams.bornRad[1] = bornRadBox[b]->open();
405  gbisParams.dEdaSum[0] = dEdaSumBox[a]->open();
406  gbisParams.dEdaSum[1] = dEdaSumBox[b]->open();
407  } else if (gbisPhase == 3) {
410  }
411 
412  //make call to calculate GBIS
415  }
416 
417  //close boxes
418  if (gbisPhase == 1) {
419  psiSumBox[0]->close(&(gbisParams.psiSum[a]));
420  psiSumBox[1]->close(&(gbisParams.psiSum[b]));
421  } else if (gbisPhase == 2) {
422  dEdaSumBox[0]->close(&(gbisParams.dEdaSum[a]));
423  dEdaSumBox[1]->close(&(gbisParams.dEdaSum[b]));
424 
425 
426  } else if (gbisPhase == 3) {
427  bornRadBox[0]->close(&(gbisParams.bornRad[a]));
428  bornRadBox[1]->close(&(gbisParams.bornRad[b]));
431  intRadBox[0]->close(&(gbisParams.intRad[a]));
432  intRadBox[1]->close(&(gbisParams.intRad[b]));
435  }
436 
437 }//end if doGBIS
438 
439  if (!patch[0]->flags.doGBIS || gbisPhase == 3) {
441  if (pressureProfileOn)
443 
444 #ifdef TRACE_COMPUTE_OBJECTS
445  traceUserBracketEvent(TRACE_COMPOBJ_IDOFFSET+cid, traceObjStartTime, CmiWallTimer());
446 #endif
447 
448 
449  reduction->submit();
450  if (pressureProfileOn)
452  }//end gbis end phase
453 
454 }//end do Force
455 
456 #ifdef NAMD_AVXTILES
457 void ComputeNonbondedPair::doForceTiles(CompAtom* p[2], CompAtomExt* pExt[2],
458  Results* r[2])
459 {
460  const int doEnergy = patch[0]->flags.doEnergy;
461  const int doVirial = patch[0]->flags.doVirial;
462  const int doFull = patch[0]->flags.doFullElectrostatics;
463  const int pid0 = tileLists.patchOrder0();
464  const int pid1 = tileLists.patchOrder1();
465 
466 #ifdef TRACE_COMPUTE_OBJECTS
467  double traceObjStartTime = CmiWallTimer();
468 #endif
469 
470  DebugM(2,"doForceTiles() called.\n");
471  DebugM(2, numAtoms[0] << " patch #1 atoms and " <<
472  numAtoms[1] << " patch #2 atoms\n");
473 
474  const Lattice &lattice = patch[pid0]->lattice;
475  for ( int i = 0; i < reductionDataSize; ++i )
476  reductionData[i] = 0;
477 
478  int doList = false;
479  if ( patch[0]->flags.savePairlists ) {
480  doList = 1;
481  pairlistsValid = 1;
484  } else if ( patch[0]->flags.usePairlists && patch[1]->flags.usePairlists ) {
485  if (!pairlistsValid || (patch[0]->flags.maxAtomMovement +
489  pairlistsValid = 0;
490  doList = 1;
491  }
492  } else
493  doList = 1;
494 
495  if (doList) {
496  double plcutoff = cutoff;
497  if (patch[0]->flags.savePairlists)
498  plcutoff += patch[0]->flags.pairlistTolerance +
500 
501  tileLists.updateBuildInfo(patch[0]->flags.step, minPart, maxPart,
502  numParts, plcutoff);
503  }
504 
505  const PatchMap* patchMap = PatchMap::Object();
506  Vector offset = patchMap->center(patchID[pid0]) -
507  patchMap->center(patchID[pid1]);
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);
513  tileLists.updateParams(lattice, offset, cutoff);
514 
515  tileLists.nbForceAVX512(doEnergy, doVirial, doFull, doList);
516 
517  reductionData[exclChecksumIndex] = tileLists.exclusionChecksum();
518  if (doEnergy) {
519  reductionData[vdwEnergyIndex] = tileLists.energyVdw();
520  reductionData[electEnergyIndex] = tileLists.energyElec();
521  if (doFull)
523  }
524  if (doVirial) {
525  reductionData[virialIndex_XX] = tileLists.virialXX();
526  reductionData[virialIndex_XY] = tileLists.virialXY();
527  reductionData[virialIndex_XZ] = tileLists.virialXZ();
528  reductionData[virialIndex_YX] = tileLists.virialXY();
529  reductionData[virialIndex_YY] = tileLists.virialYY();
530  reductionData[virialIndex_YZ] = tileLists.virialYZ();
531  reductionData[virialIndex_ZX] = tileLists.virialXZ();
532  reductionData[virialIndex_ZY] = tileLists.virialYZ();
533  reductionData[virialIndex_ZZ] = tileLists.virialZZ();
534  if (doFull) {
535  reductionData[fullElectVirialIndex_XX] = tileLists.virialSlowXX();
536  reductionData[fullElectVirialIndex_XY] = tileLists.virialSlowXY();
537  reductionData[fullElectVirialIndex_XZ] = tileLists.virialSlowXZ();
538  reductionData[fullElectVirialIndex_YX] = tileLists.virialSlowXY();
539  reductionData[fullElectVirialIndex_YY] = tileLists.virialSlowYY();
540  reductionData[fullElectVirialIndex_YZ] = tileLists.virialSlowYZ();
541  reductionData[fullElectVirialIndex_ZX] = tileLists.virialSlowXZ();
542  reductionData[fullElectVirialIndex_ZY] = tileLists.virialSlowYZ();
543  reductionData[fullElectVirialIndex_ZZ] = tileLists.virialSlowZZ();
544  }
545  }
546 
547  if (!patch[0]->flags.doGBIS || gbisPhase == 3) {
549 
550 #ifdef TRACE_COMPUTE_OBJECTS
551  traceUserBracketEvent(TRACE_COMPOBJ_IDOFFSET+cid, traceObjStartTime,
552  CmiWallTimer());
553 #endif
554 
555  reduction->submit();
556  }
557 }//end doForceTiles
558 #endif
static Node * Object()
Definition: Node.h:86
Box< Patch, GBReal > * registerDEdaSumDeposit(Compute *cid)
Definition: Patch.C:204
Pairlists * pairlists
Box< Patch, CompAtom > * registerAvgPositionPickup(Compute *cid)
Definition: Patch.C:134
int sequence(void)
Definition: Compute.h:64
void unregisterAvgPositionPickup(Compute *cid, Box< Patch, CompAtom > **const box)
Definition: Patch.C:140
BigReal solvent_dielectric
Box< Patch, Real > * dHdrPrefixBox[2]
static void submitReductionData(BigReal *, SubmitReduction *)
Parameters * parameters
int ComputeID
Definition: NamdTypes.h:183
#define TRACE_COMPOBJ_IDOFFSET
Definition: Compute.h:77
void unregisterPsiSumDeposit(Compute *cid, Box< Patch, GBReal > **const box)
Definition: Patch.C:174
CompAtom * p[2]
Lattice & lattice
Definition: Patch.h:126
Box< Patch, CompAtom > * positionBox[2]
static PatchMap * Object()
Definition: PatchMap.h:27
CompAtom * p[2]
void calcGBIS(nonbonded *params, GBISParamStruct *gbisParams)
Definition: ComputeGBIS.C:261
Definition: Vector.h:64
SimParameters * simParameters
Definition: Node.h:178
ComputeNonbondedWorkArrays *const workArrays
Box< Patch, Real > * intRadBox[2]
virtual void doForce(CompAtom *p[2], CompAtomExt *pExt[2], Results *r[2])
BigReal & item(int i)
Definition: ReductionMgr.h:312
#define DebugM(x, y)
Definition: Debug.h:59
CompAtomExt * pExt[2]
CompAtom * v[2]
BigReal gbis_gamma
BigReal z
Definition: Vector.h:66
void unregisterDHdrPrefixPickup(Compute *cid, Box< Patch, Real > **const box)
Definition: Patch.C:222
int usePairlists
Definition: PatchTypes.h:38
BigReal * reduction
static void submitPressureProfileData(BigReal *, SubmitReduction *)
static BigReal pressureProfileThickness
SubmitReduction * willSubmit(int setID, int size=-1)
Definition: ReductionMgr.C:365
Box< Patch, Real > * registerBornRadPickup(Compute *cid)
Definition: Patch.C:196
int get_table_dim() const
Definition: LJTable.h:44
if(ComputeNonbondedUtil::goMethod==2)
LDObjHandle ldObjHandle
Definition: Compute.h:44
static ReductionMgr * Object(void)
Definition: ReductionMgr.h:278
int doLoweAndersen
Definition: PatchTypes.h:26
static void(* calcMergePair)(nonbonded *)
BigReal length(void) const
Definition: Vector.h:169
Pairlists gbisStepPairlists[numGBISPairlists]
static void(* calcMergePairEnergy)(nonbonded *)
Vector origin() const
Definition: Lattice.h:262
SimParameters * simParameters
BigReal coulomb_radius_offset
static void(* calcSlowPairEnergy)(nonbonded *)
Flags flags
Definition: Patch.h:127
SubmitReduction * pressureProfileReduction
Box< Patch, CompAtom > * avgPositionBox[2]
__global__ void const int const TileList *__restrict__ tileLists
static void(* calcPair)(nonbonded *)
BigReal alpha_max
Box< Patch, GBReal > * registerPsiSumDeposit(Compute *cid)
Definition: Patch.C:164
static void(* calcSlowPair)(nonbonded *)
Box< Patch, Results > * forceBox[2]
const TableEntry * table_val(unsigned int i, unsigned int j) const
Definition: LJTable.h:35
virtual void initialize()
Pairlists * gbisStepPairlists[4]
Box< Patch, CompAtom > * velocityBox[2]
static BigReal pressureProfileMin
Box< Patch, GBReal > * psiSumBox[2]
int doEnergy
Definition: PatchTypes.h:20
int doFullElectrostatics
Definition: PatchTypes.h:23
Box< Patch, Real > * registerIntRadPickup(Compute *cid)
Definition: Patch.C:179
void skipWork(const LDObjHandle &handle)
Force * f[maxNumForces]
Definition: PatchTypes.h:67
BigReal gbis_beta
CompAtomExt * pExt[2]
BigReal x
Definition: Vector.h:66
virtual void atomUpdate()
BigReal alpha_cutoff
int PatchID
Definition: NamdTypes.h:182
Box< Patch, Real > * bornRadBox[2]
static LdbCoordinator * Object()
BigReal maxAtomMovement
Definition: PatchTypes.h:41
Box< Patch, GBReal > * dEdaSumBox[2]
void skip(void)
Definition: Box.h:63
void unregisterBornRadPickup(Compute *cid, Box< Patch, Real > **const box)
Definition: Patch.C:199
int gbisPhase
Definition: Compute.h:39
Parameters * parameters
Definition: Node.h:177
PatchID getPatchID()
Definition: Patch.h:114
static void(* calcPairEnergy)(nonbonded *)
ScaledPosition center(int pid) const
Definition: PatchMap.h:99
SubmitReduction * reduction
Vector offset(int i) const
Definition: Lattice.h:242
#define simParams
Definition: Output.C:127
Position unscale(ScaledPosition s) const
Definition: Lattice.h:77
Random * rand
Definition: Node.h:172
GBISParamStruct gbisParams
static void(* calcFullPair)(nonbonded *)
void unregisterIntRadPickup(Compute *cid, Box< Patch, Real > **const box)
Definition: Patch.C:182
Box< Patch, Real > * registerDHdrPrefixPickup(Compute *cid)
Definition: Patch.C:218
ComputeNonbondedPair(ComputeID c, PatchID pid[], int trans[], ComputeNonbondedWorkArrays *_workArrays, int minPartition=0, int maxPartition=1, int numPartitions=1)
int doVirial
Definition: PatchTypes.h:21
BigReal y
Definition: Vector.h:66
static const LJTable * ljTable
static const int numGBISPairlists
static void(* calcFullPairEnergy)(nonbonded *)
BigReal pairlistTolerance
Definition: PatchTypes.h:40
BigReal gbis_delta
void unregisterVelocityPickup(Compute *cid, Box< Patch, CompAtom > **const box)
Definition: Patch.C:154
BigReal dielectric
int doGBIS
Definition: PatchTypes.h:28
void submit(void)
Definition: ReductionMgr.h:323
Force * fullf[2]
BigReal reductionData[reductionDataSize]
BigReal * pressureProfileReduction
Data * open(void)
Definition: Box.h:39
BigReal maxGroupRadius
Definition: PatchTypes.h:42
const ComputeID cid
Definition: Compute.h:43
void unregisterDEdaSumDeposit(Compute *cid, Box< Patch, GBReal > **const box)
Definition: Patch.C:212
Box< Patch, CompAtom > * registerVelocityPickup(Compute *cid)
Definition: Patch.C:148
void close(Data **const t)
Definition: Box.h:49
Vector c() const
Definition: Lattice.h:254
ComputeNonbondedWorkArrays * workArrays
double BigReal
Definition: common.h:114
int step
Definition: PatchTypes.h:16
void register_cuda_compute_pair(ComputeID c, PatchID pid[], int t[])