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