NAMD
ComputeBondedCUDAKernel.h
Go to the documentation of this file.
1 #ifndef COMPUTEBONDEDCUDAKERNEL_H
2 #define COMPUTEBONDEDCUDAKERNEL_H
3 #include "CudaUtils.h"
4 #include "TupleTypesCUDA.h"
5 #include "CudaNonbondedTables.h"
6 
7 #if defined(NAMD_CUDA) || defined(NAMD_HIP)
8 
9 // Use Fixed point (24.40) force?
10 //#define USE_FP_FORCE
11 #ifdef USE_FP_FORCE
12 #define FORCE_TYPE long long int
13 #else
14 #define FORCE_TYPE double
15 #endif
16 #define USE_STRIDED_FORCE
17 
18 #ifndef USE_STRIDED_FORCE
19 #error "Non-USE_STRIDED_FORCE not implemented"
20 #endif
21 
22 // Use Fixed point (34.30) virial?
23 //#define USE_FP_VIRIAL
24 #ifdef USE_FP_VIRIAL
25 #define VIRIAL_TYPE long long int
26 #else
27 #define VIRIAL_TYPE double
28 #endif
29 
30 #define WRITE_FULL_VIRIALS
31 
32 // #ifdef NAMD_CUDA
33 #define USE_BONDED_FORCE_ATOMIC_STORE
34 // #endif
35 
36 // Scaling factors for 24.40 fixed point
37 #ifdef USE_FP_FORCE
38 static __constant__ const float float_to_force = (float)(1ll << 40);
39 static __constant__ const float force_to_float = (float)1.0/(float)(1ll << 40);
40 static __constant__ const double force_to_double = (double)1.0/(double)(1ll << 40);
41 #else
42 static __constant__ const float float_to_force = 1.0f;
43 static __constant__ const float force_to_float = 1.0f;
44 static __constant__ const double force_to_double = 1.0;
45 #endif
46 
47 #ifdef USE_FP_VIRIAL
48 static __constant__ const float float_to_virial = (float)(1ll << 30);
49 static __constant__ const double double_to_virial = (double)(1ll << 30);
50 static __constant__ const double virial_to_double = (double)1.0/(double)(1ll << 30);
51 static __constant__ const long long int CONVERT_TO_VIR = (1ll << 10);
52 #endif
53 
55 public:
56 
57  // Enumeration for energies_virials[]
73 
74  template <typename T>
75  struct BondedVirial {
76 #ifdef WRITE_FULL_VIRIALS
77  T xx;
78  T xy;
79  T xz;
80  T yx;
81  T yy;
82  T yz;
83  T zx;
84  T zy;
85  T zz;
86 #else
87 #error "non-WRITE_FULL_VIRIALS not implemented yet"
88  union {
89  double sforce_dp[27][3];
90  long long int sforce_fp[27][3];
91  };
92 #endif
93  };
94 
95 private:
96  const int deviceID;
97  CudaNonbondedTables& cudaNonbondedTables;
98 
99  // This stores all bonds, angles, dihedrals, and impropers in a single
100  // contigious memory array.
101  char* tupleData;
102  int tupleDataSize;
103 
104  // ---------------------------------------------------------------------------------
105  // NOTE: bonds, angles, dihedrals, impropers, etc. - pointers below are
106  // computed pointers pointing to tupleData -array
107  // DO NOT DEALLOCATE THESE!
108  int numBonds;
109  CudaBond* bonds;
110 
111  int numAngles;
112  CudaAngle* angles;
113 
114  int numDihedrals;
115  CudaDihedral* dihedrals;
116 
117  int numImpropers;
118  CudaDihedral* impropers;
119 
120  int numModifiedExclusions;
121  CudaExclusion* modifiedExclusions;
122 
123  int numExclusions;
124  CudaExclusion* exclusions;
125 
126  int numCrossterms;
127  CudaCrossterm* crossterms;
128  // ---------------------------------------------------------------------------------
129 
130  // Device memory for coordinates
131  float4* xyzq;
132  int xyzqSize;
133 
134  FORCE_TYPE* forceList;
135  int* forceListCounter;
136  int* forceListStarts;
137  int* forceListNexts;
138  int forceListSize;
139  int forceListStartsSize;
140  int forceListNextsSize;
141 
142  // Device memory for forces:
143  // [normal, nbond, slow]
144  FORCE_TYPE* forces;
145  int forcesSize;
146 
147  CudaBondValue* bondValues;
148  CudaAngleValue* angleValues;
149  CudaDihedralValue* dihedralValues;
150  CudaDihedralValue* improperValues;
151  CudaCrosstermValue* crosstermValues;
152 
153  // Accumulated energy values for every bonded type
154  double* energies_virials;
155 
156 public:
157 
158  ComputeBondedCUDAKernel(int deviceID, CudaNonbondedTables& cudaNonbondedTables);
160 
161  static int warpAlign(const int n) {return ((n + WARPSIZE - 1)/WARPSIZE)*WARPSIZE;}
162 
163  void update(
164  const int numBondsIn,
165  const int numAnglesIn,
166  const int numDihedralsIn,
167  const int numImpropersIn,
168  const int numModifiedExclusionsIn,
169  const int numExclusionsIn,
170  const int numCrosstermsIn,
171  const char* h_tupleData,
172  cudaStream_t stream);
173 
174  void setupBondValues(int numBondValues, CudaBondValue* h_bondValues);
175  void setupAngleValues(int numAngleValues, CudaAngleValue* h_angleValues);
176  void setupDihedralValues(int numDihedralValues, CudaDihedralValue* h_dihedralValues);
177  void setupImproperValues(int numImproperValues, CudaDihedralValue* h_improperValues);
178  void setupCrosstermValues(int numCrosstermValues, CudaCrosstermValue* h_crosstermValues);
179 
180  int getForceStride(const int atomStorageSize);
181  int getForceSize(const int atomStorageSize);
182  int getAllForceSize(const int atomStorageSize, const bool doSlow);
183 
184  void bondedForce(
185  const double scale14, const int atomStorageSize,
186  const bool doEnergy, const bool doVirial, const bool doSlow,
187  const float3 lata, const float3 latb, const float3 latc,
188  const float cutoff2, const float r2_delta, const int r2_delta_expc,
189  const float4* h_xyzq, FORCE_TYPE* h_forces,
190  double *h_energies,
191  cudaStream_t stream);
192 
193 };
194 
195 #endif
196 
197 #endif // COMPUTEBONDEDCUDAKERNEL_H
static __constant__ const double force_to_double
static int warpAlign(const int n)
static __constant__ const float force_to_float
int getForceStride(const int atomStorageSize)
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ const int *__restrict__ const int const float2 *__restrict__ cudaTextureObject_t const int *__restrict__ const float3 lata
int getAllForceSize(const int atomStorageSize, const bool doSlow)
void bondedForce(const double scale14, const int atomStorageSize, const bool doEnergy, const bool doVirial, const bool doSlow, const float3 lata, const float3 latb, const float3 latc, const float cutoff2, const float r2_delta, const int r2_delta_expc, const float4 *h_xyzq, FORCE_TYPE *h_forces, double *h_energies, cudaStream_t stream)
static __constant__ const float float_to_force
#define WARPSIZE
Definition: CudaUtils.h:10
void setupAngleValues(int numAngleValues, CudaAngleValue *h_angleValues)
__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 latb
__thread cudaStream_t stream
void setupBondValues(int numBondValues, CudaBondValue *h_bondValues)
int getForceSize(const int atomStorageSize)
void setupImproperValues(int numImproperValues, CudaDihedralValue *h_improperValues)
void setupCrosstermValues(int numCrosstermValues, CudaCrosstermValue *h_crosstermValues)
#define FORCE_TYPE
ComputeBondedCUDAKernel(int deviceID, CudaNonbondedTables &cudaNonbondedTables)
__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 cutoff2
void setupDihedralValues(int numDihedralValues, CudaDihedralValue *h_dihedralValues)
__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 latc
void update(const int numBondsIn, const int numAnglesIn, const int numDihedralsIn, const int numImpropersIn, const int numModifiedExclusionsIn, const int numExclusionsIn, const int numCrosstermsIn, const char *h_tupleData, cudaStream_t stream)