jianmu
/
lammps
forked from lijiext/lammps
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Remove deprecated Kokkos code

This commit is contained in:
Stan Moore 2018-03-22 11:19:17 -06:00
parent f6c76e04b8
commit a6d10ba91d
40 changed files with 241 additions and 241 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
src/KOKKOS/angle_charmm_kokkos.cpp
View File

@ -70,14 +70,14 @@ void AngleCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
if (eflag_atom) {
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
//if(k_eatom.extent(0)<maxeatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_eatom,eatom);
memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
d_eatom = k_eatom.template view<DeviceType>();
//}
}
if (vflag_atom) {
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
//if(k_vatom.extent(0)<maxvatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_vatom,vatom);
memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
d_vatom = k_vatom.template view<DeviceType>();

2
src/KOKKOS/atom_kokkos.cpp
View File

@ -119,7 +119,7 @@ void AtomKokkos::allocate_type_arrays()
{
if (avec->mass_type) {
k_mass = DAT::tdual_float_1d("Mass",ntypes+1);
mass = k_mass.h_view.ptr_on_device();
mass = k_mass.h_view.data();
mass_setflag = new int[ntypes+1];
for (int itype = 1; itype <= ntypes; itype++) mass_setflag[itype] = 0;
k_mass.modify<LMPHostType>();

20
src/KOKKOS/atom_kokkos.h
View File

@ -84,34 +84,34 @@ class SortFunctor {
Kokkos::View<typename ViewType::non_const_data_type,typename ViewType::array_type,device_type> dest;
IndexView index;
SortFunctor(ViewType src, typename Kokkos::Impl::enable_if<ViewType::dynamic_rank==1,IndexView>::type ind):source(src),index(ind){
dest = Kokkos::View<typename ViewType::non_const_data_type,typename ViewType::array_type,device_type>("",src.dimension_0());
dest = Kokkos::View<typename ViewType::non_const_data_type,typename ViewType::array_type,device_type>("",src.extent(0));
}
SortFunctor(ViewType src, typename Kokkos::Impl::enable_if<ViewType::dynamic_rank==2,IndexView>::type ind):source(src),index(ind){
dest = Kokkos::View<typename ViewType::non_const_data_type,typename ViewType::array_type,device_type>("",src.dimension_0(),src.dimension_1());
dest = Kokkos::View<typename ViewType::non_const_data_type,typename ViewType::array_type,device_type>("",src.extent(0),src.extent(1));
}
SortFunctor(ViewType src, typename Kokkos::Impl::enable_if<ViewType::dynamic_rank==3,IndexView>::type ind):source(src),index(ind){
dest = Kokkos::View<typename ViewType::non_const_data_type,typename ViewType::array_type,device_type>("",src.dimension_0(),src.dimension_1(),src.dimension_2());
dest = Kokkos::View<typename ViewType::non_const_data_type,typename ViewType::array_type,device_type>("",src.extent(0),src.extent(1),src.extent(2));
}
SortFunctor(ViewType src, typename Kokkos::Impl::enable_if<ViewType::dynamic_rank==4,IndexView>::type ind):source(src),index(ind){
dest = Kokkos::View<typename ViewType::non_const_data_type,typename ViewType::array_type,device_type>("",src.dimension_0(),src.dimension_1(),src.dimension_2(),src.dimension_3());
dest = Kokkos::View<typename ViewType::non_const_data_type,typename ViewType::array_type,device_type>("",src.extent(0),src.extent(1),src.extent(2),src.extent(3));
}
KOKKOS_INLINE_FUNCTION
void operator()(const typename Kokkos::Impl::enable_if<ViewType::rank==1, int>::type& i) {
dest(i) = source(index(i));
}
void operator()(const typename Kokkos::Impl::enable_if<ViewType::rank==2, int>::type& i) {
for(int j=0;j<source.dimension_1();j++)
for(int j=0;j<source.extent(1);j++)
dest(i,j) = source(index(i),j);
}
void operator()(const typename Kokkos::Impl::enable_if<ViewType::rank==3, int>::type& i) {
for(int j=0;j<source.dimension_1();j++)
for(int k=0;k<source.dimension_2();k++)
for(int j=0;j<source.extent(1);j++)
for(int k=0;k<source.extent(2);k++)
dest(i,j,k) = source(index(i),j,k);
}
void operator()(const typename Kokkos::Impl::enable_if<ViewType::rank==4, int>::type& i) {
for(int j=0;j<source.dimension_1();j++)
for(int k=0;k<source.dimension_2();k++)
for(int l=0;l<source.dimension_3();l++)
for(int j=0;j<source.extent(1);j++)
for(int k=0;k<source.extent(2);k++)
for(int l=0;l<source.extent(3);l++)
dest(i,j,k,l) = source(index(i),j,k,l);
}
};

20
src/KOKKOS/atom_vec_angle_kokkos.cpp
View File

@ -240,8 +240,8 @@ struct AtomVecAngleKokkos_PackComm {
_x(x.view<DeviceType>()),_list(list.view<DeviceType>()),_iswap(iswap),
_xprd(xprd),_yprd(yprd),_zprd(zprd),
_xy(xy),_xz(xz),_yz(yz) {
const size_t maxsend = (buf.view<DeviceType>().dimension_0()
*buf.view<DeviceType>().dimension_1())/3;
const size_t maxsend = (buf.view<DeviceType>().extent(0)
*buf.view<DeviceType>().extent(1))/3;
const size_t elements = 3;
buffer_view<DeviceType>(_buf,buf,maxsend,elements);
_pbc[0] = pbc[0]; _pbc[1] = pbc[1]; _pbc[2] = pbc[2];
@ -1136,8 +1136,8 @@ struct AtomVecAngleKokkos_PackExchangeFunctor {
// and angle_atom3
// 1 to store buffer length
elements = 17+atom->maxspecial+2*atom->bond_per_atom+4*atom->angle_per_atom;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*
buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*
buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}
@ -1220,10 +1220,10 @@ int AtomVecAngleKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_
X_FLOAT hi )
{
const int elements = 17+atom->maxspecial+2*atom->bond_per_atom+4*atom->angle_per_atom;
if(nsend > (int) (k_buf.view<LMPHostType>().dimension_0()*
k_buf.view<LMPHostType>().dimension_1())/elements) {
int newsize = nsend*elements/k_buf.view<LMPHostType>().dimension_1()+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().dimension_1());
if(nsend > (int) (k_buf.view<LMPHostType>().extent(0)*
k_buf.view<LMPHostType>().extent(1))/elements) {
int newsize = nsend*elements/k_buf.view<LMPHostType>().extent(1)+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().extent(1));
}
if(space == Host) {
AtomVecAngleKokkos_PackExchangeFunctor<LMPHostType>
@ -1335,8 +1335,8 @@ struct AtomVecAngleKokkos_UnpackExchangeFunctor {
_nlocal(nlocal.template view<DeviceType>()),_dim(dim),
_lo(lo),_hi(hi){
elements =17+atom->maxspecial+2*atom->bond_per_atom+4*atom->angle_per_atom;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*
buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*
buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}

10
src/KOKKOS/atom_vec_atomic_kokkos.cpp
View File

@ -506,7 +506,7 @@ struct AtomVecAtomicKokkos_PackExchangeFunctor {
_nlocal(nlocal),_dim(dim),
_lo(lo),_hi(hi){
const size_t elements = 11;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}
@ -546,9 +546,9 @@ struct AtomVecAtomicKokkos_PackExchangeFunctor {
int AtomVecAtomicKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_2d &k_buf, DAT::tdual_int_1d k_sendlist,DAT::tdual_int_1d k_copylist,ExecutionSpace space,int dim,X_FLOAT lo,X_FLOAT hi )
{
if(nsend > (int) (k_buf.view<LMPHostType>().dimension_0()*k_buf.view<LMPHostType>().dimension_1())/11) {
int newsize = nsend*11/k_buf.view<LMPHostType>().dimension_1()+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().dimension_1());
if(nsend > (int) (k_buf.view<LMPHostType>().extent(0)*k_buf.view<LMPHostType>().extent(1))/11) {
int newsize = nsend*11/k_buf.view<LMPHostType>().extent(1)+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().extent(1));
}
if(space == Host) {
AtomVecAtomicKokkos_PackExchangeFunctor<LMPHostType> f(atomKK,k_buf,k_sendlist,k_copylist,atom->nlocal,dim,lo,hi);
@ -617,7 +617,7 @@ struct AtomVecAtomicKokkos_UnpackExchangeFunctor {
_nlocal(nlocal.template view<DeviceType>()),_dim(dim),
_lo(lo),_hi(hi){
const size_t elements = 11;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}

16
src/KOKKOS/atom_vec_bond_kokkos.cpp
View File

@ -624,8 +624,8 @@ struct AtomVecBondKokkos_PackExchangeFunctor {
// maxspecial special, 1 num_bond, bond_per_atom bond_type, bond_per_atom bond_atom,
// 1 to store buffer lenght
elements = 16+atom->maxspecial+atom->bond_per_atom+atom->bond_per_atom;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*
buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*
buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}
@ -694,10 +694,10 @@ int AtomVecBondKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_2
X_FLOAT hi )
{
const int elements = 16+atomKK->maxspecial+atomKK->bond_per_atom+atomKK->bond_per_atom;
if(nsend > (int) (k_buf.view<LMPHostType>().dimension_0()*
k_buf.view<LMPHostType>().dimension_1())/elements) {
int newsize = nsend*elements/k_buf.view<LMPHostType>().dimension_1()+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().dimension_1());
if(nsend > (int) (k_buf.view<LMPHostType>().extent(0)*
k_buf.view<LMPHostType>().extent(1))/elements) {
int newsize = nsend*elements/k_buf.view<LMPHostType>().extent(1)+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().extent(1));
}
if(space == Host) {
AtomVecBondKokkos_PackExchangeFunctor<LMPHostType>
@ -794,8 +794,8 @@ struct AtomVecBondKokkos_UnpackExchangeFunctor {
_nlocal(nlocal.template view<DeviceType>()),_dim(dim),
_lo(lo),_hi(hi){
elements = 16+atom->maxspecial+atom->bond_per_atom+atom->bond_per_atom;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*
buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*
buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}

14
src/KOKKOS/atom_vec_charge_kokkos.cpp
View File

@ -169,7 +169,7 @@ struct AtomVecChargeKokkos_PackComm {
_x(x.view<DeviceType>()),_list(list.view<DeviceType>()),_iswap(iswap),
_xprd(xprd),_yprd(yprd),_zprd(zprd),
_xy(xy),_xz(xz),_yz(yz) {
const size_t maxsend = (buf.view<DeviceType>().dimension_0()*buf.view<DeviceType>().dimension_1())/3;
const size_t maxsend = (buf.view<DeviceType>().extent(0)*buf.view<DeviceType>().extent(1))/3;
const size_t elements = 3;
buffer_view<DeviceType>(_buf,buf,maxsend,elements);
_pbc[0] = pbc[0]; _pbc[1] = pbc[1]; _pbc[2] = pbc[2];
@ -620,8 +620,8 @@ struct AtomVecChargeKokkos_PackExchangeFunctor {
_nlocal(nlocal),_dim(dim),
_lo(lo),_hi(hi){
const size_t elements = 12;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*
buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*
buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}
@ -667,9 +667,9 @@ int AtomVecChargeKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat
ExecutionSpace space,int dim,
X_FLOAT lo,X_FLOAT hi )
{
if(nsend > (int) (k_buf.view<LMPHostType>().dimension_0()*k_buf.view<LMPHostType>().dimension_1())/12) {
int newsize = nsend*12/k_buf.view<LMPHostType>().dimension_1()+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().dimension_1());
if(nsend > (int) (k_buf.view<LMPHostType>().extent(0)*k_buf.view<LMPHostType>().extent(1))/12) {
int newsize = nsend*12/k_buf.view<LMPHostType>().extent(1)+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().extent(1));
}
if(space == Host) {
AtomVecChargeKokkos_PackExchangeFunctor<LMPHostType>
@ -742,7 +742,7 @@ struct AtomVecChargeKokkos_UnpackExchangeFunctor {
_nlocal(nlocal.template view<DeviceType>()),_dim(dim),
_lo(lo),_hi(hi){
const size_t elements = 12;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}

12
src/KOKKOS/atom_vec_dpd_kokkos.cpp
View File

@ -219,7 +219,7 @@ struct AtomVecDPDKokkos_PackComm {
_list(list.view<DeviceType>()),_iswap(iswap),
_xprd(xprd),_yprd(yprd),_zprd(zprd),
_xy(xy),_xz(xz),_yz(yz) {
const size_t maxsend = (buf.view<DeviceType>().dimension_0()*buf.view<DeviceType>().dimension_1())/3;
const size_t maxsend = (buf.view<DeviceType>().extent(0)*buf.view<DeviceType>().extent(1))/3;
const size_t elements = 3;
buffer_view<DeviceType>(_buf,buf,maxsend,elements);
_pbc[0] = pbc[0]; _pbc[1] = pbc[1]; _pbc[2] = pbc[2];
@ -1324,7 +1324,7 @@ struct AtomVecDPDKokkos_PackExchangeFunctor {
_nlocal(nlocal),_dim(dim),
_lo(lo),_hi(hi){
const size_t elements = 17;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}
@ -1376,9 +1376,9 @@ struct AtomVecDPDKokkos_PackExchangeFunctor {
int AtomVecDPDKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_2d &k_buf, DAT::tdual_int_1d k_sendlist,DAT::tdual_int_1d k_copylist,ExecutionSpace space,int dim,X_FLOAT lo,X_FLOAT hi )
{
if(nsend > (int) (k_buf.view<LMPHostType>().dimension_0()*k_buf.view<LMPHostType>().dimension_1())/17) {
int newsize = nsend*17/k_buf.view<LMPHostType>().dimension_1()+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().dimension_1());
if(nsend > (int) (k_buf.view<LMPHostType>().extent(0)*k_buf.view<LMPHostType>().extent(1))/17) {
int newsize = nsend*17/k_buf.view<LMPHostType>().extent(1)+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().extent(1));
}
sync(space,X_MASK | V_MASK | TAG_MASK | TYPE_MASK |
MASK_MASK | IMAGE_MASK| DPDTHETA_MASK | UCOND_MASK |
@ -1467,7 +1467,7 @@ struct AtomVecDPDKokkos_UnpackExchangeFunctor {
_nlocal(nlocal.template view<DeviceType>()),_dim(dim),
_lo(lo),_hi(hi){
const size_t elements = 17;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}

16
src/KOKKOS/atom_vec_full_kokkos.cpp
View File

@ -835,8 +835,8 @@ struct AtomVecFullKokkos_PackExchangeFunctor {
// 1 to store buffer length
elements = 20+atom->maxspecial+2*atom->bond_per_atom+4*atom->angle_per_atom+
5*atom->dihedral_per_atom + 5*atom->improper_per_atom;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*
buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*
buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}
@ -955,10 +955,10 @@ int AtomVecFullKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfloat_2
{
const int elements = 20+atom->maxspecial+2*atom->bond_per_atom+4*atom->angle_per_atom+
5*atom->dihedral_per_atom + 5*atom->improper_per_atom;
if(nsend > (int) (k_buf.view<LMPHostType>().dimension_0()*
k_buf.view<LMPHostType>().dimension_1())/elements) {
int newsize = nsend*elements/k_buf.view<LMPHostType>().dimension_1()+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().dimension_1());
if(nsend > (int) (k_buf.view<LMPHostType>().extent(0)*
k_buf.view<LMPHostType>().extent(1))/elements) {
int newsize = nsend*elements/k_buf.view<LMPHostType>().extent(1)+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().extent(1));
}
if(space == Host) {
AtomVecFullKokkos_PackExchangeFunctor<LMPHostType>
@ -1110,8 +1110,8 @@ struct AtomVecFullKokkos_UnpackExchangeFunctor {
elements = 20+atom->maxspecial+2*atom->bond_per_atom+4*atom->angle_per_atom+
5*atom->dihedral_per_atom + 5*atom->improper_per_atom;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*
buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*
buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}

4
src/KOKKOS/atom_vec_kokkos.cpp
View File

@ -51,7 +51,7 @@ struct AtomVecKokkos_PackComm {
_x(x.view<DeviceType>()),_list(list.view<DeviceType>()),_iswap(iswap),
_xprd(xprd),_yprd(yprd),_zprd(zprd),
_xy(xy),_xz(xz),_yz(yz) {
const size_t maxsend = (buf.view<DeviceType>().dimension_0()*buf.view<DeviceType>().dimension_1())/3;
const size_t maxsend = (buf.view<DeviceType>().extent(0)*buf.view<DeviceType>().extent(1))/3;
const size_t elements = 3;
buffer_view<DeviceType>(_buf,buf,maxsend,elements);
_pbc[0] = pbc[0]; _pbc[1] = pbc[1]; _pbc[2] = pbc[2];
@ -542,7 +542,7 @@ struct AtomVecKokkos_UnPackReverse {
const typename DAT::tdual_int_2d &list,
const int & iswap):
_f(f.view<DeviceType>()),_list(list.view<DeviceType>()),_iswap(iswap) {
const size_t maxsend = (buf.view<DeviceType>().dimension_0()*buf.view<DeviceType>().dimension_1())/3;
const size_t maxsend = (buf.view<DeviceType>().extent(0)*buf.view<DeviceType>().extent(1))/3;
const size_t elements = 3;
buffer_view<DeviceType>(_buf,buf,maxsend,elements);
};

32
src/KOKKOS/atom_vec_kokkos.h
View File

@ -125,14 +125,14 @@ class AtomVecKokkos : public AtomVec {
buffer_size = src.capacity();
}
return mirror_type( buffer ,
src.dimension_0() ,
src.dimension_1() ,
src.dimension_2() ,
src.dimension_3() ,
src.dimension_4() ,
src.dimension_5() ,
src.dimension_6() ,
src.dimension_7() );
src.extent(0) ,
src.extent(1) ,
src.extent(2) ,
src.extent(3) ,
src.extent(4) ,
src.extent(5) ,
src.extent(6) ,
src.extent(7) );
}
template<class ViewType>
@ -151,14 +151,14 @@ class AtomVecKokkos : public AtomVec {
buffer_size = src.capacity();
}
mirror_type tmp_view( (typename ViewType::value_type*)buffer ,
src.dimension_0() ,
src.dimension_1() ,
src.dimension_2() ,
src.dimension_3() ,
src.dimension_4() ,
src.dimension_5() ,
src.dimension_6() ,
src.dimension_7() );
src.extent(0) ,
src.extent(1) ,
src.extent(2) ,
src.extent(3) ,
src.extent(4) ,
src.extent(5) ,
src.extent(6) ,
src.extent(7) );
if(space == Device) {
Kokkos::deep_copy(LMPHostType(),tmp_view,src.h_view),
Kokkos::deep_copy(LMPHostType(),src.d_view,tmp_view);

20
src/KOKKOS/atom_vec_molecular_kokkos.cpp
View File

@ -319,8 +319,8 @@ struct AtomVecMolecularKokkos_PackComm {
_x(x.view<DeviceType>()),_list(list.view<DeviceType>()),_iswap(iswap),
_xprd(xprd),_yprd(yprd),_zprd(zprd),
_xy(xy),_xz(xz),_yz(yz) {
const size_t maxsend = (buf.view<DeviceType>().dimension_0()
*buf.view<DeviceType>().dimension_1())/3;
const size_t maxsend = (buf.view<DeviceType>().extent(0)
*buf.view<DeviceType>().extent(1))/3;
const size_t elements = 3;
buffer_view<DeviceType>(_buf,buf,maxsend,elements);
_pbc[0] = pbc[0]; _pbc[1] = pbc[1]; _pbc[2] = pbc[2];
@ -1250,8 +1250,8 @@ struct AtomVecMolecularKokkos_PackExchangeFunctor {
// 1 to store buffer length
elements = 19+atom->maxspecial+2*atom->bond_per_atom+4*atom->angle_per_atom+
5*atom->dihedral_per_atom + 5*atom->improper_per_atom;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*
buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*
buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}
@ -1368,10 +1368,10 @@ int AtomVecMolecularKokkos::pack_exchange_kokkos(const int &nsend,DAT::tdual_xfl
{
const int elements = 19+atom->maxspecial+2*atom->bond_per_atom+4*atom->angle_per_atom+
5*atom->dihedral_per_atom + 5*atom->improper_per_atom;
if(nsend > (int) (k_buf.view<LMPHostType>().dimension_0()*
k_buf.view<LMPHostType>().dimension_1())/elements) {
int newsize = nsend*elements/k_buf.view<LMPHostType>().dimension_1()+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().dimension_1());
if(nsend > (int) (k_buf.view<LMPHostType>().extent(0)*
k_buf.view<LMPHostType>().extent(1))/elements) {
int newsize = nsend*elements/k_buf.view<LMPHostType>().extent(1)+1;
k_buf.resize(newsize,k_buf.view<LMPHostType>().extent(1));
}
if(space == Host) {
AtomVecMolecularKokkos_PackExchangeFunctor<LMPHostType>
@ -1521,8 +1521,8 @@ struct AtomVecMolecularKokkos_UnpackExchangeFunctor {
elements = 19+atom->maxspecial+2*atom->bond_per_atom+4*atom->angle_per_atom+
5*atom->dihedral_per_atom + 5*atom->improper_per_atom;
const int maxsendlist = (buf.template view<DeviceType>().dimension_0()*
buf.template view<DeviceType>().dimension_1())/elements;
const int maxsendlist = (buf.template view<DeviceType>().extent(0)*
buf.template view<DeviceType>().extent(1))/elements;
buffer_view<DeviceType>(_buf,buf,maxsendlist,elements);
}

4
src/KOKKOS/bond_class2_kokkos.cpp
View File

@ -66,14 +66,14 @@ void BondClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
if (eflag_atom) {
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
//if(k_eatom.extent(0)<maxeatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_eatom,eatom);
memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
d_eatom = k_eatom.template view<DeviceType>();
//}
}
if (vflag_atom) {
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
//if(k_vatom.extent(0)<maxvatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_vatom,vatom);
memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
d_vatom = k_vatom.template view<DeviceType>();

4
src/KOKKOS/bond_harmonic_kokkos.cpp
View File

@ -67,14 +67,14 @@ void BondHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
if (eflag_atom) {
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
//if(k_eatom.extent(0)<maxeatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_eatom,eatom);
memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
d_eatom = k_eatom.template view<DeviceType>();
//}
}
if (vflag_atom) {
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
//if(k_vatom.extent(0)<maxvatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_vatom,vatom);
memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
d_vatom = k_vatom.template view<DeviceType>();

58
src/KOKKOS/comm_kokkos.cpp
View File

@ -193,8 +193,8 @@ void CommKokkos::forward_comm_device(int dummy)
if (sendproc[iswap] != me) {
if (comm_x_only) {
if (size_forward_recv[iswap]) {
buf = atomKK->k_x.view<DeviceType>().ptr_on_device() +
firstrecv[iswap]*atomKK->k_x.view<DeviceType>().dimension_1();
buf = atomKK->k_x.view<DeviceType>().data() +
firstrecv[iswap]*atomKK->k_x.view<DeviceType>().extent(1);
MPI_Irecv(buf,size_forward_recv[iswap],MPI_DOUBLE,
recvproc[iswap],0,world,&request);
}
@ -202,7 +202,7 @@ void CommKokkos::forward_comm_device(int dummy)
iswap,k_buf_send,pbc_flag[iswap],pbc[iswap]);
DeviceType::fence();
if (n) {
MPI_Send(k_buf_send.view<DeviceType>().ptr_on_device(),
MPI_Send(k_buf_send.view<DeviceType>().data(),
n,MPI_DOUBLE,sendproc[iswap],0,world);
}
@ -215,7 +215,7 @@ void CommKokkos::forward_comm_device(int dummy)
error->all(FLERR,"Ghost velocity forward comm not yet "
"implemented with Kokkos");
if (size_forward_recv[iswap])
MPI_Irecv(k_buf_recv.view<LMPHostType>().ptr_on_device(),
MPI_Irecv(k_buf_recv.view<LMPHostType>().data(),
size_forward_recv[iswap],MPI_DOUBLE,
recvproc[iswap],0,world,&request);
n = avec->pack_comm_vel(sendnum[iswap],sendlist[iswap],
@ -225,14 +225,14 @@ void CommKokkos::forward_comm_device(int dummy)
avec->unpack_comm_vel(recvnum[iswap],firstrecv[iswap],buf_recv);
} else {
if (size_forward_recv[iswap])
MPI_Irecv(k_buf_recv.view<DeviceType>().ptr_on_device(),
MPI_Irecv(k_buf_recv.view<DeviceType>().data(),
size_forward_recv[iswap],MPI_DOUBLE,
recvproc[iswap],0,world,&request);
n = avec->pack_comm_kokkos(sendnum[iswap],k_sendlist,iswap,
k_buf_send,pbc_flag[iswap],pbc[iswap]);
DeviceType::fence();
if (n)
MPI_Send(k_buf_send.view<DeviceType>().ptr_on_device(),n,
MPI_Send(k_buf_send.view<DeviceType>().data(),n,
MPI_DOUBLE,sendproc[iswap],0,world);
if (size_forward_recv[iswap]) MPI_Wait(&request,MPI_STATUS_IGNORE);
avec->unpack_comm_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv);
@ -304,11 +304,11 @@ void CommKokkos::reverse_comm_device()
if (sendproc[iswap] != me) {
if (comm_f_only) {
if (size_reverse_recv[iswap])
MPI_Irecv(k_buf_recv.view<DeviceType>().ptr_on_device(),size_reverse_recv[iswap],MPI_DOUBLE,
MPI_Irecv(k_buf_recv.view<DeviceType>().data(),size_reverse_recv[iswap],MPI_DOUBLE,
sendproc[iswap],0,world,&request);
if (size_reverse_send[iswap]) {
buf = atomKK->k_f.view<DeviceType>().ptr_on_device() +
firstrecv[iswap]*atomKK->k_f.view<DeviceType>().dimension_1();
buf = atomKK->k_f.view<DeviceType>().data() +
firstrecv[iswap]*atomKK->k_f.view<DeviceType>().extent(1);
MPI_Send(buf,size_reverse_send[iswap],MPI_DOUBLE,
recvproc[iswap],0,world);
@ -320,13 +320,13 @@ void CommKokkos::reverse_comm_device()
}
} else {
if (size_reverse_recv[iswap])
MPI_Irecv(k_buf_recv.view<DeviceType>().ptr_on_device(),
MPI_Irecv(k_buf_recv.view<DeviceType>().data(),
size_reverse_recv[iswap],MPI_DOUBLE,
sendproc[iswap],0,world,&request);
n = avec->pack_reverse_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_send);
DeviceType::fence();
if (n)
MPI_Send(k_buf_send.view<DeviceType>().ptr_on_device(),n,
MPI_Send(k_buf_send.view<DeviceType>().data(),n,
MPI_DOUBLE,recvproc[iswap],0,world);
if (size_reverse_recv[iswap]) MPI_Wait(&request,MPI_STATUS_IGNORE);
}
@ -407,10 +407,10 @@ void CommKokkos::forward_comm_pair_device(Pair *pair)
if (sendproc[iswap] != me) {
if (recvnum[iswap])
MPI_Irecv(k_buf_recv_pair.view<DeviceType>().ptr_on_device(),nsize*recvnum[iswap],MPI_DOUBLE,
MPI_Irecv(k_buf_recv_pair.view<DeviceType>().data(),nsize*recvnum[iswap],MPI_DOUBLE,
recvproc[iswap],0,world,&request);
if (sendnum[iswap])
MPI_Send(k_buf_send_pair.view<DeviceType>().ptr_on_device(),n,MPI_DOUBLE,sendproc[iswap],0,world);
MPI_Send(k_buf_send_pair.view<DeviceType>().data(),n,MPI_DOUBLE,sendproc[iswap],0,world);
if (recvnum[iswap]) MPI_Wait(&request,MPI_STATUS_IGNORE);
} else k_buf_recv_pair = k_buf_send_pair;
@ -514,7 +514,7 @@ struct BuildExchangeListFunctor {
void operator() (int i) const {
if (_x(i,_dim) < _lo || _x(i,_dim) >= _hi) {
const int mysend=Kokkos::atomic_fetch_add(&_nsend(),1);
if(mysend<_sendlist.dimension_0()) {
if(mysend<_sendlist.extent(0)) {
_sendlist(mysend) = i;
_sendflag(i) = 1;
}
@ -570,10 +570,10 @@ void CommKokkos::exchange_device()
i = nsend = 0;
if (true) {
if (k_sendflag.h_view.dimension_0()<nlocal) k_sendflag.resize(nlocal);
if (k_sendflag.h_view.extent(0)<nlocal) k_sendflag.resize(nlocal);
k_sendflag.sync<DeviceType>();
k_count.h_view() = k_exchange_sendlist.h_view.dimension_0();
while (k_count.h_view()>=k_exchange_sendlist.h_view.dimension_0()) {
k_count.h_view() = k_exchange_sendlist.h_view.extent(0);
while (k_count.h_view()>=k_exchange_sendlist.h_view.extent(0)) {
k_count.h_view() = 0;
k_count.modify<LMPHostType>();
k_count.sync<DeviceType>();
@ -587,10 +587,10 @@ void CommKokkos::exchange_device()
k_count.modify<DeviceType>();
k_count.sync<LMPHostType>();
if (k_count.h_view()>=k_exchange_sendlist.h_view.dimension_0()) {
if (k_count.h_view()>=k_exchange_sendlist.h_view.extent(0)) {
k_exchange_sendlist.resize(k_count.h_view()*1.1);
k_exchange_copylist.resize(k_count.h_view()*1.1);
k_count.h_view()=k_exchange_sendlist.h_view.dimension_0();
k_count.h_view()=k_exchange_sendlist.h_view.extent(0);
}
}
k_exchange_copylist.sync<LMPHostType>();
@ -655,18 +655,18 @@ void CommKokkos::exchange_device()
}
if (nrecv > maxrecv) grow_recv_kokkos(nrecv);
MPI_Irecv(k_buf_recv.view<DeviceType>().ptr_on_device(),nrecv1,
MPI_Irecv(k_buf_recv.view<DeviceType>().data(),nrecv1,
MPI_DOUBLE,procneigh[dim][1],0,
world,&request);
MPI_Send(k_buf_send.view<DeviceType>().ptr_on_device(),nsend,
MPI_Send(k_buf_send.view<DeviceType>().data(),nsend,
MPI_DOUBLE,procneigh[dim][0],0,world);
MPI_Wait(&request,MPI_STATUS_IGNORE);
if (procgrid[dim] > 2) {
MPI_Irecv(k_buf_recv.view<DeviceType>().ptr_on_device()+nrecv1,
MPI_Irecv(k_buf_recv.view<DeviceType>().data()+nrecv1,
nrecv2,MPI_DOUBLE,procneigh[dim][0],0,
world,&request);
MPI_Send(k_buf_send.view<DeviceType>().ptr_on_device(),nsend,
MPI_Send(k_buf_send.view<DeviceType>().data(),nsend,
MPI_DOUBLE,procneigh[dim][1],0,world);
MPI_Wait(&request,MPI_STATUS_IGNORE);
}
@ -952,10 +952,10 @@ void CommKokkos::borders_device() {
MPI_Sendrecv(&nsend,1,MPI_INT,sendproc[iswap],0,
&nrecv,1,MPI_INT,recvproc[iswap],0,world,MPI_STATUS_IGNORE);
if (nrecv*size_border > maxrecv) grow_recv_kokkos(nrecv*size_border);
if (nrecv) MPI_Irecv(k_buf_recv.view<DeviceType>().ptr_on_device(),
if (nrecv) MPI_Irecv(k_buf_recv.view<DeviceType>().data(),
nrecv*size_border,MPI_DOUBLE,
recvproc[iswap],0,world,&request);
if (n) MPI_Send(k_buf_send.view<DeviceType>().ptr_on_device(),n,
if (n) MPI_Send(k_buf_send.view<DeviceType>().data(),n,
MPI_DOUBLE,sendproc[iswap],0,world);
if (nrecv) MPI_Wait(&request,MPI_STATUS_IGNORE);
buf = buf_recv;
@ -1049,12 +1049,12 @@ void CommKokkos::grow_send_kokkos(int n, int flag, ExecutionSpace space)
k_buf_send.modify<LMPHostType>();
k_buf_send.resize(maxsend_border,atom->avec->size_border);
buf_send = k_buf_send.view<LMPHostType>().ptr_on_device();
buf_send = k_buf_send.view<LMPHostType>().data();
}
else {
k_buf_send = DAT::
tdual_xfloat_2d("comm:k_buf_send",maxsend_border,atom->avec->size_border);
buf_send = k_buf_send.view<LMPHostType>().ptr_on_device();
buf_send = k_buf_send.view<LMPHostType>().data();
}
}
@ -1068,7 +1068,7 @@ void CommKokkos::grow_recv_kokkos(int n, ExecutionSpace space)
int maxrecv_border = (maxrecv+BUFEXTRA+5)/atom->avec->size_border + 2;
k_buf_recv = DAT::
tdual_xfloat_2d("comm:k_buf_recv",maxrecv_border,atom->avec->size_border);
buf_recv = k_buf_recv.view<LMPHostType>().ptr_on_device();
buf_recv = k_buf_recv.view<LMPHostType>().data();
}
/* ----------------------------------------------------------------------
@ -1105,7 +1105,7 @@ void CommKokkos::grow_swap(int n)
}
maxswap = n;
int size = MAX(k_sendlist.d_view.dimension_1(),BUFMIN);
int size = MAX(k_sendlist.d_view.extent(1),BUFMIN);
if (exchange_comm_classic) { // force realloc on Host
k_sendlist.sync<LMPHostType>();

4
src/KOKKOS/dihedral_charmm_kokkos.cpp
View File

@ -80,7 +80,7 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
if (eflag_atom) {
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
//if(k_eatom.extent(0)<maxeatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_eatom,eatom);
memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom");
d_eatom = k_eatom.template view<DeviceType>();
@ -89,7 +89,7 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
//}
}
if (vflag_atom) {
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
//if(k_vatom.extent(0)<maxvatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_vatom,vatom);
memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom");
d_vatom = k_vatom.template view<DeviceType>();

6
src/KOKKOS/fix_eos_table_rx_kokkos.cpp
View File

@ -534,11 +534,11 @@ void FixEOStableRXKokkos<DeviceType>::create_kokkos_tables()
h_table->hi[i] = tb->hi;
h_table->invdelta[i] = tb->invdelta;
for(int j = 0; j<h_table->r.dimension_1(); j++)
for(int j = 0; j<h_table->r.extent(1); j++)
h_table->r(i,j) = tb->r[j];
for(int j = 0; j<h_table->e.dimension_1(); j++)
for(int j = 0; j<h_table->e.extent(1); j++)
h_table->e(i,j) = tb->e[j];
for(int j = 0; j<h_table->de.dimension_1(); j++)
for(int j = 0; j<h_table->de.extent(1); j++)
h_table->de(i,j) = tb->de[j];
}

2
src/KOKKOS/fix_property_atom_kokkos.cpp
View File

@ -61,7 +61,7 @@ void FixPropertyAtomKokkos::grow_arrays(int nmax)
size_t nbytes = (nmax-nmax_old) * sizeof(int);
memset(&atom->ivector[index[m]][nmax_old],0,nbytes);
} else if (style[m] == DOUBLE) {
memoryKK->grow_kokkos(atomKK->k_dvector,atomKK->dvector,atomKK->k_dvector.dimension_0(),nmax,
memoryKK->grow_kokkos(atomKK->k_dvector,atomKK->dvector,atomKK->k_dvector.extent(0),nmax,
"atom:dvector");
//memory->grow(atom->dvector[index[m]],nmax,"atom:dvector");
//size_t nbytes = (nmax-nmax_old) * sizeof(double);

12
src/KOKKOS/fix_rx_kokkos.cpp
View File

@ -1352,7 +1352,7 @@ void FixRxKokkos<DeviceType>::operator()(Tag_FixRxKokkos_solveSystems<ZERO_RATES
{
if (d_mask(i) & groupbit)
{
StridedArrayType<double,1> y( d_scratchSpace.ptr_on_device() + scratchSpaceSize * i );
StridedArrayType<double,1> y( d_scratchSpace.data() + scratchSpaceSize * i );
StridedArrayType<double,1> rwork( &y[nspecies] );
UserRHSDataKokkos<1> userData;
@ -1447,7 +1447,7 @@ void FixRxKokkos<DeviceType>::solve_reactions(const int vflag, const bool isPreF
{
const int count = nlocal + (newton_pair ? nghost : 0);
if (count > k_dpdThetaLocal.template view<DeviceType>().dimension_0()) {
if (count > k_dpdThetaLocal.template view<DeviceType>().extent(0)) {
memoryKK->destroy_kokkos (k_dpdThetaLocal, dpdThetaLocal);
memoryKK->create_kokkos (k_dpdThetaLocal, dpdThetaLocal, count, "FixRxKokkos::dpdThetaLocal");
this->d_dpdThetaLocal = k_dpdThetaLocal.template view<DeviceType>();
@ -1544,7 +1544,7 @@ void FixRxKokkos<DeviceType>::solve_reactions(const int vflag, const bool isPreF
//double *scratchSpace = new double[ scratchSpaceSize * nlocal ];
//typename ArrayTypes<DeviceType>::t_double_1d d_scratchSpace("d_scratchSpace", scratchSpaceSize * nlocal);
if (nlocal*scratchSpaceSize > d_scratchSpace.dimension_0()) {
if (nlocal*scratchSpaceSize > d_scratchSpace.extent(0)) {
memoryKK->destroy_kokkos (d_scratchSpace);
memoryKK->create_kokkos (d_scratchSpace, nlocal*scratchSpaceSize, "FixRxKokkos::d_scratchSpace");
}
@ -1560,7 +1560,7 @@ void FixRxKokkos<DeviceType>::solve_reactions(const int vflag, const bool isPreF
//StridedArrayType<double,1> _y( y );
//StridedArrayType<double,1> _rwork( rwork );
StridedArrayType<double,1> y( d_scratchSpace.ptr_on_device() + scratchSpaceSize * i );
StridedArrayType<double,1> y( d_scratchSpace.data() + scratchSpaceSize * i );
StridedArrayType<double,1> rwork( &y[nspecies] );
//UserRHSData userData;
@ -2026,7 +2026,7 @@ void FixRxKokkos<DeviceType>::computeLocalTemperature()
const int ntypes = atom->ntypes;
//memoryKK->create_kokkos (k_cutsq, h_cutsq, ntypes+1, ntypes+1, "pair:cutsq");
if (ntypes+1 > k_cutsq.dimension_0()) {
if (ntypes+1 > k_cutsq.extent(0)) {
memoryKK->destroy_kokkos (k_cutsq);
memoryKK->create_kokkos (k_cutsq, ntypes+1, ntypes+1, "FixRxKokkos::k_cutsq");
d_cutsq = k_cutsq.template view<DeviceType>();
@ -2047,7 +2047,7 @@ void FixRxKokkos<DeviceType>::computeLocalTemperature()
int sumWeightsCt = nlocal + (NEWTON_PAIR ? nghost : 0);
//memoryKK->create_kokkos (k_sumWeights, sumWeights, sumWeightsCt, "FixRxKokkos::sumWeights");
if (sumWeightsCt > k_sumWeights.template view<DeviceType>().dimension_0()) {
if (sumWeightsCt > k_sumWeights.template view<DeviceType>().extent(0)) {
memoryKK->destroy_kokkos(k_sumWeights, sumWeights);
memoryKK->create_kokkos (k_sumWeights, sumWeightsCt, "FixRxKokkos::sumWeights");
d_sumWeights = k_sumWeights.template view<DeviceType>();

6
src/KOKKOS/fix_shardlow_kokkos.cpp
View File

@ -602,9 +602,9 @@ void FixShardlowKokkos<DeviceType>::initial_integrate(int vflag)
auto h_ssa_phaseLen = np_ssa->k_ssa_phaseLen.h_view;
auto h_ssa_gphaseLen = np_ssa->k_ssa_gphaseLen.h_view;
int maxWorkItemCt = (int) ssa_itemLoc.dimension_1();
if (maxWorkItemCt < (int) ssa_gitemLoc.dimension_1()) {
maxWorkItemCt = (int) ssa_gitemLoc.dimension_1();
int maxWorkItemCt = (int) ssa_itemLoc.extent(1);
if (maxWorkItemCt < (int) ssa_gitemLoc.extent(1)) {
maxWorkItemCt = (int) ssa_gitemLoc.extent(1);
}
if (maxWorkItemCt > maxRNG) {
es_RNG_t serial_rand_state;

36
src/KOKKOS/gridcomm_kokkos.cpp
View File

@ -241,8 +241,8 @@ void GridCommKokkos<DeviceType>::setup()
maxswap += SWAPDELTA;
swap = (Swap *)
memory->srealloc(swap,maxswap*sizeof(Swap),"Commgrid:swap");
k_packlist.resize(maxswap,k_packlist.dimension_1());
k_unpacklist.resize(maxswap,k_unpacklist.dimension_1());
k_packlist.resize(maxswap,k_packlist.extent(1));
k_unpacklist.resize(maxswap,k_unpacklist.extent(1));
}
swap[nswap].sendproc = procxlo;
@ -285,8 +285,8 @@ void GridCommKokkos<DeviceType>::setup()
maxswap += 1;
swap = (Swap *)
memory->srealloc(swap,maxswap*sizeof(Swap),"Commgrid:swap");
k_packlist.resize(maxswap,k_packlist.dimension_1());
k_unpacklist.resize(maxswap,k_unpacklist.dimension_1());
k_packlist.resize(maxswap,k_packlist.extent(1));
k_unpacklist.resize(maxswap,k_unpacklist.extent(1));
}
swap[nswap].sendproc = procxhi;
@ -329,8 +329,8 @@ void GridCommKokkos<DeviceType>::setup()
maxswap += SWAPDELTA;
swap = (Swap *)
memory->srealloc(swap,maxswap*sizeof(Swap),"Commgrid:swap");
k_packlist.resize(maxswap,k_packlist.dimension_1());
k_unpacklist.resize(maxswap,k_unpacklist.dimension_1());
k_packlist.resize(maxswap,k_packlist.extent(1));
k_unpacklist.resize(maxswap,k_unpacklist.extent(1));
}
swap[nswap].sendproc = procylo;
@ -373,8 +373,8 @@ void GridCommKokkos<DeviceType>::setup()
maxswap += 1;
swap = (Swap *)
memory->srealloc(swap,maxswap*sizeof(Swap),"Commgrid:swap");
k_packlist.resize(maxswap,k_packlist.dimension_1());
k_unpacklist.resize(maxswap,k_unpacklist.dimension_1());
k_packlist.resize(maxswap,k_packlist.extent(1));
k_unpacklist.resize(maxswap,k_unpacklist.extent(1));
}
swap[nswap].sendproc = procyhi;
@ -417,8 +417,8 @@ void GridCommKokkos<DeviceType>::setup()
maxswap += SWAPDELTA;
swap = (Swap *)
memory->srealloc(swap,maxswap*sizeof(Swap),"Commgrid:swap");
k_packlist.resize(maxswap,k_packlist.dimension_1());
k_unpacklist.resize(maxswap,k_unpacklist.dimension_1());
k_packlist.resize(maxswap,k_packlist.extent(1));
k_unpacklist.resize(maxswap,k_unpacklist.extent(1));
}
swap[nswap].sendproc = proczlo;
@ -461,8 +461,8 @@ void GridCommKokkos<DeviceType>::setup()
maxswap += 1;
swap = (Swap *)
memory->srealloc(swap,maxswap*sizeof(Swap),"Commgrid:swap");
k_packlist.resize(maxswap,k_packlist.dimension_1());
k_unpacklist.resize(maxswap,k_unpacklist.dimension_1());
k_packlist.resize(maxswap,k_packlist.extent(1));
k_unpacklist.resize(maxswap,k_unpacklist.extent(1));
}
swap[nswap].sendproc = proczhi;
@ -526,9 +526,9 @@ void GridCommKokkos<DeviceType>::forward_comm(KSpace *kspace, int which)
DeviceType::fence();
if (swap[m].sendproc != me) {
MPI_Irecv(k_buf2.view<DeviceType>().ptr_on_device(),nforward*swap[m].nunpack,MPI_FFT_SCALAR,
MPI_Irecv(k_buf2.view<DeviceType>().data(),nforward*swap[m].nunpack,MPI_FFT_SCALAR,
swap[m].recvproc,0,gridcomm,&request);
MPI_Send(k_buf1.view<DeviceType>().ptr_on_device(),nforward*swap[m].npack,MPI_FFT_SCALAR,
MPI_Send(k_buf1.view<DeviceType>().data(),nforward*swap[m].npack,MPI_FFT_SCALAR,
swap[m].sendproc,0,gridcomm);
MPI_Wait(&request,MPI_STATUS_IGNORE);
}
@ -559,9 +559,9 @@ void GridCommKokkos<DeviceType>::reverse_comm(KSpace *kspace, int which)
DeviceType::fence();
if (swap[m].recvproc != me) {
MPI_Irecv(k_buf2.view<DeviceType>().ptr_on_device(),nreverse*swap[m].npack,MPI_FFT_SCALAR,
MPI_Irecv(k_buf2.view<DeviceType>().data(),nreverse*swap[m].npack,MPI_FFT_SCALAR,
swap[m].sendproc,0,gridcomm,&request);
MPI_Send(k_buf1.view<DeviceType>().ptr_on_device(),nreverse*swap[m].nunpack,MPI_FFT_SCALAR,
MPI_Send(k_buf1.view<DeviceType>().data(),nreverse*swap[m].nunpack,MPI_FFT_SCALAR,
swap[m].recvproc,0,gridcomm);
MPI_Wait(&request,MPI_STATUS_IGNORE);
}
@ -582,8 +582,8 @@ int GridCommKokkos<DeviceType>::indices(DAT::tdual_int_2d &k_list, int index,
int xlo, int xhi, int ylo, int yhi, int zlo, int zhi)
{
int nmax = (xhi-xlo+1) * (yhi-ylo+1) * (zhi-zlo+1);
if (k_list.dimension_1() < nmax)
k_list.resize(k_list.dimension_0(),nmax);
if (k_list.extent(1) < nmax)
k_list.resize(k_list.extent(0),nmax);
int nx = (outxhi_max-outxlo_max+1);
int ny = (outyhi_max-outylo_max+1);

4
src/KOKKOS/improper_class2_kokkos.cpp
View File

@ -77,14 +77,14 @@ void ImproperClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
if (eflag_atom) {
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
//if(k_eatom.extent(0)<maxeatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_eatom,eatom);
memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
d_eatom = k_eatom.template view<DeviceType>();
//}
}
if (vflag_atom) {
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
//if(k_vatom.extent(0)<maxvatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_vatom,vatom);
memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
d_vatom = k_vatom.template view<DeviceType>();

4
src/KOKKOS/improper_harmonic_kokkos.cpp
View File

@ -77,14 +77,14 @@ void ImproperHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// reallocate per-atom arrays if necessary
if (eflag_atom) {
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
//if(k_eatom.extent(0)<maxeatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_eatom,eatom);
memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
d_eatom = k_eatom.template view<DeviceType>();
//}
}
if (vflag_atom) {
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
//if(k_vatom.extent(0)<maxvatom) { // won't work without adding zero functor
memoryKK->destroy_kokkos(k_vatom,vatom);
memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
d_vatom = k_vatom.template view<DeviceType>();

4
src/KOKKOS/kokkos_type.h
View File

@ -956,7 +956,7 @@ void buffer_view(BufferView &buf, DualView &view,
const size_t n7 = 0) {
buf = BufferView(
view.template view<DeviceType>().ptr_on_device(),
view.template view<DeviceType>().data(),
n0,n1,n2,n3,n4,n5,n6,n7);
}
@ -973,7 +973,7 @@ struct MemsetZeroFunctor {
template<class ViewType>
void memset_kokkos (ViewType &view) {
static MemsetZeroFunctor<typename ViewType::execution_space> f;
f.ptr = view.ptr_on_device();
f.ptr = view.data();
#ifndef KOKKOS_USING_DEPRECATED_VIEW
Kokkos::parallel_for(view.span()*sizeof(typename ViewType::value_type)/4, f);
#else

14
src/KOKKOS/memory_kokkos.h
View File

@ -36,7 +36,7 @@ TYPE create_kokkos(TYPE &data, typename TYPE::value_type *&array,
int n1, const char *name)
{
data = TYPE(name,n1);
array = data.h_view.ptr_on_device();
array = data.h_view.data();
return data;
}
@ -51,7 +51,7 @@ template <typename TYPE, typename HTYPE>
#else
h_data = data;
#endif
array = h_data.ptr_on_device();
array = h_data.data();
return data;
}
@ -81,7 +81,7 @@ TYPE grow_kokkos(TYPE &data, typename TYPE::value_type *&array,
if (array == NULL) return create_kokkos(data,array,n1,name);
data.resize(n1);
array = data.h_view.ptr_on_device();
array = data.h_view.data();
return data;
}
@ -100,8 +100,8 @@ void destroy_kokkos(TYPE data, typename TYPE::value_type* &array)
template <typename TYPE>
TYPE destroy_kokkos(TYPE &data)
{
/*if(data.ptr_on_device()!=NULL)
free(data.ptr_on_device());*/
/*if(data.data()!=NULL)
free(data.data());*/
data = TYPE();
return data;
}
@ -251,7 +251,7 @@ TYPE create_kokkos(TYPE &data, typename TYPE::value_type **&array,
array = (typename TYPE::value_type **) smalloc(nbytes,name);
for (int i = 0; i < n1; i++)
if(data.h_view.dimension_1()==0)
if(data.h_view.extent(1)==0)
array[i] = NULL;
else
array[i] = &data.h_view(i,0);
@ -271,7 +271,7 @@ TYPE grow_kokkos(TYPE &data, typename TYPE::value_type **&array,
array = (typename TYPE::value_type **) smalloc(nbytes,name);
for (int i = 0; i < n1; i++)
if(data.h_view.dimension_1()==0)
if(data.h_view.extent(1)==0)
array[i] = NULL;
else
array[i] = &data.h_view(i,0);

8
src/KOKKOS/nbin_kokkos.cpp
View File

@ -68,14 +68,14 @@ NBinKokkos<DeviceType>::NBinKokkos(LAMMPS *lmp) : NBinStandard(lmp) {
template<class DeviceType>
void NBinKokkos<DeviceType>::bin_atoms_setup(int nall)
{
if (mbins > k_bins.d_view.dimension_0()) {
if (mbins > k_bins.d_view.extent(0)) {
k_bins = DAT::tdual_int_2d("Neighbor::d_bins",mbins,atoms_per_bin);
bins = k_bins.view<DeviceType>();
k_bincount = DAT::tdual_int_1d("Neighbor::d_bincount",mbins);
bincount = k_bincount.view<DeviceType>();
}
if (nall > k_atom2bin.d_view.dimension_0()) {
if (nall > k_atom2bin.d_view.extent(0)) {
k_atom2bin = DAT::tdual_int_1d("Neighbor::d_atom2bin",nall);
atom2bin = k_atom2bin.view<DeviceType>();
}
@ -101,7 +101,7 @@ void NBinKokkos<DeviceType>::bin_atoms()
deep_copy(d_resize, h_resize);
MemsetZeroFunctor<DeviceType> f_zero;
f_zero.ptr = (void*) k_bincount.view<DeviceType>().ptr_on_device();
f_zero.ptr = (void*) k_bincount.view<DeviceType>().data();
Kokkos::parallel_for(mbins, f_zero);
atomKK->sync(ExecutionSpaceFromDevice<DeviceType>::space,X_MASK);
@ -139,7 +139,7 @@ void NBinKokkos<DeviceType>::binatomsItem(const int &i) const
atom2bin(i) = ibin;
const int ac = Kokkos::atomic_fetch_add(&bincount[ibin], (int)1);
if(ac < bins.dimension_1()) {
if(ac < bins.extent(1)) {
bins(ibin, ac) = i;
} else {
d_resize() = 1;

12
src/KOKKOS/nbin_ssa_kokkos.cpp
View File

@ -72,7 +72,7 @@ NBinSSAKokkos<DeviceType>::NBinSSAKokkos(LAMMPS *lmp) : NBinStandard(lmp)
template<class DeviceType>
void NBinSSAKokkos<DeviceType>::bin_atoms_setup(int nall)
{
if (mbins > (int) k_bins.h_view.dimension_0()) {
if (mbins > (int) k_bins.h_view.extent(0)) {
k_bins = DAT::tdual_int_2d("NBinSSAKokkos::bins",mbins,atoms_per_bin);
bins = k_bins.view<DeviceType>();
@ -82,7 +82,7 @@ void NBinSSAKokkos<DeviceType>::bin_atoms_setup(int nall)
ghosts_per_gbin = atom->nghost / 7; // estimate needed size
if (ghosts_per_gbin > (int) k_gbins.h_view.dimension_1()) {
if (ghosts_per_gbin > (int) k_gbins.h_view.extent(1)) {
k_gbins = DAT::tdual_int_2d("NBinSSAKokkos::gbins",8,ghosts_per_gbin);
gbins = k_gbins.view<DeviceType>();
}
@ -159,7 +159,7 @@ void NBinSSAKokkos<DeviceType>::bin_atoms()
// actually bin the ghost atoms
{
if(ghosts_per_gbin > (int) gbins.dimension_1()) {
if(ghosts_per_gbin > (int) gbins.extent(1)) {
k_gbins = DAT::tdual_int_2d("gbins", 8, ghosts_per_gbin);
gbins = k_gbins.view<DeviceType>();
}
@ -188,13 +188,13 @@ void NBinSSAKokkos<DeviceType>::bin_atoms()
// actually bin the local atoms
{
if ((mbins > (int) bins.dimension_0()) ||
(atoms_per_bin > (int) bins.dimension_1())) {
if ((mbins > (int) bins.extent(0)) ||
(atoms_per_bin > (int) bins.extent(1))) {
k_bins = DAT::tdual_int_2d("bins", mbins, atoms_per_bin);
bins = k_bins.view<DeviceType>();
}
MemsetZeroFunctor<DeviceType> f_zero;
f_zero.ptr = (void*) k_bincount.view<DeviceType>().ptr_on_device();
f_zero.ptr = (void*) k_bincount.view<DeviceType>().data();
Kokkos::parallel_for(mbins, f_zero);
auto bincount_ = bincount;

2
src/KOKKOS/neighbor_kokkos.cpp
View File

@ -265,7 +265,7 @@ void NeighborKokkos::build_kokkos(int topoflag)
atomKK->sync(ExecutionSpaceFromDevice<DeviceType>::space,X_MASK);
x = atomKK->k_x;
if (includegroup) nlocal = atom->nfirst;
int maxhold_kokkos = xhold.view<DeviceType>().dimension_0();
int maxhold_kokkos = xhold.view<DeviceType>().extent(0);
if (atom->nmax > maxhold || maxhold_kokkos < maxhold) {
maxhold = atom->nmax;
xhold = DAT::tdual_x_array("neigh:xhold",maxhold);

10
src/KOKKOS/npair_kokkos.cpp
View File

@ -89,14 +89,14 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_stencil_info()
int maxstencil = ns->get_maxstencil();
if (maxstencil > k_stencil.dimension_0())
if (maxstencil > k_stencil.extent(0))
k_stencil = DAT::tdual_int_1d("neighlist:stencil",maxstencil);
for (int k = 0; k < maxstencil; k++)
k_stencil.h_view(k) = ns->stencil[k];
k_stencil.modify<LMPHostType>();
k_stencil.sync<DeviceType>();
if (GHOST) {
if (maxstencil > k_stencilxyz.dimension_0())
if (maxstencil > k_stencilxyz.extent(0))
k_stencilxyz = DAT::tdual_int_1d_3("neighlist:stencilxyz",maxstencil);
for (int k = 0; k < maxstencil; k++) {
k_stencilxyz.h_view(k,0) = ns->stencilxyz[k][0];
@ -224,7 +224,7 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::build(NeighList *list_)
if(data.h_resize()) {
deep_copy(data.h_new_maxneighs, data.new_maxneighs);
list->maxneighs = data.h_new_maxneighs() * 1.2;
list->d_neighbors = typename ArrayTypes<DeviceType>::t_neighbors_2d("neighbors", list->d_neighbors.dimension_0(), list->maxneighs);
list->d_neighbors = typename ArrayTypes<DeviceType>::t_neighbors_2d("neighbors", list->d_neighbors.extent(0), list->maxneighs);
data.neigh_list.d_neighbors = list->d_neighbors;
data.neigh_list.maxneighs = list->maxneighs;
}
@ -455,14 +455,14 @@ void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPoli
{
/* loop over atoms in i's bin,
*/
const int atoms_per_bin = c_bins.dimension_1();
const int atoms_per_bin = c_bins.extent(1);
const int BINS_PER_TEAM = dev.team_size()/atoms_per_bin<1?1:dev.team_size()/atoms_per_bin;
const int TEAMS_PER_BIN = atoms_per_bin/dev.team_size()<1?1:atoms_per_bin/dev.team_size();
const int MY_BIN = dev.team_rank()/atoms_per_bin;
const int ibin = dev.league_rank()*BINS_PER_TEAM+MY_BIN;
if(ibin >=c_bincount.dimension_0()) return;
if(ibin >=c_bincount.extent(0)) return;
X_FLOAT* other_x = sharedmem;
other_x = other_x + 5*atoms_per_bin*MY_BIN;

10
src/KOKKOS/npair_ssa_kokkos.cpp
View File

@ -144,7 +144,7 @@ void NPairSSAKokkos<DeviceType>::copy_stencil_info()
// Setup the phases of the workplan for locals
ssa_phaseCt = sz1*sy1*sx1;
if (ssa_phaseCt > (int) k_ssa_phaseLen.dimension_0()) {
if (ssa_phaseCt > (int) k_ssa_phaseLen.extent(0)) {
k_ssa_phaseLen = DAT::tdual_int_1d("NPairSSAKokkos:ssa_phaseLen",ssa_phaseCt);
ssa_phaseLen = k_ssa_phaseLen.view<DeviceType>();
k_ssa_phaseOff = DAT::tdual_int_1d_3("NPairSSAKokkos:ssa_phaseOff",ssa_phaseCt);
@ -251,8 +251,8 @@ void NPairSSAKokkos<DeviceType>::build(NeighList *list_)
int zbinCt = (lbinzhi - lbinzlo + sz1 - 1) / sz1 + 1;
int phaseLenEstimate = xbinCt*ybinCt*zbinCt;
if ((ssa_phaseCt > (int) k_ssa_itemLoc.dimension_0()) ||
(phaseLenEstimate > (int) k_ssa_itemLoc.dimension_1())) {
if ((ssa_phaseCt > (int) k_ssa_itemLoc.extent(0)) ||
(phaseLenEstimate > (int) k_ssa_itemLoc.extent(1))) {
k_ssa_itemLoc = DAT::tdual_int_2d("NPairSSAKokkos::ssa_itemLoc",ssa_phaseCt,phaseLenEstimate);
ssa_itemLoc = k_ssa_itemLoc.view<DeviceType>();
k_ssa_itemLen = DAT::tdual_int_2d("NPairSSAKokkos::ssa_itemLen",ssa_phaseCt,phaseLenEstimate);
@ -492,7 +492,7 @@ fprintf(stdout, "tota%03d total %3d could use %6d inums, expected %6d inums. inu
if(data.h_resize()) {
deep_copy(data.h_new_maxneighs, data.new_maxneighs);
list->maxneighs = data.h_new_maxneighs() * 1.2;
list->d_neighbors = typename ArrayTypes<DeviceType>::t_neighbors_2d("neighbors", list->d_neighbors.dimension_0(), list->maxneighs);
list->d_neighbors = typename ArrayTypes<DeviceType>::t_neighbors_2d("neighbors", list->d_neighbors.extent(0), list->maxneighs);
data.neigh_list.d_neighbors = list->d_neighbors;
data.neigh_list.maxneighs = list->maxneighs;
}
@ -650,7 +650,7 @@ fprintf(stdout, "Phas%03d phase %3d used %6d inums, workItems = %3d, skipped = %
// record where workPhase actually ends
if (firstTry) {
d_ssa_phaseLen(workPhase) = workItem;
while (workItem < (int) d_ssa_itemLen.dimension_1()) {
while (workItem < (int) d_ssa_itemLen.extent(1)) {
d_ssa_itemLen(workPhase,workItem++) = 0;
}
}

6
src/KOKKOS/pair_exp6_rx_kokkos.cpp
View File

@ -189,7 +189,7 @@ void PairExp6rxKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
{
const int np_total = nlocal + atom->nghost;
if (np_total > PairExp6ParamData.epsilon1.dimension_0()) {
if (np_total > PairExp6ParamData.epsilon1.extent(0)) {
PairExp6ParamData.epsilon1 = typename AT::t_float_1d("PairExp6ParamData.epsilon1" ,np_total);
PairExp6ParamData.alpha1 = typename AT::t_float_1d("PairExp6ParamData.alpha1" ,np_total);
PairExp6ParamData.rm1 = typename AT::t_float_1d("PairExp6ParamData.rm1" ,np_total);
@ -308,8 +308,8 @@ void PairExp6rxKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
#else // No atomics
num_threads = lmp->kokkos->num_threads;
int nmax = f.dimension_0();
if (nmax > t_f.dimension_1()) {
int nmax = f.extent(0);
if (nmax > t_f.extent(1)) {
t_f = t_f_array_thread("pair_exp6_rx:t_f",num_threads,nmax);
t_uCG = t_efloat_1d_thread("pair_exp6_rx:t_uCG",num_threads,nmax);
t_uCGnew = t_efloat_1d_thread("pair_exp6_rx:t_UCGnew",num_threads,nmax);

12
src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp
View File

@ -181,7 +181,7 @@ void PairMultiLucyRXKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in
{
const int ntotal = nlocal + nghost;
if (ntotal > d_mixWtSite1.dimension_0()) {
if (ntotal > d_mixWtSite1.extent(0)) {
d_mixWtSite1old = typename AT::t_float_1d("PairMultiLucyRX::mixWtSite1old",ntotal);
d_mixWtSite2old = typename AT::t_float_1d("PairMultiLucyRX::mixWtSite2old",ntotal);
d_mixWtSite1 = typename AT::t_float_1d("PairMultiLucyRX::mixWtSite1",ntotal);
@ -887,15 +887,15 @@ void PairMultiLucyRXKokkos<DeviceType>::create_kokkos_tables()
h_table->innersq[i] = tb->innersq;
h_table->invdelta[i] = tb->invdelta;
for(int j = 0; j<h_table->rsq.dimension_1(); j++)
for(int j = 0; j<h_table->rsq.extent(1); j++)
h_table->rsq(i,j) = tb->rsq[j];
for(int j = 0; j<h_table->e.dimension_1(); j++)
for(int j = 0; j<h_table->e.extent(1); j++)
h_table->e(i,j) = tb->e[j];
for(int j = 0; j<h_table->de.dimension_1(); j++)
for(int j = 0; j<h_table->de.extent(1); j++)
h_table->de(i,j) = tb->de[j];
for(int j = 0; j<h_table->f.dimension_1(); j++)
for(int j = 0; j<h_table->f.extent(1); j++)
h_table->f(i,j) = tb->f[j];
for(int j = 0; j<h_table->df.dimension_1(); j++)
for(int j = 0; j<h_table->df.extent(1); j++)
h_table->df(i,j) = tb->df[j];
}

8
src/KOKKOS/pair_sw_kokkos.cpp
View File

@ -122,13 +122,13 @@ void PairSWKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// build short neighbor list
int max_neighs = d_neighbors.dimension_1();
int max_neighs = d_neighbors.extent(1);
if ((d_neighbors_short.dimension_1() != max_neighs) ||
(d_neighbors_short.dimension_0() != ignum)) {
if ((d_neighbors_short.extent(1) != max_neighs) ||
(d_neighbors_short.extent(0) != ignum)) {
d_neighbors_short = Kokkos::View<int**,DeviceType>("SW::neighbors_short",ignum,max_neighs);
}
if (d_numneigh_short.dimension_0()!=ignum)
if (d_numneigh_short.extent(0)!=ignum)
d_numneigh_short = Kokkos::View<int*,DeviceType>("SW::numneighs_short",ignum);
Kokkos::parallel_for(Kokkos::RangePolicy<DeviceType,TagPairSWComputeShortNeigh>(0,neighflag==FULL?ignum:inum), *this);

16
src/KOKKOS/pair_table_kokkos.cpp
View File

@ -313,21 +313,21 @@ void PairTableKokkos<DeviceType>::create_kokkos_tables()
h_table->invdelta[i] = tb->invdelta;
h_table->deltasq6[i] = tb->deltasq6;
for(int j = 0; j<h_table->rsq.dimension_1(); j++)
for(int j = 0; j<h_table->rsq.extent(1); j++)
h_table->rsq(i,j) = tb->rsq[j];
for(int j = 0; j<h_table->drsq.dimension_1(); j++)
for(int j = 0; j<h_table->drsq.extent(1); j++)
h_table->drsq(i,j) = tb->drsq[j];
for(int j = 0; j<h_table->e.dimension_1(); j++)
for(int j = 0; j<h_table->e.extent(1); j++)
h_table->e(i,j) = tb->e[j];
for(int j = 0; j<h_table->de.dimension_1(); j++)
for(int j = 0; j<h_table->de.extent(1); j++)
h_table->de(i,j) = tb->de[j];
for(int j = 0; j<h_table->f.dimension_1(); j++)
for(int j = 0; j<h_table->f.extent(1); j++)
h_table->f(i,j) = tb->f[j];
for(int j = 0; j<h_table->df.dimension_1(); j++)
for(int j = 0; j<h_table->df.extent(1); j++)
h_table->df(i,j) = tb->df[j];
for(int j = 0; j<h_table->e2.dimension_1(); j++)
for(int j = 0; j<h_table->e2.extent(1); j++)
h_table->e2(i,j) = tb->e2[j];
for(int j = 0; j<h_table->f2.dimension_1(); j++)
for(int j = 0; j<h_table->f2.extent(1); j++)
h_table->f2(i,j) = tb->f2[j];
}

24
src/KOKKOS/pair_table_rx_kokkos.cpp
View File

@ -70,11 +70,11 @@ void getMixingWeights(
nTotal = 0.0;
nTotalOld = 0.0;
assert(id >= 0);
assert(id < dvector.dimension_1());
assert(id < dvector.extent(1));
for (int ispecies = 0; ispecies < nspecies; ++ispecies){
assert(ispecies < dvector.dimension_0());
assert(ispecies < dvector.extent(0));
nTotal += dvector(ispecies,id);
assert(ispecies+nspecies < dvector.dimension_0());
assert(ispecies+nspecies < dvector.extent(0));
nTotalOld += dvector(ispecies+nspecies,id);
}
@ -653,7 +653,7 @@ void PairTableRXKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in)
// loop over neighbors of my atoms
const int ntotal = atom->nlocal + atom->nghost;
if (ntotal > mixWtSite1.dimension_0()) {
if (ntotal > mixWtSite1.extent(0)) {
mixWtSite1old = Kokkos::View<double*, DeviceType>("PairTableRXKokkos::mixWtSite1old", ntotal);
mixWtSite2old = Kokkos::View<double*, DeviceType>("PairTableRXKokkos::mixWtSite2old", ntotal);
mixWtSite1 = Kokkos::View<double*, DeviceType>("PairTableRXKokkos::mixWtSite1", ntotal);
@ -848,21 +848,21 @@ void PairTableRXKokkos<DeviceType>::create_kokkos_tables()
h_table->invdelta[i] = tb->invdelta;
h_table->deltasq6[i] = tb->deltasq6;
for(int j = 0; j<h_table->rsq.dimension_1(); j++)
for(int j = 0; j<h_table->rsq.extent(1); j++)
h_table->rsq(i,j) = tb->rsq[j];
for(int j = 0; j<h_table->drsq.dimension_1(); j++)
for(int j = 0; j<h_table->drsq.extent(1); j++)
h_table->drsq(i,j) = tb->drsq[j];
for(int j = 0; j<h_table->e.dimension_1(); j++)
for(int j = 0; j<h_table->e.extent(1); j++)
h_table->e(i,j) = tb->e[j];
for(int j = 0; j<h_table->de.dimension_1(); j++)
for(int j = 0; j<h_table->de.extent(1); j++)
h_table->de(i,j) = tb->de[j];
for(int j = 0; j<h_table->f.dimension_1(); j++)
for(int j = 0; j<h_table->f.extent(1); j++)
h_table->f(i,j) = tb->f[j];
for(int j = 0; j<h_table->df.dimension_1(); j++)
for(int j = 0; j<h_table->df.extent(1); j++)
h_table->df(i,j) = tb->df[j];
for(int j = 0; j<h_table->e2.dimension_1(); j++)
for(int j = 0; j<h_table->e2.extent(1); j++)
h_table->e2(i,j) = tb->e2[j];
for(int j = 0; j<h_table->f2.dimension_1(); j++)
for(int j = 0; j<h_table->f2.extent(1); j++)
h_table->f2(i,j) = tb->f2[j];
}

8
src/KOKKOS/pair_tersoff_kokkos.cpp
View File

@ -207,13 +207,13 @@ void PairTersoffKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// build short neighbor list
int max_neighs = d_neighbors.dimension_1();
int max_neighs = d_neighbors.extent(1);
if ((d_neighbors_short.dimension_1() != max_neighs) ||
(d_neighbors_short.dimension_0() != ignum)) {
if ((d_neighbors_short.extent(1) != max_neighs) ||
(d_neighbors_short.extent(0) != ignum)) {
d_neighbors_short = Kokkos::View<int**,DeviceType>("Tersoff::neighbors_short",ignum,max_neighs);
}
if (d_numneigh_short.dimension_0()!=ignum)
if (d_numneigh_short.extent(0)!=ignum)
d_numneigh_short = Kokkos::View<int*,DeviceType>("Tersoff::numneighs_short",ignum);
Kokkos::parallel_for(Kokkos::RangePolicy<DeviceType,TagPairTersoffComputeShortNeigh>(0,neighflag==FULL?ignum:inum), *this);

8
src/KOKKOS/pair_tersoff_mod_kokkos.cpp
View File

@ -207,13 +207,13 @@ void PairTersoffMODKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// build short neighbor list
int max_neighs = d_neighbors.dimension_1();
int max_neighs = d_neighbors.extent(1);
if ((d_neighbors_short.dimension_1() != max_neighs) ||
(d_neighbors_short.dimension_0() != ignum)) {
if ((d_neighbors_short.extent(1) != max_neighs) ||
(d_neighbors_short.extent(0) != ignum)) {
d_neighbors_short = Kokkos::View<int**,DeviceType>("Tersoff::neighbors_short",ignum,max_neighs);
}
if (d_numneigh_short.dimension_0()!=ignum)
if (d_numneigh_short.extent(0)!=ignum)
d_numneigh_short = Kokkos::View<int*,DeviceType>("Tersoff::numneighs_short",ignum);
Kokkos::parallel_for(Kokkos::RangePolicy<DeviceType,TagPairTersoffMODComputeShortNeigh>(0,neighflag==FULL?ignum:inum), *this);

8
src/KOKKOS/pair_tersoff_zbl_kokkos.cpp
View File

@ -221,13 +221,13 @@ void PairTersoffZBLKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
// build short neighbor list
int max_neighs = d_neighbors.dimension_1();
int max_neighs = d_neighbors.extent(1);
if ((d_neighbors_short.dimension_1() != max_neighs) ||
(d_neighbors_short.dimension_0() != ignum)) {
if ((d_neighbors_short.extent(1) != max_neighs) ||
(d_neighbors_short.extent(0) != ignum)) {
d_neighbors_short = Kokkos::View<int**,DeviceType>("Tersoff::neighbors_short",ignum,max_neighs);
}
if (d_numneigh_short.dimension_0()!=ignum)
if (d_numneigh_short.extent(0)!=ignum)
d_numneigh_short = Kokkos::View<int*,DeviceType>("Tersoff::numneighs_short",ignum);
Kokkos::parallel_for(Kokkos::RangePolicy<DeviceType,TagPairTersoffZBLComputeShortNeigh>(0,neighflag==FULL?ignum:inum), *this);

14
src/KOKKOS/pair_vashishta_kokkos.cpp
View File

@ -119,22 +119,22 @@ void PairVashishtaKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
EV_FLOAT ev;
EV_FLOAT ev_all;
int max_neighs = d_neighbors.dimension_1();
int max_neighs = d_neighbors.extent(1);
if ((d_neighbors_short_2body.dimension_1() != max_neighs) ||
(d_neighbors_short_2body.dimension_0() != ignum)) {
if ((d_neighbors_short_2body.extent(1) != max_neighs) ||
(d_neighbors_short_2body.extent(0) != ignum)) {
d_neighbors_short_2body = Kokkos::View<int**,DeviceType>("Vashishta::neighbors_short_2body",ignum,max_neighs);
}
if (d_numneigh_short_2body.dimension_0()!=ignum) {
if (d_numneigh_short_2body.extent(0)!=ignum) {
d_numneigh_short_2body = Kokkos::View<int*,DeviceType>("Vashishta::numneighs_short_2body",ignum);
}
if ((d_neighbors_short_3body.dimension_1() != max_neighs) ||
(d_neighbors_short_3body.dimension_0() != ignum)) {
if ((d_neighbors_short_3body.extent(1) != max_neighs) ||
(d_neighbors_short_3body.extent(0) != ignum)) {
d_neighbors_short_3body = Kokkos::View<int**,DeviceType>("Vashishta::neighbors_short_3body",ignum,max_neighs);
}
if (d_numneigh_short_3body.dimension_0()!=ignum) {
if (d_numneigh_short_3body.extent(0)!=ignum) {
d_numneigh_short_3body = Kokkos::View<int*,DeviceType>("Vashishta::numneighs_short_3body",ignum);
}

8
src/KOKKOS/verlet_kokkos.cpp
View File

@ -304,7 +304,7 @@ void VerletKokkos::run(int n)
if (atomKK->sortfreq > 0) sortflag = 1;
else sortflag = 0;
f_merge_copy = DAT::t_f_array("VerletKokkos::f_merge_copy",atomKK->k_f.dimension_0());
f_merge_copy = DAT::t_f_array("VerletKokkos::f_merge_copy",atomKK->k_f.extent(0));
atomKK->sync(Device,ALL_MASK);
//static double time = 0.0;
@ -514,12 +514,12 @@ void VerletKokkos::run(int n)
}
if(execute_on_host && !std::is_same<LMPHostType,LMPDeviceType>::value) {
if(f_merge_copy.dimension_0()<atomKK->k_f.dimension_0()) {
f_merge_copy = DAT::t_f_array("VerletKokkos::f_merge_copy",atomKK->k_f.dimension_0());
if(f_merge_copy.extent(0)<atomKK->k_f.extent(0)) {
f_merge_copy = DAT::t_f_array("VerletKokkos::f_merge_copy",atomKK->k_f.extent(0));
}
f = atomKK->k_f.d_view;
Kokkos::deep_copy(LMPHostType(),f_merge_copy,atomKK->k_f.h_view);
Kokkos::parallel_for(atomKK->k_f.dimension_0(),
Kokkos::parallel_for(atomKK->k_f.extent(0),
ForceAdder<DAT::t_f_array,DAT::t_f_array>(atomKK->k_f.d_view,f_merge_copy));
atomKK->k_f.modified_host() = 0; // special case
atomKK->k_f.modify<LMPDeviceType>();