From cfe818a17564d063cc2e954b8490812d0155f3cb Mon Sep 17 00:00:00 2001 From: Dan Ibanez Date: Thu, 26 Jan 2017 08:37:56 -0700 Subject: [PATCH 01/25] remove unused variables from fix_cmap --- src/MOLECULE/fix_cmap.cpp | 11 ++++------- 1 file changed, 4 insertions(+), 7 deletions(-) diff --git a/src/MOLECULE/fix_cmap.cpp b/src/MOLECULE/fix_cmap.cpp index 26aaef60d7..8534a8bf33 100644 --- a/src/MOLECULE/fix_cmap.cpp +++ b/src/MOLECULE/fix_cmap.cpp @@ -230,7 +230,7 @@ void FixCMAP::min_setup(int vflag) void FixCMAP::pre_neighbor() { - int i,m,itype,atom1,atom2,atom3,atom4,atom5; + int i,m,atom1,atom2,atom3,atom4,atom5; // guesstimate initial length of local crossterm list // if ncmap was not set (due to read_restart, no read_data), @@ -842,7 +842,7 @@ void FixCMAP::set_map_derivatives(double **map, double **d1yo, double **d2yo, // use the bicubic spline to calculate the derivatives int i, j, k, ii, jj, xm, p; - double phi, psi, y, d1y, d2y, d12y, tyyk,tdyk; + double phi, psi, d1y, d2y, d12y, tyyk,tdyk; double *tmp_y, *tmp_dy, *tmp_ddy, **tmap, **tddmap; int ix; double a,b,a1,b1,a2,b2; @@ -850,7 +850,6 @@ void FixCMAP::set_map_derivatives(double **map, double **d1yo, double **d2yo, xm = CMAPDIM/2; p = CMAPDIM; - y = 0.; d1y = 0.; d2y = 0.; d12y = 0.; @@ -907,8 +906,6 @@ void FixCMAP::set_map_derivatives(double **map, double **d1yo, double **d2yo, b1 = b*b*b-b; a2 = 3.0*a*a-1.0; b2 = 3.0*b*b-1.0; - y = a*tmp_y[ix]+b*tmp_y[ix+1]+ - (a1*tmp_ddy[ix]+b1*tmp_ddy[ix+1])*(CMAPDX*CMAPDX)/6.0; d1y = (tmp_y[ix+1]-tmp_y[ix])/CMAPDX- a2/6.0*CMAPDX*tmp_ddy[ix]+b2/6.0*CMAPDX*tmp_ddy[ix+1]; spline(tmp_dy,tmp_ddy,CMAPDIM+xm+xm); @@ -1015,8 +1012,8 @@ void FixCMAP::bc_interpol(double x1, double x2, int low1, int low2, double *gs, // gradients and cross-derivatives // calculate the interpolated value of the point of interest (POI) - int i, p=12; - double t, u, fac, gs1l, gs2l, gs1u, gs2u; + int i; + double t, u, gs1l, gs2l; // set the interpolation coefficients From f0a4ed615dc5df1d481eeeb767afde7c1d92919b Mon Sep 17 00:00:00 2001 From: Dan Ibanez Date: Thu, 26 Jan 2017 08:46:03 -0700 Subject: [PATCH 02/25] add missing KOKKOS_INLINE_FUNCTION for params --- src/KOKKOS/pair_coul_long_kokkos.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/src/KOKKOS/pair_coul_long_kokkos.h b/src/KOKKOS/pair_coul_long_kokkos.h index 19789fbd86..92d08ddaba 100644 --- a/src/KOKKOS/pair_coul_long_kokkos.h +++ b/src/KOKKOS/pair_coul_long_kokkos.h @@ -45,7 +45,9 @@ class PairCoulLongKokkos : public PairCoulLong { double init_one(int, int); struct params_coul{ + KOKKOS_INLINE_FUNCTION params_coul(){cut_coulsq=0;}; + KOKKOS_INLINE_FUNCTION params_coul(int i){cut_coulsq=0;}; F_FLOAT cut_coulsq; }; From 6a74a81da023992be4eb90ee1b106204b092a505 Mon Sep 17 00:00:00 2001 From: Dan Ibanez Date: Thu, 26 Jan 2017 08:46:28 -0700 Subject: [PATCH 03/25] consistent #ifdefs for fft3d variable this variable is only used when FFTW3 is enabled, so its declaration and initialization should be protected under the same conditions to avoid compiler warnings --- src/KSPACE/fft3d.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/src/KSPACE/fft3d.cpp b/src/KSPACE/fft3d.cpp index a0896ed22d..907ab0c6ee 100644 --- a/src/KSPACE/fft3d.cpp +++ b/src/KSPACE/fft3d.cpp @@ -67,7 +67,10 @@ void fft_3d(FFT_DATA *in, FFT_DATA *out, int flag, struct fft_plan_3d *plan) { int i,total,length,offset,num; - FFT_SCALAR norm, *out_ptr; + FFT_SCALAR norm; +#if defined(FFT_FFTW3) + FFT_SCALAR *out_ptr; +#endif FFT_DATA *data,*copy; // system specific constants @@ -296,7 +299,9 @@ void fft_3d(FFT_DATA *in, FFT_DATA *out, int flag, struct fft_plan_3d *plan) if (flag == 1 && plan->scaled) { norm = plan->norm; num = plan->normnum; +#if defined(FFT_FFTW3) out_ptr = (FFT_SCALAR *)out; +#endif for (i = 0; i < num; i++) { #if defined(FFT_FFTW3) *(out_ptr++) *= norm; From cf95ea070973cf6718e574e0e33a2b89ac5b404b Mon Sep 17 00:00:00 2001 From: Dan Ibanez Date: Thu, 26 Jan 2017 08:50:18 -0700 Subject: [PATCH 04/25] fft3d: only declare variables when used avoids compiler warnings --- src/KSPACE/fft3d.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/src/KSPACE/fft3d.cpp b/src/KSPACE/fft3d.cpp index 907ab0c6ee..45bbf17507 100644 --- a/src/KSPACE/fft3d.cpp +++ b/src/KSPACE/fft3d.cpp @@ -366,13 +366,12 @@ struct fft_plan_3d *fft_3d_create_plan( { struct fft_plan_3d *plan; int me,nprocs; - int i,num,flag,remapflag,fftflag; + int flag,remapflag; int first_ilo,first_ihi,first_jlo,first_jhi,first_klo,first_khi; int second_ilo,second_ihi,second_jlo,second_jhi,second_klo,second_khi; int third_ilo,third_ihi,third_jlo,third_jhi,third_klo,third_khi; int out_size,first_size,second_size,third_size,copy_size,scratch_size; int np1,np2,ip1,ip2; - int list[50]; // system specific variables @@ -695,6 +694,9 @@ struct fft_plan_3d *fft_3d_create_plan( flag = 0; + int i,num,fftflag; + int list[50]; + num = 0; factor(nfast,&num,list); for (i = 0; i < num; i++) From 5b18421dd204a3ba6fb1fdb47c2419054b264c2d Mon Sep 17 00:00:00 2001 From: Dan Ibanez Date: Thu, 26 Jan 2017 08:54:59 -0700 Subject: [PATCH 05/25] fft3d : remove unused variables --- src/KSPACE/fft3d.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/KSPACE/fft3d.cpp b/src/KSPACE/fft3d.cpp index 45bbf17507..7936051f6d 100644 --- a/src/KSPACE/fft3d.cpp +++ b/src/KSPACE/fft3d.cpp @@ -1093,7 +1093,7 @@ void bifactor(int n, int *factor1, int *factor2) void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan) { - int i,total,length,offset,num; + int i,offset,num; FFT_SCALAR norm, *data_ptr; // system specific constants From 5b29515849adb590a408cbedabe6b374255ec4a2 Mon Sep 17 00:00:00 2001 From: Dan Ibanez Date: Thu, 26 Jan 2017 08:57:12 -0700 Subject: [PATCH 06/25] fft3d: use C++ loop declarations the variable (offset) is only used in a subset of numerous scenarios with #ifdef, it seems better just to have each loop declare it as needed. (avoids compiler warnings) --- src/KSPACE/fft3d.cpp | 50 ++++++++++++++++++++++---------------------- 1 file changed, 25 insertions(+), 25 deletions(-) diff --git a/src/KSPACE/fft3d.cpp b/src/KSPACE/fft3d.cpp index 7936051f6d..9751655f97 100644 --- a/src/KSPACE/fft3d.cpp +++ b/src/KSPACE/fft3d.cpp @@ -1093,7 +1093,7 @@ void bifactor(int n, int *factor1, int *factor2) void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan) { - int i,offset,num; + int i,num; FFT_SCALAR norm, *data_ptr; // system specific constants @@ -1139,20 +1139,20 @@ void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan) // data is just an array of 0.0 #ifdef FFT_SGI - for (offset = 0; offset < total1; offset += length1) + for (int offset = 0; offset < total1; offset += length1) FFT_1D(flag,length1,&data[offset],1,plan->coeff1); - for (offset = 0; offset < total2; offset += length2) + for (int offset = 0; offset < total2; offset += length2) FFT_1D(flag,length2,&data[offset],1,plan->coeff2); - for (offset = 0; offset < total3; offset += length3) + for (int offset = 0; offset < total3; offset += length3) FFT_1D(flag,length3,&data[offset],1,plan->coeff3); #elif defined(FFT_SCSL) - for (offset = 0; offset < total1; offset += length1) + for (int offset = 0; offset < total1; offset += length1) FFT_1D(flag,length1,scalef,&data[offset],&data[offset],plan->coeff1, plan->work1,&isys); - for (offset = 0; offset < total2; offset += length2) + for (int offset = 0; offset < total2; offset += length2) FFT_1D(flag,length2,scalef,&data[offset],&data[offset],plan->coeff2, plan->work2,&isys); - for (offset = 0; offset < total3; offset += length3) + for (int offset = 0; offset < total3; offset += length3) FFT_1D(flag,length3,scalef,&data[offset],&data[offset],plan->coeff3, plan->work3,&isys); #elif defined(FFT_ACML) @@ -1164,11 +1164,11 @@ void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan) num=total3/length3; FFT_1D(&flag,&num,&length3,data,plan->coeff3,&info); #elif defined(FFT_INTEL) - for (offset = 0; offset < total1; offset += length1) + for (int offset = 0; offset < total1; offset += length1) FFT_1D(&data[offset],&length1,&flag,plan->coeff1); - for (offset = 0; offset < total2; offset += length2) + for (int offset = 0; offset < total2; offset += length2) FFT_1D(&data[offset],&length2,&flag,plan->coeff2); - for (offset = 0; offset < total3; offset += length3) + for (int offset = 0; offset < total3; offset += length3) FFT_1D(&data[offset],&length3,&flag,plan->coeff3); #elif defined(FFT_MKL) if (flag == -1) { @@ -1182,28 +1182,28 @@ void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan) } #elif defined(FFT_DEC) if (flag == -1) { - for (offset = 0; offset < total1; offset += length1) + for (int offset = 0; offset < total1; offset += length1) FFT_1D(&c,&c,&f,&data[offset],&data[offset],&length1,&one); - for (offset = 0; offset < total2; offset += length2) + for (int offset = 0; offset < total2; offset += length2) FFT_1D(&c,&c,&f,&data[offset],&data[offset],&length2,&one); - for (offset = 0; offset < total3; offset += length3) + for (int offset = 0; offset < total3; offset += length3) FFT_1D(&c,&c,&f,&data[offset],&data[offset],&length3,&one); } else { - for (offset = 0; offset < total1; offset += length1) + for (int offset = 0; offset < total1; offset += length1) FFT_1D(&c,&c,&b,&data[offset],&data[offset],&length1,&one); - for (offset = 0; offset < total2; offset += length2) + for (int offset = 0; offset < total2; offset += length2) FFT_1D(&c,&c,&b,&data[offset],&data[offset],&length2,&one); - for (offset = 0; offset < total3; offset += length3) + for (int offset = 0; offset < total3; offset += length3) FFT_1D(&c,&c,&b,&data[offset],&data[offset],&length3,&one); } #elif defined(FFT_T3E) - for (offset = 0; offset < total1; offset += length1) + for (int offset = 0; offset < total1; offset += length1) FFT_1D(&flag,&length1,&scalef,&data[offset],&data[offset],plan->coeff1, plan->work1,&isys); - for (offset = 0; offset < total2; offset += length2) + for (int offset = 0; offset < total2; offset += length2) FFT_1D(&flag,&length2,&scalef,&data[offset],&data[offset],plan->coeff2, plan->work2,&isys); - for (offset = 0; offset < total3; offset += length3) + for (int offset = 0; offset < total3; offset += length3) FFT_1D(&flag,&length3,&scalef,&data[offset],&data[offset],plan->coeff3, plan->work3,&isys); #elif defined(FFT_FFTW2) @@ -1235,18 +1235,18 @@ void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan) FFTW_API(execute_dft)(theplan,data,data); #else if (flag == -1) { - for (offset = 0; offset < total1; offset += length1) + for (int offset = 0; offset < total1; offset += length1) kiss_fft(plan->cfg_fast_forward,&data[offset],&data[offset]); - for (offset = 0; offset < total2; offset += length2) + for (int offset = 0; offset < total2; offset += length2) kiss_fft(plan->cfg_mid_forward,&data[offset],&data[offset]); - for (offset = 0; offset < total3; offset += length3) + for (int offset = 0; offset < total3; offset += length3) kiss_fft(plan->cfg_slow_forward,&data[offset],&data[offset]); } else { - for (offset = 0; offset < total1; offset += length1) + for (int offset = 0; offset < total1; offset += length1) kiss_fft(plan->cfg_fast_backward,&data[offset],&data[offset]); - for (offset = 0; offset < total2; offset += length2) + for (int offset = 0; offset < total2; offset += length2) kiss_fft(plan->cfg_mid_backward,&data[offset],&data[offset]); - for (offset = 0; offset < total3; offset += length3) + for (int offset = 0; offset < total3; offset += length3) kiss_fft(plan->cfg_slow_backward,&data[offset],&data[offset]); } #endif From 2023fa28e01736802c5563cb6d31bfb7b3d9cb53 Mon Sep 17 00:00:00 2001 From: Dan Ibanez Date: Thu, 26 Jan 2017 10:03:54 -0700 Subject: [PATCH 07/25] consistent #ifdefs for fft3d variable (2) this variable is only used when FFTW3 is enabled, so its declaration and initialization should be protected under the same conditions to avoid compiler warnings --- src/KSPACE/fft3d.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/src/KSPACE/fft3d.cpp b/src/KSPACE/fft3d.cpp index 9751655f97..947d5dbe7f 100644 --- a/src/KSPACE/fft3d.cpp +++ b/src/KSPACE/fft3d.cpp @@ -1094,7 +1094,10 @@ void bifactor(int n, int *factor1, int *factor2) void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan) { int i,num; - FFT_SCALAR norm, *data_ptr; + FFT_SCALAR norm; +#if defined(FFT_FFTW3) + FFT_SCALAR *data_ptr; +#endif // system specific constants @@ -1258,7 +1261,9 @@ void fft_1d_only(FFT_DATA *data, int nsize, int flag, struct fft_plan_3d *plan) if (flag == 1 && plan->scaled) { norm = plan->norm; num = MIN(plan->normnum,nsize); +#if defined(FFT_FFTW3) data_ptr = (FFT_SCALAR *)data; +#endif for (i = 0; i < num; i++) { #if defined(FFT_FFTW3) *(data_ptr++) *= norm; From b5d622c6a381b95b6dfac063b6da896425c0d30f Mon Sep 17 00:00:00 2001 From: Dan Ibanez Date: Thu, 26 Jan 2017 10:18:45 -0700 Subject: [PATCH 08/25] pppm_kokkos: remove unused variables --- src/KOKKOS/pppm_kokkos.cpp | 24 +++--------------------- 1 file changed, 3 insertions(+), 21 deletions(-) diff --git a/src/KOKKOS/pppm_kokkos.cpp b/src/KOKKOS/pppm_kokkos.cpp index de9c0ae630..f923eb35c5 100644 --- a/src/KOKKOS/pppm_kokkos.cpp +++ b/src/KOKKOS/pppm_kokkos.cpp @@ -374,7 +374,6 @@ void PPPMKokkos::setup() error->all(FLERR,"Incorrect boundaries with slab PPPM"); } - int i,j,k,n; double *prd; // volume-dependent factors @@ -417,10 +416,6 @@ void PPPMKokkos::setup() DeviceType::fence(); copymode = 0; - // virial coefficients - - double sqk,vterm; - // merge three outer loops into one for better threading numz_fft = nzhi_fft-nzlo_fft + 1; @@ -615,7 +610,7 @@ void PPPMKokkos::setup_grid() template void PPPMKokkos::compute(int eflag, int vflag) { - int i,j; + int i; // set energy/virial flags // invoke allocate_peratom() if needed for first time @@ -977,7 +972,6 @@ void PPPMKokkos::set_grid_global() // fluid-occupied volume used to estimate real-space error // zprd used rather than zprd_slab - double h; bigint natoms = atomKK->natoms; if (!gewaldflag) { @@ -1639,9 +1633,6 @@ void PPPMKokkos::operator()(TagPPPM_particle_map, const int &i) cons template void PPPMKokkos::make_rho() { - int l,m,n,nx,ny,nz,mx,my,mz; - FFT_SCALAR dx,dy,dz,x0,y0,z0; - // clear 3d density array //memset(&(density_brick(nzlo_out,nylo_out,nxlo_out)),0, @@ -1848,8 +1839,7 @@ void PPPMKokkos::poisson() template void PPPMKokkos::poisson_ik() { - int i,j,k,n; - double eng; + int j; // transform charge density (r -> k) @@ -1880,7 +1870,6 @@ void PPPMKokkos::poisson_ik() for (j = 0; j < 6; j++) virial[j] += ev.v[j]; energy += ev.ecoul; } else { - n = 0; copymode = 1; Kokkos::parallel_reduce(Kokkos::RangePolicy(0,nfft),*this,ev); DeviceType::fence(); @@ -2217,8 +2206,6 @@ void PPPMKokkos::operator()(TagPPPM_poisson_ik_triclinic6, const int template void PPPMKokkos::poisson_peratom() { - int i,j,k,n; - // merge three outer loops into one for better threading numz_inout = (nzhi_in-nzlo_out)-(nzlo_in-nzlo_out) + 1; @@ -2570,22 +2557,17 @@ KOKKOS_INLINE_FUNCTION void PPPMKokkos::operator()(TagPPPM_fieldforce_ik, const int &i) const { int l,m,n,nx,ny,nz,mx,my,mz; - FFT_SCALAR dx,dy,dz,x0,y0,z0; + FFT_SCALAR x0,y0,z0; FFT_SCALAR ekx,eky,ekz; nx = d_part2grid(i,0); ny = d_part2grid(i,1); nz = d_part2grid(i,2); - dx = nx+shiftone - (x(i,0)-boxlo[0])*delxinv; - dy = ny+shiftone - (x(i,1)-boxlo[1])*delyinv; - dz = nz+shiftone - (x(i,2)-boxlo[2])*delzinv; nz -= nzlo_out; ny -= nylo_out; nx -= nxlo_out; - //compute_rho1d(i,dx,dy,dz); // hasn't changed from make_rho - ekx = eky = ekz = ZEROF; for (n = nlower; n <= nupper; n++) { mz = n+nz; From 9daf7fb6509e04b952719b8c464113bf02f6d6a5 Mon Sep 17 00:00:00 2001 From: Dan Ibanez Date: Thu, 26 Jan 2017 10:19:06 -0700 Subject: [PATCH 09/25] pppm_kokkos: don't shadow member variables --- src/KOKKOS/pppm_kokkos.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/KOKKOS/pppm_kokkos.cpp b/src/KOKKOS/pppm_kokkos.cpp index f923eb35c5..df1949f977 100644 --- a/src/KOKKOS/pppm_kokkos.cpp +++ b/src/KOKKOS/pppm_kokkos.cpp @@ -2827,8 +2827,8 @@ void PPPMKokkos::unpack_reverse_kokkos(int flag, Kokkos::DualView(); - int nx = (nxhi_out-nxlo_out+1); - int ny = (nyhi_out-nylo_out+1); + nx = (nxhi_out-nxlo_out+1); + ny = (nyhi_out-nylo_out+1); copymode = 1; Kokkos::parallel_for(Kokkos::RangePolicy(0,nlist),*this); From e5d0bde78339800f88990142fd6051e744f22e95 Mon Sep 17 00:00:00 2001 From: Dan Ibanez Date: Thu, 26 Jan 2017 10:19:44 -0700 Subject: [PATCH 10/25] pppm_kokkos: remove useless statement --- src/KOKKOS/pppm_kokkos.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/src/KOKKOS/pppm_kokkos.cpp b/src/KOKKOS/pppm_kokkos.cpp index df1949f977..56cd43f1be 100644 --- a/src/KOKKOS/pppm_kokkos.cpp +++ b/src/KOKKOS/pppm_kokkos.cpp @@ -2997,7 +2997,6 @@ void PPPMKokkos::slabcorr() // sum local contributions to get global dipole moment - dipole_all; MPI_Allreduce(&dipole,&dipole_all,1,MPI_DOUBLE,MPI_SUM,world); // need to make non-neutral systems and/or From 51a0b6b445e543c269eccc15ea0784911a6a0ade Mon Sep 17 00:00:00 2001 From: Richard Berger Date: Sat, 28 Jan 2017 07:49:08 -0500 Subject: [PATCH 11/25] Fix data type of molecule array in npair_kokkos.h This showed up when trying to compile with -DLAMMPS_BIGBIG. Fixes issue #365 --- src/KOKKOS/npair_kokkos.h | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/src/KOKKOS/npair_kokkos.h b/src/KOKKOS/npair_kokkos.h index 4b77175191..9ce361c057 100644 --- a/src/KOKKOS/npair_kokkos.h +++ b/src/KOKKOS/npair_kokkos.h @@ -147,7 +147,8 @@ class NeighborKokkosExecute // data from Atom class const typename AT::t_x_array_randomread x; - const typename AT::t_int_1d_const type,mask,molecule; + const typename AT::t_int_1d_const type,mask; + const typename AT::t_tagint_1d_const molecule; const typename AT::t_tagint_1d_const tag; const typename AT::t_tagint_2d_const special; const typename AT::t_int_2d_const nspecial; @@ -184,7 +185,7 @@ class NeighborKokkosExecute const typename AT::t_x_array_randomread &_x, const typename AT::t_int_1d_const &_type, const typename AT::t_int_1d_const &_mask, - const typename AT::t_int_1d_const &_molecule, + const typename AT::t_tagint_1d_const &_molecule, const typename AT::t_tagint_1d_const &_tag, const typename AT::t_tagint_2d_const &_special, const typename AT::t_int_2d_const &_nspecial, From fbe30b56835cb01cbaffce2d8f817ecbfa1ac2d0 Mon Sep 17 00:00:00 2001 From: Axel Kohlmeyer Date: Tue, 31 Jan 2017 18:13:44 -0500 Subject: [PATCH 12/25] correct issue with compiling for -DLAMMPS_BIGBIG in fix qeq/reax/kk --- src/KOKKOS/fix_qeq_reax_kokkos.cpp | 10 +++++----- src/KOKKOS/fix_qeq_reax_kokkos.h | 3 ++- 2 files changed, 7 insertions(+), 6 deletions(-) diff --git a/src/KOKKOS/fix_qeq_reax_kokkos.cpp b/src/KOKKOS/fix_qeq_reax_kokkos.cpp index ab25e557e4..8e155b4158 100644 --- a/src/KOKKOS/fix_qeq_reax_kokkos.cpp +++ b/src/KOKKOS/fix_qeq_reax_kokkos.cpp @@ -387,7 +387,7 @@ KOKKOS_INLINE_FUNCTION void FixQEqReaxKokkos::compute_h_item(int ii, int &m_fill, const bool &final) const { const int i = d_ilist[ii]; - int j,jj,jtag,jtype,flag; + int j,jj,jtype,flag; if (mask[i] & groupbit) { @@ -395,7 +395,7 @@ void FixQEqReaxKokkos::compute_h_item(int ii, int &m_fill, const boo const X_FLOAT ytmp = x(i,1); const X_FLOAT ztmp = x(i,2); const int itype = type(i); - const int itag = tag(i); + const tagint itag = tag(i); const int jnum = d_numneigh[i]; if (final) d_firstnbr[i] = m_fill; @@ -403,7 +403,6 @@ void FixQEqReaxKokkos::compute_h_item(int ii, int &m_fill, const boo for (jj = 0; jj < jnum; jj++) { j = d_neighbors(i,jj); j &= NEIGHMASK; - jtype = type(j); const X_FLOAT delx = x(j,0) - xtmp; @@ -411,10 +410,11 @@ void FixQEqReaxKokkos::compute_h_item(int ii, int &m_fill, const boo const X_FLOAT delz = x(j,2) - ztmp; if (neighflag != FULL) { + const tagint jtag = tag(j); flag = 0; if (j < nlocal) flag = 1; - else if (tag[i] < tag[j]) flag = 1; - else if (tag[i] == tag[j]) { + else if (itag < jtag) flag = 1; + else if (itag == jtag) { if (delz > SMALL) flag = 1; else if (fabs(delz) < SMALL) { if (dely > SMALL) flag = 1; diff --git a/src/KOKKOS/fix_qeq_reax_kokkos.h b/src/KOKKOS/fix_qeq_reax_kokkos.h index eca0d761b7..9014e20b8e 100644 --- a/src/KOKKOS/fix_qeq_reax_kokkos.h +++ b/src/KOKKOS/fix_qeq_reax_kokkos.h @@ -159,7 +159,8 @@ class FixQEqReaxKokkos : public FixQEqReax { //typename ArrayTypes::t_float_1d_randomread mass, q; typename ArrayTypes::t_float_1d_randomread mass; typename ArrayTypes::t_float_1d q; - typename ArrayTypes::t_int_1d type, tag, mask; + typename ArrayTypes::t_int_1d type, mask; + typename ArrayTypes::t_tagint_1d tag; DAT::tdual_float_1d k_q; typename AT::t_float_1d d_q; From bcb1d94b9a1240437fd0bfe5d5985976668cc5a8 Mon Sep 17 00:00:00 2001 From: Axel Kohlmeyer Date: Tue, 31 Jan 2017 18:28:04 -0500 Subject: [PATCH 13/25] silence compiler warning about dead code --- src/KOKKOS/fix_reaxc_bonds_kokkos.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/src/KOKKOS/fix_reaxc_bonds_kokkos.cpp b/src/KOKKOS/fix_reaxc_bonds_kokkos.cpp index 6fefb6b667..7688d6745a 100644 --- a/src/KOKKOS/fix_reaxc_bonds_kokkos.cpp +++ b/src/KOKKOS/fix_reaxc_bonds_kokkos.cpp @@ -76,7 +76,6 @@ void FixReaxCBondsKokkos::init() void FixReaxCBondsKokkos::Output_ReaxC_Bonds(bigint ntimestep, FILE *fp) { - int i, j; int nbuf_local; int nlocal_max, numbonds, numbonds_max; double *buf; From ed59193d13f8ae75b349a00b8356084b5316b170 Mon Sep 17 00:00:00 2001 From: Stefan Paquay Date: Wed, 1 Feb 2017 17:39:06 +0100 Subject: [PATCH 14/25] Removed traces of pair morse/kk --- src/pair_morse.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/pair_morse.h b/src/pair_morse.h index 83cf3276e3..ca6cca62a6 100644 --- a/src/pair_morse.h +++ b/src/pair_morse.h @@ -28,8 +28,8 @@ class PairMorse : public Pair { public: PairMorse(class LAMMPS *); virtual ~PairMorse(); - virtual void compute(int, int); + void settings(int, char **); void coeff(int, char **); double init_one(int, int); From 6232073d3b76439a53ddc66086732780f11d85d8 Mon Sep 17 00:00:00 2001 From: Stefan Paquay Date: Wed, 1 Feb 2017 17:39:37 +0100 Subject: [PATCH 15/25] Removed traces of pair morse/kk --- src/pair_morse.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/pair_morse.h b/src/pair_morse.h index ca6cca62a6..83cf3276e3 100644 --- a/src/pair_morse.h +++ b/src/pair_morse.h @@ -28,8 +28,8 @@ class PairMorse : public Pair { public: PairMorse(class LAMMPS *); virtual ~PairMorse(); + virtual void compute(int, int); - void settings(int, char **); void coeff(int, char **); double init_one(int, int); From d4c9e2500be82a30a8a5a4db268896f8e0a3f11f Mon Sep 17 00:00:00 2001 From: Stefan Paquay Date: Wed, 1 Feb 2017 17:45:21 +0100 Subject: [PATCH 16/25] Ported Morse to KOKKOS --- src/KOKKOS/Install.sh | 2 + src/KOKKOS/pair_morse_kokkos.cpp | 306 +++++++++++++++++++++++++++++++ src/KOKKOS/pair_morse_kokkos.h | 138 ++++++++++++++ src/pair_morse.h | 4 +- 4 files changed, 448 insertions(+), 2 deletions(-) create mode 100644 src/KOKKOS/pair_morse_kokkos.cpp create mode 100644 src/KOKKOS/pair_morse_kokkos.h diff --git a/src/KOKKOS/Install.sh b/src/KOKKOS/Install.sh index 4f7249d67c..691ebbed3e 100644 --- a/src/KOKKOS/Install.sh +++ b/src/KOKKOS/Install.sh @@ -173,6 +173,8 @@ action pair_lj_gromacs_kokkos.cpp action pair_lj_gromacs_kokkos.h action pair_lj_sdk_kokkos.cpp pair_lj_sdk.cpp action pair_lj_sdk_kokkos.h pair_lj_sdk.h +action pair_morse_kokkos.cpp +action pair_morse_kokkos.h action pair_reax_c_kokkos.cpp pair_reax_c.cpp action pair_reax_c_kokkos.h pair_reax_c.h action pair_sw_kokkos.cpp pair_sw.cpp diff --git a/src/KOKKOS/pair_morse_kokkos.cpp b/src/KOKKOS/pair_morse_kokkos.cpp new file mode 100644 index 0000000000..bd5ab22703 --- /dev/null +++ b/src/KOKKOS/pair_morse_kokkos.cpp @@ -0,0 +1,306 @@ +/* ---------------------------------------------------------------------- + LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator + http://lammps.sandia.gov, Sandia National Laboratories + Steve Plimpton, sjplimp@sandia.gov + + Copyright (2003) Sandia Corporation. Under the terms of Contract + DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains + certain rights in this software. This software is distributed under + the GNU General Public License. + + See the README file in the top-level LAMMPS directory. +------------------------------------------------------------------------- */ + +#include +#include +#include +#include +#include "pair_morse_kokkos.h" +#include "kokkos.h" +#include "atom_kokkos.h" +#include "comm.h" +#include "force.h" +#include "neighbor.h" +#include "neigh_list.h" +#include "neigh_request.h" +#include "update.h" +#include "integrate.h" +#include "respa.h" +#include "math_const.h" +#include "memory.h" +#include "error.h" +#include "atom_masks.h" + +using namespace LAMMPS_NS; +using namespace MathConst; + +#define KOKKOS_CUDA_MAX_THREADS 256 +#define KOKKOS_CUDA_MIN_BLOCKS 8 + +/* ---------------------------------------------------------------------- */ + +template +PairMorseKokkos::PairMorseKokkos(LAMMPS *lmp) : PairMorse(lmp) +{ + respa_enable = 0; + + atomKK = (AtomKokkos *) atom; + execution_space = ExecutionSpaceFromDevice::space; + datamask_read = X_MASK | F_MASK | TYPE_MASK | ENERGY_MASK | VIRIAL_MASK; + datamask_modify = F_MASK | ENERGY_MASK | VIRIAL_MASK; + cutsq = NULL; +} + +/* ---------------------------------------------------------------------- */ + +template +PairMorseKokkos::~PairMorseKokkos() +{ + if (allocated) { + memory->destroy_kokkos(k_eatom,eatom); + memory->destroy_kokkos(k_vatom,vatom); + k_cutsq = DAT::tdual_ffloat_2d(); + memory->sfree(cutsq); + eatom = NULL; + vatom = NULL; + cutsq = NULL; + } +} + +/* ---------------------------------------------------------------------- */ + +template +void PairMorseKokkos::cleanup_copy() { + // WHY needed: this prevents parent copy from deallocating any arrays + allocated = 0; + cutsq = NULL; + eatom = NULL; + vatom = NULL; +} + +/* ---------------------------------------------------------------------- */ + +template +void PairMorseKokkos::compute(int eflag_in, int vflag_in) +{ + eflag = eflag_in; + vflag = vflag_in; + + + if (neighflag == FULL) no_virial_fdotr_compute = 1; + + if (eflag || vflag) ev_setup(eflag,vflag); + else evflag = vflag_fdotr = 0; + + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + + atomKK->sync(execution_space,datamask_read); + k_cutsq.template sync(); + k_params.template sync(); + if (eflag || vflag) atomKK->modified(execution_space,datamask_modify); + else atomKK->modified(execution_space,F_MASK); + + x = atomKK->k_x.view(); + c_x = atomKK->k_x.view(); + f = atomKK->k_f.view(); + type = atomKK->k_type.view(); + tag = atomKK->k_tag.view(); + nlocal = atom->nlocal; + nall = atom->nlocal + atom->nghost; + newton_pair = force->newton_pair; + special_lj[0] = force->special_lj[0]; + special_lj[1] = force->special_lj[1]; + special_lj[2] = force->special_lj[2]; + special_lj[3] = force->special_lj[3]; + + // loop over neighbors of my atoms + + EV_FLOAT ev = pair_compute,void >(this,(NeighListKokkos*)list); + + if (eflag_global) eng_vdwl += ev.evdwl; + if (vflag_global) { + virial[0] += ev.v[0]; + virial[1] += ev.v[1]; + virial[2] += ev.v[2]; + virial[3] += ev.v[3]; + virial[4] += ev.v[4]; + virial[5] += ev.v[5]; + } + + if (vflag_fdotr) pair_virial_fdotr_compute(this); + + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } +} + +template +template +KOKKOS_INLINE_FUNCTION +F_FLOAT PairMorseKokkos:: +compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const { + (void) i; + (void) j; + const F_FLOAT rr = sqrt(rsq); + const F_FLOAT r0 = STACKPARAMS ? m_params[itype][jtype].r0 : params(itype,jtype).r0; + const F_FLOAT d0 = STACKPARAMS ? m_params[itype][jtype].d0 : params(itype,jtype).d0; + const F_FLOAT aa = STACKPARAMS ? m_params[itype][jtype].alpha : params(itype,jtype).alpha; + const F_FLOAT dr = rr - r0; + + // U = d0 * [ exp( -2*a*(x-r0)) - 2*exp(-a*(x-r0)) ] + // f = -2*a*d0*[ -exp( -2*a*(x-r0) ) + exp( -a*(x-r0) ) ] * grad(r) + // = +2*a*d0*[ exp( -2*a*(x-r0) ) - exp( -a*(x-r0) ) ] * grad(r) + const F_FLOAT dexp = exp( -aa*dr ); + const F_FLOAT forcelj = 2*aa*d0*dexp*(dexp-1.0); + + return forcelj / rr; +} + +template +template +KOKKOS_INLINE_FUNCTION +F_FLOAT PairMorseKokkos:: +compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const { + (void) i; + (void) j; + const F_FLOAT rr = sqrt(rsq); + const F_FLOAT r0 = STACKPARAMS ? m_params[itype][jtype].r0 : params(itype,jtype).r0; + const F_FLOAT d0 = STACKPARAMS ? m_params[itype][jtype].d0 : params(itype,jtype).d0; + const F_FLOAT aa = STACKPARAMS ? m_params[itype][jtype].alpha : params(itype,jtype).alpha; + const F_FLOAT dr = rr - r0; + + // U = d0 * [ exp( -2*a*(x-r0)) - 2*exp(-a*(x-r0)) ] + // f = -2*a*d0*[ -exp( -2*a*(x-r0) ) + exp( -a*(x-r0) ) ] * grad(r) + // = +2*a*d0*[ exp( -2*a*(x-r0) ) - exp( -a*(x-r0) ) ] * grad(r) + const F_FLOAT dexp = exp( -aa*dr ); + + return d0 * dexp * ( dexp - 2.0 ); +} + +/* ---------------------------------------------------------------------- + allocate all arrays +------------------------------------------------------------------------- */ + +template +void PairMorseKokkos::allocate() +{ + PairMorse::allocate(); + + int n = atom->ntypes; + memory->destroy(cutsq); + memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + d_cutsq = k_cutsq.template view(); + k_params = Kokkos::DualView("PairMorse::params",n+1,n+1); + params = k_params.d_view; +} + +/* ---------------------------------------------------------------------- + global settings +------------------------------------------------------------------------- */ + +template +void PairMorseKokkos::settings(int narg, char **arg) +{ + if (narg > 2) error->all(FLERR,"Illegal pair_style command"); + + PairMorse::settings(1,arg); +} + +/* ---------------------------------------------------------------------- + init specific to this pair style +------------------------------------------------------------------------- */ + +template +void PairMorseKokkos::init_style() +{ + PairMorse::init_style(); + + // error if rRESPA with inner levels + + if (update->whichflag == 1 && strstr(update->integrate_style,"respa")) { + int respa = 0; + if (((Respa *) update->integrate)->level_inner >= 0) respa = 1; + if (((Respa *) update->integrate)->level_middle >= 0) respa = 2; + if (respa) + error->all(FLERR,"Cannot use Kokkos pair style with rRESPA inner/middle"); + } + + // irequest = neigh request made by parent class + + neighflag = lmp->kokkos->neighflag; + int irequest = neighbor->nrequest - 1; + + neighbor->requests[irequest]-> + kokkos_host = Kokkos::Impl::is_same::value && + !Kokkos::Impl::is_same::value; + neighbor->requests[irequest]-> + kokkos_device = Kokkos::Impl::is_same::value; + + if (neighflag == FULL) { + neighbor->requests[irequest]->full = 1; + neighbor->requests[irequest]->half = 0; + } else if (neighflag == HALF || neighflag == HALFTHREAD) { + neighbor->requests[irequest]->full = 0; + neighbor->requests[irequest]->half = 1; + } else if (neighflag == N2) { + neighbor->requests[irequest]->full = 0; + neighbor->requests[irequest]->half = 0; + } else { + error->all(FLERR,"Cannot use chosen neighbor list style with morse/kk"); + } +} + +/* ---------------------------------------------------------------------- + init for one type pair i,j and corresponding j,i +------------------------------------------------------------------------- */ +// Rewrite this. +template +double PairMorseKokkos::init_one(int i, int j) +{ + double cutone = PairMorse::init_one(i,j); + + k_params.h_view(i,j).d0 = d0[i][j]; + k_params.h_view(i,j).alpha = alpha[i][j]; + k_params.h_view(i,j).r0 = r0[i][j]; + k_params.h_view(i,j).offset = offset[i][j]; + k_params.h_view(i,j).cutsq = cutone*cutone; + k_params.h_view(j,i) = k_params.h_view(i,j); + + if(i(); + k_params.template modify(); + + return cutone; +} + + + +namespace LAMMPS_NS { +template class PairMorseKokkos; +#ifdef KOKKOS_HAVE_CUDA +template class PairMorseKokkos; +#endif +} + diff --git a/src/KOKKOS/pair_morse_kokkos.h b/src/KOKKOS/pair_morse_kokkos.h new file mode 100644 index 0000000000..b62b59df54 --- /dev/null +++ b/src/KOKKOS/pair_morse_kokkos.h @@ -0,0 +1,138 @@ +/* -*- c++ -*- ---------------------------------------------------------- + LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator + http://lammps.sandia.gov, Sandia National Laboratories + Steve Plimpton, sjplimp@sandia.gov + + Copyright (2003) Sandia Corporation. Under the terms of Contract + DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains + certain rights in this software. This software is distributed under + the GNU General Public License. + + See the README file in the top-level LAMMPS directory. +------------------------------------------------------------------------- + Contributed by Stefan Paquay @ Eindhoven University of Technology + Based on pair lj/cut/kokkos. +------------------------------------------------------------------------- */ + +#ifdef PAIR_CLASS + +PairStyle(morse/kk,PairMorseKokkos) +PairStyle(morse/kk/device,PairMorseKokkos) +PairStyle(morse/kk/host,PairMorseKokkos) + +#else + +#ifndef LMP_PAIR_MORSE_KOKKOS_H +#define LMP_PAIR_MORSE_KOKKOS_H + +#include "pair_kokkos.h" +#include "pair_morse.h" +#include "neigh_list_kokkos.h" + +namespace LAMMPS_NS { + +template +class PairMorseKokkos : public PairMorse { + public: + enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2}; + enum {COUL_FLAG=0}; + typedef DeviceType device_type; + PairMorseKokkos(class LAMMPS *); + virtual ~PairMorseKokkos(); + + void compute(int, int); + + void settings(int, char **); + void init_style(); + double init_one(int, int); + + struct params_morse{ + KOKKOS_INLINE_FUNCTION + params_morse(){cutsq=0,d0=0;alpha=0;r0=0;offset=0;} + KOKKOS_INLINE_FUNCTION + params_morse(int i){cutsq=0,d0=0;alpha=0;r0=0;offset=0;} + F_FLOAT cutsq,d0,alpha,r0,offset; + }; + + protected: + void cleanup_copy(); + + template + KOKKOS_INLINE_FUNCTION + F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; + + template + KOKKOS_INLINE_FUNCTION + F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; + + template + KOKKOS_INLINE_FUNCTION + F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const { + return 0; + } + + + Kokkos::DualView k_params; + typename Kokkos::DualView::t_dev_const_um params; + params_morse m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; + F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; + typename ArrayTypes::t_x_array_randomread x; + typename ArrayTypes::t_x_array c_x; + typename ArrayTypes::t_f_array f; + typename ArrayTypes::t_int_1d_randomread type; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; + typename ArrayTypes::t_tagint_1d tag; + + int newton_pair; + double special_lj[4]; + + typename ArrayTypes::tdual_ffloat_2d k_cutsq; + typename ArrayTypes::t_ffloat_2d d_cutsq; + + + int neighflag; + int nlocal,nall,eflag,vflag; + + void allocate(); + friend class PairComputeFunctor; + friend class PairComputeFunctor; + friend class PairComputeFunctor; + friend class PairComputeFunctor; + friend class PairComputeFunctor; + friend class PairComputeFunctor; + friend class PairComputeFunctor; + friend class PairComputeFunctor; + friend EV_FLOAT pair_compute_neighlist(PairMorseKokkos*,NeighListKokkos*); + friend EV_FLOAT pair_compute_neighlist(PairMorseKokkos*,NeighListKokkos*); + friend EV_FLOAT pair_compute_neighlist(PairMorseKokkos*,NeighListKokkos*); + friend EV_FLOAT pair_compute_neighlist(PairMorseKokkos*,NeighListKokkos*); + friend EV_FLOAT pair_compute(PairMorseKokkos*,NeighListKokkos*); + friend void pair_virial_fdotr_compute(PairMorseKokkos*); +}; + +} + +#endif +#endif + +/* ERROR/WARNING messages: + +E: Illegal ... command + +Self-explanatory. Check the input script syntax and compare to the +documentation for the command. You can use -echo screen as a +command-line option when running LAMMPS to see the offending line. + +E: Cannot use Kokkos pair style with rRESPA inner/middle + +Self-explanatory. + +E: Cannot use chosen neighbor list style with morse/kk + +That style is not supported by Kokkos. + +*/ diff --git a/src/pair_morse.h b/src/pair_morse.h index 83cf3276e3..d1282f8bf1 100644 --- a/src/pair_morse.h +++ b/src/pair_morse.h @@ -28,8 +28,8 @@ class PairMorse : public Pair { public: PairMorse(class LAMMPS *); virtual ~PairMorse(); - virtual void compute(int, int); + void settings(int, char **); void coeff(int, char **); double init_one(int, int); @@ -49,7 +49,7 @@ class PairMorse : public Pair { double **morse1; double **offset; - void allocate(); + virtual void allocate(); }; } From 7ee45ec5f3546e7772913bedb08dedd84da9a6fb Mon Sep 17 00:00:00 2001 From: Stan Moore Date: Wed, 1 Feb 2017 11:52:27 -0700 Subject: [PATCH 17/25] Fixing tagint and imageint issues in Kokkos package --- src/KOKKOS/atom_vec_angle_kokkos.cpp | 4 +-- src/KOKKOS/atom_vec_atomic_kokkos.cpp | 2 +- src/KOKKOS/atom_vec_bond_kokkos.cpp | 4 +-- src/KOKKOS/atom_vec_charge_kokkos.cpp | 2 +- src/KOKKOS/atom_vec_molecular_kokkos.cpp | 4 +-- src/KOKKOS/domain_kokkos.cpp | 4 +-- src/KOKKOS/pair_reax_c_kokkos.cpp | 38 ++++++++++++------------ src/KOKKOS/pair_tersoff_kokkos.cpp | 4 +-- src/KOKKOS/pair_tersoff_mod_kokkos.cpp | 4 +-- src/KOKKOS/pair_tersoff_zbl_kokkos.cpp | 4 +-- 10 files changed, 35 insertions(+), 35 deletions(-) diff --git a/src/KOKKOS/atom_vec_angle_kokkos.cpp b/src/KOKKOS/atom_vec_angle_kokkos.cpp index 221a5e738a..48fc3a352c 100644 --- a/src/KOKKOS/atom_vec_angle_kokkos.cpp +++ b/src/KOKKOS/atom_vec_angle_kokkos.cpp @@ -957,10 +957,10 @@ struct AtomVecAngleKokkos_UnpackBorder { _x(i+_first,0) = _buf(i,0); _x(i+_first,1) = _buf(i,1); _x(i+_first,2) = _buf(i,2); - _tag(i+_first) = static_cast (_buf(i,3)); + _tag(i+_first) = static_cast (_buf(i,3)); _type(i+_first) = static_cast (_buf(i,4)); _mask(i+_first) = static_cast (_buf(i,5)); - _molecule(i+_first) = static_cast (_buf(i,6)); + _molecule(i+_first) = static_cast (_buf(i,6)); } }; diff --git a/src/KOKKOS/atom_vec_atomic_kokkos.cpp b/src/KOKKOS/atom_vec_atomic_kokkos.cpp index 0a03cf9f6a..dc254e6a7e 100644 --- a/src/KOKKOS/atom_vec_atomic_kokkos.cpp +++ b/src/KOKKOS/atom_vec_atomic_kokkos.cpp @@ -836,7 +836,7 @@ struct AtomVecAtomicKokkos_UnpackBorder { _x(i+_first,0) = _buf(i,0); _x(i+_first,1) = _buf(i,1); _x(i+_first,2) = _buf(i,2); - _tag(i+_first) = static_cast (_buf(i,3)); + _tag(i+_first) = static_cast (_buf(i,3)); _type(i+_first) = static_cast (_buf(i,4)); _mask(i+_first) = static_cast (_buf(i,5)); // printf("%i %i %lf %lf %lf %i BORDER\n",_tag(i+_first),i+_first,_x(i+_first,0),_x(i+_first,1),_x(i+_first,2),_type(i+_first)); diff --git a/src/KOKKOS/atom_vec_bond_kokkos.cpp b/src/KOKKOS/atom_vec_bond_kokkos.cpp index f7bbb4dad5..f10decac28 100644 --- a/src/KOKKOS/atom_vec_bond_kokkos.cpp +++ b/src/KOKKOS/atom_vec_bond_kokkos.cpp @@ -905,10 +905,10 @@ struct AtomVecBondKokkos_UnpackBorder { _x(i+_first,0) = _buf(i,0); _x(i+_first,1) = _buf(i,1); _x(i+_first,2) = _buf(i,2); - _tag(i+_first) = static_cast (_buf(i,3)); + _tag(i+_first) = static_cast (_buf(i,3)); _type(i+_first) = static_cast (_buf(i,4)); _mask(i+_first) = static_cast (_buf(i,5)); - _molecule(i+_first) = static_cast (_buf(i,6)); + _molecule(i+_first) = static_cast (_buf(i,6)); } }; diff --git a/src/KOKKOS/atom_vec_charge_kokkos.cpp b/src/KOKKOS/atom_vec_charge_kokkos.cpp index b7b5cbac7d..f6952f127c 100644 --- a/src/KOKKOS/atom_vec_charge_kokkos.cpp +++ b/src/KOKKOS/atom_vec_charge_kokkos.cpp @@ -872,7 +872,7 @@ struct AtomVecChargeKokkos_UnpackBorder { _x(i+_first,0) = _buf(i,0); _x(i+_first,1) = _buf(i,1); _x(i+_first,2) = _buf(i,2); - _tag(i+_first) = static_cast (_buf(i,3)); + _tag(i+_first) = static_cast (_buf(i,3)); _type(i+_first) = static_cast (_buf(i,4)); _mask(i+_first) = static_cast (_buf(i,5)); _q(i+_first) = _buf(i,6); diff --git a/src/KOKKOS/atom_vec_molecular_kokkos.cpp b/src/KOKKOS/atom_vec_molecular_kokkos.cpp index b4370e239d..4fd8114376 100644 --- a/src/KOKKOS/atom_vec_molecular_kokkos.cpp +++ b/src/KOKKOS/atom_vec_molecular_kokkos.cpp @@ -1029,10 +1029,10 @@ struct AtomVecMolecularKokkos_UnpackBorder { _x(i+_first,0) = _buf(i,0); _x(i+_first,1) = _buf(i,1); _x(i+_first,2) = _buf(i,2); - _tag(i+_first) = static_cast (_buf(i,3)); + _tag(i+_first) = static_cast (_buf(i,3)); _type(i+_first) = static_cast (_buf(i,4)); _mask(i+_first) = static_cast (_buf(i,5)); - _molecule(i+_first) = static_cast (_buf(i,6)); + _molecule(i+_first) = static_cast (_buf(i,6)); } }; diff --git a/src/KOKKOS/domain_kokkos.cpp b/src/KOKKOS/domain_kokkos.cpp index 9742dc36b0..4bf8dc9841 100644 --- a/src/KOKKOS/domain_kokkos.cpp +++ b/src/KOKKOS/domain_kokkos.cpp @@ -250,7 +250,7 @@ struct DomainPBCFunctor { x(i,0) += period[0]; if (DEFORM_VREMAP && (mask[i] & deform_groupbit)) v(i,0) += h_rate[0]; imageint idim = image[i] & IMGMASK; - const int otherdims = image[i] ^ idim; + const imageint otherdims = image[i] ^ idim; idim--; idim &= IMGMASK; image[i] = otherdims | idim; @@ -260,7 +260,7 @@ struct DomainPBCFunctor { x(i,0) = MAX(x(i,0),lo[0]); if (DEFORM_VREMAP && (mask[i] & deform_groupbit)) v(i,0) -= h_rate[0]; imageint idim = image[i] & IMGMASK; - const int otherdims = image[i] ^ idim; + const imageint otherdims = image[i] ^ idim; idim++; idim &= IMGMASK; image[i] = otherdims | idim; diff --git a/src/KOKKOS/pair_reax_c_kokkos.cpp b/src/KOKKOS/pair_reax_c_kokkos.cpp index 0fbf579a92..5490268e56 100644 --- a/src/KOKKOS/pair_reax_c_kokkos.cpp +++ b/src/KOKKOS/pair_reax_c_kokkos.cpp @@ -1046,7 +1046,7 @@ void PairReaxCKokkos::operator()(PairReaxComputeLJCoulomb::operator()(PairReaxComputeLJCoulomb::operator()(PairReaxComputeTabulatedLJCoulomb::operator()(PairReaxComputeTabulatedLJCoulomb::operator()(PairReaxBuildListsHalf, const X_FLOAT ytmp = x(i,1); const X_FLOAT ztmp = x(i,2); const int itype = type(i); - const int itag = tag(i); + const tagint itag = tag(i); const int jnum = d_numneigh[i]; F_FLOAT C12, C34, C56, BO_s, BO_pi, BO_pi2, BO, delij[3], dBOp_i[3], dln_BOp_pi_i[3], dln_BOp_pi2_i[3]; @@ -1605,7 +1605,7 @@ void PairReaxCKokkos::operator()(PairReaxBuildListsHalf, for (int jj = 0; jj < jnum; jj++) { int j = d_neighbors(i,jj); j &= NEIGHMASK; - const int jtag = tag(j); + const tagint jtag = tag(j); d_bo_first[j] = j*maxbo; d_hb_first[j] = j*maxhb; @@ -1802,7 +1802,7 @@ void PairReaxCKokkos::operator()(PairReaxBuildListsHalf_LessAtomics< const X_FLOAT ytmp = x(i,1); const X_FLOAT ztmp = x(i,2); const int itype = type(i); - const int itag = tag(i); + const tagint itag = tag(i); const int jnum = d_numneigh[i]; F_FLOAT C12, C34, C56, BO_s, BO_pi, BO_pi2, BO, delij[3]; @@ -1826,7 +1826,7 @@ void PairReaxCKokkos::operator()(PairReaxBuildListsHalf_LessAtomics< for (int jj = 0; jj < jnum; jj++) { int j = d_neighbors(i,jj); j &= NEIGHMASK; - const int jtag = tag(j); + const tagint jtag = tag(j); d_bo_first[j] = j*maxbo; d_hb_first[j] = j*maxhb; @@ -2752,7 +2752,7 @@ void PairReaxCKokkos::operator()(PairReaxComputeTorsion::operator()(PairReaxComputeTorsion::operator()(PairReaxComputeHydrogen::operator()(PairReaxComputeHydrogen::operator()(PairReaxComputeHydrogen::operator()(PairReaxUpdateBond, cons Kokkos::View::value> > a_Cdbopi2 = d_Cdbopi2; const int i = d_ilist[ii]; - const int itag = tag(i); + const tagint itag = tag(i); const int j_start = d_bo_first[i]; const int j_end = j_start + d_bo_num[i]; for (int jj = j_start; jj < j_end; jj++) { int j = d_bo_list[jj]; j &= NEIGHMASK; - const int jtag = tag(j); + const tagint jtag = tag(j); const int j_index = jj - j_start; const F_FLOAT Cdbo_i = d_Cdbo(i,j_index); const F_FLOAT Cdbopi_i = d_Cdbopi(i,j_index); @@ -3302,7 +3302,7 @@ void PairReaxCKokkos::operator()(PairReaxComputeBond1::operator()(PairReaxComputeBond1 jtag) { if ((itag+jtag) % 2 == 0) continue; @@ -3438,7 +3438,7 @@ void PairReaxCKokkos::operator()(PairReaxComputeBond2::operator()(PairReaxComputeBond2 jtag) { if ((itag+jtag) % 2 == 0) continue; diff --git a/src/KOKKOS/pair_tersoff_kokkos.cpp b/src/KOKKOS/pair_tersoff_kokkos.cpp index 66faa956c6..c94ae5c6f7 100644 --- a/src/KOKKOS/pair_tersoff_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_kokkos.cpp @@ -314,7 +314,7 @@ void PairTersoffKokkos::operator()(TagPairTersoffComputeHalf::operator()(TagPairTersoffComputeHalf jtag) { if ((itag+jtag) % 2 == 0) continue; diff --git a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp index d89b411110..25d949d58a 100644 --- a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp @@ -314,7 +314,7 @@ void PairTersoffMODKokkos::operator()(TagPairTersoffMODComputeHalf::operator()(TagPairTersoffMODComputeHalf jtag) { if ((itag+jtag) % 2 == 0) continue; diff --git a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp index 04195d4e45..fc67b8e953 100644 --- a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp @@ -328,7 +328,7 @@ void PairTersoffZBLKokkos::operator()(TagPairTersoffZBLComputeHalf::operator()(TagPairTersoffZBLComputeHalf jtag) { if ((itag+jtag) % 2 == 0) continue; From 0a87b7443a7509c3b47da36129e31ce23628224a Mon Sep 17 00:00:00 2001 From: Stefan Paquay Date: Thu, 2 Feb 2017 13:42:47 +0100 Subject: [PATCH 18/25] Updated contributing authors and docs --- doc/src/pair_morse.txt | 1 + src/KOKKOS/pair_morse_kokkos.cpp | 4 ++++ src/KOKKOS/pair_morse_kokkos.h | 3 --- 3 files changed, 5 insertions(+), 3 deletions(-) diff --git a/doc/src/pair_morse.txt b/doc/src/pair_morse.txt index a22e6ea3c6..73bc1a314f 100644 --- a/doc/src/pair_morse.txt +++ b/doc/src/pair_morse.txt @@ -13,6 +13,7 @@ pair_style morse/opt command :h3 pair_style morse/smooth/linear command :h3 pair_style morse/smooth/linear/omp command :h3 pair_style morse/soft command :h3 +pair_style morse/kk command :h3 [Syntax:] diff --git a/src/KOKKOS/pair_morse_kokkos.cpp b/src/KOKKOS/pair_morse_kokkos.cpp index bd5ab22703..e0685f1376 100644 --- a/src/KOKKOS/pair_morse_kokkos.cpp +++ b/src/KOKKOS/pair_morse_kokkos.cpp @@ -11,6 +11,10 @@ See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ +/* ---------------------------------------------------------------------- + Contributing authors: Stefan Paquay (Eindhoven University of Technology) +------------------------------------------------------------------------- */ + #include #include #include diff --git a/src/KOKKOS/pair_morse_kokkos.h b/src/KOKKOS/pair_morse_kokkos.h index b62b59df54..a671f2a713 100644 --- a/src/KOKKOS/pair_morse_kokkos.h +++ b/src/KOKKOS/pair_morse_kokkos.h @@ -9,9 +9,6 @@ the GNU General Public License. See the README file in the top-level LAMMPS directory. -------------------------------------------------------------------------- - Contributed by Stefan Paquay @ Eindhoven University of Technology - Based on pair lj/cut/kokkos. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS From daf9f95381c240c47a4ce1aa0cf40161c1af72d9 Mon Sep 17 00:00:00 2001 From: Stan Moore Date: Thu, 2 Feb 2017 13:09:52 -0700 Subject: [PATCH 19/25] Fixing Kokkos per-atom e/v issue --- src/KOKKOS/pair_buck_coul_cut_kokkos.cpp | 2 +- src/KOKKOS/pair_buck_coul_cut_kokkos.h | 27 ++++++----- src/KOKKOS/pair_buck_coul_long_kokkos.cpp | 13 +++++ src/KOKKOS/pair_buck_coul_long_kokkos.h | 29 +++++------ src/KOKKOS/pair_buck_kokkos.h | 19 ++++---- src/KOKKOS/pair_coul_cut_kokkos.cpp | 23 +++++++++ src/KOKKOS/pair_coul_cut_kokkos.h | 30 +++++++----- src/KOKKOS/pair_coul_debye_kokkos.cpp | 25 +++++++++- src/KOKKOS/pair_coul_debye_kokkos.h | 30 +++++++----- src/KOKKOS/pair_coul_dsf_kokkos.cpp | 4 +- src/KOKKOS/pair_coul_dsf_kokkos.h | 17 +++---- src/KOKKOS/pair_coul_long_kokkos.cpp | 23 +++++++++ src/KOKKOS/pair_coul_long_kokkos.h | 29 +++++------ src/KOKKOS/pair_coul_wolf_kokkos.cpp | 4 +- src/KOKKOS/pair_coul_wolf_kokkos.h | 13 ++--- src/KOKKOS/pair_eam_alloy_kokkos.cpp | 4 +- src/KOKKOS/pair_eam_alloy_kokkos.h | 6 +-- src/KOKKOS/pair_eam_fs_kokkos.h | 6 +-- src/KOKKOS/pair_eam_kokkos.h | 6 +-- ..._lj_charmm_coul_charmm_implicit_kokkos.cpp | 23 +++++++++ ...ir_lj_charmm_coul_charmm_implicit_kokkos.h | 29 +++++------ .../pair_lj_charmm_coul_charmm_kokkos.cpp | 23 +++++++++ .../pair_lj_charmm_coul_charmm_kokkos.h | 29 +++++------ .../pair_lj_charmm_coul_long_kokkos.cpp | 4 +- src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h | 29 +++++------ src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp | 25 +++++++++- src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h | 30 +++++++----- .../pair_lj_class2_coul_long_kokkos.cpp | 23 +++++++++ src/KOKKOS/pair_lj_class2_coul_long_kokkos.h | 32 +++++++------ src/KOKKOS/pair_lj_class2_kokkos.cpp | 24 ++++++++++ src/KOKKOS/pair_lj_class2_kokkos.h | 22 +++++---- src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp | 26 +++++++++- src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h | 30 +++++++----- src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp | 25 +++++++++- src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h | 29 ++++++----- src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp | 25 +++++++++- src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h | 30 +++++++----- src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp | 5 +- src/KOKKOS/pair_lj_cut_coul_long_kokkos.h | 29 +++++------ src/KOKKOS/pair_lj_cut_kokkos.cpp | 5 +- src/KOKKOS/pair_lj_cut_kokkos.h | 19 ++++---- src/KOKKOS/pair_lj_expand_kokkos.cpp | 23 +++++++++ src/KOKKOS/pair_lj_expand_kokkos.h | 21 ++++---- .../pair_lj_gromacs_coul_gromacs_kokkos.cpp | 23 +++++++++ .../pair_lj_gromacs_coul_gromacs_kokkos.h | 29 +++++------ src/KOKKOS/pair_lj_gromacs_kokkos.cpp | 23 +++++++++ src/KOKKOS/pair_lj_gromacs_kokkos.h | 29 +++++------ src/KOKKOS/pair_lj_sdk_kokkos.cpp | 24 ++++++++++ src/KOKKOS/pair_lj_sdk_kokkos.h | 22 +++++---- src/KOKKOS/pair_reax_c_kokkos.h | 6 +-- src/KOKKOS/pair_sw_kokkos.cpp | 4 +- src/KOKKOS/pair_sw_kokkos.h | 14 +++--- src/KOKKOS/pair_table_kokkos.cpp | 26 +++++++++- src/KOKKOS/pair_table_kokkos.h | 48 ++++++++++--------- src/KOKKOS/pair_tersoff_kokkos.cpp | 4 +- src/KOKKOS/pair_tersoff_kokkos.h | 6 +-- src/KOKKOS/pair_tersoff_mod_kokkos.cpp | 4 +- src/KOKKOS/pair_tersoff_mod_kokkos.h | 6 +-- src/KOKKOS/pair_tersoff_zbl_kokkos.cpp | 4 +- src/KOKKOS/pair_vashishta_kokkos.cpp | 7 ++- src/KOKKOS/pair_vashishta_kokkos.h | 14 +++--- 61 files changed, 802 insertions(+), 361 deletions(-) diff --git a/src/KOKKOS/pair_buck_coul_cut_kokkos.cpp b/src/KOKKOS/pair_buck_coul_cut_kokkos.cpp index a176ca2be4..afe2d14fcc 100644 --- a/src/KOKKOS/pair_buck_coul_cut_kokkos.cpp +++ b/src/KOKKOS/pair_buck_coul_cut_kokkos.cpp @@ -154,7 +154,7 @@ void PairBuckCoulCutKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) virial_fdotr_compute(); + if (vflag_fdotr) pair_virial_fdotr_compute(this); if (eflag_atom) { k_eatom.template modify(); diff --git a/src/KOKKOS/pair_buck_coul_cut_kokkos.h b/src/KOKKOS/pair_buck_coul_cut_kokkos.h index 0b6aba5e92..fb0034d767 100644 --- a/src/KOKKOS/pair_buck_coul_cut_kokkos.h +++ b/src/KOKKOS/pair_buck_coul_cut_kokkos.h @@ -34,6 +34,7 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairBuckCoulCutKokkos(class LAMMPS *); ~PairBuckCoulCutKokkos(); @@ -83,25 +84,25 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; int neighflag; diff --git a/src/KOKKOS/pair_buck_coul_long_kokkos.cpp b/src/KOKKOS/pair_buck_coul_long_kokkos.cpp index 413f38370d..fa05c99995 100644 --- a/src/KOKKOS/pair_buck_coul_long_kokkos.cpp +++ b/src/KOKKOS/pair_buck_coul_long_kokkos.cpp @@ -114,6 +114,19 @@ void PairBuckCoulLongKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); diff --git a/src/KOKKOS/pair_buck_coul_long_kokkos.h b/src/KOKKOS/pair_buck_coul_long_kokkos.h index fe63a2c124..f48258c2e5 100644 --- a/src/KOKKOS/pair_buck_coul_long_kokkos.h +++ b/src/KOKKOS/pair_buck_coul_long_kokkos.h @@ -34,6 +34,7 @@ class PairBuckCoulLongKokkos : public PairBuckCoulLong { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairBuckCoulLongKokkos(class LAMMPS *); ~PairBuckCoulLongKokkos(); @@ -84,27 +85,27 @@ class PairBuckCoulLongKokkos : public PairBuckCoulLong { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; - typename ArrayTypes::t_ffloat_1d_randomread + typename AT::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; diff --git a/src/KOKKOS/pair_buck_kokkos.h b/src/KOKKOS/pair_buck_kokkos.h index 0b0bbf94cf..d57e320e99 100644 --- a/src/KOKKOS/pair_buck_kokkos.h +++ b/src/KOKKOS/pair_buck_kokkos.h @@ -34,6 +34,7 @@ class PairBuckKokkos : public PairBuck { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2}; enum {COUL_FLAG=0}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairBuckKokkos(class LAMMPS *); ~PairBuckKokkos(); @@ -72,22 +73,22 @@ class PairBuckKokkos : public PairBuck { typename Kokkos::DualView::t_dev_const_um params; params_buck m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_tagint_1d tag; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_tagint_1d tag; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; double special_lj[4]; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; int neighflag; diff --git a/src/KOKKOS/pair_coul_cut_kokkos.cpp b/src/KOKKOS/pair_coul_cut_kokkos.cpp index 19d4306317..a770d55be9 100644 --- a/src/KOKKOS/pair_coul_cut_kokkos.cpp +++ b/src/KOKKOS/pair_coul_cut_kokkos.cpp @@ -83,6 +83,19 @@ void PairCoulCutKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); @@ -124,6 +137,16 @@ void PairCoulCutKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); } diff --git a/src/KOKKOS/pair_coul_cut_kokkos.h b/src/KOKKOS/pair_coul_cut_kokkos.h index 2049385c3f..78d0b192f8 100644 --- a/src/KOKKOS/pair_coul_cut_kokkos.h +++ b/src/KOKKOS/pair_coul_cut_kokkos.h @@ -34,6 +34,7 @@ class PairCoulCutKokkos : public PairCoulCut { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairCoulCutKokkos(class LAMMPS *); ~PairCoulCutKokkos(); @@ -87,22 +88,25 @@ class PairCoulCutKokkos : public PairCoulCut { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_float_1d_randomread q; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_float_1d_randomread q; + typename AT::t_int_1d_randomread type; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; int neighflag; diff --git a/src/KOKKOS/pair_coul_debye_kokkos.cpp b/src/KOKKOS/pair_coul_debye_kokkos.cpp index 9a6e1b8020..e9314c2a0a 100644 --- a/src/KOKKOS/pair_coul_debye_kokkos.cpp +++ b/src/KOKKOS/pair_coul_debye_kokkos.cpp @@ -90,6 +90,19 @@ void PairCoulDebyeKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); @@ -136,7 +149,17 @@ void PairCoulDebyeKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) virial_fdotr_compute(); + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; } diff --git a/src/KOKKOS/pair_coul_debye_kokkos.h b/src/KOKKOS/pair_coul_debye_kokkos.h index 4aadcbe4e1..143b1c2354 100644 --- a/src/KOKKOS/pair_coul_debye_kokkos.h +++ b/src/KOKKOS/pair_coul_debye_kokkos.h @@ -34,6 +34,7 @@ class PairCoulDebyeKokkos : public PairCoulDebye { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairCoulDebyeKokkos(class LAMMPS *); ~PairCoulDebyeKokkos(); @@ -87,22 +88,25 @@ class PairCoulDebyeKokkos : public PairCoulDebye { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; int neighflag; diff --git a/src/KOKKOS/pair_coul_dsf_kokkos.cpp b/src/KOKKOS/pair_coul_dsf_kokkos.cpp index e689754d0a..3ca1fb8ea4 100644 --- a/src/KOKKOS/pair_coul_dsf_kokkos.cpp +++ b/src/KOKKOS/pair_coul_dsf_kokkos.cpp @@ -183,8 +183,6 @@ void PairCoulDSFKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -195,6 +193,8 @@ void PairCoulDSFKokkos::compute(int eflag_in, int vflag_in) k_vatom.template sync(); } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + copymode = 0; } diff --git a/src/KOKKOS/pair_coul_dsf_kokkos.h b/src/KOKKOS/pair_coul_dsf_kokkos.h index e50f92c370..33740aa916 100644 --- a/src/KOKKOS/pair_coul_dsf_kokkos.h +++ b/src/KOKKOS/pair_coul_dsf_kokkos.h @@ -37,6 +37,7 @@ class PairCoulDSFKokkos : public PairCoulDSF { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; typedef EV_FLOAT value_type; PairCoulDSFKokkos(class LAMMPS *); ~PairCoulDSFKokkos(); @@ -63,14 +64,14 @@ class PairCoulDSFKokkos : public PairCoulDSF { protected: - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_f_array f; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int neighflag,newton_pair; @@ -79,9 +80,9 @@ class PairCoulDSFKokkos : public PairCoulDSF { double special_coul[4]; double qqrd2e; - typename ArrayTypes::t_neighbors_2d d_neighbors; - typename ArrayTypes::t_int_1d_randomread d_ilist; - typename ArrayTypes::t_int_1d_randomread d_numneigh; + typename AT::t_neighbors_2d d_neighbors; + typename AT::t_int_1d_randomread d_ilist; + typename AT::t_int_1d_randomread d_numneigh; //NeighListKokkos k_list; friend void pair_virial_fdotr_compute(PairCoulDSFKokkos*); diff --git a/src/KOKKOS/pair_coul_long_kokkos.cpp b/src/KOKKOS/pair_coul_long_kokkos.cpp index 7536549bf4..dcf4943104 100644 --- a/src/KOKKOS/pair_coul_long_kokkos.cpp +++ b/src/KOKKOS/pair_coul_long_kokkos.cpp @@ -107,6 +107,19 @@ void PairCoulLongKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_coulsq.template sync(); @@ -158,6 +171,16 @@ void PairCoulLongKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; diff --git a/src/KOKKOS/pair_coul_long_kokkos.h b/src/KOKKOS/pair_coul_long_kokkos.h index 19789fbd86..5547f0651f 100644 --- a/src/KOKKOS/pair_coul_long_kokkos.h +++ b/src/KOKKOS/pair_coul_long_kokkos.h @@ -34,6 +34,7 @@ class PairCoulLongKokkos : public PairCoulLong { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairCoulLongKokkos(class LAMMPS *); ~PairCoulLongKokkos(); @@ -86,27 +87,27 @@ class PairCoulLongKokkos : public PairCoulLong { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; - typename ArrayTypes::t_ffloat_1d_randomread + typename AT::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; diff --git a/src/KOKKOS/pair_coul_wolf_kokkos.cpp b/src/KOKKOS/pair_coul_wolf_kokkos.cpp index 1785ba2731..22a11c4a18 100644 --- a/src/KOKKOS/pair_coul_wolf_kokkos.cpp +++ b/src/KOKKOS/pair_coul_wolf_kokkos.cpp @@ -184,8 +184,6 @@ void PairCoulWolfKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -196,6 +194,8 @@ void PairCoulWolfKokkos::compute(int eflag_in, int vflag_in) k_vatom.template sync(); } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + copymode = 0; } diff --git a/src/KOKKOS/pair_coul_wolf_kokkos.h b/src/KOKKOS/pair_coul_wolf_kokkos.h index bde26c0c3e..cedafe7a21 100644 --- a/src/KOKKOS/pair_coul_wolf_kokkos.h +++ b/src/KOKKOS/pair_coul_wolf_kokkos.h @@ -37,6 +37,7 @@ class PairCoulWolfKokkos : public PairCoulWolf { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; typedef EV_FLOAT value_type; PairCoulWolfKokkos(class LAMMPS *); ~PairCoulWolfKokkos(); @@ -63,9 +64,9 @@ class PairCoulWolfKokkos : public PairCoulWolf { protected: - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_f_array f; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; @@ -81,9 +82,9 @@ class PairCoulWolfKokkos : public PairCoulWolf { double special_coul[4]; double qqrd2e; - typename ArrayTypes::t_neighbors_2d d_neighbors; - typename ArrayTypes::t_int_1d_randomread d_ilist; - typename ArrayTypes::t_int_1d_randomread d_numneigh; + typename AT::t_neighbors_2d d_neighbors; + typename AT::t_int_1d_randomread d_ilist; + typename AT::t_int_1d_randomread d_numneigh; //NeighListKokkos k_list; friend void pair_virial_fdotr_compute(PairCoulWolfKokkos*); diff --git a/src/KOKKOS/pair_eam_alloy_kokkos.cpp b/src/KOKKOS/pair_eam_alloy_kokkos.cpp index f3b7c36106..866fe2cfa1 100644 --- a/src/KOKKOS/pair_eam_alloy_kokkos.cpp +++ b/src/KOKKOS/pair_eam_alloy_kokkos.cpp @@ -246,8 +246,6 @@ void PairEAMAlloyKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -258,6 +256,8 @@ void PairEAMAlloyKokkos::compute(int eflag_in, int vflag_in) k_vatom.template sync(); } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + copymode = 0; } diff --git a/src/KOKKOS/pair_eam_alloy_kokkos.h b/src/KOKKOS/pair_eam_alloy_kokkos.h index 391349d9cb..44420f2d9e 100644 --- a/src/KOKKOS/pair_eam_alloy_kokkos.h +++ b/src/KOKKOS/pair_eam_alloy_kokkos.h @@ -154,9 +154,9 @@ class PairEAMAlloyKokkos : public PairEAM { void interpolate(int, double, double *, t_host_ffloat_2d_n7, int); void read_file(char *); - typename ArrayTypes::t_neighbors_2d d_neighbors; - typename ArrayTypes::t_int_1d_randomread d_ilist; - typename ArrayTypes::t_int_1d_randomread d_numneigh; + typename AT::t_neighbors_2d d_neighbors; + typename AT::t_int_1d_randomread d_ilist; + typename AT::t_int_1d_randomread d_numneigh; //NeighListKokkos k_list; int iswap; diff --git a/src/KOKKOS/pair_eam_fs_kokkos.h b/src/KOKKOS/pair_eam_fs_kokkos.h index c3b0f32457..7bb4d0abfb 100644 --- a/src/KOKKOS/pair_eam_fs_kokkos.h +++ b/src/KOKKOS/pair_eam_fs_kokkos.h @@ -154,9 +154,9 @@ class PairEAMFSKokkos : public PairEAM { void interpolate(int, double, double *, t_host_ffloat_2d_n7, int); void read_file(char *); - typename ArrayTypes::t_neighbors_2d d_neighbors; - typename ArrayTypes::t_int_1d_randomread d_ilist; - typename ArrayTypes::t_int_1d_randomread d_numneigh; + typename AT::t_neighbors_2d d_neighbors; + typename AT::t_int_1d_randomread d_ilist; + typename AT::t_int_1d_randomread d_numneigh; //NeighListKokkos k_list; int iswap; diff --git a/src/KOKKOS/pair_eam_kokkos.h b/src/KOKKOS/pair_eam_kokkos.h index f933a4a0d7..d36ef0f95a 100644 --- a/src/KOKKOS/pair_eam_kokkos.h +++ b/src/KOKKOS/pair_eam_kokkos.h @@ -149,9 +149,9 @@ class PairEAMKokkos : public PairEAM { virtual void file2array(); void array2spline(); - typename ArrayTypes::t_neighbors_2d d_neighbors; - typename ArrayTypes::t_int_1d_randomread d_ilist; - typename ArrayTypes::t_int_1d_randomread d_numneigh; + typename AT::t_neighbors_2d d_neighbors; + typename AT::t_int_1d_randomread d_ilist; + typename AT::t_int_1d_randomread d_numneigh; //NeighListKokkos k_list; int iswap; diff --git a/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.cpp b/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.cpp index 1cf2dfe054..a4e0fa963d 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.cpp +++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.cpp @@ -115,6 +115,19 @@ void PairLJCharmmCoulCharmmImplicitKokkos::compute(int eflag_in, int if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); @@ -167,6 +180,16 @@ void PairLJCharmmCoulCharmmImplicitKokkos::compute(int eflag_in, int virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; diff --git a/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h b/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h index 048a7dab60..9405092716 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h +++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h @@ -34,6 +34,7 @@ class PairLJCharmmCoulCharmmImplicitKokkos : public PairLJCharmmCoulCharmmImplic enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJCharmmCoulCharmmImplicitKokkos(class LAMMPS *); ~PairLJCharmmCoulCharmmImplicitKokkos(); @@ -77,27 +78,27 @@ class PairLJCharmmCoulCharmmImplicitKokkos : public PairLJCharmmCoulCharmmImplic F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; - typename ArrayTypes::t_ffloat_1d_randomread + typename AT::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; diff --git a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp index a4d0316ca7..1e5509c39c 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp +++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp @@ -115,6 +115,19 @@ void PairLJCharmmCoulCharmmKokkos::compute(int eflag_in, int vflag_i if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); @@ -167,6 +180,16 @@ void PairLJCharmmCoulCharmmKokkos::compute(int eflag_in, int vflag_i virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; diff --git a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h index db0b14a84e..4e00040fd9 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h +++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h @@ -34,6 +34,7 @@ class PairLJCharmmCoulCharmmKokkos : public PairLJCharmmCoulCharmm { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJCharmmCoulCharmmKokkos(class LAMMPS *); ~PairLJCharmmCoulCharmmKokkos(); @@ -77,27 +78,27 @@ class PairLJCharmmCoulCharmmKokkos : public PairLJCharmmCoulCharmm { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; - typename ArrayTypes::t_ffloat_1d_randomread + typename AT::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; diff --git a/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.cpp b/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.cpp index 5b1176278a..2f94c47bae 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.cpp +++ b/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.cpp @@ -180,8 +180,6 @@ void PairLJCharmmCoulLongKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -192,6 +190,8 @@ void PairLJCharmmCoulLongKokkos::compute(int eflag_in, int vflag_in) k_vatom.template sync(); } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + copymode = 0; } diff --git a/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h b/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h index 0969d11b0e..b030783401 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h +++ b/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h @@ -34,6 +34,7 @@ class PairLJCharmmCoulLongKokkos : public PairLJCharmmCoulLong { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJCharmmCoulLongKokkos(class LAMMPS *); ~PairLJCharmmCoulLongKokkos(); @@ -75,27 +76,27 @@ class PairLJCharmmCoulLongKokkos : public PairLJCharmmCoulLong { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; - typename ArrayTypes::t_ffloat_1d_randomread + typename AT::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; diff --git a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp index 0e33958a22..4283cfa981 100644 --- a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp +++ b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp @@ -92,6 +92,19 @@ void PairLJClass2CoulCutKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); @@ -138,7 +151,17 @@ void PairLJClass2CoulCutKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) virial_fdotr_compute(); + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; } diff --git a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h index c3492666de..1679460dbe 100644 --- a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h +++ b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h @@ -34,6 +34,7 @@ class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJClass2CoulCutKokkos(class LAMMPS *); ~PairLJClass2CoulCutKokkos(); @@ -76,22 +77,25 @@ class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; int neighflag; int nlocal,nall,eflag,vflag; diff --git a/src/KOKKOS/pair_lj_class2_coul_long_kokkos.cpp b/src/KOKKOS/pair_lj_class2_coul_long_kokkos.cpp index 00becd0ec6..adbffff50f 100644 --- a/src/KOKKOS/pair_lj_class2_coul_long_kokkos.cpp +++ b/src/KOKKOS/pair_lj_class2_coul_long_kokkos.cpp @@ -100,6 +100,19 @@ void PairLJClass2CoulLongKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); @@ -152,6 +165,16 @@ void PairLJClass2CoulLongKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; diff --git a/src/KOKKOS/pair_lj_class2_coul_long_kokkos.h b/src/KOKKOS/pair_lj_class2_coul_long_kokkos.h index c5c46ed2d5..8a2e993353 100644 --- a/src/KOKKOS/pair_lj_class2_coul_long_kokkos.h +++ b/src/KOKKOS/pair_lj_class2_coul_long_kokkos.h @@ -34,6 +34,7 @@ class PairLJClass2CoulLongKokkos : public PairLJClass2CoulLong { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJClass2CoulLongKokkos(class LAMMPS *); ~PairLJClass2CoulLongKokkos(); @@ -76,24 +77,27 @@ class PairLJClass2CoulLongKokkos : public PairLJClass2CoulLong { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; - typename ArrayTypes::t_ffloat_1d_randomread + typename AT::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; diff --git a/src/KOKKOS/pair_lj_class2_kokkos.cpp b/src/KOKKOS/pair_lj_class2_kokkos.cpp index 9f9e3c72a5..c26aa4f467 100644 --- a/src/KOKKOS/pair_lj_class2_kokkos.cpp +++ b/src/KOKKOS/pair_lj_class2_kokkos.cpp @@ -92,6 +92,19 @@ void PairLJClass2Kokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_params.template sync(); @@ -125,7 +138,18 @@ void PairLJClass2Kokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + } template diff --git a/src/KOKKOS/pair_lj_class2_kokkos.h b/src/KOKKOS/pair_lj_class2_kokkos.h index ccff6821f3..ae0676c8a5 100644 --- a/src/KOKKOS/pair_lj_class2_kokkos.h +++ b/src/KOKKOS/pair_lj_class2_kokkos.h @@ -34,6 +34,7 @@ class PairLJClass2Kokkos : public PairLJClass2 { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2}; enum {COUL_FLAG=0}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJClass2Kokkos(class LAMMPS *); ~PairLJClass2Kokkos(); @@ -79,19 +80,22 @@ class PairLJClass2Kokkos : public PairLJClass2 { typename Kokkos::DualView::t_dev_const_um params; params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; - typename ArrayTypes::t_tagint_1d tag; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_tagint_1d tag; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; double special_lj[4]; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; int neighflag; int nlocal,nall,eflag,vflag; diff --git a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp index f1a7950521..e9f587d0cc 100644 --- a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp +++ b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp @@ -92,6 +92,19 @@ void PairLJCutCoulCutKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); @@ -136,7 +149,18 @@ void PairLJCutCoulCutKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) virial_fdotr_compute(); + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + + if (vflag_fdotr) pair_virial_fdotr_compute(this); + } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h index 5891371d14..7dcee07400 100644 --- a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h @@ -34,6 +34,7 @@ class PairLJCutCoulCutKokkos : public PairLJCutCoulCut { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJCutCoulCutKokkos(class LAMMPS *); ~PairLJCutCoulCutKokkos(); @@ -75,22 +76,25 @@ class PairLJCutCoulCutKokkos : public PairLJCutCoulCut { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; int neighflag; diff --git a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp index 3eb793c52a..0af781affd 100644 --- a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp +++ b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp @@ -96,6 +96,19 @@ void PairLJCutCoulDebyeKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); @@ -142,7 +155,17 @@ void PairLJCutCoulDebyeKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) virial_fdotr_compute(); + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; } diff --git a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h index d507f76a3a..e0ab231359 100644 --- a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h @@ -75,22 +75,25 @@ class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; int neighflag; diff --git a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp index 6507bb0272..5ddf035b62 100644 --- a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp +++ b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp @@ -104,6 +104,19 @@ void PairLJCutCoulDSFKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); @@ -160,7 +173,17 @@ void PairLJCutCoulDSFKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) virial_fdotr_compute(); + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; } diff --git a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h index 3e378757c0..13030fbf7c 100644 --- a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h @@ -34,6 +34,7 @@ class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJCutCoulDSFKokkos(class LAMMPS *); ~PairLJCutCoulDSFKokkos(); @@ -74,22 +75,25 @@ class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; int neighflag; diff --git a/src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp b/src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp index 11095bf4de..e3754cdb5d 100644 --- a/src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp +++ b/src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp @@ -168,8 +168,6 @@ void PairLJCutCoulLongKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -179,6 +177,9 @@ void PairLJCutCoulLongKokkos::compute(int eflag_in, int vflag_in) k_vatom.template modify(); k_vatom.template sync(); } + + if (vflag_fdotr) pair_virial_fdotr_compute(this); + } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/pair_lj_cut_coul_long_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_long_kokkos.h index 7328329232..79e77a7508 100644 --- a/src/KOKKOS/pair_lj_cut_coul_long_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_long_kokkos.h @@ -34,6 +34,7 @@ class PairLJCutCoulLongKokkos : public PairLJCutCoulLong { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJCutCoulLongKokkos(class LAMMPS *); ~PairLJCutCoulLongKokkos(); @@ -76,27 +77,27 @@ class PairLJCutCoulLongKokkos : public PairLJCutCoulLong { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; - typename ArrayTypes::t_ffloat_1d_randomread + typename AT::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; diff --git a/src/KOKKOS/pair_lj_cut_kokkos.cpp b/src/KOKKOS/pair_lj_cut_kokkos.cpp index 066961cc89..9a2a7e23a9 100644 --- a/src/KOKKOS/pair_lj_cut_kokkos.cpp +++ b/src/KOKKOS/pair_lj_cut_kokkos.cpp @@ -138,8 +138,6 @@ void PairLJCutKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -149,6 +147,9 @@ void PairLJCutKokkos::compute(int eflag_in, int vflag_in) k_vatom.template modify(); k_vatom.template sync(); } + + if (vflag_fdotr) pair_virial_fdotr_compute(this); + } template diff --git a/src/KOKKOS/pair_lj_cut_kokkos.h b/src/KOKKOS/pair_lj_cut_kokkos.h index 81c25c20d8..7e2b8fd91a 100644 --- a/src/KOKKOS/pair_lj_cut_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_kokkos.h @@ -34,6 +34,7 @@ class PairLJCutKokkos : public PairLJCut { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2}; enum {COUL_FLAG=0}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJCutKokkos(class LAMMPS *); ~PairLJCutKokkos(); @@ -73,22 +74,22 @@ class PairLJCutKokkos : public PairLJCut { typename Kokkos::DualView::t_dev_const_um params; params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; - typename ArrayTypes::t_tagint_1d tag; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; + typename AT::t_tagint_1d tag; int newton_pair; double special_lj[4]; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; int neighflag; diff --git a/src/KOKKOS/pair_lj_expand_kokkos.cpp b/src/KOKKOS/pair_lj_expand_kokkos.cpp index d7daef2a66..cdede8bd52 100644 --- a/src/KOKKOS/pair_lj_expand_kokkos.cpp +++ b/src/KOKKOS/pair_lj_expand_kokkos.cpp @@ -91,6 +91,19 @@ void PairLJExpandKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_params.template sync(); @@ -126,6 +139,16 @@ void PairLJExpandKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; diff --git a/src/KOKKOS/pair_lj_expand_kokkos.h b/src/KOKKOS/pair_lj_expand_kokkos.h index ee110bab80..3b0067bcd8 100644 --- a/src/KOKKOS/pair_lj_expand_kokkos.h +++ b/src/KOKKOS/pair_lj_expand_kokkos.h @@ -79,19 +79,22 @@ class PairLJExpandKokkos : public PairLJExpand { typename Kokkos::DualView::t_dev_const_um params; params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; - typename ArrayTypes::t_tagint_1d tag; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_tagint_1d tag; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; double special_lj[4]; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; int neighflag; diff --git a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp index b636f36499..0ca8375657 100644 --- a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp +++ b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp @@ -106,6 +106,19 @@ void PairLJGromacsCoulGromacsKokkos::compute(int eflag_in, int vflag if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); @@ -158,6 +171,16 @@ void PairLJGromacsCoulGromacsKokkos::compute(int eflag_in, int vflag virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; diff --git a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h index bbf5c50a62..51fc0b71b7 100644 --- a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h +++ b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h @@ -34,6 +34,7 @@ class PairLJGromacsCoulGromacsKokkos : public PairLJGromacsCoulGromacs { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJGromacsCoulGromacsKokkos(class LAMMPS *); ~PairLJGromacsCoulGromacsKokkos(); @@ -84,27 +85,27 @@ class PairLJGromacsCoulGromacsKokkos : public PairLJGromacsCoulGromacs { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; - typename ArrayTypes::t_ffloat_2d d_cut_ljsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; - typename ArrayTypes::t_ffloat_2d d_cut_coulsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_ljsq; + typename AT::t_ffloat_2d d_cut_ljsq; + typename AT::tdual_ffloat_2d k_cut_coulsq; + typename AT::t_ffloat_2d d_cut_coulsq; - typename ArrayTypes::t_ffloat_1d_randomread + typename AT::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; diff --git a/src/KOKKOS/pair_lj_gromacs_kokkos.cpp b/src/KOKKOS/pair_lj_gromacs_kokkos.cpp index a72fb7dbf9..e6606d30cf 100644 --- a/src/KOKKOS/pair_lj_gromacs_kokkos.cpp +++ b/src/KOKKOS/pair_lj_gromacs_kokkos.cpp @@ -103,6 +103,19 @@ void PairLJGromacsKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_inner.template sync(); @@ -145,6 +158,16 @@ void PairLJGromacsKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; diff --git a/src/KOKKOS/pair_lj_gromacs_kokkos.h b/src/KOKKOS/pair_lj_gromacs_kokkos.h index 1f16409dc5..c8c792038d 100644 --- a/src/KOKKOS/pair_lj_gromacs_kokkos.h +++ b/src/KOKKOS/pair_lj_gromacs_kokkos.h @@ -34,6 +34,7 @@ class PairLJGromacsKokkos : public PairLJGromacs { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=0}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJGromacsKokkos(class LAMMPS *); ~PairLJGromacsKokkos(); @@ -87,27 +88,27 @@ class PairLJGromacsKokkos : public PairLJGromacs { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_inner[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_inner_sq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_float_1d_randomread q; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; - typename ArrayTypes::tdual_ffloat_2d k_cut_inner; - typename ArrayTypes::t_ffloat_2d d_cut_inner; - typename ArrayTypes::tdual_ffloat_2d k_cut_inner_sq; - typename ArrayTypes::t_ffloat_2d d_cut_inner_sq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cut_inner; + typename AT::t_ffloat_2d d_cut_inner; + typename AT::tdual_ffloat_2d k_cut_inner_sq; + typename AT::t_ffloat_2d d_cut_inner_sq; - typename ArrayTypes::t_ffloat_1d_randomread + typename AT::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; diff --git a/src/KOKKOS/pair_lj_sdk_kokkos.cpp b/src/KOKKOS/pair_lj_sdk_kokkos.cpp index 1acf35f18d..2d74e37192 100644 --- a/src/KOKKOS/pair_lj_sdk_kokkos.cpp +++ b/src/KOKKOS/pair_lj_sdk_kokkos.cpp @@ -91,6 +91,19 @@ void PairLJSDKKokkos::compute(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_params.template sync(); @@ -124,7 +137,18 @@ void PairLJSDKKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + } template diff --git a/src/KOKKOS/pair_lj_sdk_kokkos.h b/src/KOKKOS/pair_lj_sdk_kokkos.h index 0b9056a23e..f313c4342f 100644 --- a/src/KOKKOS/pair_lj_sdk_kokkos.h +++ b/src/KOKKOS/pair_lj_sdk_kokkos.h @@ -34,6 +34,7 @@ class PairLJSDKKokkos : public PairLJSDK { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2}; enum {COUL_FLAG=0}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJSDKKokkos(class LAMMPS *); ~PairLJSDKKokkos(); @@ -74,19 +75,22 @@ class PairLJSDKKokkos : public PairLJSDK { typename Kokkos::DualView::t_dev_const_um params; params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; - typename ArrayTypes::t_tagint_1d tag; + typename AT::t_x_array_randomread x; + typename AT::t_x_array c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + typename AT::t_tagint_1d tag; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; int newton_pair; double special_lj[4]; - typename ArrayTypes::tdual_ffloat_2d k_cutsq; - typename ArrayTypes::t_ffloat_2d d_cutsq; + typename AT::tdual_ffloat_2d k_cutsq; + typename AT::t_ffloat_2d d_cutsq; int neighflag; diff --git a/src/KOKKOS/pair_reax_c_kokkos.h b/src/KOKKOS/pair_reax_c_kokkos.h index 2d746dee0d..38ae3bd5b5 100644 --- a/src/KOKKOS/pair_reax_c_kokkos.h +++ b/src/KOKKOS/pair_reax_c_kokkos.h @@ -414,9 +414,9 @@ class PairReaxCKokkos : public PairReaxC { typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread; typedef typename tdual_ffloat_2d_n7::t_host t_host_ffloat_2d_n7; - typename ArrayTypes::t_neighbors_2d d_neighbors; - typename ArrayTypes::t_int_1d_randomread d_ilist; - typename ArrayTypes::t_int_1d_randomread d_numneigh; + typename AT::t_neighbors_2d d_neighbors; + typename AT::t_int_1d_randomread d_ilist; + typename AT::t_int_1d_randomread d_numneigh; typename AT::t_int_1d d_bo_first, d_bo_num, d_bo_list, d_hb_first, d_hb_num, d_hb_list; diff --git a/src/KOKKOS/pair_sw_kokkos.cpp b/src/KOKKOS/pair_sw_kokkos.cpp index 8d0f2fcfc3..1190374b30 100644 --- a/src/KOKKOS/pair_sw_kokkos.cpp +++ b/src/KOKKOS/pair_sw_kokkos.cpp @@ -171,8 +171,6 @@ void PairSWKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev_all.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -183,6 +181,8 @@ void PairSWKokkos::compute(int eflag_in, int vflag_in) k_vatom.template sync(); } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + copymode = 0; } diff --git a/src/KOKKOS/pair_sw_kokkos.h b/src/KOKKOS/pair_sw_kokkos.h index c722d9d52c..b5bd74d19b 100644 --- a/src/KOKKOS/pair_sw_kokkos.h +++ b/src/KOKKOS/pair_sw_kokkos.h @@ -118,10 +118,10 @@ class PairSWKokkos : public PairSW { void threebodyj(const Param&, const Param&, const Param&, const F_FLOAT&, const F_FLOAT&, F_FLOAT *, F_FLOAT *, F_FLOAT *) const; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_tagint_1d tag; - typename ArrayTypes::t_int_1d_randomread type; + typename AT::t_x_array_randomread x; + typename AT::t_f_array f; + typename AT::t_tagint_1d tag; + typename AT::t_int_1d_randomread type; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; @@ -132,9 +132,9 @@ class PairSWKokkos : public PairSW { DAT::t_int_2d_randomread d_type2rhor; DAT::t_int_2d_randomread d_type2z2r; - typename ArrayTypes::t_neighbors_2d d_neighbors; - typename ArrayTypes::t_int_1d_randomread d_ilist; - typename ArrayTypes::t_int_1d_randomread d_numneigh; + typename AT::t_neighbors_2d d_neighbors; + typename AT::t_int_1d_randomread d_ilist; + typename AT::t_int_1d_randomread d_numneigh; //NeighListKokkos k_list; int neighflag,newton_pair; diff --git a/src/KOKKOS/pair_table_kokkos.cpp b/src/KOKKOS/pair_table_kokkos.cpp index 88773af7b3..93556635c2 100644 --- a/src/KOKKOS/pair_table_kokkos.cpp +++ b/src/KOKKOS/pair_table_kokkos.cpp @@ -12,7 +12,7 @@ ------------------------------------------------------------------------- */ /* ---------------------------------------------------------------------- - Contributing author: Paul Crozier (SNL) + Contributing author: Christian Trott (SNL) ------------------------------------------------------------------------- */ #include @@ -89,6 +89,19 @@ void PairTableKokkos::compute_style(int eflag_in, int vflag_in) if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; + // reallocate per-atom arrays if necessary + + if (eflag_atom) { + memory->destroy_kokkos(k_eatom,eatom); + memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + d_eatom = k_eatom.view(); + } + if (vflag_atom) { + memory->destroy_kokkos(k_vatom,vatom); + memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + d_vatom = k_vatom.view(); + } + atomKK->sync(execution_space,datamask_read); //k_cutsq.template sync(); //k_params.template sync(); @@ -165,7 +178,18 @@ void PairTableKokkos::compute_style(int eflag_in, int vflag_in) virial[5] += ev.v[5]; } + if (eflag_atom) { + k_eatom.template modify(); + k_eatom.template sync(); + } + + if (vflag_atom) { + k_vatom.template modify(); + k_vatom.template sync(); + } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + } template diff --git a/src/KOKKOS/pair_table_kokkos.h b/src/KOKKOS/pair_table_kokkos.h index ce8b2e3ee5..9e785b68fd 100644 --- a/src/KOKKOS/pair_table_kokkos.h +++ b/src/KOKKOS/pair_table_kokkos.h @@ -44,6 +44,7 @@ class PairTableKokkos : public PairTable { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2}; enum {COUL_FLAG=0}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairTableKokkos(class LAMMPS *); virtual ~PairTableKokkos(); @@ -62,27 +63,27 @@ class PairTableKokkos : public PairTable { protected: /*struct TableDeviceConst { - typename ArrayTypes::t_ffloat_2d_randomread cutsq; - typename ArrayTypes::t_int_2d_randomread tabindex; - typename ArrayTypes::t_int_1d_randomread nshiftbits,nmask; - typename ArrayTypes::t_ffloat_1d_randomread innersq,invdelta,deltasq6; - typename ArrayTypes::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2; + typename AT::t_ffloat_2d_randomread cutsq; + typename AT::t_int_2d_randomread tabindex; + typename AT::t_int_1d_randomread nshiftbits,nmask; + typename AT::t_ffloat_1d_randomread innersq,invdelta,deltasq6; + typename AT::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2; };*/ //Its faster not to use texture fetch if the number of tables is less than 32! struct TableDeviceConst { - typename ArrayTypes::t_ffloat_2d cutsq; - typename ArrayTypes::t_int_2d tabindex; - typename ArrayTypes::t_int_1d nshiftbits,nmask; - typename ArrayTypes::t_ffloat_1d innersq,invdelta,deltasq6; - typename ArrayTypes::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2; + typename AT::t_ffloat_2d cutsq; + typename AT::t_int_2d tabindex; + typename AT::t_int_1d nshiftbits,nmask; + typename AT::t_ffloat_1d innersq,invdelta,deltasq6; + typename AT::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2; }; struct TableDevice { - typename ArrayTypes::t_ffloat_2d cutsq; - typename ArrayTypes::t_int_2d tabindex; - typename ArrayTypes::t_int_1d nshiftbits,nmask; - typename ArrayTypes::t_ffloat_1d innersq,invdelta,deltasq6; - typename ArrayTypes::t_ffloat_2d rsq,drsq,e,de,f,df,e2,f2; + typename AT::t_ffloat_2d cutsq; + typename AT::t_int_2d tabindex; + typename AT::t_int_1d nshiftbits,nmask; + typename AT::t_ffloat_1d innersq,invdelta,deltasq6; + typename AT::t_ffloat_2d rsq,drsq,e,de,f,df,e2,f2; }; struct TableHost { @@ -99,17 +100,20 @@ class PairTableKokkos : public PairTable { F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; - typename ArrayTypes::t_ffloat_2d d_cutsq; + typename AT::t_ffloat_2d d_cutsq; virtual void allocate(); void compute_table(Table *); - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_x_array_const c_x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_int_1d_randomread type; - typename ArrayTypes::t_efloat_1d d_eatom; - typename ArrayTypes::t_virial_array d_vatom; + typename AT::t_x_array_randomread x; + typename AT::t_x_array_const c_x; + typename AT::t_f_array f; + typename AT::t_int_1d_randomread type; + + DAT::tdual_efloat_1d k_eatom; + DAT::tdual_virial_array k_vatom; + typename AT::t_efloat_1d d_eatom; + typename AT::t_virial_array d_vatom; protected: int nlocal,nall,eflag,vflag,neighflag,newton_pair; diff --git a/src/KOKKOS/pair_tersoff_kokkos.cpp b/src/KOKKOS/pair_tersoff_kokkos.cpp index 66faa956c6..63f34b2786 100644 --- a/src/KOKKOS/pair_tersoff_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_kokkos.cpp @@ -254,8 +254,6 @@ void PairTersoffKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev_all.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -266,6 +264,8 @@ void PairTersoffKokkos::compute(int eflag_in, int vflag_in) k_vatom.template sync(); } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + copymode = 0; } diff --git a/src/KOKKOS/pair_tersoff_kokkos.h b/src/KOKKOS/pair_tersoff_kokkos.h index 9a3c152604..cc8c0c8ec7 100644 --- a/src/KOKKOS/pair_tersoff_kokkos.h +++ b/src/KOKKOS/pair_tersoff_kokkos.h @@ -204,9 +204,9 @@ class PairTersoffKokkos : public PairTersoff { typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread; typedef typename tdual_ffloat_2d_n7::t_host t_host_ffloat_2d_n7; - typename ArrayTypes::t_neighbors_2d d_neighbors; - typename ArrayTypes::t_int_1d_randomread d_ilist; - typename ArrayTypes::t_int_1d_randomread d_numneigh; + typename AT::t_neighbors_2d d_neighbors; + typename AT::t_int_1d_randomread d_ilist; + typename AT::t_int_1d_randomread d_numneigh; //NeighListKokkos k_list; int neighflag,newton_pair; diff --git a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp index d89b411110..bba13e16ac 100644 --- a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp @@ -254,8 +254,6 @@ void PairTersoffMODKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev_all.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -266,6 +264,8 @@ void PairTersoffMODKokkos::compute(int eflag_in, int vflag_in) k_vatom.template sync(); } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + copymode = 0; } diff --git a/src/KOKKOS/pair_tersoff_mod_kokkos.h b/src/KOKKOS/pair_tersoff_mod_kokkos.h index 38320ca1be..aaa4e51dd9 100644 --- a/src/KOKKOS/pair_tersoff_mod_kokkos.h +++ b/src/KOKKOS/pair_tersoff_mod_kokkos.h @@ -204,9 +204,9 @@ class PairTersoffMODKokkos : public PairTersoffMOD { typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread; typedef typename tdual_ffloat_2d_n7::t_host t_host_ffloat_2d_n7; - typename ArrayTypes::t_neighbors_2d d_neighbors; - typename ArrayTypes::t_int_1d_randomread d_ilist; - typename ArrayTypes::t_int_1d_randomread d_numneigh; + typename AT::t_neighbors_2d d_neighbors; + typename AT::t_int_1d_randomread d_ilist; + typename AT::t_int_1d_randomread d_numneigh; //NeighListKokkos k_list; int neighflag,newton_pair; diff --git a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp index 04195d4e45..b52b3686af 100644 --- a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp @@ -268,8 +268,6 @@ void PairTersoffZBLKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev_all.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -280,6 +278,8 @@ void PairTersoffZBLKokkos::compute(int eflag_in, int vflag_in) k_vatom.template sync(); } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + copymode = 0; } diff --git a/src/KOKKOS/pair_vashishta_kokkos.cpp b/src/KOKKOS/pair_vashishta_kokkos.cpp index 73e4e04f98..a87766bfb2 100644 --- a/src/KOKKOS/pair_vashishta_kokkos.cpp +++ b/src/KOKKOS/pair_vashishta_kokkos.cpp @@ -12,8 +12,7 @@ ------------------------------------------------------------------------- */ /* ---------------------------------------------------------------------- - Contributing author: Stan Moore (SNL) - Anders Hafreager (UiO), andershaf@gmail.com + Contributing author: Anders Hafreager (UiO), andershaf@gmail.com ------------------------------------------------------------------------- */ #include @@ -182,8 +181,6 @@ void PairVashishtaKokkos::compute(int eflag_in, int vflag_in) virial[5] += ev_all.v[5]; } - if (vflag_fdotr) pair_virial_fdotr_compute(this); - if (eflag_atom) { k_eatom.template modify(); k_eatom.template sync(); @@ -194,6 +191,8 @@ void PairVashishtaKokkos::compute(int eflag_in, int vflag_in) k_vatom.template sync(); } + if (vflag_fdotr) pair_virial_fdotr_compute(this); + copymode = 0; } diff --git a/src/KOKKOS/pair_vashishta_kokkos.h b/src/KOKKOS/pair_vashishta_kokkos.h index 8bc4d621c2..1793f0ee5e 100644 --- a/src/KOKKOS/pair_vashishta_kokkos.h +++ b/src/KOKKOS/pair_vashishta_kokkos.h @@ -118,10 +118,10 @@ class PairVashishtaKokkos : public PairVashishta { void threebodyj(const Param&, const Param&, const Param&, const F_FLOAT&, const F_FLOAT&, F_FLOAT *, F_FLOAT *, F_FLOAT *) const; - typename ArrayTypes::t_x_array_randomread x; - typename ArrayTypes::t_f_array f; - typename ArrayTypes::t_tagint_1d tag; - typename ArrayTypes::t_int_1d_randomread type; + typename AT::t_x_array_randomread x; + typename AT::t_f_array f; + typename AT::t_tagint_1d tag; + typename AT::t_int_1d_randomread type; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; @@ -132,9 +132,9 @@ class PairVashishtaKokkos : public PairVashishta { DAT::t_int_2d_randomread d_type2rhor; DAT::t_int_2d_randomread d_type2z2r; - typename ArrayTypes::t_neighbors_2d d_neighbors; - typename ArrayTypes::t_int_1d_randomread d_ilist; - typename ArrayTypes::t_int_1d_randomread d_numneigh; + typename AT::t_neighbors_2d d_neighbors; + typename AT::t_int_1d_randomread d_ilist; + typename AT::t_int_1d_randomread d_numneigh; //NeighListKokkos k_list; int neighflag,newton_pair; From 38e0e4bb69ff7859f6e0fa7db8c774f488d5c866 Mon Sep 17 00:00:00 2001 From: Stan Moore Date: Thu, 2 Feb 2017 13:24:05 -0700 Subject: [PATCH 20/25] Add missing typedef in Kokkos pair styles --- src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h | 1 + src/KOKKOS/pair_lj_expand_kokkos.h | 1 + 2 files changed, 2 insertions(+) diff --git a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h index e0ab231359..2f13db57c7 100644 --- a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h @@ -34,6 +34,7 @@ class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJCutCoulDebyeKokkos(class LAMMPS *); ~PairLJCutCoulDebyeKokkos(); diff --git a/src/KOKKOS/pair_lj_expand_kokkos.h b/src/KOKKOS/pair_lj_expand_kokkos.h index 3b0067bcd8..093031f7d7 100644 --- a/src/KOKKOS/pair_lj_expand_kokkos.h +++ b/src/KOKKOS/pair_lj_expand_kokkos.h @@ -34,6 +34,7 @@ class PairLJExpandKokkos : public PairLJExpand { enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2}; enum {COUL_FLAG=0}; typedef DeviceType device_type; + typedef ArrayTypes AT; PairLJExpandKokkos(class LAMMPS *); ~PairLJExpandKokkos(); From 6f92429602d6946ca754f44046e0cf6375dcd66e Mon Sep 17 00:00:00 2001 From: Stan Moore Date: Thu, 2 Feb 2017 13:34:27 -0700 Subject: [PATCH 21/25] Fixing per-atom ev issue --- src/KOKKOS/angle_charmm_kokkos.cpp | 8 ++++---- src/KOKKOS/bond_harmonic_kokkos.cpp | 8 ++++---- src/KOKKOS/dihedral_charmm_kokkos.cpp | 8 ++++---- src/KOKKOS/improper_harmonic_kokkos.cpp | 8 ++++---- 4 files changed, 16 insertions(+), 16 deletions(-) diff --git a/src/KOKKOS/angle_charmm_kokkos.cpp b/src/KOKKOS/angle_charmm_kokkos.cpp index dcd5d91e05..b4e2d47f36 100644 --- a/src/KOKKOS/angle_charmm_kokkos.cpp +++ b/src/KOKKOS/angle_charmm_kokkos.cpp @@ -70,18 +70,18 @@ void AngleCharmmKokkos::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - if(k_eatom.dimension_0()destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.d_view; - } + //} } if (vflag_atom) { - if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.d_view; - } + //} } x = atomKK->k_x.view(); diff --git a/src/KOKKOS/bond_harmonic_kokkos.cpp b/src/KOKKOS/bond_harmonic_kokkos.cpp index 9b34786ec0..d07923901f 100644 --- a/src/KOKKOS/bond_harmonic_kokkos.cpp +++ b/src/KOKKOS/bond_harmonic_kokkos.cpp @@ -67,18 +67,18 @@ void BondHarmonicKokkos::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - if(k_eatom.dimension_0()destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.d_view; - } + //} } if (vflag_atom) { - if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.d_view; - } + //} } // if (eflag || vflag) atomKK->modified(execution_space,datamask_modify); diff --git a/src/KOKKOS/dihedral_charmm_kokkos.cpp b/src/KOKKOS/dihedral_charmm_kokkos.cpp index b0dd1a2172..e623b70857 100644 --- a/src/KOKKOS/dihedral_charmm_kokkos.cpp +++ b/src/KOKKOS/dihedral_charmm_kokkos.cpp @@ -80,22 +80,22 @@ void DihedralCharmmKokkos::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - if(k_eatom.dimension_0()destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); d_eatom = k_eatom.d_view; k_eatom_pair = Kokkos::DualView("dihedral:eatom_pair",maxeatom); d_eatom_pair = k_eatom.d_view; - } + //} } if (vflag_atom) { - if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom"); d_vatom = k_vatom.d_view; k_vatom_pair = Kokkos::DualView("dihedral:vatom_pair",maxvatom); d_vatom_pair = k_vatom.d_view; - } + //} } x = atomKK->k_x.view(); diff --git a/src/KOKKOS/improper_harmonic_kokkos.cpp b/src/KOKKOS/improper_harmonic_kokkos.cpp index 34d3d437d6..0bf729cdfd 100644 --- a/src/KOKKOS/improper_harmonic_kokkos.cpp +++ b/src/KOKKOS/improper_harmonic_kokkos.cpp @@ -77,18 +77,18 @@ void ImproperHarmonicKokkos::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - if(k_eatom.dimension_0()destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.d_view; - } + //} } if (vflag_atom) { - if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.d_view; - } + //} } //atomKK->sync(execution_space,datamask_read); From 5da0d393921df2212831d2a6af7ae8500b7218bd Mon Sep 17 00:00:00 2001 From: Stan Moore Date: Thu, 2 Feb 2017 13:35:51 -0700 Subject: [PATCH 22/25] Fixing fdotr in pair_buck_coul_cut_kokkos --- src/KOKKOS/pair_buck_coul_cut_kokkos.h | 1 + 1 file changed, 1 insertion(+) diff --git a/src/KOKKOS/pair_buck_coul_cut_kokkos.h b/src/KOKKOS/pair_buck_coul_cut_kokkos.h index fb0034d767..f758fe3206 100644 --- a/src/KOKKOS/pair_buck_coul_cut_kokkos.h +++ b/src/KOKKOS/pair_buck_coul_cut_kokkos.h @@ -124,6 +124,7 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut { friend EV_FLOAT pair_compute_neighlist(PairBuckCoulCutKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairBuckCoulCutKokkos*, NeighListKokkos*); + friend void pair_virial_fdotr_compute(PairBuckCoulLongKokkos*); }; From 6a4918b39abef87e9e3e0fc03fe4fa2896b2202c Mon Sep 17 00:00:00 2001 From: Stan Moore Date: Thu, 2 Feb 2017 13:43:21 -0700 Subject: [PATCH 23/25] Fixing typo in pair_buck_coul_cut_kokkos --- src/KOKKOS/pair_buck_coul_cut_kokkos.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/KOKKOS/pair_buck_coul_cut_kokkos.h b/src/KOKKOS/pair_buck_coul_cut_kokkos.h index f758fe3206..d026da458f 100644 --- a/src/KOKKOS/pair_buck_coul_cut_kokkos.h +++ b/src/KOKKOS/pair_buck_coul_cut_kokkos.h @@ -124,7 +124,7 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut { friend EV_FLOAT pair_compute_neighlist(PairBuckCoulCutKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairBuckCoulCutKokkos*, NeighListKokkos*); - friend void pair_virial_fdotr_compute(PairBuckCoulLongKokkos*); + friend void pair_virial_fdotr_compute(PairBuckCoulCutKokkos*); }; From 1834a5e46c3a482df1740b3a5c37a07ffe277b67 Mon Sep 17 00:00:00 2001 From: Stan Moore Date: Thu, 2 Feb 2017 14:58:04 -0700 Subject: [PATCH 24/25] Fixing more Kokkos per-atom and fdotr issues --- src/KOKKOS/angle_charmm_kokkos.cpp | 4 ++-- src/KOKKOS/bond_fene_kokkos.cpp | 4 ++-- src/KOKKOS/bond_fene_kokkos.h | 4 ++-- src/KOKKOS/bond_harmonic_kokkos.cpp | 4 ++-- src/KOKKOS/dihedral_charmm_kokkos.cpp | 8 ++++---- src/KOKKOS/dihedral_opls_kokkos.cpp | 4 ++-- src/KOKKOS/dihedral_opls_kokkos.h | 4 ++-- src/KOKKOS/improper_harmonic_kokkos.cpp | 4 ++-- src/KOKKOS/pair_coul_debye_kokkos.h | 1 + src/KOKKOS/pair_coul_dsf_kokkos.cpp | 4 ++-- src/KOKKOS/pair_coul_dsf_kokkos.h | 4 ++-- src/KOKKOS/pair_coul_wolf_kokkos.cpp | 4 ++-- src/KOKKOS/pair_coul_wolf_kokkos.h | 4 ++-- src/KOKKOS/pair_eam_alloy_kokkos.cpp | 4 ++-- src/KOKKOS/pair_eam_alloy_kokkos.h | 4 ++-- src/KOKKOS/pair_eam_fs_kokkos.cpp | 4 ++-- src/KOKKOS/pair_eam_fs_kokkos.h | 4 ++-- src/KOKKOS/pair_eam_kokkos.cpp | 4 ++-- src/KOKKOS/pair_eam_kokkos.h | 4 ++-- src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h | 1 + src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h | 1 + src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h | 1 + src/KOKKOS/pair_reax_c_kokkos.h | 2 +- src/KOKKOS/pair_sw_kokkos.cpp | 4 ++-- src/KOKKOS/pair_sw_kokkos.h | 4 ++-- src/KOKKOS/pair_tersoff_kokkos.cpp | 4 ++-- src/KOKKOS/pair_tersoff_kokkos.h | 4 ++-- src/KOKKOS/pair_tersoff_mod_kokkos.cpp | 4 ++-- src/KOKKOS/pair_tersoff_mod_kokkos.h | 4 ++-- src/KOKKOS/pair_tersoff_zbl_kokkos.cpp | 4 ++-- src/KOKKOS/pair_tersoff_zbl_kokkos.h | 4 ++-- src/KOKKOS/pair_vashishta_kokkos.cpp | 4 ++-- src/KOKKOS/pair_vashishta_kokkos.h | 4 ++-- 33 files changed, 63 insertions(+), 59 deletions(-) diff --git a/src/KOKKOS/angle_charmm_kokkos.cpp b/src/KOKKOS/angle_charmm_kokkos.cpp index b4e2d47f36..e2a605b802 100644 --- a/src/KOKKOS/angle_charmm_kokkos.cpp +++ b/src/KOKKOS/angle_charmm_kokkos.cpp @@ -73,14 +73,14 @@ void AngleCharmmKokkos::compute(int eflag_in, int vflag_in) //if(k_eatom.dimension_0()destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.template view(); //} } if (vflag_atom) { //if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.template view(); //} } diff --git a/src/KOKKOS/bond_fene_kokkos.cpp b/src/KOKKOS/bond_fene_kokkos.cpp index b8b0295182..b10469b45e 100644 --- a/src/KOKKOS/bond_fene_kokkos.cpp +++ b/src/KOKKOS/bond_fene_kokkos.cpp @@ -77,12 +77,12 @@ void BondFENEKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"bond:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"bond:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } k_k.template sync(); diff --git a/src/KOKKOS/bond_fene_kokkos.h b/src/KOKKOS/bond_fene_kokkos.h index ef37992925..c630abcafc 100644 --- a/src/KOKKOS/bond_fene_kokkos.h +++ b/src/KOKKOS/bond_fene_kokkos.h @@ -66,8 +66,8 @@ class BondFENEKokkos : public BondFENE { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; DAT::tdual_int_scalar k_warning_flag; typename AT::t_int_scalar d_warning_flag; diff --git a/src/KOKKOS/bond_harmonic_kokkos.cpp b/src/KOKKOS/bond_harmonic_kokkos.cpp index d07923901f..aead602747 100644 --- a/src/KOKKOS/bond_harmonic_kokkos.cpp +++ b/src/KOKKOS/bond_harmonic_kokkos.cpp @@ -70,14 +70,14 @@ void BondHarmonicKokkos::compute(int eflag_in, int vflag_in) //if(k_eatom.dimension_0()destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.template view(); //} } if (vflag_atom) { //if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.template view(); //} } diff --git a/src/KOKKOS/dihedral_charmm_kokkos.cpp b/src/KOKKOS/dihedral_charmm_kokkos.cpp index e623b70857..0af9fc4e83 100644 --- a/src/KOKKOS/dihedral_charmm_kokkos.cpp +++ b/src/KOKKOS/dihedral_charmm_kokkos.cpp @@ -83,18 +83,18 @@ void DihedralCharmmKokkos::compute(int eflag_in, int vflag_in) //if(k_eatom.dimension_0()destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.template view(); k_eatom_pair = Kokkos::DualView("dihedral:eatom_pair",maxeatom); - d_eatom_pair = k_eatom.d_view; + d_eatom_pair = k_eatom.template view(); //} } if (vflag_atom) { //if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.template view(); k_vatom_pair = Kokkos::DualView("dihedral:vatom_pair",maxvatom); - d_vatom_pair = k_vatom.d_view; + d_vatom_pair = k_vatom.template view(); //} } diff --git a/src/KOKKOS/dihedral_opls_kokkos.cpp b/src/KOKKOS/dihedral_opls_kokkos.cpp index ed4a5b259c..5406a6daaf 100644 --- a/src/KOKKOS/dihedral_opls_kokkos.cpp +++ b/src/KOKKOS/dihedral_opls_kokkos.cpp @@ -77,12 +77,12 @@ void DihedralOPLSKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } k_k1.template sync(); diff --git a/src/KOKKOS/dihedral_opls_kokkos.h b/src/KOKKOS/dihedral_opls_kokkos.h index 5bf2552d9b..0a8860c87c 100644 --- a/src/KOKKOS/dihedral_opls_kokkos.h +++ b/src/KOKKOS/dihedral_opls_kokkos.h @@ -68,8 +68,8 @@ class DihedralOPLSKokkos : public DihedralOPLS { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; int nlocal,newton_bond; int eflag,vflag; diff --git a/src/KOKKOS/improper_harmonic_kokkos.cpp b/src/KOKKOS/improper_harmonic_kokkos.cpp index 0bf729cdfd..cde7d60e8e 100644 --- a/src/KOKKOS/improper_harmonic_kokkos.cpp +++ b/src/KOKKOS/improper_harmonic_kokkos.cpp @@ -80,14 +80,14 @@ void ImproperHarmonicKokkos::compute(int eflag_in, int vflag_in) //if(k_eatom.dimension_0()destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.template view(); //} } if (vflag_atom) { //if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.template view(); //} } diff --git a/src/KOKKOS/pair_coul_debye_kokkos.h b/src/KOKKOS/pair_coul_debye_kokkos.h index 143b1c2354..b6e87b55f7 100644 --- a/src/KOKKOS/pair_coul_debye_kokkos.h +++ b/src/KOKKOS/pair_coul_debye_kokkos.h @@ -128,6 +128,7 @@ class PairCoulDebyeKokkos : public PairCoulDebye { friend EV_FLOAT pair_compute_neighlist(PairCoulDebyeKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairCoulDebyeKokkos*, NeighListKokkos*); + friend void pair_virial_fdotr_compute(PairCoulDebyeKokkos*); }; diff --git a/src/KOKKOS/pair_coul_dsf_kokkos.cpp b/src/KOKKOS/pair_coul_dsf_kokkos.cpp index 3ca1fb8ea4..948a234a4a 100644 --- a/src/KOKKOS/pair_coul_dsf_kokkos.cpp +++ b/src/KOKKOS/pair_coul_dsf_kokkos.cpp @@ -88,12 +88,12 @@ void PairCoulDSFKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } atomKK->sync(execution_space,datamask_read); diff --git a/src/KOKKOS/pair_coul_dsf_kokkos.h b/src/KOKKOS/pair_coul_dsf_kokkos.h index 33740aa916..49b1806e3c 100644 --- a/src/KOKKOS/pair_coul_dsf_kokkos.h +++ b/src/KOKKOS/pair_coul_dsf_kokkos.h @@ -62,12 +62,12 @@ class PairCoulDSFKokkos : public PairCoulDSF { KOKKOS_INLINE_FUNCTION int sbmask(const int& j) const; - protected: - typename AT::t_x_array_randomread x; typename AT::t_f_array f; typename AT::t_float_1d_randomread q; + protected: + DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; typename AT::t_efloat_1d d_eatom; diff --git a/src/KOKKOS/pair_coul_wolf_kokkos.cpp b/src/KOKKOS/pair_coul_wolf_kokkos.cpp index 22a11c4a18..55b19958a3 100644 --- a/src/KOKKOS/pair_coul_wolf_kokkos.cpp +++ b/src/KOKKOS/pair_coul_wolf_kokkos.cpp @@ -83,12 +83,12 @@ void PairCoulWolfKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } atomKK->sync(execution_space,datamask_read); diff --git a/src/KOKKOS/pair_coul_wolf_kokkos.h b/src/KOKKOS/pair_coul_wolf_kokkos.h index cedafe7a21..f6da745de3 100644 --- a/src/KOKKOS/pair_coul_wolf_kokkos.h +++ b/src/KOKKOS/pair_coul_wolf_kokkos.h @@ -70,8 +70,8 @@ class PairCoulWolfKokkos : public PairCoulWolf { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; int neighflag,newton_pair; diff --git a/src/KOKKOS/pair_eam_alloy_kokkos.cpp b/src/KOKKOS/pair_eam_alloy_kokkos.cpp index 866fe2cfa1..491ec10f66 100644 --- a/src/KOKKOS/pair_eam_alloy_kokkos.cpp +++ b/src/KOKKOS/pair_eam_alloy_kokkos.cpp @@ -82,12 +82,12 @@ void PairEAMAlloyKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } atomKK->sync(execution_space,datamask_read); diff --git a/src/KOKKOS/pair_eam_alloy_kokkos.h b/src/KOKKOS/pair_eam_alloy_kokkos.h index 44420f2d9e..d8c745b230 100644 --- a/src/KOKKOS/pair_eam_alloy_kokkos.h +++ b/src/KOKKOS/pair_eam_alloy_kokkos.h @@ -125,8 +125,8 @@ class PairEAMAlloyKokkos : public PairEAM { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; DAT::tdual_ffloat_1d k_rho; DAT::tdual_ffloat_1d k_fp; diff --git a/src/KOKKOS/pair_eam_fs_kokkos.cpp b/src/KOKKOS/pair_eam_fs_kokkos.cpp index ba450b0872..006db6abc2 100644 --- a/src/KOKKOS/pair_eam_fs_kokkos.cpp +++ b/src/KOKKOS/pair_eam_fs_kokkos.cpp @@ -82,12 +82,12 @@ void PairEAMFSKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } atomKK->sync(execution_space,datamask_read); diff --git a/src/KOKKOS/pair_eam_fs_kokkos.h b/src/KOKKOS/pair_eam_fs_kokkos.h index 7bb4d0abfb..40375af066 100644 --- a/src/KOKKOS/pair_eam_fs_kokkos.h +++ b/src/KOKKOS/pair_eam_fs_kokkos.h @@ -125,8 +125,8 @@ class PairEAMFSKokkos : public PairEAM { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; DAT::tdual_ffloat_1d k_rho; DAT::tdual_ffloat_1d k_fp; diff --git a/src/KOKKOS/pair_eam_kokkos.cpp b/src/KOKKOS/pair_eam_kokkos.cpp index 3d8223ed66..656dd11c62 100644 --- a/src/KOKKOS/pair_eam_kokkos.cpp +++ b/src/KOKKOS/pair_eam_kokkos.cpp @@ -77,12 +77,12 @@ void PairEAMKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } atomKK->sync(execution_space,datamask_read); diff --git a/src/KOKKOS/pair_eam_kokkos.h b/src/KOKKOS/pair_eam_kokkos.h index d36ef0f95a..229b16471d 100644 --- a/src/KOKKOS/pair_eam_kokkos.h +++ b/src/KOKKOS/pair_eam_kokkos.h @@ -122,8 +122,8 @@ class PairEAMKokkos : public PairEAM { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; DAT::tdual_ffloat_1d k_rho; DAT::tdual_ffloat_1d k_fp; diff --git a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h index 7dcee07400..4dd9379821 100644 --- a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h @@ -116,6 +116,7 @@ class PairLJCutCoulCutKokkos : public PairLJCutCoulCut { friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulCutKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairLJCutCoulCutKokkos*, NeighListKokkos*); + friend void pair_virial_fdotr_compute(PairLJCutCoulCutKokkos*); }; diff --git a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h index 2f13db57c7..315969bc62 100644 --- a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h @@ -116,6 +116,7 @@ class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye { friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulDebyeKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairLJCutCoulDebyeKokkos*, NeighListKokkos*); + friend void pair_virial_fdotr_compute(PairLJCutCoulDebyeKokkos*); }; diff --git a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h index 13030fbf7c..558d40e174 100644 --- a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h @@ -115,6 +115,7 @@ class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF { friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulDSFKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairLJCutCoulDSFKokkos*, NeighListKokkos*); + friend void pair_virial_fdotr_compute(PairLJCutCoulDSFKokkos*); }; diff --git a/src/KOKKOS/pair_reax_c_kokkos.h b/src/KOKKOS/pair_reax_c_kokkos.h index 38ae3bd5b5..8a0c08b660 100644 --- a/src/KOKKOS/pair_reax_c_kokkos.h +++ b/src/KOKKOS/pair_reax_c_kokkos.h @@ -392,7 +392,7 @@ class PairReaxCKokkos : public PairReaxC { typename AT::t_efloat_1d v_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_virial_array d_vatom; typename AT::t_virial_array v_vatom; HAT::t_virial_array h_vatom; diff --git a/src/KOKKOS/pair_sw_kokkos.cpp b/src/KOKKOS/pair_sw_kokkos.cpp index 1190374b30..6b2c1ca04b 100644 --- a/src/KOKKOS/pair_sw_kokkos.cpp +++ b/src/KOKKOS/pair_sw_kokkos.cpp @@ -88,12 +88,12 @@ void PairSWKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } atomKK->sync(execution_space,datamask_read); diff --git a/src/KOKKOS/pair_sw_kokkos.h b/src/KOKKOS/pair_sw_kokkos.h index b5bd74d19b..d899edfc1b 100644 --- a/src/KOKKOS/pair_sw_kokkos.h +++ b/src/KOKKOS/pair_sw_kokkos.h @@ -125,8 +125,8 @@ class PairSWKokkos : public PairSW { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; DAT::t_int_1d_randomread d_type2frho; DAT::t_int_2d_randomread d_type2rhor; diff --git a/src/KOKKOS/pair_tersoff_kokkos.cpp b/src/KOKKOS/pair_tersoff_kokkos.cpp index 63f34b2786..6e1981a31e 100644 --- a/src/KOKKOS/pair_tersoff_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_kokkos.cpp @@ -172,12 +172,12 @@ void PairTersoffKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } atomKK->sync(execution_space,datamask_read); diff --git a/src/KOKKOS/pair_tersoff_kokkos.h b/src/KOKKOS/pair_tersoff_kokkos.h index cc8c0c8ec7..d34cbe0783 100644 --- a/src/KOKKOS/pair_tersoff_kokkos.h +++ b/src/KOKKOS/pair_tersoff_kokkos.h @@ -197,8 +197,8 @@ class PairTersoffKokkos : public PairTersoff { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; typedef Kokkos::DualView tdual_ffloat_2d_n7; typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread; diff --git a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp index bba13e16ac..0be7a7c05b 100644 --- a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp @@ -172,12 +172,12 @@ void PairTersoffMODKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } atomKK->sync(execution_space,datamask_read); diff --git a/src/KOKKOS/pair_tersoff_mod_kokkos.h b/src/KOKKOS/pair_tersoff_mod_kokkos.h index aaa4e51dd9..540aad659e 100644 --- a/src/KOKKOS/pair_tersoff_mod_kokkos.h +++ b/src/KOKKOS/pair_tersoff_mod_kokkos.h @@ -197,8 +197,8 @@ class PairTersoffMODKokkos : public PairTersoffMOD { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; typedef Kokkos::DualView tdual_ffloat_2d_n7; typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread; diff --git a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp index b52b3686af..6d89226324 100644 --- a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp @@ -186,12 +186,12 @@ void PairTersoffZBLKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } atomKK->sync(execution_space,datamask_read); diff --git a/src/KOKKOS/pair_tersoff_zbl_kokkos.h b/src/KOKKOS/pair_tersoff_zbl_kokkos.h index 31c5912ff6..58d21ecae3 100644 --- a/src/KOKKOS/pair_tersoff_zbl_kokkos.h +++ b/src/KOKKOS/pair_tersoff_zbl_kokkos.h @@ -202,8 +202,8 @@ class PairTersoffZBLKokkos : public PairTersoffZBL { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; typedef Kokkos::DualView tdual_ffloat_2d_n7; typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread; diff --git a/src/KOKKOS/pair_vashishta_kokkos.cpp b/src/KOKKOS/pair_vashishta_kokkos.cpp index a87766bfb2..2e73bd6eb4 100644 --- a/src/KOKKOS/pair_vashishta_kokkos.cpp +++ b/src/KOKKOS/pair_vashishta_kokkos.cpp @@ -87,12 +87,12 @@ void PairVashishtaKokkos::compute(int eflag_in, int vflag_in) if (eflag_atom) { memory->destroy_kokkos(k_eatom,eatom); memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); - d_eatom = k_eatom.d_view; + d_eatom = k_eatom.view(); } if (vflag_atom) { memory->destroy_kokkos(k_vatom,vatom); memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); - d_vatom = k_vatom.d_view; + d_vatom = k_vatom.view(); } atomKK->sync(execution_space,datamask_read); diff --git a/src/KOKKOS/pair_vashishta_kokkos.h b/src/KOKKOS/pair_vashishta_kokkos.h index 1793f0ee5e..49c936185d 100644 --- a/src/KOKKOS/pair_vashishta_kokkos.h +++ b/src/KOKKOS/pair_vashishta_kokkos.h @@ -125,8 +125,8 @@ class PairVashishtaKokkos : public PairVashishta { DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; - DAT::t_efloat_1d d_eatom; - DAT::t_virial_array d_vatom; + typename ArrayTypes::t_efloat_1d d_eatom; + typename ArrayTypes::t_virial_array d_vatom; DAT::t_int_1d_randomread d_type2frho; DAT::t_int_2d_randomread d_type2rhor; From 067119f6c6420174da8b8c61b08a7a2dae027a24 Mon Sep 17 00:00:00 2001 From: Stan Moore Date: Thu, 2 Feb 2017 15:13:51 -0700 Subject: [PATCH 25/25] Adding missing friend statement to pair_lj_class2_coul_cut_kokkos --- src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h | 1 + 1 file changed, 1 insertion(+) diff --git a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h index 1679460dbe..991d7c9626 100644 --- a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h +++ b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h @@ -116,6 +116,7 @@ class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut { friend EV_FLOAT pair_compute_neighlist(PairLJClass2CoulCutKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairLJClass2CoulCutKokkos*, NeighListKokkos*); + friend void pair_virial_fdotr_compute(PairLJClass2CoulCutKokkos*); };