1 #if defined(NAMD_CUDA) || defined(NAMD_HIP)
39 #define PATCH_PAIR_SIZE (sizeof(patch_pair)/4)
53 #define COPY_ATOM( DEST, SOURCE ) { \
54 DEST.position.x = SOURCE.position.x; \
55 DEST.position.y = SOURCE.position.y; \
56 DEST.position.z = SOURCE.position.z; \
57 DEST.charge = SOURCE.charge; \
60 #define COPY_PARAM( DEST, SOURCE ) { \
61 DEST.sqrt_epsilon = SOURCE.sqrt_epsilon; \
62 DEST.half_sigma = SOURCE.half_sigma; \
63 DEST.index = SOURCE.index; \
64 DEST.excl_index = SOURCE.excl_index; \
65 DEST.excl_maxdiff = SOURCE.excl_maxdiff; \
68 #define COPY_ATOM_TO_SHARED( ATOM, PARAM, SHARED ) { \
69 COPY_ATOM( SHARED, ATOM ) \
70 COPY_PARAM( SHARED, PARAM ) \
73 #define COPY_ATOM_FROM_SHARED( ATOM, PARAM, SHARED ) { \
74 COPY_ATOM( ATOM, SHARED ) \
75 COPY_PARAM( PARAM, SHARED ) \
80 #define MAX_EXCLUSIONS (1<<27)
81 #define MAX_CONST_EXCLUSIONS 2048 // cache size is 8k
91 #define FORCE_TABLE_SIZE 4096
98 int npatches,
int natoms,
int nexclmask,
int plist_len);
111 int cbegin,
int ccount,
int ctotal,
112 int doSlow,
int doEnergy,
int usePairlists,
int savePairlists,
113 int doStreaming,
int saveOrder, cudaStream_t &strm);
void cuda_bind_force_table(const float4 *t, const float4 *et)
void cuda_bind_forces(float4 *f, float4 *f_slow)
int cuda_stream_finished()
void cuda_bind_exclusions(const unsigned int *t, int n)
void cuda_bind_atoms(const atom *a)
__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
static __thread int plist_size
void cuda_bind_GBIS_energy(float *e)
void cuda_bind_GBIS_dEdaSum(GBReal *dEdaSumH)
__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
__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 cudaTextureObject_t cudaTextureObject_t float plcutoff2
void cuda_nonbonded_forces(float3 lata, float3 latb, float3 latc, float cutoff2, float plcutoff2, int cbegin, int ccount, int ctotal, int doSlow, int doEnergy, int usePairlists, int savePairlists, int doStreaming, int saveOrder, cudaStream_t &strm)
static __thread float * intRadSH
void cuda_bind_patch_pairs(patch_pair *h_patch_pairs, int npatch_pairs, int npatches, int natoms, int plist_len, int nexclmask)
void cuda_bind_vdw_types(const int *t)
void cuda_bind_lj_table(const float2 *t, int _lj_table_size)
void cuda_bind_GBIS_dHdrPrefix(float *dHdrPrefixH)
void cuda_bind_GBIS_bornRad(float *bornRadH)
void cuda_bind_virials(float *v, int *queue, int *blockorder)
void cuda_bind_GBIS_psiSum(GBReal *psiSumH)
void cuda_GBIS_P3(int cbegin, int ccount, int pbegin, int pcount, float a_cut, float rho_0, float scaling, float3 lata, float3 latb, float3 latc, cudaStream_t &strm)
void cuda_GBIS_P1(int cbegin, int ccount, int pbegin, int pcount, float a_cut, float rho_0, float3 lata, float3 latb, float3 latc, cudaStream_t &strm)
__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
void cuda_errcheck(const char *msg)
void cuda_bind_GBIS_intRad(float *intRad0H, float *intRadSH)
static __thread float * energy_gbis
struct __align__(16) patch_pair
k< npairi;++k){TABENERGY(const int numtypes=simParams->tableNumTypes;const float table_spacing=simParams->tableSpacing;const int npertype=(int)(namdnearbyint(simParams->tableMaxDist/simParams->tableSpacing)+1);) int table_i=(r2iilist[2 *k] >> 14)+r2_delta_expc;const int j=pairlisti[k];#define p_j BigReal diffa=r2list[k]-r2_table[table_i];#define table_four_i TABENERGY(register const int tabtype=-1-(lj_pars->A< 0?lj_pars->A:0);) BigReal kqq=kq_i *p_j-> charge
void cuda_GBIS_P2(int cbegin, int ccount, int pbegin, int pcount, float a_cut, float r_cut, float scaling, float kappa, float smoothDist, float epsilon_p, float epsilon_s, float3 lata, float3 latb, float3 latc, int doEnergy, int doFullElec, cudaStream_t &strm)
__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 cuda_bind_atom_params(const atom_param *t)
static __thread float * intRad0H