NAMD
CudaComputeNonbondedKernel.h
Go to the documentation of this file.
1 #ifndef CUDACOMPUTENONBONDEDKERNEL_H
2 #define CUDACOMPUTENONBONDEDKERNEL_H
3 #include "CudaUtils.h"
4 #include "CudaTileListKernel.h"
5 #include "CudaNonbondedTables.h"
6 #include "HipDefines.h"
7 #if defined(NAMD_CUDA) || defined(NAMD_HIP)
8 
10 private:
11 
12  const int deviceID;
13  CudaNonbondedTables& cudaNonbondedTables;
14  const bool doStreaming;
15 
16  // Exclusions
17  unsigned int* overflowExclusions;
18  int overflowExclusionsSize;
19 
20  int2* exclIndexMaxDiff;
21  int exclIndexMaxDiffSize;
22 
23  // Atom indices
24  int* atomIndex;
25  int atomIndexSize;
26 
27  // VdW types
28  int* vdwTypes;
29  int vdwTypesSize;
30 
31  unsigned int* patchNumCount;
32  int patchNumCountSize;
33 
34  int* patchReadyQueue;
35  int patchReadyQueueSize;
36 
37  float *force_x, *force_y, *force_z, *force_w;
38  int forceSize;
39  float *forceSlow_x, *forceSlow_y, *forceSlow_z, *forceSlow_w;
40  int forceSlowSize;
41 public:
42  CudaComputeNonbondedKernel(int deviceID, CudaNonbondedTables& cudaNonbondedTables, bool doStreaming);
44 
45  void updateVdwTypesExcl(const int atomStorageSize, const int* h_vdwTypes,
46  const int2* h_exclIndexMaxDiff, const int* h_atomIndex, cudaStream_t stream);
47 
48  void nonbondedForce(CudaTileListKernel& tlKernel,
49  const int atomStorageSize, const bool doPairlist,
50  const bool doEnergy, const bool doVirial, const bool doSlow,
51  const float3 lata, const float3 latb, const float3 latc,
52  const float4* h_xyzq, const float cutoff2,
53  float4* d_forces, float4* d_forcesSlow,
54  float4* h_forces, float4* h_forcesSlow,
55  cudaStream_t stream);
56 
58  const int atomStorageSize, const bool doEnergy, const bool doVirial, const bool doSlow, const bool doGBIS,
59  float4* d_forces, float4* d_forcesSlow,
60  VirialEnergy* d_virialEnergy, cudaStream_t stream);
61 
62  void getVirialEnergy(VirialEnergy* h_virialEnergy, cudaStream_t stream);
63 
64  void bindExclusions(int numExclusions, unsigned int* exclusion_bits);
65 
66  int* getPatchReadyQueue();
67 
68  void reallocate_forceSOA(int atomStorageSize);
69 };
70 
71 #endif // NAMD_CUDA
72 #endif // CUDACOMPUTENONBONDEDKERNEL_H
void nonbondedForce(CudaTileListKernel &tlKernel, const int atomStorageSize, const bool doPairlist, const bool doEnergy, const bool doVirial, const bool doSlow, const float3 lata, const float3 latb, const float3 latc, const float4 *h_xyzq, const float cutoff2, float4 *d_forces, float4 *d_forcesSlow, float4 *h_forces, float4 *h_forcesSlow, 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)
__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
void reallocate_forceSOA(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 const float3 latb
__thread cudaStream_t stream
void bindExclusions(int numExclusions, unsigned int *exclusion_bits)
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 getVirialEnergy(VirialEnergy *h_virialEnergy, cudaStream_t stream)
__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
__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