37f22c8627
- Optimizations for molecular systems - Improved kernel performance and greater CPU overlap - Reduced GPU to CPU communications for discrete devices - Switch classic Intel makefiles to use LLVM-based compilers - Prefetch optimizations supported for OpenCL - Optimized data repack for quaternions
823 lines
35 KiB
Plaintext
823 lines
35 KiB
Plaintext
// **************************************************************************
|
|
// sw.cu
|
|
// -------------------
|
|
// W. Michael Brown (ORNL)
|
|
//
|
|
// Device code for acceleration of the sw pair style
|
|
//
|
|
// __________________________________________________________________________
|
|
// This file is part of the LAMMPS Accelerator Library (LAMMPS_AL)
|
|
// __________________________________________________________________________
|
|
//
|
|
// begin : Tue March 26, 2013
|
|
// email : brownw@ornl.gov
|
|
// ***************************************************************************
|
|
|
|
#if defined(NV_KERNEL) || defined(USE_HIP)
|
|
#include "lal_aux_fun1.h"
|
|
|
|
#ifndef _DOUBLE_DOUBLE
|
|
_texture( pos_tex,float4);
|
|
_texture( sw1_tex,float4);
|
|
_texture( sw2_tex,float4);
|
|
_texture( sw3_tex,float4);
|
|
#else
|
|
_texture_2d( pos_tex,int4);
|
|
_texture( sw1_tex,int4);
|
|
_texture( sw2_tex,int4);
|
|
_texture( sw3_tex,int4);
|
|
#endif
|
|
|
|
#else
|
|
#define pos_tex x_
|
|
#define sw1_tex sw1
|
|
#define sw2_tex sw2
|
|
#define sw3_tex sw3
|
|
#endif
|
|
|
|
#define THIRD (numtyp)0.66666666666666666667
|
|
|
|
//#define THREE_CONCURRENT
|
|
|
|
#if (SHUFFLE_AVAIL == 0)
|
|
|
|
#define store_answers_p(f, energy, virial, ii, inum, tid, t_per_atom, \
|
|
offset, eflag, vflag, ans, engv, ev_stride) \
|
|
if (t_per_atom>1) { \
|
|
simd_reduce_add3(t_per_atom, red_acc, offset, tid, f.x, f.y, f.z); \
|
|
if (EVFLAG && (vflag==2 || eflag==2)) { \
|
|
if (eflag) { \
|
|
simdsync(); \
|
|
simd_reduce_add1(t_per_atom, red_acc, offset, tid, energy); \
|
|
} \
|
|
if (vflag) { \
|
|
simdsync(); \
|
|
simd_reduce_arr(6, t_per_atom, red_acc, offset, tid, virial); \
|
|
} \
|
|
} \
|
|
} \
|
|
if (offset==0 && ii<inum) { \
|
|
acctyp3 old=ans[ii]; \
|
|
old.x+=f.x; \
|
|
old.y+=f.y; \
|
|
old.z+=f.z; \
|
|
ans[ii]=old; \
|
|
} \
|
|
if (EVFLAG && (eflag || vflag)) { \
|
|
int ei=BLOCK_ID_X; \
|
|
if (eflag!=2 && vflag!=2) { \
|
|
if (eflag) { \
|
|
simdsync(); \
|
|
block_reduce_add1(simd_size(), red_acc, tid, energy); \
|
|
if (vflag) __syncthreads(); \
|
|
if (tid==0) { \
|
|
engv[ei]+=energy*(acctyp)0.5; \
|
|
ei+=ev_stride; \
|
|
} \
|
|
} \
|
|
if (vflag) { \
|
|
simdsync(); \
|
|
block_reduce_arr(6, simd_size(), red_acc, tid, virial); \
|
|
if (tid==0) { \
|
|
for (int r=0; r<6; r++) { \
|
|
engv[ei]+=virial[r]*(acctyp)0.5; \
|
|
ei+=ev_stride; \
|
|
} \
|
|
} \
|
|
} \
|
|
} else if (offset==0 && ii<inum) { \
|
|
int ei=ii; \
|
|
if (EVFLAG && eflag) { \
|
|
engv[ei]+=energy*(acctyp)0.5; \
|
|
ei+=inum; \
|
|
} \
|
|
if (EVFLAG && vflag) { \
|
|
for (int i=0; i<6; i++) { \
|
|
engv[ei]+=virial[i]*(acctyp)0.5; \
|
|
ei+=inum; \
|
|
} \
|
|
} \
|
|
} \
|
|
}
|
|
|
|
#else
|
|
|
|
#if (EVFLAG == 1)
|
|
|
|
#define store_answers_p(f, energy, virial, ii, inum, tid, t_per_atom, \
|
|
offset, eflag, vflag, ans, engv, ev_stride) \
|
|
if (t_per_atom>1) { \
|
|
simd_reduce_add3(t_per_atom, f.x, f.y, f.z); \
|
|
if (vflag==2 || eflag==2) { \
|
|
if (eflag) \
|
|
simd_reduce_add1(t_per_atom,energy); \
|
|
if (vflag) \
|
|
simd_reduce_arr(6, t_per_atom,virial); \
|
|
} \
|
|
} \
|
|
if (offset==0 && ii<inum) { \
|
|
acctyp3 old=ans[ii]; \
|
|
old.x+=f.x; \
|
|
old.y+=f.y; \
|
|
old.z+=f.z; \
|
|
ans[ii]=old; \
|
|
} \
|
|
if (eflag || vflag) { \
|
|
if (eflag!=2 && vflag!=2) { \
|
|
const int vwidth = simd_size(); \
|
|
const int voffset = tid & (simd_size() - 1); \
|
|
const int bnum = tid/simd_size(); \
|
|
int active_subgs = BLOCK_SIZE_X/simd_size(); \
|
|
for ( ; active_subgs > 1; active_subgs /= vwidth) { \
|
|
if (active_subgs < BLOCK_SIZE_X/simd_size()) __syncthreads(); \
|
|
if (bnum < active_subgs) { \
|
|
if (eflag) { \
|
|
simd_reduce_add1(vwidth, energy); \
|
|
if (voffset==0) red_acc[6][bnum] = energy; \
|
|
} \
|
|
if (vflag) { \
|
|
simd_reduce_arr(6, vwidth, virial); \
|
|
if (voffset==0) \
|
|
for (int r=0; r<6; r++) red_acc[r][bnum]=virial[r]; \
|
|
} \
|
|
} \
|
|
\
|
|
__syncthreads(); \
|
|
if (tid < active_subgs) { \
|
|
if (eflag) energy = red_acc[6][tid]; \
|
|
if (vflag) \
|
|
for (int r = 0; r < 6; r++) virial[r] = red_acc[r][tid]; \
|
|
} else { \
|
|
if (eflag) energy = (acctyp)0; \
|
|
if (vflag) for (int r = 0; r < 6; r++) virial[r] = (acctyp)0; \
|
|
} \
|
|
} \
|
|
\
|
|
if (bnum == 0) { \
|
|
int ei=BLOCK_ID_X; \
|
|
if (eflag) { \
|
|
simd_reduce_add1(vwidth, energy); \
|
|
if (tid==0) { \
|
|
engv[ei]+=energy*(acctyp)0.5; \
|
|
ei+=ev_stride; \
|
|
} \
|
|
} \
|
|
if (vflag) { \
|
|
simd_reduce_arr(6, vwidth, virial); \
|
|
if (tid==0) { \
|
|
for (int r=0; r<6; r++) { \
|
|
engv[ei]+=virial[r]*(acctyp)0.5; \
|
|
ei+=ev_stride; \
|
|
} \
|
|
} \
|
|
} \
|
|
} \
|
|
} else if (offset==0 && ii<inum) { \
|
|
int ei=ii; \
|
|
if (eflag) { \
|
|
engv[ei]+=energy*(acctyp)0.5; \
|
|
ei+=inum; \
|
|
} \
|
|
if (vflag) { \
|
|
for (int i=0; i<6; i++) { \
|
|
engv[ei]+=virial[i]*(acctyp)0.5; \
|
|
ei+=inum; \
|
|
} \
|
|
} \
|
|
} \
|
|
}
|
|
|
|
#else
|
|
|
|
#define store_answers_p(f, energy, virial, ii, inum, tid, t_per_atom, \
|
|
offset, eflag, vflag, ans, engv, ev_stride) \
|
|
if (t_per_atom>1) \
|
|
simd_reduce_add3(t_per_atom, f.x, f.y, f.z); \
|
|
if (offset==0 && ii<inum) { \
|
|
acctyp3 old=ans[ii]; \
|
|
old.x+=f.x; \
|
|
old.y+=f.y; \
|
|
old.z+=f.z; \
|
|
ans[ii]=old; \
|
|
}
|
|
|
|
#endif
|
|
#endif
|
|
|
|
__kernel void k_sw_short_nbor(const __global numtyp4 *restrict x_,
|
|
const __global numtyp * restrict cutsq,
|
|
const int ntypes, __global int * dev_nbor,
|
|
const __global int * dev_packed,
|
|
const int inum, const int nbor_pitch,
|
|
const int t_per_atom) {
|
|
const int ii=GLOBAL_ID_X;
|
|
|
|
#ifdef ONETYPE
|
|
const numtyp sw_cutsq=cutsq[ONETYPE];
|
|
#endif
|
|
|
|
if (ii<inum) {
|
|
const int i=dev_packed[ii];
|
|
int nbor=ii+nbor_pitch;
|
|
const int numj=dev_packed[nbor];
|
|
nbor+=nbor_pitch;
|
|
const int nbor_end=nbor+fast_mul(numj,nbor_pitch);
|
|
|
|
numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
|
|
#ifndef ONETYPE
|
|
const int itype=ix.w*ntypes;
|
|
#endif
|
|
int newj=0;
|
|
|
|
__global int *out_list=dev_nbor+2*nbor_pitch+ii*t_per_atom;
|
|
const int out_stride=nbor_pitch*t_per_atom-t_per_atom;
|
|
|
|
for ( ; nbor<nbor_end; nbor+=nbor_pitch) {
|
|
int sj=dev_packed[nbor];
|
|
int j = sj & NEIGHMASK;
|
|
numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
|
|
|
|
#ifndef ONETYPE
|
|
const int mtype=jx.w+itype;
|
|
const numtyp sw_cutsq=cutsq[mtype];
|
|
#endif
|
|
|
|
// Compute r12
|
|
numtyp delx = ix.x-jx.x;
|
|
numtyp dely = ix.y-jx.y;
|
|
numtyp delz = ix.z-jx.z;
|
|
numtyp rsq = delx*delx+dely*dely+delz*delz;
|
|
|
|
if (rsq<sw_cutsq) {
|
|
*out_list=sj;
|
|
out_list++;
|
|
newj++;
|
|
if ((newj & (t_per_atom-1))==0)
|
|
out_list+=out_stride;
|
|
}
|
|
} // for nbor
|
|
dev_nbor[ii+nbor_pitch]=newj;
|
|
} // if ii
|
|
}
|
|
|
|
__kernel void k_sw(const __global numtyp4 *restrict x_,
|
|
const __global numtyp4 * restrict sw_pre,
|
|
const __global numtyp4 * restrict c_14,
|
|
const __global numtyp2 * restrict c_56,
|
|
const int ntypes, const __global int * dev_nbor,
|
|
__global acctyp3 *restrict ans,
|
|
__global acctyp *restrict engv,
|
|
const int eflag, const int vflag, const int inum,
|
|
const int nbor_pitch, const int t_per_atom,
|
|
const int ev_stride) {
|
|
int n_stride;
|
|
int tid, ii, offset;
|
|
atom_info(t_per_atom,ii,tid,offset);
|
|
local_allocate_store_pair();
|
|
|
|
#ifdef ONETYPE
|
|
const numtyp4 pre_sw=sw_pre[ONETYPE];
|
|
const numtyp4 pre_sw_c14=c_14[ONETYPE];
|
|
numtyp2 pre_sw_c56;
|
|
if (EVFLAG && eflag) pre_sw_c56=c_56[ONETYPE];
|
|
#endif
|
|
|
|
acctyp3 f;
|
|
f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0;
|
|
acctyp energy, virial[6];
|
|
if (EVFLAG) {
|
|
energy=(acctyp)0;
|
|
for (int i=0; i<6; i++) virial[i]=(acctyp)0;
|
|
}
|
|
|
|
if (ii<inum) {
|
|
int nbor, nbor_end, i, numj;
|
|
nbor_info_p(dev_nbor,nbor_pitch,t_per_atom,ii,offset,i,numj,
|
|
n_stride,nbor_end,nbor);
|
|
|
|
numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
|
|
#ifndef ONETYPE
|
|
int itype=ix.w;
|
|
itype*=ntypes;
|
|
#endif
|
|
|
|
for ( ; nbor<nbor_end; nbor+=n_stride) {
|
|
|
|
int j=dev_nbor[nbor];
|
|
j &= NEIGHMASK;
|
|
|
|
numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
|
|
#ifndef ONETYPE
|
|
int mtype=jx.w;
|
|
mtype+=itype;
|
|
#endif
|
|
|
|
// Compute r12
|
|
numtyp delx = ix.x-jx.x;
|
|
numtyp dely = ix.y-jx.y;
|
|
numtyp delz = ix.z-jx.z;
|
|
numtyp rsq = delx*delx+dely*dely+delz*delz;
|
|
|
|
#ifndef ONETYPE
|
|
numtyp4 pre_sw=sw_pre[mtype];
|
|
numtyp4 pre_sw_c14=c_14[mtype];
|
|
#endif
|
|
numtyp r=ucl_sqrt(rsq);
|
|
#ifdef SPQ
|
|
numtyp rp=r*r;
|
|
rp=ucl_recip(rp*rp);
|
|
numtyp rq=(numtyp)1.0;
|
|
#else
|
|
numtyp rp=ucl_powr(r,-pre_sw.z);
|
|
numtyp rq=ucl_powr(r,-pre_sw.w);
|
|
#endif
|
|
numtyp rainv=ucl_recip(r-pre_sw.x);
|
|
numtyp expsrainv=ucl_exp(pre_sw.y*rainv);
|
|
rainv*=rainv*r;
|
|
numtyp force = (pre_sw_c14.x*rp-pre_sw_c14.y*rq +
|
|
(pre_sw_c14.z*rp-pre_sw_c14.w*rq) * rainv)*
|
|
expsrainv*ucl_recip(rsq);
|
|
|
|
f.x+=delx*force;
|
|
f.y+=dely*force;
|
|
f.z+=delz*force;
|
|
|
|
if (EVFLAG && eflag) {
|
|
#ifndef ONETYPE
|
|
numtyp2 pre_sw_c56=c_56[mtype];
|
|
#endif
|
|
energy+=(pre_sw_c56.x*rp - pre_sw_c56.y*rq) * expsrainv;
|
|
}
|
|
|
|
if (EVFLAG && vflag) {
|
|
virial[0] += delx*delx*force;
|
|
virial[1] += dely*dely*force;
|
|
virial[2] += delz*delz*force;
|
|
virial[3] += delx*dely*force;
|
|
virial[4] += delx*delz*force;
|
|
virial[5] += dely*delz*force;
|
|
}
|
|
} // for nbor
|
|
} // if ii
|
|
store_answers_p(f,energy,virial,ii,inum,tid,t_per_atom,offset,eflag,vflag,
|
|
ans,engv,ev_stride);
|
|
}
|
|
|
|
#define threebody(delr1x,delr1y,delr1z,delr2x,delr2y,delr2z, eflag, energy) \
|
|
{ \
|
|
numtyp r1 = ucl_sqrt(rsq1); \
|
|
numtyp rinvsq1 = ucl_recip(rsq1); \
|
|
numtyp rainv1 = ucl_recip(r1 - sw_cut_ij); \
|
|
numtyp gsrainv1 = sw_sigma_gamma_ij * rainv1; \
|
|
numtyp gsrainvsq1 = gsrainv1*rainv1/r1; \
|
|
numtyp expgsrainv1 = ucl_exp(gsrainv1); \
|
|
\
|
|
numtyp r2 = ucl_sqrt(rsq2); \
|
|
numtyp rinvsq2 = ucl_recip(rsq2); \
|
|
numtyp rainv2 = ucl_recip(r2 - sw_cut_ik); \
|
|
numtyp gsrainv2 = sw_sigma_gamma_ik * rainv2; \
|
|
numtyp gsrainvsq2 = gsrainv2*rainv2/r2; \
|
|
numtyp expgsrainv2 = ucl_exp(gsrainv2); \
|
|
\
|
|
numtyp rinv12 = ucl_recip(r1*r2); \
|
|
numtyp cs = (delr1x*delr2x + delr1y*delr2y + delr1z*delr2z) * rinv12; \
|
|
numtyp delcs = cs - sw_costheta_ijk; \
|
|
numtyp delcssq = delcs*delcs; \
|
|
\
|
|
numtyp facexp = expgsrainv1*expgsrainv2; \
|
|
\
|
|
numtyp facrad = sw_lambda_epsilon_ijk * facexp*delcssq; \
|
|
numtyp frad1 = facrad*gsrainvsq1; \
|
|
numtyp frad2 = facrad*gsrainvsq2; \
|
|
numtyp facang = (numtyp)2.0 * sw_lambda_epsilon_ijk * facexp*delcs; \
|
|
numtyp facang12 = rinv12*facang; \
|
|
numtyp csfacang = cs*facang; \
|
|
numtyp csfac1 = rinvsq1*csfacang; \
|
|
\
|
|
fjx = delr1x*(frad1+csfac1)-delr2x*facang12; \
|
|
fjy = delr1y*(frad1+csfac1)-delr2y*facang12; \
|
|
fjz = delr1z*(frad1+csfac1)-delr2z*facang12; \
|
|
\
|
|
numtyp csfac2 = rinvsq2*csfacang; \
|
|
\
|
|
fkx = delr2x*(frad2+csfac2)-delr1x*facang12; \
|
|
fky = delr2y*(frad2+csfac2)-delr1y*facang12; \
|
|
fkz = delr2z*(frad2+csfac2)-delr1z*facang12; \
|
|
\
|
|
if (EVFLAG && eflag) \
|
|
energy+=facrad; \
|
|
if (EVFLAG && vflag) { \
|
|
virial[0] += delr1x*fjx + delr2x*fkx; \
|
|
virial[1] += delr1y*fjy + delr2y*fky; \
|
|
virial[2] += delr1z*fjz + delr2z*fkz; \
|
|
virial[3] += delr1x*fjy + delr2x*fky; \
|
|
virial[4] += delr1x*fjz + delr2x*fkz; \
|
|
virial[5] += delr1y*fjz + delr2y*fkz; \
|
|
} \
|
|
}
|
|
|
|
#define threebody_half(delr1x, delr1y, delr1z, delr2x, delr2y, delr2z) \
|
|
{ \
|
|
numtyp r1 = ucl_sqrt(rsq1); \
|
|
numtyp rinvsq1 = ucl_recip(rsq1); \
|
|
numtyp rainv1 = ucl_recip(r1 - sw_cut_ij); \
|
|
numtyp gsrainv1 = sw_sigma_gamma_ij * rainv1; \
|
|
numtyp gsrainvsq1 = gsrainv1*rainv1/r1; \
|
|
numtyp expgsrainv1 = ucl_exp(gsrainv1); \
|
|
\
|
|
numtyp r2 = ucl_sqrt(rsq2); \
|
|
numtyp rainv2 = ucl_recip(r2 - sw_cut_ik); \
|
|
numtyp gsrainv2 = sw_sigma_gamma_ik * rainv2; \
|
|
numtyp expgsrainv2 = ucl_exp(gsrainv2); \
|
|
\
|
|
numtyp rinv12 = ucl_recip(r1*r2); \
|
|
numtyp cs = (delr1x*delr2x + delr1y*delr2y + delr1z*delr2z) * rinv12; \
|
|
numtyp delcs = cs - sw_costheta_ijk; \
|
|
numtyp delcssq = delcs*delcs; \
|
|
\
|
|
numtyp facexp = expgsrainv1*expgsrainv2; \
|
|
\
|
|
numtyp facrad = sw_lambda_epsilon_ijk * facexp*delcssq; \
|
|
numtyp frad1 = facrad*gsrainvsq1; \
|
|
numtyp facang = (numtyp)2.0 * sw_lambda_epsilon_ijk * facexp*delcs; \
|
|
numtyp facang12 = rinv12*facang; \
|
|
numtyp csfacang = cs*facang; \
|
|
numtyp csfac1 = rinvsq1*csfacang; \
|
|
\
|
|
fjx = delr1x*(frad1+csfac1)-delr2x*facang12; \
|
|
fjy = delr1y*(frad1+csfac1)-delr2y*facang12; \
|
|
fjz = delr1z*(frad1+csfac1)-delr2z*facang12; \
|
|
}
|
|
|
|
#ifdef ONETYPE
|
|
#define sw_cut_ij sw_cut
|
|
#define sw_cut_ik sw_cut
|
|
#define sw_sigma_gamma_ij sw_sigma_gamma
|
|
#define sw_sigma_gamma_ik sw_sigma_gamma
|
|
#endif
|
|
|
|
__kernel void k_sw_three_center(const __global numtyp4 *restrict x_,
|
|
const __global numtyp2 *restrict cut_sig_gamma,
|
|
const __global numtyp2 *restrict sw_pre3,
|
|
const int ntypes,
|
|
const __global int * dev_nbor,
|
|
__global acctyp3 *restrict ans,
|
|
__global acctyp *restrict engv,
|
|
const int eflag, const int vflag,
|
|
const int inum, const int nbor_pitch,
|
|
const int t_per_atom, const int evatom) {
|
|
int n_stride;
|
|
const int tpa_sq=fast_mul(t_per_atom,t_per_atom);
|
|
local_allocate_store_three();
|
|
|
|
int tid, ii, offset;
|
|
atom_info(tpa_sq,ii,tid,offset);
|
|
|
|
#ifdef ONETYPE
|
|
const numtyp sw_cut=cut_sig_gamma[ONETYPE].x;
|
|
const numtyp sw_sigma_gamma=cut_sig_gamma[ONETYPE].y;
|
|
const numtyp sw_lambda_epsilon_ijk=sw_pre3[ONETYPE3].x;
|
|
const numtyp sw_costheta_ijk=sw_pre3[ONETYPE3].y;
|
|
#endif
|
|
|
|
acctyp3 f;
|
|
f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0;
|
|
acctyp energy, virial[6];
|
|
if (EVFLAG) {
|
|
energy=(acctyp)0;
|
|
for (int i=0; i<6; i++) virial[i]=(acctyp)0;
|
|
}
|
|
|
|
if (ii<inum) {
|
|
int i, numj, nbor_j, nbor_end;
|
|
int offset_j=offset/t_per_atom;
|
|
nbor_info_p(dev_nbor,nbor_pitch,t_per_atom,ii,offset_j,i,numj,
|
|
n_stride,nbor_end,nbor_j);
|
|
int offset_k=tid & (t_per_atom-1);
|
|
|
|
numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
|
|
#ifndef ONETYPE
|
|
int itype=ix.w;
|
|
itype*=ntypes;
|
|
#endif
|
|
|
|
for ( ; nbor_j<nbor_end; nbor_j+=n_stride) {
|
|
|
|
int j=dev_nbor[nbor_j];
|
|
j &= NEIGHMASK;
|
|
|
|
numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
|
|
#ifndef ONETYPE
|
|
int mtypej=jx.w;
|
|
mtypej+=itype;
|
|
#endif
|
|
|
|
// Compute r12
|
|
numtyp delr1x = jx.x-ix.x;
|
|
numtyp delr1y = jx.y-ix.y;
|
|
numtyp delr1z = jx.z-ix.z;
|
|
numtyp rsq1 = delr1x*delr1x+delr1y*delr1y+delr1z*delr1z;
|
|
|
|
#ifndef ONETYPE
|
|
const numtyp sw_cut_ij=cut_sig_gamma[mtypej].x;
|
|
const numtyp sw_sigma_gamma_ij=cut_sig_gamma[mtypej].y;
|
|
#endif
|
|
|
|
int nbor_k;
|
|
nbor_k = nbor_j-offset_j+offset_k;
|
|
if (nbor_k<=nbor_j) nbor_k += n_stride;
|
|
|
|
for ( ; nbor_k<nbor_end; nbor_k+=n_stride) {
|
|
int k=dev_nbor[nbor_k];
|
|
k &= NEIGHMASK;
|
|
|
|
numtyp4 kx; fetch4(kx,k,pos_tex);
|
|
#ifndef ONETYPE
|
|
const int ktype=kx.w;
|
|
const int mtypek=itype+ktype;
|
|
#endif
|
|
|
|
numtyp delr2x = kx.x-ix.x;
|
|
numtyp delr2y = kx.y-ix.y;
|
|
numtyp delr2z = kx.z-ix.z;
|
|
numtyp rsq2 = delr2x*delr2x + delr2y*delr2y + delr2z*delr2z;
|
|
#ifndef ONETYPE
|
|
const numtyp sw_cut_ik=cut_sig_gamma[mtypek].x;
|
|
const numtyp sw_sigma_gamma_ik=cut_sig_gamma[mtypek].y;
|
|
const int mtypejk=ntypes*mtypej+ktype;
|
|
const numtyp sw_lambda_epsilon_ijk=sw_pre3[mtypejk].x;
|
|
const numtyp sw_costheta_ijk=sw_pre3[mtypejk].y;
|
|
#endif
|
|
|
|
numtyp fjx, fjy, fjz, fkx, fky, fkz;
|
|
threebody(delr1x,delr1y,delr1z,delr2x,delr2y,delr2z,eflag,energy);
|
|
|
|
f.x -= fjx + fkx;
|
|
f.y -= fjy + fky;
|
|
f.z -= fjz + fkz;
|
|
}
|
|
} // for nbor
|
|
|
|
if (EVFLAG) {
|
|
numtyp pre;
|
|
if (evatom==1)
|
|
pre=THIRD;
|
|
else
|
|
pre=(numtyp)2.0;
|
|
energy*=pre;
|
|
if (vflag)
|
|
for (int i=0; i<6; i++)
|
|
virial[i]*=pre;
|
|
}
|
|
} // if ii
|
|
store_answers(f,energy,virial,ii,inum,tid,tpa_sq,offset,
|
|
eflag,vflag,ans,engv);
|
|
}
|
|
|
|
__kernel void k_sw_three_end(const __global numtyp4 *restrict x_,
|
|
const __global numtyp2 *restrict cut_sig_gamma,
|
|
const __global numtyp2 *restrict sw_pre3,
|
|
const int ntypes, const __global int * dev_nbor,
|
|
const __global int * dev_ilist,
|
|
__global acctyp3 *restrict ans,
|
|
__global acctyp *restrict engv,
|
|
const int eflag, const int vflag,
|
|
const int inum, const int nbor_pitch,
|
|
const int t_per_atom, const int gpu_nbor) {
|
|
int n_stride;
|
|
const int tpa_sq=fast_mul(t_per_atom,t_per_atom);
|
|
int tid, ii, offset;
|
|
atom_info(tpa_sq,ii,tid,offset);
|
|
|
|
local_allocate_store_three();
|
|
|
|
#ifdef ONETYPE
|
|
const numtyp sw_cut=cut_sig_gamma[ONETYPE].x;
|
|
const numtyp sw_sigma_gamma=cut_sig_gamma[ONETYPE].y;
|
|
const numtyp sw_lambda_epsilon_ijk=sw_pre3[ONETYPE3].x;
|
|
const numtyp sw_costheta_ijk=sw_pre3[ONETYPE3].y;
|
|
#endif
|
|
|
|
acctyp3 f;
|
|
f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0;
|
|
acctyp energy, virial[6];
|
|
if (EVFLAG) {
|
|
energy=(acctyp)0;
|
|
for (int i=0; i<6; i++) virial[i]=(acctyp)0;
|
|
}
|
|
|
|
if (ii<inum) {
|
|
int i, numj, nbor_j, nbor_end, k_end;
|
|
int offset_j=offset/t_per_atom;
|
|
nbor_info_p(dev_nbor,nbor_pitch,t_per_atom,ii,offset_j,i,numj,
|
|
n_stride,nbor_end,nbor_j);
|
|
int offset_k=tid & (t_per_atom-1);
|
|
|
|
numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
|
|
#ifndef ONETYPE
|
|
int itype=ix.w;
|
|
itype*=ntypes;
|
|
#endif
|
|
|
|
for ( ; nbor_j<nbor_end; nbor_j+=n_stride) {
|
|
int j=dev_nbor[nbor_j];
|
|
j &= NEIGHMASK;
|
|
|
|
numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
|
|
#ifndef ONETYPE
|
|
const int jtype=jx.w;
|
|
const int mtypej=itype+jtype;
|
|
#endif
|
|
|
|
// Compute r12
|
|
numtyp delr1x = ix.x-jx.x;
|
|
numtyp delr1y = ix.y-jx.y;
|
|
numtyp delr1z = ix.z-jx.z;
|
|
numtyp rsq1 = delr1x*delr1x+delr1y*delr1y+delr1z*delr1z;
|
|
|
|
#ifndef ONETYPE
|
|
const numtyp sw_cut_ij=cut_sig_gamma[mtypej].x;
|
|
const numtyp sw_sigma_gamma_ij=cut_sig_gamma[mtypej].y;
|
|
#endif
|
|
|
|
int nbor_k;
|
|
if (gpu_nbor) nbor_k=j+nbor_pitch;
|
|
else nbor_k=dev_ilist[j]+nbor_pitch;
|
|
const int numk=dev_nbor[nbor_k];
|
|
nbor_k+=nbor_pitch+fast_mul(j,t_per_atom-1);
|
|
k_end=nbor_k+fast_mul(numk/t_per_atom,n_stride)+(numk&(t_per_atom-1));
|
|
nbor_k+=offset_k;
|
|
|
|
for ( ; nbor_k<k_end; nbor_k+=n_stride) {
|
|
int k=dev_nbor[nbor_k];
|
|
k &= NEIGHMASK;
|
|
|
|
if (k == i) continue;
|
|
|
|
numtyp4 kx; fetch4(kx,k,pos_tex);
|
|
#ifndef ONETYPE
|
|
const int ktype=kx.w;
|
|
const int mtypek=jtype*ntypes+ktype;
|
|
const numtyp sw_cut_ik=cut_sig_gamma[mtypek].x;
|
|
#endif
|
|
|
|
numtyp delr2x = kx.x - jx.x;
|
|
numtyp delr2y = kx.y - jx.y;
|
|
numtyp delr2z = kx.z - jx.z;
|
|
numtyp rsq2 = delr2x*delr2x + delr2y*delr2y + delr2z*delr2z;
|
|
|
|
// for neigh no within pair hybrid: still need the check below
|
|
// this loop apparently iterates over neighbors of j that are not intended
|
|
const numtyp cutsq_jk=sw_cut_ik*sw_cut_ik;
|
|
if (!(rsq2>(numtyp)0 && rsq2<cutsq_jk)) continue;
|
|
|
|
#ifndef ONETYPE
|
|
const numtyp sw_sigma_gamma_ik=cut_sig_gamma[mtypek].y;
|
|
const int mtypejik=jtype*ntypes*ntypes+itype+ktype;
|
|
const numtyp sw_lambda_epsilon_ijk=sw_pre3[mtypejik].x;
|
|
const numtyp sw_costheta_ijk=sw_pre3[mtypejik].y;
|
|
#endif
|
|
|
|
numtyp fjx, fjy, fjz;
|
|
threebody_half(delr1x,delr1y,delr1z,delr2x,delr2y,delr2z);
|
|
|
|
f.x += fjx;
|
|
f.y += fjy;
|
|
f.z += fjz;
|
|
}
|
|
} // for nbor
|
|
} // if ii
|
|
#ifdef THREE_CONCURRENT
|
|
store_answers(f,energy,virial,ii,inum,tid,tpa_sq,offset,
|
|
eflag,vflag,ans,engv);
|
|
#else
|
|
store_answers_p(f,energy,virial,ii,inum,tid,tpa_sq,offset,
|
|
eflag,vflag,ans,engv,NUM_BLOCKS_X);
|
|
#endif
|
|
}
|
|
|
|
__kernel void k_sw_three_end_vatom(const __global numtyp4 *restrict x_,
|
|
const __global numtyp2 *restrict cut_sig_gamma,
|
|
const __global numtyp2 *restrict sw_pre3,
|
|
const int ntypes, const __global int * dev_nbor,
|
|
const __global int * dev_ilist,
|
|
__global acctyp3 *restrict ans,
|
|
__global acctyp *restrict engv,
|
|
const int eflag, const int vflag,
|
|
const int inum, const int nbor_pitch,
|
|
const int t_per_atom, const int gpu_nbor) {
|
|
int n_stride;
|
|
const int tpa_sq=fast_mul(t_per_atom,t_per_atom);
|
|
int tid, ii, offset;
|
|
atom_info(tpa_sq,ii,tid,offset);
|
|
|
|
local_allocate_store_three();
|
|
|
|
#ifdef ONETYPE
|
|
const numtyp sw_cut=cut_sig_gamma[ONETYPE].x;
|
|
const numtyp sw_sigma_gamma=cut_sig_gamma[ONETYPE].y;
|
|
const numtyp sw_lambda_epsilon_ijk=sw_pre3[ONETYPE3].x;
|
|
const numtyp sw_costheta_ijk=sw_pre3[ONETYPE3].y;
|
|
#endif
|
|
|
|
acctyp3 f;
|
|
f.x=(acctyp)0; f.y=(acctyp)0; f.z=(acctyp)0;
|
|
acctyp energy, virial[6];
|
|
if (EVFLAG) {
|
|
energy=(acctyp)0;
|
|
for (int i=0; i<6; i++) virial[i]=(acctyp)0;
|
|
}
|
|
|
|
if (ii<inum) {
|
|
int i, numj, nbor_j, nbor_end, k_end;
|
|
int offset_j=offset/t_per_atom;
|
|
nbor_info_p(dev_nbor,nbor_pitch,t_per_atom,ii,offset_j,i,numj,
|
|
n_stride,nbor_end,nbor_j);
|
|
int offset_k=tid & (t_per_atom-1);
|
|
|
|
numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
|
|
#ifndef ONETYPE
|
|
int itype=ix.w;
|
|
itype*=ntypes;
|
|
#endif
|
|
|
|
for ( ; nbor_j<nbor_end; nbor_j+=n_stride) {
|
|
int j=dev_nbor[nbor_j];
|
|
j &= NEIGHMASK;
|
|
|
|
numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
|
|
#ifndef ONETYPE
|
|
const int jtype=jx.w;
|
|
const int mtypej=itype+jtype;
|
|
#endif
|
|
|
|
// Compute r12
|
|
numtyp delr1x = ix.x-jx.x;
|
|
numtyp delr1y = ix.y-jx.y;
|
|
numtyp delr1z = ix.z-jx.z;
|
|
numtyp rsq1 = delr1x*delr1x+delr1y*delr1y+delr1z*delr1z;
|
|
|
|
#ifndef ONETYPE
|
|
const numtyp sw_cut_ij=cut_sig_gamma[mtypej].x;
|
|
const numtyp sw_sigma_gamma_ij=cut_sig_gamma[mtypej].y;
|
|
#endif
|
|
|
|
int nbor_k;
|
|
if (gpu_nbor) nbor_k=j+nbor_pitch;
|
|
else nbor_k=dev_ilist[j]+nbor_pitch;
|
|
const int numk=dev_nbor[nbor_k];
|
|
nbor_k+=nbor_pitch+fast_mul(j,t_per_atom-1);
|
|
k_end=nbor_k+fast_mul(numk/t_per_atom,n_stride)+(numk&(t_per_atom-1));
|
|
nbor_k+=offset_k;
|
|
|
|
for ( ; nbor_k<k_end; nbor_k+=n_stride) {
|
|
int k=dev_nbor[nbor_k];
|
|
k &= NEIGHMASK;
|
|
|
|
if (k == i) continue;
|
|
|
|
numtyp4 kx; fetch4(kx,k,pos_tex);
|
|
#ifndef ONETYPE
|
|
const int ktype=kx.w;
|
|
const int mtypek=jtype*ntypes+ktype;
|
|
const numtyp sw_cut_ik=cut_sig_gamma[mtypek].x;
|
|
#endif
|
|
|
|
numtyp delr2x = kx.x - jx.x;
|
|
numtyp delr2y = kx.y - jx.y;
|
|
numtyp delr2z = kx.z - jx.z;
|
|
numtyp rsq2 = delr2x*delr2x + delr2y*delr2y + delr2z*delr2z;
|
|
|
|
// for neigh no within pair hybrid: still need the check below
|
|
// this loop apparently iterates over neighbors of j that are not intended
|
|
const numtyp cutsq_jk=sw_cut_ik*sw_cut_ik;
|
|
if (!(rsq2>(numtyp)0 && rsq2<cutsq_jk)) continue;
|
|
|
|
#ifndef ONETYPE
|
|
const numtyp sw_sigma_gamma_ik=cut_sig_gamma[mtypek].y;
|
|
const int mtypejik=jtype*ntypes*ntypes+itype+ktype;
|
|
const numtyp sw_lambda_epsilon_ijk=sw_pre3[mtypejik].x;
|
|
const numtyp sw_costheta_ijk=sw_pre3[mtypejik].y;
|
|
#endif
|
|
|
|
numtyp fjx, fjy, fjz, fkx, fky, fkz;
|
|
threebody(delr1x,delr1y,delr1z,delr2x,delr2y,delr2z,eflag,energy);
|
|
|
|
f.x += fjx;
|
|
f.y += fjy;
|
|
f.z += fjz;
|
|
}
|
|
} // for nbor
|
|
energy*=THIRD;
|
|
for (int i=0; i<6; i++)
|
|
virial[i]*=THIRD;
|
|
} // if ii
|
|
#ifdef THREE_CONCURRENT
|
|
store_answers(f,energy,virial,ii,inum,tid,tpa_sq,offset,
|
|
eflag,vflag,ans,engv);
|
|
#else
|
|
store_answers_p(f,energy,virial,ii,inum,tid,tpa_sq,offset,
|
|
eflag,vflag,ans,engv,NUM_BLOCKS_X);
|
|
#endif
|
|
}
|