NAMD
AVXTiles.h
Go to the documentation of this file.
1 #ifndef AVXTILES_H
2 #define AVXTILES_H
3 
4 #include "common.h"
5 #include "NamdTypes.h"
6 
7 #ifdef NAMD_AVXTILES
8 
9 #if defined(NAMD_CUDA) || defined(NAMD_HIP)
10 #error Cannot build with GPU Acceleration (NAMD_CUDA or NAMD_HIP) and NAMD_AVXTILES defined
11 #endif
12 
13 // Each patch is associtated with an AVXTiles object with
14 // - pointer to CUDA single precision {x,y,z,q} data
15 // - bounding box data for atoms as divided into tiles
16 // - arrays for (slow) forces
17 // - local storage of sort order, LJ types, and exclusion data
18 class AVXTiles {
19  public:
20  struct AVXTilesAtom { float x, y, z, q; };
21  struct AVXTilesForce { float x, y, z, w; };
22 
23  AVXTiles();
24  ~AVXTiles();
25 
26  inline int numAtoms() const { return _numAtoms; }
27  inline int numFreeAtoms() const { return _numFreeAtoms; }
28  inline int numTiles() const { return _numTiles; }
29 
30  // Signal that force data has been touched and should be reduced with Patch
31  // data and zeroed on subsequent step
32  inline void touch() { _touched = true; }
33 
34  inline void realloc(const int numAtoms, CudaAtom *ac) {
35  atoms = (AVXTilesAtom *)ac;
36  if (numAtoms != _numAtoms) {
37  if (numAtoms > _numAtoms) _touched = true;
38  _numAtoms = numAtoms;
39  _numTiles = ((numAtoms - 1) >> 4) + 1;
40  if (_numTiles > _numTilesAlloc) _realloc();
41  }
42  }
43 
44  inline void zeroForces(const int doSlow) {
45  if (_touched) {
46  memset(forces, 0, _numAtoms * sizeof(AVXTilesForce));
47  if (doSlow) memset(forcesSlow, 0, _numAtoms * sizeof(AVXTilesForce));
48  _touched = false;
49  }
50  }
51 
52  void atomUpdate(const CompAtom *compAtom,
53  const CompAtomExt *compAtomExt);
54 
55  // Needed on steps where we will a rebuild of the tile list is required
56  // - if multiple Computes use the patch, only calculate once per timestep
57  inline void buildBoundingBoxes(const int step) {
58  if (step != _lastBuild) _buildBoundingBoxes(step);
59  }
60 
61  void nativeForceVirialUpdate(const int doSlow, const int doVirial,
62  const CompAtom *p, const Vector &center,
63  Force * __restrict__ natForces,
64  Force * __restrict__ natForcesSlow,
65  const Force * __restrict__ natForcesVirial,
66  const Force * __restrict__ natForcesSlowVirial,
67  double virial[6], double virialSlow[6]);
68 
69  // ----------------------- Per Atom Data --------------------------
70  AVXTilesAtom *atoms;
71  AVXTilesForce *forces;
72  AVXTilesForce *forcesSlow;
73 
74  int *vdwTypes;
75  int *atomIndex;
76 #ifdef MEM_OPT_VERSION
77  int *atomExclIndex;
78 #endif
79  int *exclIndexStart, *exclIndexMaxDiff;
80  int *reverseOrder;
81 
82  // ----------------------- Per Tile Data --------------------------
83 
84  float *bbox_x, *bbox_y, *bbox_z, *bbox_wx, *bbox_wy, *bbox_wz;
85 
86  private:
87  int _numAtoms, _numFreeAtoms, _numTiles, _numTilesAlloc, _lastBuild;
88  bool _touched;
89  void _realloc();
90  void _buildBoundingBoxes(const int step);
91  template <int doSlow, int doVirial, int touched>
92  void _nativeForceVirialUpdate(const CompAtom *p, const Vector &center,
93  Force * __restrict__ natForces,
94  Force * __restrict__ natForcesSlow,
95  const Force * __restrict__ natForcesVirial,
96  const Force * __restrict__ natForcesSlowVirial,
97  double virial[6], double virialSlow[6]);
98 };
99 
100 #endif // NAMD_AVXTILES
101 #endif // AVXTILES_H
Definition: Vector.h:64
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ const int *__restrict__ const int const float2 *__restrict__ cudaTextureObject_t const int *__restrict__ vdwTypes
static __thread atom * atoms
static __thread float4 * forces
gridSize z
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ const int *__restrict__ const int const float2 *__restrict__ cudaTextureObject_t const int *__restrict__ const float3 const float3 const float3 const float4 *__restrict__ const float cudaTextureObject_t cudaTextureObject_t float const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ exclIndexMaxDiff
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ const int *__restrict__ const int const float2 *__restrict__ cudaTextureObject_t const int *__restrict__ const float3 const float3 const float3 const float4 *__restrict__ const float cudaTextureObject_t cudaTextureObject_t float const PatchPairRecord *__restrict__ const int *__restrict__ atomIndex
gridSize y
gridSize x