NAMD
|
#include <math.h>
#include <cuda.h>
#include <namd_cub/cub.cuh>
#include "ComputeBondedCUDAKernel.h"
#include "DeviceCUDA.h"
Go to the source code of this file.
Macros | |
#define | BONDEDFORCESKERNEL_NUM_WARP 4 |
#define | NONBONDEDTABLES |
#define | CALL(DOENERGY, DOVIRIAL) |
#define | CALL(DOENERGY, DOVIRIAL, DOELECT) |
Functions | |
__device__ __forceinline__ float4 | sampleTableTex (cudaTextureObject_t tex, float k) |
template<typename T > | |
__forceinline__ __device__ void | convertForces (const float x, const float y, const float z, T &Tx, T &Ty, T &Tz) |
template<typename T > | |
__forceinline__ __device__ void | calcComponentForces (float fij, const float dx, const float dy, const float dz, T &fxij, T &fyij, T &fzij) |
template<typename T > | |
__forceinline__ __device__ void | calcComponentForces (float fij1, const float dx1, const float dy1, const float dz1, float fij2, const float dx2, const float dy2, const float dz2, T &fxij, T &fyij, T &fzij) |
__forceinline__ __device__ int | warpAggregatingAtomicInc (int *counter) |
template<typename T > | |
__forceinline__ __device__ void | storeForces (const T fx, const T fy, const T fz, const int ind, const int stride, T *force, T *forceList, int *forceListCounter, int *forceListStarts, int *forceListNexts) |
template<typename T , bool doEnergy, bool doVirial> | |
__device__ void | bondForce (const int index, const CudaBond *__restrict__ bondList, const CudaBondValue *__restrict__ bondValues, const float4 *__restrict__ xyzq, const int stride, const float3 lata, const float3 latb, const float3 latc, T *__restrict__ force, double &energy, T *__restrict__ forceList, int *forceListCounter, int *forceListStarts, int *forceListNexts, ComputeBondedCUDAKernel::BondedVirial< double > &virial) |
template<typename T , bool doEnergy, bool doVirial, bool doElect> | |
__device__ void | modifiedExclusionForce (const int index, const CudaExclusion *__restrict__ exclusionList, const bool doSlow, const float one_scale14, const int vdwCoefTableWidth, cudaTextureObject_t vdwCoefTableTex, cudaTextureObject_t forceTableTex, cudaTextureObject_t energyTableTex, const float4 *__restrict__ xyzq, const int stride, const float3 lata, const float3 latb, const float3 latc, const float cutoff2, double &energyVdw, T *__restrict__ forceNbond, double &energyNbond, T *__restrict__ forceSlow, double &energySlow, T *__restrict__ forceList, int *forceListCounter, int *forceListStartsNbond, int *forceListStartsSlow, int *forceListNexts, ComputeBondedCUDAKernel::BondedVirial< double > &virialNbond, ComputeBondedCUDAKernel::BondedVirial< double > &virialSlow) |
template<typename T , bool doEnergy, bool doVirial> | |
__device__ void | exclusionForce (const int index, const CudaExclusion *__restrict__ exclusionList, const float r2_delta, const int r2_delta_expc, cudaTextureObject_t r2_table_tex, cudaTextureObject_t exclusionTableTex, const float4 *__restrict__ xyzq, const int stride, const float3 lata, const float3 latb, const float3 latc, const float cutoff2, T *__restrict__ forceSlow, double &energySlow, T *__restrict__ forceList, int *forceListCounter, int *forceListStartsSlow, int *forceListNexts, ComputeBondedCUDAKernel::BondedVirial< double > &virialSlow) |
template<typename T , bool doEnergy, bool doVirial> | |
__device__ void | angleForce (const int index, const CudaAngle *__restrict__ angleList, const CudaAngleValue *__restrict__ angleValues, const float4 *__restrict__ xyzq, const int stride, const float3 lata, const float3 latb, const float3 latc, T *__restrict__ force, double &energy, T *__restrict__ forceList, int *forceListCounter, int *forceListStarts, int *forceListNexts, ComputeBondedCUDAKernel::BondedVirial< double > &virial) |
template<bool doEnergy> | |
__forceinline__ __device__ void | diheComp (const CudaDihedralValue *dihedralValues, int ic, const float sin_phi, const float cos_phi, const float scale, float &df, double &e) |
template<typename T , bool doEnergy, bool doVirial> | |
__device__ void | diheForce (const int index, const CudaDihedral *__restrict__ diheList, const CudaDihedralValue *__restrict__ dihedralValues, const float4 *__restrict__ xyzq, const int stride, const float3 lata, const float3 latb, const float3 latc, T *__restrict__ force, double &energy, T *__restrict__ forceList, int *forceListCounter, int *forceListStarts, int *forceListNexts, ComputeBondedCUDAKernel::BondedVirial< double > &virial) |
__device__ __forceinline__ float3 | cross (const float3 v1, const float3 v2) |
__device__ __forceinline__ float | dot (const float3 v1, const float3 v2) |
template<typename T , bool doEnergy, bool doVirial> | |
__device__ void | crosstermForce (const int index, const CudaCrossterm *__restrict__ crosstermList, const CudaCrosstermValue *__restrict__ crosstermValues, const float4 *__restrict__ xyzq, const int stride, const float3 lata, const float3 latb, const float3 latc, T *__restrict__ force, double &energy, T *__restrict__ forceList, int *forceListCounter, int *forceListStarts, int *forceListNexts, ComputeBondedCUDAKernel::BondedVirial< double > &virial) |
template<typename T , bool doEnergy, bool doVirial> | |
__global__ void | bondedForcesKernel (const int start, const int numBonds, const CudaBond *__restrict__ bonds, const CudaBondValue *__restrict__ bondValues, const int numAngles, const CudaAngle *__restrict__ angles, const CudaAngleValue *__restrict__ angleValues, const int numDihedrals, const CudaDihedral *__restrict__ dihedrals, const CudaDihedralValue *__restrict__ dihedralValues, const int numImpropers, const CudaDihedral *__restrict__ impropers, const CudaDihedralValue *__restrict__ improperValues, const int numExclusions, const CudaExclusion *__restrict__ exclusions, const int numCrossterms, const CudaCrossterm *__restrict__ crossterms, const CudaCrosstermValue *__restrict__ crosstermValues, const float cutoff2, const float r2_delta, const int r2_delta_expc, const float *__restrict__ r2_table, const float4 *__restrict__ exclusionTable, cudaTextureObject_t r2_table_tex, cudaTextureObject_t exclusionTableTex, const float4 *__restrict__ xyzq, const int stride, const float3 lata, const float3 latb, const float3 latc, T *__restrict__ force, T *__restrict__ forceSlow, T *__restrict__ forceList, int *forceListCounter, int *forceListStarts, int *forceListStartsSlow, int *forceListNexts, double *__restrict__ energies_virials) |
template<typename T , bool doEnergy, bool doVirial, bool doElect> | |
__global__ void | modifiedExclusionForcesKernel (const int start, const int numModifiedExclusions, const CudaExclusion *__restrict__ modifiedExclusions, const bool doSlow, const float one_scale14, const float cutoff2, const int vdwCoefTableWidth, const float2 *__restrict__ vdwCoefTable, cudaTextureObject_t vdwCoefTableTex, cudaTextureObject_t modifiedExclusionForceTableTex, cudaTextureObject_t modifiedExclusionEnergyTableTex, const float4 *__restrict__ xyzq, const int stride, const float3 lata, const float3 latb, const float3 latc, T *__restrict__ forceNbond, T *__restrict__ forceSlow, T *__restrict__ forceList, int *forceListCounter, int *forceListStartsNbond, int *forceListStartsSlow, int *forceListNexts, double *__restrict__ energies_virials) |
template<typename T > | |
__global__ void | gatherBondedForcesKernel (const int start, const int atomStorageSize, const int stride, const T *__restrict__ forceList, const int *__restrict__ forceListStarts, const int *__restrict__ forceListNexts, T *__restrict__ force) |
__global__ void | reduceBondedBinsKernel (double *energies_virials) |
Variables | |
__thread DeviceCUDA * | deviceCUDA |
#define BONDEDFORCESKERNEL_NUM_WARP 4 |
Definition at line 1188 of file ComputeBondedCUDAKernel.cu.
Referenced by ComputeBondedCUDAKernel::bondedForce(), bondedForcesKernel(), and modifiedExclusionForcesKernel().
#define CALL | ( | DOENERGY, | |
DOVIRIAL | |||
) |
Referenced by ComputeBondedCUDAKernel::bondedForce(), cuda_nonbonded_forces(), CudaComputeGBISKernel::GBISphase2(), and CudaComputeNonbondedKernel::nonbondedForce().
#define CALL | ( | DOENERGY, | |
DOVIRIAL, | |||
DOELECT | |||
) |
#define NONBONDEDTABLES |
__device__ void angleForce | ( | const int | index, |
const CudaAngle *__restrict__ | angleList, | ||
const CudaAngleValue *__restrict__ | angleValues, | ||
const float4 *__restrict__ | xyzq, | ||
const int | stride, | ||
const float3 | lata, | ||
const float3 | latb, | ||
const float3 | latc, | ||
T *__restrict__ | force, | ||
double & | energy, | ||
T *__restrict__ | forceList, | ||
int * | forceListCounter, | ||
int * | forceListStarts, | ||
int * | forceListNexts, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virial | ||
) |
Definition at line 588 of file ComputeBondedCUDAKernel.cu.
References force_to_float, CudaAngle::i, CudaAngle::ioffsetXYZ, CudaAngle::itype, CudaAngle::j, CudaAngle::k, CudaAngleValue::k, CudaAngleValue::k_ub, CudaAngle::koffsetXYZ, CudaAngleValue::normal, CudaAngleValue::r_ub, CudaAngle::scale, and CudaAngleValue::theta0.
__global__ void bondedForcesKernel | ( | const int | start, |
const int | numBonds, | ||
const CudaBond *__restrict__ | bonds, | ||
const CudaBondValue *__restrict__ | bondValues, | ||
const int | numAngles, | ||
const CudaAngle *__restrict__ | angles, | ||
const CudaAngleValue *__restrict__ | angleValues, | ||
const int | numDihedrals, | ||
const CudaDihedral *__restrict__ | dihedrals, | ||
const CudaDihedralValue *__restrict__ | dihedralValues, | ||
const int | numImpropers, | ||
const CudaDihedral *__restrict__ | impropers, | ||
const CudaDihedralValue *__restrict__ | improperValues, | ||
const int | numExclusions, | ||
const CudaExclusion *__restrict__ | exclusions, | ||
const int | numCrossterms, | ||
const CudaCrossterm *__restrict__ | crossterms, | ||
const CudaCrosstermValue *__restrict__ | crosstermValues, | ||
const float | cutoff2, | ||
const float | r2_delta, | ||
const int | r2_delta_expc, | ||
const float *__restrict__ | r2_table, | ||
const float4 *__restrict__ | exclusionTable, | ||
cudaTextureObject_t | r2_table_tex, | ||
cudaTextureObject_t | exclusionTableTex, | ||
const float4 *__restrict__ | xyzq, | ||
const int | stride, | ||
const float3 | lata, | ||
const float3 | latb, | ||
const float3 | latc, | ||
T *__restrict__ | force, | ||
T *__restrict__ | forceSlow, | ||
T *__restrict__ | forceList, | ||
int * | forceListCounter, | ||
int * | forceListStarts, | ||
int * | forceListStartsSlow, | ||
int * | forceListNexts, | ||
double *__restrict__ | energies_virials | ||
) |
Definition at line 1195 of file ComputeBondedCUDAKernel.cu.
References ComputeBondedCUDAKernel::amdDiheVirialIndex_XX, ATOMIC_BINS, BLOCK_SYNC, BONDEDFORCESKERNEL_NUM_WARP, cutoff2, ComputeBondedCUDAKernel::energies_virials_SIZE, ComputeBondedCUDAKernel::energyIndex_ANGLE, ComputeBondedCUDAKernel::energyIndex_BOND, ComputeBondedCUDAKernel::energyIndex_CROSSTERM, ComputeBondedCUDAKernel::energyIndex_DIHEDRAL, ComputeBondedCUDAKernel::energyIndex_ELECT_SLOW, ComputeBondedCUDAKernel::energyIndex_IMPROPER, exclusions, lata, latb, latc, ComputeBondedCUDAKernel::normalVirialIndex_XX, ComputeBondedCUDAKernel::slowVirialIndex_XX, WARP_FULL_MASK, WARP_SHUFFLE_XOR, WARPSIZE, ComputeBondedCUDAKernel::BondedVirial< T >::xx, ComputeBondedCUDAKernel::BondedVirial< T >::xy, xyzq, ComputeBondedCUDAKernel::BondedVirial< T >::xz, ComputeBondedCUDAKernel::BondedVirial< T >::yx, ComputeBondedCUDAKernel::BondedVirial< T >::yy, ComputeBondedCUDAKernel::BondedVirial< T >::yz, ComputeBondedCUDAKernel::BondedVirial< T >::zx, ComputeBondedCUDAKernel::BondedVirial< T >::zy, and ComputeBondedCUDAKernel::BondedVirial< T >::zz.
__device__ void bondForce | ( | const int | index, |
const CudaBond *__restrict__ | bondList, | ||
const CudaBondValue *__restrict__ | bondValues, | ||
const float4 *__restrict__ | xyzq, | ||
const int | stride, | ||
const float3 | lata, | ||
const float3 | latb, | ||
const float3 | latc, | ||
T *__restrict__ | force, | ||
double & | energy, | ||
T *__restrict__ | forceList, | ||
int * | forceListCounter, | ||
int * | forceListStarts, | ||
int * | forceListNexts, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virial | ||
) |
Definition at line 254 of file ComputeBondedCUDAKernel.cu.
References CudaBond::i, CudaBond::ioffsetXYZ, CudaBond::itype, CudaBond::j, CudaBondValue::k, CudaBond::scale, CudaBondValue::x0, and CudaBondValue::x1.
__forceinline__ __device__ void calcComponentForces | ( | float | fij, |
const float | dx, | ||
const float | dy, | ||
const float | dz, | ||
T & | fxij, | ||
T & | fyij, | ||
T & | fzij | ||
) |
Definition at line 99 of file ComputeBondedCUDAKernel.cu.
__forceinline__ __device__ void calcComponentForces | ( | float | fij1, |
const float | dx1, | ||
const float | dy1, | ||
const float | dz1, | ||
float | fij2, | ||
const float | dx2, | ||
const float | dy2, | ||
const float | dz2, | ||
T & | fxij, | ||
T & | fyij, | ||
T & | fzij | ||
) |
Definition at line 130 of file ComputeBondedCUDAKernel.cu.
__forceinline__ __device__ void convertForces | ( | const float | x, |
const float | y, | ||
const float | z, | ||
T & | Tx, | ||
T & | Ty, | ||
T & | Tz | ||
) |
Definition at line 77 of file ComputeBondedCUDAKernel.cu.
__device__ __forceinline__ float3 cross | ( | const float3 | v1, |
const float3 | v2 | ||
) |
Definition at line 910 of file ComputeBondedCUDAKernel.cu.
Referenced by Sequencer::addRotDragToPosition(), PmeKSpace::compute_energy(), CrosstermElem::computeForce(), DihedralElem::computeForce(), ImproperElem::computeForce(), crosstermForce(), ComputeConsTorque::doForce(), ComputeStir::doForce(), GridforceFullMainGrid::initialize(), ComputeMsmMgr::initialize(), proc_dihedralgrad(), proc_getdihedral(), Lattice::set(), settle1(), Sequencer::submitReductions(), and Lattice::volume().
__device__ void crosstermForce | ( | const int | index, |
const CudaCrossterm *__restrict__ | crosstermList, | ||
const CudaCrosstermValue *__restrict__ | crosstermValues, | ||
const float4 *__restrict__ | xyzq, | ||
const int | stride, | ||
const float3 | lata, | ||
const float3 | latb, | ||
const float3 | latc, | ||
T *__restrict__ | force, | ||
double & | energy, | ||
T *__restrict__ | forceList, | ||
int * | forceListCounter, | ||
int * | forceListStarts, | ||
int * | forceListNexts, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virial | ||
) |
Definition at line 926 of file ComputeBondedCUDAKernel.cu.
References A, B, CudaCrosstermValue::c, cross(), CudaCrosstermValue::dim, dot(), CudaCrossterm::i1, CudaCrossterm::i2, CudaCrossterm::i3, CudaCrossterm::i4, CudaCrossterm::i5, CudaCrossterm::i6, CudaCrossterm::i7, CudaCrossterm::i8, CudaCrossterm::itype, M_PI, CudaCrossterm::offset12XYZ, CudaCrossterm::offset23XYZ, CudaCrossterm::offset34XYZ, CudaCrossterm::offset56XYZ, CudaCrossterm::offset67XYZ, CudaCrossterm::offset78XYZ, and CudaCrossterm::scale.
__forceinline__ __device__ void diheComp | ( | const CudaDihedralValue * | dihedralValues, |
int | ic, | ||
const float | sin_phi, | ||
const float | cos_phi, | ||
const float | scale, | ||
float & | df, | ||
double & | e | ||
) |
Definition at line 728 of file ComputeBondedCUDAKernel.cu.
References CudaDihedralValue::delta, CudaDihedralValue::k, M_PI, and CudaDihedralValue::n.
__device__ void diheForce | ( | const int | index, |
const CudaDihedral *__restrict__ | diheList, | ||
const CudaDihedralValue *__restrict__ | dihedralValues, | ||
const float4 *__restrict__ | xyzq, | ||
const int | stride, | ||
const float3 | lata, | ||
const float3 | latb, | ||
const float3 | latc, | ||
T *__restrict__ | force, | ||
double & | energy, | ||
T *__restrict__ | forceList, | ||
int * | forceListCounter, | ||
int * | forceListStarts, | ||
int * | forceListNexts, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virial | ||
) |
Definition at line 772 of file ComputeBondedCUDAKernel.cu.
References force_to_float, CudaDihedral::i, CudaDihedral::ioffsetXYZ, CudaDihedral::itype, CudaDihedral::j, CudaDihedral::joffsetXYZ, CudaDihedral::k, CudaDihedral::l, CudaDihedral::loffsetXYZ, and CudaDihedral::scale.
__device__ __forceinline__ float dot | ( | const float3 | v1, |
const float3 | v2 | ||
) |
Definition at line 918 of file ComputeBondedCUDAKernel.cu.
Referenced by crosstermForce(), ARestraint::GetDihe(), and AmberParm7Reader::parse_fortran_format().
__device__ void exclusionForce | ( | const int | index, |
const CudaExclusion *__restrict__ | exclusionList, | ||
const float | r2_delta, | ||
const int | r2_delta_expc, | ||
cudaTextureObject_t | r2_table_tex, | ||
cudaTextureObject_t | exclusionTableTex, | ||
const float4 *__restrict__ | xyzq, | ||
const int | stride, | ||
const float3 | lata, | ||
const float3 | latb, | ||
const float3 | latc, | ||
const float | cutoff2, | ||
T *__restrict__ | forceSlow, | ||
double & | energySlow, | ||
T *__restrict__ | forceList, | ||
int * | forceListCounter, | ||
int * | forceListStartsSlow, | ||
int * | forceListNexts, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virialSlow | ||
) |
Definition at line 492 of file ComputeBondedCUDAKernel.cu.
References __ldg, CudaExclusion::i, CudaExclusion::ioffsetXYZ, and CudaExclusion::j.
__global__ void gatherBondedForcesKernel | ( | const int | start, |
const int | atomStorageSize, | ||
const int | stride, | ||
const T *__restrict__ | forceList, | ||
const int *__restrict__ | forceListStarts, | ||
const int *__restrict__ | forceListNexts, | ||
T *__restrict__ | force | ||
) |
__device__ void modifiedExclusionForce | ( | const int | index, |
const CudaExclusion *__restrict__ | exclusionList, | ||
const bool | doSlow, | ||
const float | one_scale14, | ||
const int | vdwCoefTableWidth, | ||
cudaTextureObject_t | vdwCoefTableTex, | ||
cudaTextureObject_t | forceTableTex, | ||
cudaTextureObject_t | energyTableTex, | ||
const float4 *__restrict__ | xyzq, | ||
const int | stride, | ||
const float3 | lata, | ||
const float3 | latb, | ||
const float3 | latc, | ||
const float | cutoff2, | ||
double & | energyVdw, | ||
T *__restrict__ | forceNbond, | ||
double & | energyNbond, | ||
T *__restrict__ | forceSlow, | ||
double & | energySlow, | ||
T *__restrict__ | forceList, | ||
int * | forceListCounter, | ||
int * | forceListStartsNbond, | ||
int * | forceListStartsSlow, | ||
int * | forceListNexts, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virialNbond, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virialSlow | ||
) |
Definition at line 352 of file ComputeBondedCUDAKernel.cu.
References __ldg, energyTableTex, forceTableTex, CudaExclusion::i, CudaExclusion::ioffsetXYZ, CudaExclusion::j, sampleTableTex(), vdwCoefTableTex, vdwCoefTableWidth, CudaExclusion::vdwtypei, CudaExclusion::vdwtypej, float2::x, ComputeBondedCUDAKernel::BondedVirial< T >::xx, ComputeBondedCUDAKernel::BondedVirial< T >::xy, ComputeBondedCUDAKernel::BondedVirial< T >::xz, float2::y, ComputeBondedCUDAKernel::BondedVirial< T >::yx, ComputeBondedCUDAKernel::BondedVirial< T >::yy, ComputeBondedCUDAKernel::BondedVirial< T >::yz, ComputeBondedCUDAKernel::BondedVirial< T >::zx, ComputeBondedCUDAKernel::BondedVirial< T >::zy, and ComputeBondedCUDAKernel::BondedVirial< T >::zz.
__global__ void modifiedExclusionForcesKernel | ( | const int | start, |
const int | numModifiedExclusions, | ||
const CudaExclusion *__restrict__ | modifiedExclusions, | ||
const bool | doSlow, | ||
const float | one_scale14, | ||
const float | cutoff2, | ||
const int | vdwCoefTableWidth, | ||
const float2 *__restrict__ | vdwCoefTable, | ||
cudaTextureObject_t | vdwCoefTableTex, | ||
cudaTextureObject_t | modifiedExclusionForceTableTex, | ||
cudaTextureObject_t | modifiedExclusionEnergyTableTex, | ||
const float4 *__restrict__ | xyzq, | ||
const int | stride, | ||
const float3 | lata, | ||
const float3 | latb, | ||
const float3 | latc, | ||
T *__restrict__ | forceNbond, | ||
T *__restrict__ | forceSlow, | ||
T *__restrict__ | forceList, | ||
int * | forceListCounter, | ||
int * | forceListStartsNbond, | ||
int * | forceListStartsSlow, | ||
int * | forceListNexts, | ||
double *__restrict__ | energies_virials | ||
) |
Definition at line 1504 of file ComputeBondedCUDAKernel.cu.
References ATOMIC_BINS, BLOCK_SYNC, BONDEDFORCESKERNEL_NUM_WARP, cutoff2, ComputeBondedCUDAKernel::energies_virials_SIZE, ComputeBondedCUDAKernel::energyIndex_ELECT, ComputeBondedCUDAKernel::energyIndex_ELECT_SLOW, ComputeBondedCUDAKernel::energyIndex_LJ, if(), lata, latb, latc, ComputeBondedCUDAKernel::nbondVirialIndex_XX, ComputeBondedCUDAKernel::nbondVirialIndex_XY, ComputeBondedCUDAKernel::nbondVirialIndex_XZ, ComputeBondedCUDAKernel::nbondVirialIndex_YX, ComputeBondedCUDAKernel::nbondVirialIndex_YY, ComputeBondedCUDAKernel::nbondVirialIndex_YZ, ComputeBondedCUDAKernel::nbondVirialIndex_ZX, ComputeBondedCUDAKernel::nbondVirialIndex_ZY, ComputeBondedCUDAKernel::nbondVirialIndex_ZZ, ComputeBondedCUDAKernel::slowVirialIndex_XX, ComputeBondedCUDAKernel::slowVirialIndex_XY, ComputeBondedCUDAKernel::slowVirialIndex_XZ, ComputeBondedCUDAKernel::slowVirialIndex_YX, ComputeBondedCUDAKernel::slowVirialIndex_YY, ComputeBondedCUDAKernel::slowVirialIndex_YZ, ComputeBondedCUDAKernel::slowVirialIndex_ZX, ComputeBondedCUDAKernel::slowVirialIndex_ZY, ComputeBondedCUDAKernel::slowVirialIndex_ZZ, vdwCoefTable, vdwCoefTableTex, vdwCoefTableWidth, WARP_FULL_MASK, WARP_SHUFFLE_XOR, WARPSIZE, ComputeBondedCUDAKernel::BondedVirial< T >::xx, ComputeBondedCUDAKernel::BondedVirial< T >::xy, xyzq, ComputeBondedCUDAKernel::BondedVirial< T >::xz, ComputeBondedCUDAKernel::BondedVirial< T >::yx, ComputeBondedCUDAKernel::BondedVirial< T >::yy, ComputeBondedCUDAKernel::BondedVirial< T >::yz, ComputeBondedCUDAKernel::BondedVirial< T >::zx, ComputeBondedCUDAKernel::BondedVirial< T >::zy, and ComputeBondedCUDAKernel::BondedVirial< T >::zz.
__global__ void reduceBondedBinsKernel | ( | double * | energies_virials | ) |
Definition at line 1807 of file ComputeBondedCUDAKernel.cu.
References ATOMIC_BINS, ComputeBondedCUDAKernel::energies_virials_SIZE, and tempStorage.
__device__ __forceinline__ float4 sampleTableTex | ( | cudaTextureObject_t | tex, |
float | k | ||
) |
Definition at line 30 of file ComputeBondedCUDAKernel.cu.
References FORCE_ENERGY_TABLE_SIZE, and x.
Referenced by calcForceEnergy(), and modifiedExclusionForce().
__forceinline__ __device__ void storeForces | ( | const T | fx, |
const T | fy, | ||
const T | fz, | ||
const int | ind, | ||
const int | stride, | ||
T * | force, | ||
T * | forceList, | ||
int * | forceListCounter, | ||
int * | forceListStarts, | ||
int * | forceListNexts | ||
) |
Definition at line 189 of file ComputeBondedCUDAKernel.cu.
References WARP_ALL, WARP_BALLOT, WARP_FULL_MASK, WARP_SHUFFLE, warpAggregatingAtomicInc(), and WARPSIZE.
__forceinline__ __device__ int warpAggregatingAtomicInc | ( | int * | counter | ) |
Definition at line 167 of file ComputeBondedCUDAKernel.cu.
References WARP_SHUFFLE, and WARPSIZE.
Referenced by storeForces().
__thread DeviceCUDA* deviceCUDA |
Definition at line 22 of file DeviceCUDA.C.
Referenced by ComputeNonbondedCUDA::assignPatches(), CudaComputeNonbonded::assignPatches(), ComputeBondedCUDAKernel::bondedForce(), build_cuda_exclusions(), build_cuda_force_table(), CudaTileListKernel::buildTileLists(), ComputePmeMgr::chargeGridSubmitted(), ComputeNonbondedCUDA::ComputeNonbondedCUDA(), ComputeMgr::createComputes(), ComputeCUDAMgr::createCudaComputeNonbonded(), cuda_check_local_calc(), cuda_check_remote_calc(), cuda_check_remote_progress(), cuda_initialize(), ComputePme::doWork(), ComputeNonbondedCUDA::doWork(), ComputeNonbondedCUDA::finishWork(), CudaComputeGBISKernel::GBISphase1(), CudaComputeGBISKernel::GBISphase2(), CudaComputeGBISKernel::GBISphase3(), ComputeCUDAMgr::getCudaComputeNonbonded(), ComputeCUDAMgr::initialize(), ComputePmeMgr::initialize(), ComputePmeMgr::initialize_computes(), ComputePmeMgr::initialize_pencils(), ComputePmeCUDAMgr::initialize_pencils(), ComputePmeCUDAMgr::isPmeDevice(), CudaComputeNonbondedKernel::nonbondedForce(), ComputeNonbondedCUDA::recvYieldDevice(), CudaComputeNonbondedKernel::reduceVirialEnergy(), ComputeNonbondedCUDA::requirePatch(), ComputePmeCUDAMgr::setupPencils(), ComputePmeMgr::ungridCalc(), and ComputeCUDAMgr::update().