NAMD
CudaComputeNonbondedKernel.hip.h
Go to the documentation of this file.
1 #ifndef CUDACOMPUTENONBONDEDKERNEL_HIP_H
2 #define CUDACOMPUTENONBONDEDKERNEL_HIP_H
3 #include "CudaUtils.h"
4 #include "CudaRecord.h"
6 #include "CudaNonbondedTables.h"
7 #include "HipDefines.h"
8 #if defined(NAMD_HIP)
9 
10 
19 struct AlchData{
20  float scaling;
21  float switchdist2;
22  float cutoff2;
23  float switchfactor;
24  float alchVdwShiftCoeff;
25  float alchLambda;
26  float lambdaUp;
27  float lambdaDown;
28  float elecLambdaUp;
29  float elecLambdaDown;
30  float vdwLambdaUp;
31  float vdwLambdaDown;
32 
33  float lambda2Up;
34  float lambda2Down;
35  float elecLambda2Up;
36  float elecLambda2Down;
37  float vdwLambda2Up;
38  float vdwLambda2Down;
39 
40  float vdwShiftUp;
41  float vdwShift2Up;
42  float vdwShiftDown;
43  float vdwShift2Down;
44  bool alchDecouple;
45 };
46 
48 private:
49 
50  const int deviceID;
51  CudaNonbondedTables& cudaNonbondedTables;
52  const bool doStreaming;
53 
54  // Exclusions
55  // JM: Seems like nvcc makes this point to NULL by default but clang doesn't
56  // we rely on it being set to NULL so we can initialize it on setExclusiveByAtom() laters
57  int2 *d_exclusionsByAtom = NULL;
58  unsigned int* overflowExclusions;
59  size_t overflowExclusionsSize;
60 
61  int2* exclIndexMaxDiff;
62  size_t exclIndexMaxDiffSize;
63 
64  // Atom indices
65  int* atomIndex;
66  size_t atomIndexSize;
67 
68  // VdW types
69  int* vdwTypes;
70  size_t vdwTypesSize;
71 
72  unsigned int* patchNumCount;
73  size_t patchNumCountSize;
74 
75  int* patchReadyQueue;
76  size_t patchReadyQueueSize;
77 
78  float *force_x, *force_y, *force_z, *force_w;
79  size_t forceSize;
80  float *forceSlow_x, *forceSlow_y, *forceSlow_z, *forceSlow_w;
81  size_t forceSlowSize;
82 public:
83  CudaComputeNonbondedKernel(int deviceID, CudaNonbondedTables& cudaNonbondedTables, bool doStreaming);
85 
86  static __device__ __host__ __forceinline__ int
87  computeNumTiles(const int numAtoms, const int tilesize = BOUNDINGBOXSIZE) {
88  return (numAtoms+tilesize-1)/tilesize;
89  }
90 
91  static __device__ __host__ __forceinline__ int
92  computeAtomPad(const int numAtoms, const int tilesize = BOUNDINGBOXSIZE) {
93  return computeNumTiles(numAtoms, tilesize)*tilesize;
94  }
95 
96  void updateVdwTypesExcl(const int atomStorageSize, const int* h_vdwTypes,
97  const int2* h_exclIndexMaxDiff, const int* h_atomIndex, cudaStream_t stream);
98 
100  const int numPatches, const int atomStorageSize, const bool alchOn,
101  CudaLocalRecord* localRecords,
102  const int* d_vdwTypes, const int* d_id, const int* d_sortOrder,
103  const int* d_partition, cudaStream_t stream);
104 
105  void nonbondedForce(CudaTileListKernel& tlKernel,
106  const int atomStorageSize, const bool atomsChanged,
107  const bool doMinimize, const bool doPairlist,
108  const bool doEnergy, const bool doVirial, const bool doSlow, const bool doAlch,
109  const bool doAlchVdwForceSwitching,
110  const bool doFEP, const bool doTI, const bool doTable,
111  const float3 lata, const float3 latb, const float3 latc,
112  const float4* h_xyzq, const float cutoff2,
113  const CudaNBConstants nbConstants,
114  float4* d_forces, float4* d_forcesSlow,
115  float4* h_forces, float4* h_forcesSlow, AlchData *srcFlags,
116  bool lambdaWindowUpdated, char *part,
117  bool CUDASOAintegratorOn, bool useDeviceMigration,
118  cudaStream_t stream);
119 
120  void reduceVirialEnergy(CudaTileListKernel& tlKernel,
121  const int atomStorageSize, const bool doEnergy, const bool doVirial, const bool doSlow, const bool doGBIS,
122  float4* d_forces, float4* d_forcesSlow,
123  VirialEnergy* d_virialEnergy, cudaStream_t stream);
124 
125  void getVirialEnergy(VirialEnergy* h_virialEnergy, cudaStream_t stream);
126 
127  void bindExclusions(int numExclusions, unsigned int* exclusion_bits);
128 
129  int* getPatchReadyQueue();
130 
131  void reallocate_forceSOA(int atomStorageSize);
132 
133  void setExclusionsByAtom(int2* h_data, const int num_atoms);
134 
135 };
136 
137 #endif // NAMD_HIP
138 #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)
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)
#define BOUNDINGBOXSIZE
Definition: CudaUtils.h:18
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)