Built 3-body short neighbor list for the 3-body kernels using per-pair cutoffs... (cdac5f49) · Commits · 郑智淋 / lammps

lib/gpu/lal_vashishta.cpp

+2 −1

Original line number	Diff line number	Diff line
		@@ -233,7 +233,8 @@ void VashishtaT::loop(const bool _eflag, const bool _vflag, const int evatom) {
		(BX/this->_threads_per_atom)));

		this->k_short_nbor.set_size(GX,BX);
		this->k_short_nbor.run(&this->atom->x, &_cutshortsq,
		this->k_short_nbor.run(&this->atom->x, &param4, &map,
		&elem2param, &_nelements, &_nparams,
		&this->nbor->dev_nbor, &this->_nbor_data->begin(),
		&this->dev_short_nbor, &ainum,
		&nbor_pitch, &this->_threads_per_atom);

lib/gpu/lal_vashishta.cu

+17 −5

Original line number	Diff line number	Diff line
		@@ -137,7 +137,10 @@ texture<int4> param5_tex;
		#endif

		__kernel void k_vashishta_short_nbor(const __global numtyp4 *restrict x_,
		const numtyp cutshortsq,
		const __global numtyp4 *restrict param4,
		const __global int *restrict map,
		const __global int *restrict elem2param,
		const int nelements, const int nparams,
		const __global int * dev_nbor,
		const __global int * dev_packed,
		__global int * dev_short_nbor,
		@@ -154,6 +157,8 @@ __kernel void k_vashishta_short_nbor(const __global numtyp4 *restrict x_,
		n_stride,nbor_end,nbor);

		numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
		int itype=ix.w;
		itype=map[itype];

		int ncount = 0;
		int m = nbor;
		@@ -167,6 +172,9 @@ __kernel void k_vashishta_short_nbor(const __global numtyp4 *restrict x_,
		j &= NEIGHMASK;

		numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
		int jtype=jx.w;
		jtype=map[jtype];
		int ijparam=elem2param[itypenelementsnelements+jtype*nelements+jtype];

		// Compute r12
		numtyp delx = ix.x-jx.x;
		@@ -174,7 +182,7 @@ __kernel void k_vashishta_short_nbor(const __global numtyp4 *restrict x_,
		numtyp delz = ix.z-jx.z;
		numtyp rsq = delxdelx+delydely+delz*delz;

		if (rsq<cutshortsq) {
		if (rsq<param4[ijparam].x) { //param4[ijparam].x = r0sq; //param4[ijparam].z=cutsq
		dev_short_nbor[nbor_short] = nj;
		nbor_short += n_stride;
		ncount++;
		@@ -461,7 +469,9 @@ __kernel void k_vashishta_three_center(const __global numtyp4 *restrict x_,

		numtyp4 param4_ijparam; fetch4(param4_ijparam,ijparam,param4_tex);
		param_r0sq_ij=param4_ijparam.x;
		if (rsq1 > param_r0sq_ij) continue;

		// if (rsq1 > param_r0sq_ij) continue;

		param_gamma_ij=param4_ijparam.y;
		param_r0_ij=param4_ijparam.w;

		@@ -592,7 +602,8 @@ __kernel void k_vashishta_three_end(const __global numtyp4 *restrict x_,
		int ijparam=elem2param[itypenelementsnelements+jtype*nelements+jtype];
		numtyp4 param4_ijparam; fetch4(param4_ijparam,ijparam,param4_tex);
		param_r0sq_ij = param4_ijparam.x;
		if (rsq1 > param_r0sq_ij) continue;

		// if (rsq1 > param_r0sq_ij) continue;

		param_gamma_ij=param4_ijparam.y;
		param_r0_ij = param4_ijparam.w;
		@@ -742,7 +753,8 @@ __kernel void k_vashishta_three_end_vatom(const __global numtyp4 *restrict x_,
		int ijparam=elem2param[itypenelementsnelements+jtype*nelements+jtype];
		numtyp4 param4_ijparam; fetch4(param4_ijparam,ijparam,param4_tex);
		param_r0sq_ij=param4_ijparam.x;
		if (rsq1 > param_r0sq_ij) continue;

		// if (rsq1 > param_r0sq_ij) continue;

		param_gamma_ij=param4_ijparam.y;
		param_r0_ij=param4_ijparam.w;

Admin message