NAMD
CudaComputeGBISKernel.h
Go to the documentation of this file.
1 #ifndef CUDACOMPUTEGBISKERNEL_H
2 #define CUDACOMPUTEGBISKERNEL_H
3 #if defined(NAMD_CUDA) || defined(NAMD_HIP)
4 
6 private:
7 
8  int deviceID;
9 
10  float* intRad0;
11  int intRad0Size;
12 
13  float* intRadS;
14  int intRadSSize;
15 
16  float* psiSum;
17  int psiSumSize;
18 
19  float* bornRad;
20  int bornRadSize;
21 
22  float* dEdaSum;
23  int dEdaSumSize;
24 
25  float* dHdrPrefix;
26  int dHdrPrefixSize;
27 
28 public:
29  CudaComputeGBISKernel(int deviceID);
31 
32  void updateIntRad(const int atomStorageSize, float* intRad0H, float* intRadSH,
33  cudaStream_t stream);
34 
35  void updateBornRad(const int atomStorageSize, float* bornRadH, cudaStream_t stream);
36 
37  void update_dHdrPrefix(const int atomStorageSize, float* dHdrPrefixH, cudaStream_t stream);
38 
39  void GBISphase1(CudaTileListKernel& tlKernel, const int atomStorageSize,
40  const float3 lata, const float3 latb, const float3 latc, const float a_cut, float* h_psiSum,
41  cudaStream_t stream);
42 
43  void GBISphase2(CudaTileListKernel& tlKernel, const int atomStorageSize,
44  const bool doEnergy, const bool doSlow,
45  const float3 lata, const float3 latb, const float3 latc,
46  const float r_cut, const float scaling, const float kappa, const float smoothDist,
47  const float epsilon_p, const float epsilon_s,
48  float4* d_forces,
49  float* h_dEdaSum, cudaStream_t stream);
50 
51  void GBISphase3(CudaTileListKernel& tlKernel, const int atomStorageSize,
52  const float3 lata, const float3 latb, const float3 latc, const float a_cut,
53  float4* d_forces,
54  cudaStream_t stream);
55 
56 };
57 
58 #endif // NAMD_CUDA
59 #endif //CUDACOMPUTEGBISKERNEL_H
void updateIntRad(const int atomStorageSize, float *intRad0H, float *intRadSH, cudaStream_t stream)
void GBISphase2(CudaTileListKernel &tlKernel, const int atomStorageSize, const bool doEnergy, const bool doSlow, const float3 lata, const float3 latb, const float3 latc, const float r_cut, const float scaling, const float kappa, const float smoothDist, const float epsilon_p, const float epsilon_s, float4 *d_forces, float *h_dEdaSum, cudaStream_t stream)
void GBISphase3(CudaTileListKernel &tlKernel, const int atomStorageSize, const float3 lata, const float3 latb, const float3 latc, const float a_cut, float4 *d_forces, 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
static __thread float * bornRadH
static __thread float * dHdrPrefixH
void updateBornRad(const int atomStorageSize, float *bornRadH, 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 latb
__thread cudaStream_t stream
static __thread float * intRadSH
void update_dHdrPrefix(const int atomStorageSize, float *dHdrPrefixH, 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 latc
void GBISphase1(CudaTileListKernel &tlKernel, const int atomStorageSize, const float3 lata, const float3 latb, const float3 latc, const float a_cut, float *h_psiSum, cudaStream_t stream)
static __thread float * intRad0H