1 #ifndef COMPUTEBONDEDCUDAKERNEL_H
2 #define COMPUTEBONDEDCUDAKERNEL_H
7 #if defined(NAMD_CUDA) || defined(NAMD_HIP)
12 #define FORCE_TYPE long long int
14 #define FORCE_TYPE double
16 #define USE_STRIDED_FORCE
18 #ifndef USE_STRIDED_FORCE
19 #error "Non-USE_STRIDED_FORCE not implemented"
25 #define VIRIAL_TYPE long long int
27 #define VIRIAL_TYPE double
30 #define WRITE_FULL_VIRIALS
33 #define USE_BONDED_FORCE_ATOMIC_STORE
38 static __constant__
const float float_to_force = (float)(1ll << 40);
39 static __constant__
const float force_to_float = (float)1.0/(
float)(1ll << 40);
40 static __constant__
const double force_to_double = (double)1.0/(
double)(1ll << 40);
43 static __constant__
const float force_to_float = 1.0f;
44 static __constant__
const double force_to_double = 1.0;
48 static __constant__
const float float_to_virial = (float)(1ll << 30);
49 static __constant__
const double double_to_virial = (double)(1ll << 30);
50 static __constant__
const double virial_to_double = (double)1.0/(
double)(1ll << 30);
51 static __constant__
const long long int CONVERT_TO_VIR = (1ll << 10);
76 #ifdef WRITE_FULL_VIRIALS
87 #error "non-WRITE_FULL_VIRIALS not implemented yet"
89 double sforce_dp[27][3];
90 long long int sforce_fp[27][3];
120 int numModifiedExclusions;
135 int* forceListCounter;
136 int* forceListStarts;
139 int forceListStartsSize;
140 int forceListNextsSize;
154 double* energies_virials;
164 const int numBondsIn,
165 const int numAnglesIn,
166 const int numDihedralsIn,
167 const int numImpropersIn,
168 const int numModifiedExclusionsIn,
169 const int numExclusionsIn,
170 const int numCrosstermsIn,
171 const char* h_tupleData,
185 const double scale14,
const int atomStorageSize,
186 const bool doEnergy,
const bool doVirial,
const bool doSlow,
187 const float3
lata,
const float3
latb,
const float3
latc,
188 const float cutoff2,
const float r2_delta,
const int r2_delta_expc,
197 #endif // COMPUTEBONDEDCUDAKERNEL_H
static __constant__ const double force_to_double
~ComputeBondedCUDAKernel()
static int warpAlign(const int n)
static __constant__ const float force_to_float
int getForceStride(const 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 lata
int getAllForceSize(const int atomStorageSize, const bool doSlow)
void bondedForce(const double scale14, const int atomStorageSize, const bool doEnergy, const bool doVirial, const bool doSlow, const float3 lata, const float3 latb, const float3 latc, const float cutoff2, const float r2_delta, const int r2_delta_expc, const float4 *h_xyzq, FORCE_TYPE *h_forces, double *h_energies, cudaStream_t stream)
static __constant__ const float float_to_force
void setupAngleValues(int numAngleValues, CudaAngleValue *h_angleValues)
__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 setupBondValues(int numBondValues, CudaBondValue *h_bondValues)
int getForceSize(const int atomStorageSize)
void setupImproperValues(int numImproperValues, CudaDihedralValue *h_improperValues)
void setupCrosstermValues(int numCrosstermValues, CudaCrosstermValue *h_crosstermValues)
ComputeBondedCUDAKernel(int deviceID, CudaNonbondedTables &cudaNonbondedTables)
__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 setupDihedralValues(int numDihedralValues, CudaDihedralValue *h_dihedralValues)
__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 update(const int numBondsIn, const int numAnglesIn, const int numDihedralsIn, const int numImpropersIn, const int numModifiedExclusionsIn, const int numExclusionsIn, const int numCrosstermsIn, const char *h_tupleData, cudaStream_t stream)