Add copyright and fix style (4febc7f7) · Commits · 郑智淋 / lammps

doc/src/pair_lj.rst

+3 −0

Original line number	Diff line number	Diff line
		@@ -102,6 +102,9 @@ pair\_style lj/cut/tip4p/long/omp command
		pair\_style lj/cut/tip4p/long/opt command
		=========================================

		pair\_style lj/cut/tip4p/long/gpu command
		=====================================

		Syntax
		""""""

lib/gpu/lal_lj_tip4p_long.cpp

+72 −57

Original line number	Diff line number	Diff line
		/**************************************************************************
		lj_tip4p_long.cpp
		-------------------
		V. Nikolskiy (HSE)

		Class for acceleration of the lj/tip4p/long pair style

		__________________________________________________________________________
		This file is part of the LAMMPS Accelerator Library (LAMMPS_AL)
		__________________________________________________________________________

		begin :
		email : thevsevak@gmail.com
		***************************************************************************/

		#if defined(USE_OPENCL)
		#include "lj_tip4p_long_cl.h"
		#elif defined(USE_CUDART)
		@@ -91,7 +106,7 @@ int LJ_TIP4PLong<numtyp, acctyp>::init(const int ntypes,
		}
		ucl_copy(sp_lj,host_write,8,false);

		force_comp.alloc(7272, (this->ucl_device), UCL_READ_WRITE);
		//force_comp.alloc(7272, (this->ucl_device), UCL_READ_WRITE);

		_qqrd2e=qqrd2e;
		_g_ewald=g_ewald;
		@@ -103,18 +118,18 @@ int LJ_TIP4PLong<numtyp, acctyp>::init(const int ntypes,
		ansO.alloc(nall,*(this->ucl_device), UCL_READ_WRITE);

		// Allocate a host write buffer for data initialization
		UCL_H_Vec<int> host_tag_write(nall,*(this->ucl_device),UCL_WRITE_ONLY);
		this->tag.alloc(nall,*(this->ucl_device), UCL_READ_WRITE);
		UCL_H_Vec<int> host_tag_write(nall,*(this->ucl_device),UCL_READ_WRITE);
		this->tag.alloc(nall,*(this->ucl_device), UCL_READ_ONLY);
		for(int i=0; i<nall; ++i) host_tag_write[i] = tag[i];
		ucl_copy(this->tag, host_tag_write, nall, false);

		//if(max_same>host_tag_write.cols()) host_tag_write.resize(max_same);
		this->atom_sametag.alloc(nall, *(this->ucl_device), UCL_READ_WRITE);
		this->atom_sametag.alloc(nall, *(this->ucl_device), UCL_READ_ONLY);
		for(int i=0; i<nall; ++i) host_tag_write[i] = sametag[i];
		ucl_copy(this->atom_sametag, host_tag_write, nall, false);

		if(map_size>host_tag_write.cols()) host_tag_write.resize(map_size);
		this->map_array.alloc(map_size,*(this->ucl_device), UCL_READ_WRITE);
		host_tag_write.resize_ib(map_size);
		this->map_array.alloc(map_size,*(this->ucl_device), UCL_READ_ONLY);
		for(int i=0; i<map_size; ++i) host_tag_write[i] = map_array[i];
		ucl_copy(this->map_array, host_tag_write, map_size, false);

		@@ -143,7 +158,7 @@ void LJ_TIP4PLong<numtyp, acctyp>::clear() {
		atom_sametag.clear();
		map_array.clear();
		ansO.clear();
		force_comp.clear();
		//force_comp.clear();

		k_pair_distrib.clear();

		@@ -211,10 +226,10 @@ void LJ_TIP4PLong<numtyp, acctyp>::copy_relations_data(int hn, double newsit
		int* tag, int map_array, int map_size, int sametag, int max_same, int ago){
		int nall = n;
		const int hn_sz = n4; // matrix size = col size col number
		hneight.resize_ib(hn_sz+1);
		hneight.resize_ib(hn_sz);
		if (ago == 0)
		hneight.zero();
		m.resize_ib(n+1);
		m.resize_ib(n);
		m.zero();

		UCL_H_Vec<int> host_tag_write(nall,*(this->ucl_device),UCL_WRITE_ONLY);
		@@ -222,12 +237,12 @@ void LJ_TIP4PLong<numtyp, acctyp>::copy_relations_data(int hn, double newsit
		for(int i=0; i<nall; ++i) host_tag_write[i] = tag[i];
		ucl_copy(this->tag, host_tag_write, nall, false);

		if(max_same>host_tag_write.cols()) host_tag_write.resize(max_same);
		host_tag_write.resize_ib(max_same);
		this->atom_sametag.resize_ib(nall);
		for(int i=0; i<nall; ++i) host_tag_write[i] = sametag[i];
		ucl_copy(this->atom_sametag, host_tag_write, nall, false);

		if(map_size>host_tag_write.cols()) host_tag_write.resize(map_size);
		host_tag_write.resize_ib(map_size);
		this->map_array.resize_ib(map_size);
		for(int i=0; i<map_size; ++i) host_tag_write[i] = map_array[i];
		ucl_copy(this->map_array, host_tag_write, map_size, false);

lib/gpu/lal_lj_tip4p_long.cu

+485 −465

Original line number	Diff line number	Diff line
		// **************************************************************************
		// lj_tip4p_long.cu
		// -------------------
		// V. Nikolskiy (HSE)
		//
		// Device code for acceleration of the lj/tip4p/long pair style
		//
		// __________________________________________________________________________
		// This file is part of the LAMMPS Accelerator Library (LAMMPS_AL)
		// __________________________________________________________________________
		//
		// begin :
		// email : thevsevak@gmail.com
		// ***************************************************************************/

		#ifdef NV_KERNEL

		#include "lal_aux_fun1.h"
		@@ -14,6 +29,8 @@ texture<int2> q_tex;
		#define q_tex q_
		#endif

		#include <float.h>

		ucl_inline int atom_mapping(const __global int *map, int glob){
		return map[glob];
		}
		@@ -213,7 +230,7 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
		iH1 = hneigh[i*4 ];
		iH2 = hneigh[i*4+1];
		}
		if(m[iO].w == 0) {
		if(fabs(m[iO].w) <= FLT_EPSILON) {
		compute_newsite(iO,iH1,iH2, &m[iO], alpha, x_);
		m[iO].w = qtmp;
		}
		@@ -223,12 +240,12 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
		int iI, iH;
		iI = atom_mapping(map,tag[i] - 1);
		numtyp4 iIx; fetch4(iIx,iI,pos_tex); //x_[iI];
		if (iIx.w == typeH) {
		if ((int)iIx.w == typeH) {
		iO = atom_mapping(map,tag[i] - 2);
		iO = closest_image(i, iO, sametag, x_);
		iH1 = closest_image(i, iI, sametag, x_);
		iH2 = i;
		} else if (iIx.w == typeO) {
		} else { //if ((int)iIx.w == typeO)
		iH = atom_mapping(map, tag[i] + 1);
		iO = closest_image(i,iI,sametag, x_);
		iH1 = i;
		@@ -253,7 +270,7 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
		}
		if (iO >= inum) {
		non_local_oxy = 1;
		if(m[iO].w == 0) {
		if(fabs(m[iO].w) <= FLT_EPSILON) {
		compute_newsite(iO,iH1,iH2, &m[iO], alpha, x_);
		numtyp qO; fetch(qO,iO,q_tex);
		m[iO].w = qO;
		@@ -329,7 +346,7 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
		jH1 = hneigh[j*4 ];
		jH2 = hneigh[j*4+1];
		}
		if (m[j].w == 0) {
		if (fabs(m[j].w) <= FLT_EPSILON) {
		compute_newsite(j, jH1, jH2, &m[j], alpha, x_);
		m[j].w = qj;
		}
		@@ -356,11 +373,12 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
		f.x += delx * force_coul;
		f.y += dely * force_coul;
		f.z += delz * force_coul;
		f.w += 0;
		} else {
		fO.x += delx * force_coul;
		fO.y += dely * force_coul;
		fO.z += delz * force_coul;
		fO.w += -1;
		fO.w += 0;
		}
		if (eflag>0) {
		e_coul += prefactor*(_erfc-factor_coul);
		@@ -387,6 +405,7 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
		vdj.x = xjO.xcO + xjH1.xcH + xjH2.x*cH;
		vdj.y = xjO.ycO + xjH1.ycH + xjH2.y*cH;
		vdj.z = xjO.zcO + xjH1.zcH + xjH2.z*cH;
		vdj.w = vdj.w;
		virial[0] += (ix.x - vdj.x)*fd.x;
		virial[1] += (ix.y - vdj.y)*fd.y;
		virial[2] += (ix.z - vdj.z)*fd.z;
		@@ -403,6 +422,7 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
		vdi.x = xO.xcO + xH1.xcH + xH2.x*cH;
		vdi.y = xO.ycO + xH1.ycH + xH2.y*cH;
		vdi.z = xO.zcO + xH1.zcH + xH2.z*cH;
		vdi.w = vdi.w;
		if (jtype != typeH){
		numtyp4 xjH1; fetch4(xjH1,jH1,pos_tex);
		numtyp4 xjH2; fetch4(xjH2,jH2,pos_tex);
		@@ -410,6 +430,7 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
		vdj.x = xjO.xcO + xjH1.xcH + xjH2.x*cH;
		vdj.y = xjO.ycO + xjH1.ycH + xjH2.y*cH;
		vdj.z = xjO.zcO + xjH1.zcH + xjH2.z*cH;
		vdj.w = vdj.w;
		} else vdj = jx;
		vO[0] += 0.5(vdi.x - vdj.x)fd.x;
		vO[1] += 0.5(vdi.y - vdj.y)fd.y;
		@@ -529,5 +550,4 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
		} // if ii
		}


		__kernel void k_lj_tip4p_long_fast(){}

lib/gpu/lal_lj_tip4p_long.h

+25 −10

Original line number	Diff line number	Diff line
		/**************************************************************************
		lj_tip4p_long.h
		-------------------
		V. Nikolskiy (HSE)

		Class for acceleration of the lj/tip4p/long pair style

		__________________________________________________________________________
		This file is part of the LAMMPS Accelerator Library (LAMMPS_AL)
		__________________________________________________________________________

		begin :
		email : thevsevak@gmail.com
		***************************************************************************/

		#ifndef LAL_LJ_TIP4P_LONG_H
		#define LAL_LJ_TIP4P_LONG_H

		@@ -77,7 +92,7 @@ public:
		UCL_D_Vec<int> hneight;
		UCL_D_Vec<numtyp4> m; // position and charge of virtual particle
		UCL_D_Vec<acctyp4> ansO; // force applied to virtual particle
		UCL_D_Vec<acctyp4> force_comp;
		// UCL_D_Vec<acctyp4> force_comp;

		UCL_D_Vec<int> tag;
		UCL_D_Vec<int> map_array;

lib/gpu/lal_lj_tip4p_long_ext.cpp

+48 −35

Original line number	Diff line number	Diff line
		/***************************************************************************
		lj_tip4p_long_ext.cpp
		-------------------
		V. Nikolskiy (HSE)

		Functions for LAMMPS access to lj/tip4p/long acceleration functions

		__________________________________________________________________________
		This file is part of the LAMMPS Accelerator Library (LAMMPS_AL)
		__________________________________________________________________________

		begin :
		email : thevsevak@gmail.com
		***************************************************************************/

		#include <iostream>
		#include <cassert>
		#include <cmath>
		@@ -91,8 +106,6 @@ int ljtip4p_long_gpu_init(const int ntypes, double cutsq, double host_lj1,
		return init_ok;
		}



		void ljtip4p_long_gpu_clear() {
		LJTIP4PLMF.clear();
		}

Admin message