NAMD
CudaComputeNonbondedKernel.h
Go to the documentation of this file.
1 #ifndef CUDACOMPUTENONBONDEDKERNEL_H
2 #define CUDACOMPUTENONBONDEDKERNEL_H
3 #include "CudaUtils.h"
4 #include "CudaRecord.h"
5 #include "CudaTileListKernel.h"
6 #include "CudaNonbondedTables.h"
7 #ifdef NAMD_CUDA
8 
9 
24 struct AlchData{
25  float scaling;
26  float switchdist2;
27  float cutoff2;
28  float switchfactor;
30  // float alchLambda;
31  float lambdaUp;
32  float lambdaDown;
33  float elecLambdaUp;
35  float vdwLambdaUp;
37 
38  float lambda2Up;
39  float lambda2Down;
42  float vdwLambda2Up;
44 
45  float vdwShiftUp;
46  float vdwShift2Up;
47  float vdwShiftDown;
50 };
51 
53 private:
54 
55  const int deviceID;
56  CudaNonbondedTables& cudaNonbondedTables;
57  const bool doStreaming;
58 
59  // Exclusions
60  int2 *d_exclusionsByAtom;
61  unsigned int* overflowExclusions;
62  size_t overflowExclusionsSize;
63 
64  int2* exclIndexMaxDiff;
65  size_t exclIndexMaxDiffSize;
66 
67  // Atom indices
68  int* atomIndex;
69  size_t atomIndexSize;
70 
71  // VdW types
72  int* vdwTypes;
73  size_t vdwTypesSize;
74 
75  unsigned int* patchNumCount;
76  size_t patchNumCountSize;
77 
78  int* patchReadyQueue;
79  size_t patchReadyQueueSize;
80 
81  float *force_x, *force_y, *force_z, *force_w;
82  size_t forceSize;
83  float *forceSlow_x, *forceSlow_y, *forceSlow_z, *forceSlow_w;
84  size_t forceSlowSize;
85 public:
86  CudaComputeNonbondedKernel(int deviceID, CudaNonbondedTables& cudaNonbondedTables, bool doStreaming);
88 
89  static __device__ __host__ __forceinline__ int
90  computeNumTiles(const int numAtoms, const int tilesize = WARPSIZE) {
91  return (numAtoms+tilesize-1)/tilesize;
92  }
93 
94  static __device__ __host__ __forceinline__ int
95  computeAtomPad(const int numAtoms, const int tilesize = WARPSIZE) {
96  return computeNumTiles(numAtoms, tilesize)*tilesize;
97  }
98 
99  void updateVdwTypesExcl(const int atomStorageSize, const int* h_vdwTypes,
100  const int2* h_exclIndexMaxDiff, const int* h_atomIndex, cudaStream_t stream);
101 
103  const int numPatches, const int atomStorageSize, const bool alchOn,
104  CudaLocalRecord* localRecords,
105  const int* d_vdwTypes, const int* d_id, const int* d_sortOrder,
106  const int* d_partition, cudaStream_t stream);
107 
108  void nonbondedForce(CudaTileListKernel& tlKernel,
109  const int atomStorageSize, const bool atomsChanged, const bool doMinimize,
110  const bool doPairlist, const bool doEnergy, const bool doVirial,
111  const bool doSlow, const bool doAlch, const bool doAlchVdwForceSwitching,
112  const bool doFEP, const bool doTI, const bool doTable,
113  const float3 lata, const float3 latb, const float3 latc,
114  const float4* h_xyzq, const float cutoff2,
115  const CudaNBConstants nbConstants,
116  float4* d_forces, float4* d_forcesSlow,
117  float4* h_forces, float4* h_forcesSlow, AlchData *fepFlags,
118  bool lambdaWindowUpdated,
119  char *part, bool CUDASOAintegratorOn, bool useDeviceMigration,
120  cudaStream_t stream);
121 
122  void reduceVirialEnergy(CudaTileListKernel& tlKernel,
123  const int atomStorageSize, const bool doEnergy, const bool doVirial, const bool doSlow, const bool doGBIS,
124  float4* d_forces, float4* d_forcesSlow,
125  VirialEnergy* d_virialEnergy, cudaStream_t stream);
126 
127  void getVirialEnergy(VirialEnergy* h_virialEnergy, cudaStream_t stream);
128 
129  void bindExclusions(int numExclusions, unsigned int* exclusion_bits);
130 
131  int* getPatchReadyQueue();
132 
133  void reallocate_forceSOA(int atomStorageSize);
134 
135  void setExclusionsByAtom(int2* h_data, const int num_atoms);
136 };
137 
138 #endif // NAMD_CUDA
139 #endif // CUDACOMPUTENONBONDEDKERNEL_H
Alchemical datastructure that holds the lambda-relevant paramenters for FEP/TI.
void nonbondedForce(CudaTileListKernel &tlKernel, const int atomStorageSize, const bool atomsChanged, const bool doMinimize, const bool doPairlist, const bool doEnergy, const bool doVirial, const bool doSlow, const bool doAlch, const bool doAlchVdwForceSwitching, const bool doFEP, const bool doTI, const bool doTable, const float3 lata, const float3 latb, const float3 latc, const float4 *h_xyzq, const float cutoff2, const CudaNBConstants nbConstants, float4 *d_forces, float4 *d_forcesSlow, float4 *h_forces, float4 *h_forcesSlow, AlchData *fepFlags, bool lambdaWindowUpdated, char *part, bool CUDASOAintegratorOn, bool useDeviceMigration, cudaStream_t stream)
CudaComputeNonbondedKernel(int deviceID, CudaNonbondedTables &cudaNonbondedTables, bool doStreaming)
void updateVdwTypesExcl(const int atomStorageSize, const int *h_vdwTypes, const int2 *h_exclIndexMaxDiff, const int *h_atomIndex, cudaStream_t stream)
#define WARPSIZE
Definition: CudaUtils.h:17
void reallocate_forceSOA(int atomStorageSize)
void bindExclusions(int numExclusions, unsigned int *exclusion_bits)
static __device__ __host__ __forceinline__ int computeAtomPad(const int numAtoms, const int tilesize=WARPSIZE)
void reduceVirialEnergy(CudaTileListKernel &tlKernel, const int atomStorageSize, const bool doEnergy, const bool doVirial, const bool doSlow, const bool doGBIS, float4 *d_forces, float4 *d_forcesSlow, VirialEnergy *d_virialEnergy, cudaStream_t stream)
void setExclusionsByAtom(int2 *h_data, const int num_atoms)
void updateVdwTypesExclOnGPU(CudaTileListKernel &tlKernel, const int numPatches, const int atomStorageSize, const bool alchOn, CudaLocalRecord *localRecords, const int *d_vdwTypes, const int *d_id, const int *d_sortOrder, const int *d_partition, cudaStream_t stream)
void getVirialEnergy(VirialEnergy *h_virialEnergy, cudaStream_t stream)
static __device__ __host__ __forceinline__ int computeNumTiles(const int numAtoms, const int tilesize=WARPSIZE)