NAMD
|
#include <cuda.h>
#include <namd_cub/cub.cuh>
#include "CudaUtils.h"
#include "CudaTileListKernel.h"
#include "CudaComputeGBISKernel.h"
#include "ComputeGBIS.inl"
#include "DeviceCUDA.h"
Go to the source code of this file.
Classes | |
struct | GBISParam< phase > |
struct | GBISInput< phase > |
struct | GBISResults< phase > |
struct | GBISParam< 1 > |
struct | GBISInput< 1 > |
struct | GBISResults< 1 > |
struct | GBISParam< 2 > |
struct | GBISInput< 2 > |
struct | GBISResults< 2 > |
struct | GBISParam< 3 > |
struct | GBISInput< 3 > |
struct | GBISResults< 3 > |
Macros | |
#define | GBIS_CUDA |
#define | LARGE_FLOAT (float)(1.0e10) |
#define | GBISKERNEL_NUM_WARP 4 |
#define | CALL(DOENERGY, DOSLOW) |
Functions | |
__device__ __forceinline__ void | shuffleNext (float &w) |
template<bool doEnergy, bool doSlow> | |
__device__ __forceinline__ void | calcGBISPhase (const float r2, const float dx, const float dy, const float dz, const float cutoff2, const GBISParam< 1 > param, const GBISInput< 1 > inp, GBISResults< 1 > &resI, GBISResults< 1 > &resJ, float &energy) |
__device__ __forceinline__ void | writeResult (const int i, const GBISResults< 1 > res, float *out, float4 *forces) |
template<bool doEnergy, bool doSlow> | |
__device__ __forceinline__ void | calcGBISPhase (const float r2, const float dx, const float dy, const float dz, const float cutoff2, const GBISParam< 2 > param, const GBISInput< 2 > inp, GBISResults< 2 > &resI, GBISResults< 2 > &resJ, float &energyT) |
__device__ __forceinline__ void | writeResult (const int i, const GBISResults< 2 > res, float *out, float4 *forces) |
template<bool doEnergy, bool doSlow> | |
__device__ __forceinline__ void | calcGBISPhase (const float r2, const float dx, const float dy, const float dz, const float cutoff2, const GBISParam< 3 > param, const GBISInput< 3 > inp, GBISResults< 3 > &resI, GBISResults< 3 > &resJ, float &energy) |
__device__ __forceinline__ void | writeResult (const int i, const GBISResults< 3 > res, float *out, float4 *forces) |
template<bool doEnergy, bool doSlow, int phase> | |
__global__ void | GBIS_Kernel (const int numTileLists, const TileList *__restrict__ tileLists, const int *__restrict__ tileJatomStart, const PatchPairRecord *__restrict__ patchPairs, const float3 lata, const float3 latb, const float3 latc, const float4 *__restrict__ xyzq, const float cutoff2, const GBISParam< phase > param, const float *__restrict__ inp1, const float *__restrict__ inp2, const float *__restrict__ inp3, float *__restrict__ out, float4 *__restrict__ forces, TileListVirialEnergy *__restrict__ virialEnergy) |
Variables | |
__thread DeviceCUDA * | deviceCUDA |
#define CALL | ( | DOENERGY, | |
DOSLOW | |||
) |
#define GBIS_CUDA |
Definition at line 18 of file CudaComputeGBISKernel.cu.
#define GBISKERNEL_NUM_WARP 4 |
Definition at line 28 of file CudaComputeGBISKernel.cu.
Referenced by GBIS_Kernel(), CudaComputeGBISKernel::GBISphase1(), CudaComputeGBISKernel::GBISphase2(), and CudaComputeGBISKernel::GBISphase3().
#define LARGE_FLOAT (float)(1.0e10) |
Definition at line 26 of file CudaComputeGBISKernel.cu.
__device__ __forceinline__ void calcGBISPhase | ( | const float | r2, |
const float | dx, | ||
const float | dy, | ||
const float | dz, | ||
const float | cutoff2, | ||
const GBISParam< 1 > | param, | ||
const GBISInput< 1 > | inp, | ||
GBISResults< 1 > & | resI, | ||
GBISResults< 1 > & | resJ, | ||
float & | energy | ||
) |
Definition at line 78 of file CudaComputeGBISKernel.cu.
References GBISParam< 1 >::a_cut, CalcH(), GBISInput< 1 >::intRad0j, GBISInput< 1 >::intRadSi, GBISResults< 1 >::psiSum, GBISInput< 1 >::qi, and GBISInput< 1 >::qj.
__device__ __forceinline__ void calcGBISPhase | ( | const float | r2, |
const float | dx, | ||
const float | dy, | ||
const float | dz, | ||
const float | cutoff2, | ||
const GBISParam< 2 > | param, | ||
const GBISInput< 2 > | inp, | ||
GBISResults< 2 > & | resI, | ||
GBISResults< 2 > & | resJ, | ||
float & | energyT | ||
) |
Definition at line 147 of file CudaComputeGBISKernel.cu.
References GBISInput< 2 >::bornRadI, GBISInput< 2 >::bornRadJ, GBISResults< 2 >::dEdaSum, GBISParam< 2 >::epsilon_p_i, GBISParam< 2 >::epsilon_s_i, GBISResults< 2 >::force, if(), GBISParam< 2 >::kappa, GBISInput< 2 >::qi, GBISInput< 2 >::qj, GBISParam< 2 >::r_cut2, GBISParam< 2 >::r_cut_2, GBISParam< 2 >::r_cut_4, GBISParam< 2 >::scaling, and GBISParam< 2 >::smoothDist.
__device__ __forceinline__ void calcGBISPhase | ( | const float | r2, |
const float | dx, | ||
const float | dy, | ||
const float | dz, | ||
const float | cutoff2, | ||
const GBISParam< 3 > | param, | ||
const GBISInput< 3 > | inp, | ||
GBISResults< 3 > & | resI, | ||
GBISResults< 3 > & | resJ, | ||
float & | energy | ||
) |
Definition at line 274 of file CudaComputeGBISKernel.cu.
References GBISParam< 3 >::a_cut, CalcDH(), GBISInput< 3 >::dHdrPrefixI, GBISInput< 3 >::dHdrPrefixJ, GBISResults< 3 >::force, GBISInput< 3 >::intRadIS, GBISInput< 3 >::intRadJ0, GBISInput< 3 >::intRadSJ, and GBISInput< 3 >::qi.
__global__ void GBIS_Kernel | ( | const int | numTileLists, |
const TileList *__restrict__ | tileLists, | ||
const int *__restrict__ | tileJatomStart, | ||
const PatchPairRecord *__restrict__ | patchPairs, | ||
const float3 | lata, | ||
const float3 | latb, | ||
const float3 | latc, | ||
const float4 *__restrict__ | xyzq, | ||
const float | cutoff2, | ||
const GBISParam< phase > | param, | ||
const float *__restrict__ | inp1, | ||
const float *__restrict__ | inp2, | ||
const float *__restrict__ | inp3, | ||
float *__restrict__ | out, | ||
float4 *__restrict__ | forces, | ||
TileListVirialEnergy *__restrict__ | virialEnergy | ||
) |
Definition at line 314 of file CudaComputeGBISKernel.cu.
References cutoff2, GBISKERNEL_NUM_WARP, PatchPairRecord::iatomSize, TileList::iatomStart, itileList, PatchPairRecord::jatomSize, TileList::jtileEnd, TileList::jtileStart, TileList::offsetXYZ, tempStorage, WARP_FULL_MASK, WARP_SHUFFLE, WARPSIZE, and writeResult().
__device__ __forceinline__ void shuffleNext | ( | float & | w | ) |
Definition at line 31 of file CudaComputeGBISKernel.cu.
References WARP_FULL_MASK, WARP_SHUFFLE, and WARPSIZE.
__device__ __forceinline__ void writeResult | ( | const int | i, |
const GBISResults< 1 > | res, | ||
float * | out, | ||
float4 * | forces | ||
) |
Definition at line 97 of file CudaComputeGBISKernel.cu.
References GBISResults< 1 >::psiSum.
Referenced by GBIS_Kernel().
__device__ __forceinline__ void writeResult | ( | const int | i, |
const GBISResults< 2 > | res, | ||
float * | out, | ||
float4 * | forces | ||
) |
Definition at line 224 of file CudaComputeGBISKernel.cu.
References GBISResults< 2 >::dEdaSum, GBISResults< 2 >::force, x, y, and z.
__device__ __forceinline__ void writeResult | ( | const int | i, |
const GBISResults< 3 > | res, | ||
float * | out, | ||
float4 * | forces | ||
) |
Definition at line 300 of file CudaComputeGBISKernel.cu.
References GBISResults< 3 >::force, x, y, and z.
__thread DeviceCUDA* deviceCUDA |
Definition at line 22 of file DeviceCUDA.C.