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