1 #ifndef CUDACOMPUTENONBONDEDKERNEL_H
2 #define CUDACOMPUTENONBONDEDKERNEL_H
7 #if defined(NAMD_CUDA) || defined(NAMD_HIP)
14 const bool doStreaming;
17 unsigned int* overflowExclusions;
18 int overflowExclusionsSize;
20 int2* exclIndexMaxDiff;
21 int exclIndexMaxDiffSize;
31 unsigned int* patchNumCount;
32 int patchNumCountSize;
35 int patchReadyQueueSize;
37 float *force_x, *force_y, *force_z, *force_w;
39 float *forceSlow_x, *forceSlow_y, *forceSlow_z, *forceSlow_w;
46 const int2* h_exclIndexMaxDiff,
const int* h_atomIndex, cudaStream_t
stream);
49 const int atomStorageSize,
const bool doPairlist,
50 const bool doEnergy,
const bool doVirial,
const bool doSlow,
52 const float4* h_xyzq,
const float cutoff2,
53 float4* d_forces, float4* d_forcesSlow,
54 float4* h_forces, float4* h_forcesSlow,
58 const int atomStorageSize,
const bool doEnergy,
const bool doVirial,
const bool doSlow,
const bool doGBIS,
59 float4* d_forces, float4* d_forcesSlow,
64 void bindExclusions(
int numExclusions,
unsigned int* exclusion_bits);
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
~CudaComputeNonbondedKernel()
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)
int * getPatchReadyQueue()
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