NAMD
|
#include <cuda.h>
#include <namd_cub/cub.cuh>
#include "CudaComputeNonbondedKernel.h"
#include "CudaTileListKernel.h"
#include "DeviceCUDA.h"
Go to the source code of this file.
Macros | |
#define | NONBONDKERNEL_NUM_WARP 4 |
#define | REDUCENONBONDEDVIRIALKERNEL_NUM_WARP 32 |
#define | REDUCEVIRIALENERGYKERNEL_NUM_WARP 32 |
#define | REDUCEGBISENERGYKERNEL_NUM_WARP 32 |
#define | OVERALLOC 1.2f |
#define | MAX_CONST_EXCLUSIONS 2048 |
#define | LARGE_FLOAT (float)(1.0e10) |
#define | TABLE_PARAMS cudaNonbondedTables.getForceTableTex(), cudaNonbondedTables.getEnergyTableTex() |
#define | CALL(DOENERGY, DOVIRIAL, DOSLOW, DOPAIRLIST, DOSTREAMING) |
Functions | |
void | NAMD_die (const char *) |
__device__ __forceinline__ float4 | sampleTableTex (cudaTextureObject_t tex, float k) |
template<bool doEnergy, bool doSlow> | |
__device__ __forceinline__ void | calcForceEnergy (const float r2, const float qi, const float qj, const float dx, const float dy, const float dz, const int vdwtypei, const int vdwtypej, const float2 *__restrict__ vdwCoefTable, cudaTextureObject_t vdwCoefTableTex, cudaTextureObject_t forceTableTex, cudaTextureObject_t energyTableTex, float3 &iforce, float3 &iforceSlow, float3 &jforce, float3 &jforceSlow, float &energyVdw, float &energyElec, float &energySlow) |
template<bool doSlow> | |
__device__ __forceinline__ void | storeForces (const int pos, const float3 force, const float3 forceSlow, float *__restrict__ devForces_x, float *__restrict__ devForces_y, float *__restrict__ devForces_z, float *__restrict__ devForcesSlow_x, float *__restrict__ devForcesSlow_y, float *__restrict__ devForcesSlow_z) |
template<bool doPairlist> | |
__device__ __forceinline__ void | shuffleNext (float &xyzq_j_w, int &vdwtypej, int &jatomIndex, int &jexclMaxdiff, int &jexclIndex) |
template<bool doPairlist> | |
__device__ __forceinline__ void | shuffleNext (float &xyzq_j_w, int &vdwtypej, int &jatomIndex) |
template<bool doSlow> | |
__device__ __forceinline__ void | shuffleNext (float3 &jforce, float3 &jforceSlow) |
__device__ __forceinline__ float | distsq (const BoundingBox a, const float4 b) |
template<bool doEnergy, bool doVirial, bool doSlow, bool doPairlist, bool doStreaming> | |
__global__ void | __launch_bounds__ (WARPSIZE *NONBONDKERNEL_NUM_WARP, doPairlist?(10):(doEnergy?(10):(10))) nonbondedForceKernel(const int start |
if (itileList< numTileLists) | |
__global__ void | reduceNonbondedVirialKernel (const bool doSlow, const int atomStorageSize, const float4 *__restrict__ xyzq, const float4 *__restrict__ devForces, const float4 *__restrict__ devForcesSlow, VirialEnergy *__restrict__ virialEnergy) |
__global__ void | reduceVirialEnergyKernel (const bool doEnergy, const bool doVirial, const bool doSlow, const int numTileLists, const TileListVirialEnergy *__restrict__ tileListVirialEnergy, VirialEnergy *__restrict__ virialEnergy) |
__global__ void | reduceGBISEnergyKernel (const int numTileLists, const TileListVirialEnergy *__restrict__ tileListVirialEnergy, VirialEnergy *__restrict__ virialEnergy) |
__global__ void | reduceNonbondedBinsKernel (const bool doVirial, const bool doEnergy, const bool doSlow, const bool doGBIS, VirialEnergy *__restrict__ virialEnergy) |
template<int doSlow> | |
__global__ void | transposeForcesKernel (float4 *f, float4 *fSlow, float *fx, float *fy, float *fz, float *fw, float *fSlowx, float *fSlowy, float *fSlowz, float *fSloww, int n) |
Variables | |
__thread DeviceCUDA * | deviceCUDA |
__constant__ unsigned int | constExclusions [MAX_CONST_EXCLUSIONS] |
__global__ void const int | numTileLists |
__global__ void const int const TileList *__restrict__ | tileLists |
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ | tileExcls |
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ const int *__restrict__ | tileJatomStart |
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ const int *__restrict__ const int | vdwCoefTableWidth |
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ const int *__restrict__ const int const float2 *__restrict__ | vdwCoefTable |
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ const int *__restrict__ const int const float2 *__restrict__ cudaTextureObject_t | vdwCoefTableTex |
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ const int *__restrict__ const int const float2 *__restrict__ cudaTextureObject_t const int *__restrict__ | vdwTypes |
__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 |
__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 | latc |
__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__ | xyzq |
__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 const float4 *__restrict__ const float cudaTextureObject_t | forceTableTex |
__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 | energyTableTex |
__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 |
__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 const PatchPairRecord *__restrict__ | patchPairs |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ | atomIndex |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ | exclIndexMaxDiff |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ | overflowExclusions |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ | tileListDepth |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ | tileListOrder |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ | jtiles |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ | tileListStat |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ | boundingBoxes |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ | devForce_x |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ | devForce_y |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ | devForce_z |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ | devForce_w |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ | devForceSlow_x |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ | devForceSlow_y |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ | devForceSlow_z |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ | devForceSlow_w |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ const int | numPatches |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ const int unsigned int *__restrict__ | patchNumCount |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ const int unsigned int *__restrict__ const CudaPatchRecord *__restrict__ | cudaPatches |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ const int unsigned int *__restrict__ const CudaPatchRecord *__restrict__ float4 *__restrict__ | mapForces |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ const int unsigned int *__restrict__ const CudaPatchRecord *__restrict__ float4 *__restrict__ float4 *__restrict__ | mapForcesSlow |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ const int unsigned int *__restrict__ const CudaPatchRecord *__restrict__ float4 *__restrict__ float4 *__restrict__ int *__restrict__ | mapPatchReadyQueue |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ const int unsigned int *__restrict__ const CudaPatchRecord *__restrict__ float4 *__restrict__ float4 *__restrict__ int *__restrict__ int *__restrict__ | outputOrder |
__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 const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ const int unsigned int *__restrict__ const CudaPatchRecord *__restrict__ float4 *__restrict__ float4 *__restrict__ int *__restrict__ int *__restrict__ TileListVirialEnergy *__restrict__ virialEnergy int | itileList = start + (NONBONDKERNEL_NUM_WARP == 1 ? blockIdx.x : (threadIdx.x/WARPSIZE + NONBONDKERNEL_NUM_WARP*blockIdx.x)) |
#define CALL | ( | DOENERGY, | |
DOVIRIAL, | |||
DOSLOW, | |||
DOPAIRLIST, | |||
DOSTREAMING | |||
) |
#define LARGE_FLOAT (float)(1.0e10) |
Definition at line 238 of file CudaComputeNonbondedKernel.cu.
#define MAX_CONST_EXCLUSIONS 2048 |
Definition at line 39 of file CudaComputeNonbondedKernel.cu.
Referenced by CudaComputeNonbondedKernel::bindExclusions().
#define NONBONDKERNEL_NUM_WARP 4 |
Definition at line 24 of file CudaComputeNonbondedKernel.cu.
Referenced by CudaComputeNonbondedKernel::nonbondedForce().
#define OVERALLOC 1.2f |
Definition at line 35 of file CudaComputeNonbondedKernel.cu.
Referenced by CudaComputeNonbondedKernel::updateVdwTypesExcl().
#define REDUCEGBISENERGYKERNEL_NUM_WARP 32 |
Definition at line 27 of file CudaComputeNonbondedKernel.cu.
Referenced by CudaComputeNonbondedKernel::reduceVirialEnergy().
#define REDUCENONBONDEDVIRIALKERNEL_NUM_WARP 32 |
Definition at line 25 of file CudaComputeNonbondedKernel.cu.
Referenced by CudaComputeNonbondedKernel::reduceVirialEnergy().
#define REDUCEVIRIALENERGYKERNEL_NUM_WARP 32 |
Definition at line 26 of file CudaComputeNonbondedKernel.cu.
Referenced by CudaComputeNonbondedKernel::reduceVirialEnergy().
#define TABLE_PARAMS cudaNonbondedTables.getForceTableTex(), cudaNonbondedTables.getEnergyTableTex() |
__global__ void __launch_bounds__ | ( | WARPSIZE * | NONBONDKERNEL_NUM_WARP, |
doPairlist? | 10):(doEnergy?(10):(10) | ||
) | const |
__device__ __forceinline__ void calcForceEnergy | ( | const float | r2, |
const float | qi, | ||
const float | qj, | ||
const float | dx, | ||
const float | dy, | ||
const float | dz, | ||
const int | vdwtypei, | ||
const int | vdwtypej, | ||
const float2 *__restrict__ | vdwCoefTable, | ||
cudaTextureObject_t | vdwCoefTableTex, | ||
cudaTextureObject_t | forceTableTex, | ||
cudaTextureObject_t | energyTableTex, | ||
float3 & | iforce, | ||
float3 & | iforceSlow, | ||
float3 & | jforce, | ||
float3 & | jforceSlow, | ||
float & | energyVdw, | ||
float & | energyElec, | ||
float & | energySlow | ||
) |
Definition at line 91 of file CudaComputeNonbondedKernel.cu.
References __ldg, energyTableTex, forceTableTex, sampleTableTex(), vdwCoefTableTex, float2::x, and float2::y.
__device__ __forceinline__ float distsq | ( | const BoundingBox | a, |
const float4 | b | ||
) |
Definition at line 230 of file CudaComputeNonbondedKernel.cu.
References BoundingBox::wx, BoundingBox::wy, BoundingBox::wz, BoundingBox::x, BoundingBox::y, and BoundingBox::z.
Referenced by buildTileListsBBKernel().
if | ( | ) |
Definition at line 300 of file CudaComputeNonbondedKernel.cu.
__global__ void reduceGBISEnergyKernel | ( | const int | numTileLists, |
const TileListVirialEnergy *__restrict__ | tileListVirialEnergy, | ||
VirialEnergy *__restrict__ | virialEnergy | ||
) |
Definition at line 1049 of file CudaComputeNonbondedKernel.cu.
References ATOMIC_BINS, BLOCK_SYNC, itileList, and tempStorage.
__global__ void reduceNonbondedBinsKernel | ( | const bool | doVirial, |
const bool | doEnergy, | ||
const bool | doSlow, | ||
const bool | doGBIS, | ||
VirialEnergy *__restrict__ | virialEnergy | ||
) |
Definition at line 1071 of file CudaComputeNonbondedKernel.cu.
References ATOMIC_BINS, and tempStorage.
__global__ void reduceNonbondedVirialKernel | ( | const bool | doSlow, |
const int | atomStorageSize, | ||
const float4 *__restrict__ | xyzq, | ||
const float4 *__restrict__ | devForces, | ||
const float4 *__restrict__ | devForcesSlow, | ||
VirialEnergy *__restrict__ | virialEnergy | ||
) |
Definition at line 804 of file CudaComputeNonbondedKernel.cu.
References ATOMIC_BINS, BLOCK_SYNC, and tempStorage.
__global__ void reduceVirialEnergyKernel | ( | const bool | doEnergy, |
const bool | doVirial, | ||
const bool | doSlow, | ||
const int | numTileLists, | ||
const TileListVirialEnergy *__restrict__ | tileListVirialEnergy, | ||
VirialEnergy *__restrict__ | virialEnergy | ||
) |
Definition at line 923 of file CudaComputeNonbondedKernel.cu.
References ATOMIC_BINS, BLOCK_SYNC, TileListVirialEnergy::energyElec, TileListVirialEnergy::energySlow, TileListVirialEnergy::energyVdw, TileListVirialEnergy::forceSlowx, TileListVirialEnergy::forceSlowy, TileListVirialEnergy::forceSlowz, TileListVirialEnergy::forcex, TileListVirialEnergy::forcey, TileListVirialEnergy::forcez, itileList, TileListVirialEnergy::shx, TileListVirialEnergy::shy, TileListVirialEnergy::shz, and tempStorage.
__device__ __forceinline__ float4 sampleTableTex | ( | cudaTextureObject_t | tex, |
float | k | ||
) |
Definition at line 44 of file CudaComputeNonbondedKernel.cu.
References FORCE_ENERGY_TABLE_SIZE, and x.
__device__ __forceinline__ void shuffleNext | ( | float & | xyzq_j_w, |
int & | vdwtypej, | ||
int & | jatomIndex, | ||
int & | jexclMaxdiff, | ||
int & | jexclIndex | ||
) |
Definition at line 192 of file CudaComputeNonbondedKernel.cu.
References WARP_FULL_MASK, WARP_SHUFFLE, and WARPSIZE.
__device__ __forceinline__ void shuffleNext | ( | float & | xyzq_j_w, |
int & | vdwtypej, | ||
int & | jatomIndex | ||
) |
Definition at line 204 of file CudaComputeNonbondedKernel.cu.
References WARP_FULL_MASK, WARP_SHUFFLE, and WARPSIZE.
__device__ __forceinline__ void shuffleNext | ( | float3 & | jforce, |
float3 & | jforceSlow | ||
) |
Definition at line 214 of file CudaComputeNonbondedKernel.cu.
References WARP_FULL_MASK, WARP_SHUFFLE, and WARPSIZE.
__device__ __forceinline__ void storeForces | ( | const int | pos, |
const float3 | force, | ||
const float3 | forceSlow, | ||
float *__restrict__ | devForces_x, | ||
float *__restrict__ | devForces_y, | ||
float *__restrict__ | devForces_z, | ||
float *__restrict__ | devForcesSlow_x, | ||
float *__restrict__ | devForcesSlow_y, | ||
float *__restrict__ | devForcesSlow_z | ||
) |
Definition at line 156 of file CudaComputeNonbondedKernel.cu.
__global__ void transposeForcesKernel | ( | float4 * | f, |
float4 * | fSlow, | ||
float * | fx, | ||
float * | fy, | ||
float * | fz, | ||
float * | fw, | ||
float * | fSlowx, | ||
float * | fSlowy, | ||
float * | fSlowz, | ||
float * | fSloww, | ||
int | n | ||
) |
Definition at line 1234 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ atomIndex |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by ResidueLookupElem::append(), and HomePatch::doAtomMigration().
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ boundingBoxes |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__constant__ unsigned int constExclusions[MAX_CONST_EXCLUSIONS] |
Definition at line 40 of file CudaComputeNonbondedKernel.cu.
Referenced by CudaComputeNonbondedKernel::bindExclusions().
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ const int unsigned int* __restrict__ const CudaPatchRecord* __restrict__ cudaPatches |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by bondedForcesKernel(), cuda_nonbonded_forces(), GBIS_Kernel(), CudaComputeGBISKernel::GBISphase1(), CudaComputeGBISKernel::GBISphase3(), modifiedExclusionForcesKernel(), vdw_fswitch_energy(), vdw_fswitch_forceandenergy(), and vdw_fswitch_shift().
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ devForce_w |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ devForce_x |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ devForce_y |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ devForce_z |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ devForceSlow_w |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ devForceSlow_x |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ devForceSlow_y |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ devForceSlow_z |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__thread DeviceCUDA* deviceCUDA |
Definition at line 22 of file DeviceCUDA.C.
__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 energyTableTex |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by calcForceEnergy(), and modifiedExclusionForce().
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ exclIndexMaxDiff |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 forceTableTex |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by calcForceEnergy(), and modifiedExclusionForce().
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ const int unsigned int* __restrict__ const CudaPatchRecord* __restrict__ float4* __restrict__ float4* __restrict__ int* __restrict__ int* __restrict__ TileListVirialEnergy* __restrict__ virialEnergy int itileList = start + (NONBONDKERNEL_NUM_WARP == 1 ? blockIdx.x : (threadIdx.x/WARPSIZE + NONBONDKERNEL_NUM_WARP*blockIdx.x)) |
Definition at line 300 of file CudaComputeNonbondedKernel.cu.
Referenced by bitshiftTileListDepth(), buildRemoveZerosSortKey(), buildTileListsBBKernel(), GBIS_Kernel(), reduceGBISEnergyKernel(), reduceVirialEnergyKernel(), and setupSortKey().
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ jtiles |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by bondedForcesKernel(), CudaTileListKernel::buildTileLists(), cuda_GBIS_P1(), cuda_GBIS_P2(), cuda_GBIS_P3(), ComputeNonbondedCUDA::doWork(), CudaComputeGBISKernel::GBISphase1(), CudaComputeGBISKernel::GBISphase3(), modifiedExclusionForcesKernel(), and ComputeNonbondedCUDA::recvYieldDevice().
__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 |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by bondedForcesKernel(), CudaTileListKernel::buildTileLists(), cuda_GBIS_P1(), cuda_GBIS_P2(), cuda_GBIS_P3(), ComputeNonbondedCUDA::doWork(), CudaComputeGBISKernel::GBISphase1(), CudaComputeGBISKernel::GBISphase3(), modifiedExclusionForcesKernel(), and ComputeNonbondedCUDA::recvYieldDevice().
__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 |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by bondedForcesKernel(), CudaTileListKernel::buildTileLists(), cuda_GBIS_P1(), cuda_GBIS_P2(), cuda_GBIS_P3(), ComputeNonbondedCUDA::doWork(), CudaComputeGBISKernel::GBISphase1(), CudaComputeGBISKernel::GBISphase3(), modifiedExclusionForcesKernel(), and ComputeNonbondedCUDA::recvYieldDevice().
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ const int unsigned int* __restrict__ const CudaPatchRecord* __restrict__ float4* __restrict__ mapForces |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ const int unsigned int* __restrict__ const CudaPatchRecord* __restrict__ float4* __restrict__ float4* __restrict__ mapForcesSlow |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ const int unsigned int* __restrict__ const CudaPatchRecord* __restrict__ float4* __restrict__ float4* __restrict__ int* __restrict__ mapPatchReadyQueue |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ const int numPatches |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by ProxyMgr::buildSpanningTree0(), ParallelIOMgr::calcAtomsInEachPatch(), ComputeGlobal::ComputeGlobal(), Controller::Controller(), WorkDistrib::createAtomLists(), WorkDistrib::createHomePatches(), ProxyMgr::createProxies(), NodeProxyMgr::createProxyInfo(), dumpbench(), ComputePmeMgr::initialize(), ComputePmeMgr::initialize_pencils(), NodeProxyMgr::NodeProxyMgr(), CudaComputeNonbondedKernel::nonbondedForce(), Node::outputPatchComputeMaps(), WorkDistrib::patchMapInit(), PatchProxyListMsg::PatchProxyListMsg(), GlobalMasterServer::recvData(), ComputePmeCUDADevice::recvForcesFromNeighbor(), WorkDistrib::reinitAtoms(), WorkDistrib::savePatchMap(), ParallelIOMgr::sendAtomsToHomePatchProcs(), WorkDistrib::sendPatchMap(), ProxyMgr::sendSpanningTrees(), Node::startup(), NamdCentLB::Strategy(), and NodeProxyMgr::~NodeProxyMgr().
__global__ void const int numTileLists |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by buildTileListsBBKernel().
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ const int unsigned int* __restrict__ const CudaPatchRecord* __restrict__ float4* __restrict__ float4* __restrict__ int* __restrict__ int* __restrict__ outputOrder |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ overflowExclusions |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ const BoundingBox* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ float* __restrict__ const int unsigned int* __restrict__ patchNumCount |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ patchPairs |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__global__ void const int const TileList* __restrict__ TileExcl* __restrict__ const int* __restrict__ tileJatomStart |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ tileListDepth |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ tileListOrder |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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 const PatchPairRecord* __restrict__ const int* __restrict__ const int2* __restrict__ const unsigned int* __restrict__ unsigned int* __restrict__ int* __restrict__ int* __restrict__ TileListStat* __restrict__ tileListStat |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__global__ void const int const TileList* __restrict__ TileExcl* __restrict__ const int* __restrict__ const int const float2* __restrict__ vdwCoefTable |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by modifiedExclusionForcesKernel().
__global__ void const int const TileList* __restrict__ TileExcl* __restrict__ const int* __restrict__ const int const float2* __restrict__ cudaTextureObject_t vdwCoefTableTex |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by calcForceEnergy(), modifiedExclusionForce(), and modifiedExclusionForcesKernel().
__global__ void const int const TileList* __restrict__ TileExcl* __restrict__ const int* __restrict__ const int vdwCoefTableWidth |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by modifiedExclusionForce(), and modifiedExclusionForcesKernel().
__global__ void const int const TileList* __restrict__ TileExcl* __restrict__ const int* __restrict__ const int const float2* __restrict__ cudaTextureObject_t const int* __restrict__ vdwTypes |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
__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__ xyzq |
Definition at line 254 of file CudaComputeNonbondedKernel.cu.
Referenced by bondedForcesKernel(), and modifiedExclusionForcesKernel().