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 | CALL(DOENERGY, DOVIRIAL) |
#define | CALL(DOENERGY, DOVIRIAL, DOELECT) |
Functions | |
__device__ long long int | lliroundf (float f) |
__device__ unsigned long long int | llitoulli (long long int l) |
template<typename T > | |
__forceinline__ __device__ void | convertForces (const float x, const float y, const float z, T &Tx, T &Ty, T &Tz) |
template<> | |
__forceinline__ __device__ void | convertForces< long long int > (const float x, const float y, const float z, long long int &Tx, long long int &Ty, long long int &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<> | |
__forceinline__ __device__ void | calcComponentForces< long long int > (float fij, const float dx, const float dy, const float dz, long long int &fxij, long long int &fyij, long long int &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) |
template<> | |
__forceinline__ __device__ void | calcComponentForces< long long int > (float fij1, const float dx1, const float dy1, const float dz1, float fij2, const float dx2, const float dy2, const float dz2, long long int &fxij, long long int &fyij, long long int &fzij) |
template<typename T > | |
__forceinline__ __device__ void | storeForces (const T fx, const T fy, const T fz, const int ind, const int stride, T *force) |
template<> | |
__forceinline__ __device__ void | storeForces< long long int > (const long long int fx, const long long int fy, const long long int fz, const int ind, const int stride, long long int *force) |
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, 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, 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, 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, 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, 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, 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, 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, double *__restrict__ energies_virials) |
Variables | |
__thread DeviceCUDA * | deviceCUDA |
#define BONDEDFORCESKERNEL_NUM_WARP 4 |
Definition at line 1032 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 | |||
) |
__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, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virial | ||
) |
Definition at line 440 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, | ||
double *__restrict__ | energies_virials | ||
) |
Definition at line 1037 of file ComputeBondedCUDAKernel.cu.
References ComputeBondedCUDAKernel::amdDiheVirialIndex_XX, BLOCK_SYNC, BONDEDFORCESKERNEL_NUM_WARP, cutoff2, double_to_virial, ComputeBondedCUDAKernel::energyIndex_ANGLE, ComputeBondedCUDAKernel::energyIndex_BOND, ComputeBondedCUDAKernel::energyIndex_CROSSTERM, ComputeBondedCUDAKernel::energyIndex_DIHEDRAL, ComputeBondedCUDAKernel::energyIndex_ELECT_SLOW, ComputeBondedCUDAKernel::energyIndex_IMPROPER, exclusions, lata, latb, latc, llitoulli(), 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, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virial | ||
) |
Definition at line 145 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 57 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 86 of file ComputeBondedCUDAKernel.cu.
__forceinline__ __device__ void calcComponentForces< long long int > | ( | float | fij, |
const float | dx, | ||
const float | dy, | ||
const float | dz, | ||
long long int & | fxij, | ||
long long int & | fyij, | ||
long long int & | fzij | ||
) |
Definition at line 70 of file ComputeBondedCUDAKernel.cu.
References float_to_force, and lliroundf().
__forceinline__ __device__ void calcComponentForces< long long int > | ( | float | fij1, |
const float | dx1, | ||
const float | dy1, | ||
const float | dz1, | ||
float | fij2, | ||
const float | dx2, | ||
const float | dy2, | ||
const float | dz2, | ||
long long int & | fxij, | ||
long long int & | fyij, | ||
long long int & | fzij | ||
) |
Definition at line 100 of file ComputeBondedCUDAKernel.cu.
References float_to_force, and lliroundf().
__forceinline__ __device__ void convertForces | ( | const float | x, |
const float | y, | ||
const float | z, | ||
T & | Tx, | ||
T & | Ty, | ||
T & | Tz | ||
) |
Definition at line 37 of file ComputeBondedCUDAKernel.cu.
__forceinline__ __device__ void convertForces< long long int > | ( | const float | x, |
const float | y, | ||
const float | z, | ||
long long int & | Tx, | ||
long long int & | Ty, | ||
long long int & | Tz | ||
) |
Definition at line 47 of file ComputeBondedCUDAKernel.cu.
References float_to_force, lliroundf(), x, y, and z.
__device__ __forceinline__ float3 cross | ( | const float3 | v1, |
const float3 | v2 | ||
) |
Definition at line 759 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, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virial | ||
) |
Definition at line 775 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 578 of file ComputeBondedCUDAKernel.cu.
References CudaDihedralValue::delta, CudaDihedralValue::k, M_PI, CudaDihedralValue::n, and x.
__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, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virial | ||
) |
Definition at line 622 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 767 of file ComputeBondedCUDAKernel.cu.
Referenced by crosstermForce(), and ARestraint::GetDihe().
__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, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virialSlow | ||
) |
Definition at line 350 of file ComputeBondedCUDAKernel.cu.
References __ldg, CudaExclusion::i, CudaExclusion::ioffsetXYZ, and CudaExclusion::j.
|
inline |
Definition at line 21 of file ComputeBondedCUDAKernel.cu.
Referenced by calcComponentForces< long long int >(), and convertForces< long long int >().
|
inline |
Definition at line 28 of file ComputeBondedCUDAKernel.cu.
Referenced by bondedForcesKernel(), modifiedExclusionForcesKernel(), and storeForces< long long int >().
__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, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virialNbond, | ||
ComputeBondedCUDAKernel::BondedVirial< double > & | virialSlow | ||
) |
Definition at line 220 of file ComputeBondedCUDAKernel.cu.
References __ldg, energyTableTex, forceTableTex, CudaExclusion::i, CudaExclusion::ioffsetXYZ, CudaExclusion::j, 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, | ||
double *__restrict__ | energies_virials | ||
) |
Definition at line 1324 of file ComputeBondedCUDAKernel.cu.
References BLOCK_SYNC, BONDEDFORCESKERNEL_NUM_WARP, cutoff2, double_to_virial, ComputeBondedCUDAKernel::energyIndex_ELECT, ComputeBondedCUDAKernel::energyIndex_ELECT_SLOW, ComputeBondedCUDAKernel::energyIndex_LJ, if(), lata, latb, latc, llitoulli(), 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.
__forceinline__ __device__ void storeForces | ( | const T | fx, |
const T | fy, | ||
const T | fz, | ||
const int | ind, | ||
const int | stride, | ||
T * | force | ||
) |
Definition at line 122 of file ComputeBondedCUDAKernel.cu.
__forceinline__ __device__ void storeForces< long long int > | ( | const long long int | fx, |
const long long int | fy, | ||
const long long int | fz, | ||
const int | ind, | ||
const int | stride, | ||
long long int * | force | ||
) |
Definition at line 131 of file ComputeBondedCUDAKernel.cu.
References llitoulli().
__thread DeviceCUDA* deviceCUDA |
Definition at line 18 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().