diff --git a/src/KOKKOS/angle_charmm_kokkos.cpp b/src/KOKKOS/angle_charmm_kokkos.cpp index 59a20c25df..f048738489 100644 --- a/src/KOKKOS/angle_charmm_kokkos.cpp +++ b/src/KOKKOS/angle_charmm_kokkos.cpp @@ -70,14 +70,14 @@ 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); memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view(); //} } if (vflag_atom) { - //if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.template view(); diff --git a/src/KOKKOS/atom_kokkos.cpp b/src/KOKKOS/atom_kokkos.cpp index 4ecead5b1d..b54719e852 100644 --- a/src/KOKKOS/atom_kokkos.cpp +++ b/src/KOKKOS/atom_kokkos.cpp @@ -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(); diff --git a/src/KOKKOS/atom_kokkos.h b/src/KOKKOS/atom_kokkos.h index 2245023189..a83b299ebd 100644 --- a/src/KOKKOS/atom_kokkos.h +++ b/src/KOKKOS/atom_kokkos.h @@ -84,34 +84,34 @@ class SortFunctor { Kokkos::View dest; IndexView index; SortFunctor(ViewType src, typename Kokkos::Impl::enable_if::type ind):source(src),index(ind){ - dest = Kokkos::View("",src.dimension_0()); + dest = Kokkos::View("",src.extent(0)); } SortFunctor(ViewType src, typename Kokkos::Impl::enable_if::type ind):source(src),index(ind){ - dest = Kokkos::View("",src.dimension_0(),src.dimension_1()); + dest = Kokkos::View("",src.extent(0),src.extent(1)); } SortFunctor(ViewType src, typename Kokkos::Impl::enable_if::type ind):source(src),index(ind){ - dest = Kokkos::View("",src.dimension_0(),src.dimension_1(),src.dimension_2()); + dest = Kokkos::View("",src.extent(0),src.extent(1),src.extent(2)); } SortFunctor(ViewType src, typename Kokkos::Impl::enable_if::type ind):source(src),index(ind){ - dest = Kokkos::View("",src.dimension_0(),src.dimension_1(),src.dimension_2(),src.dimension_3()); + dest = Kokkos::View("",src.extent(0),src.extent(1),src.extent(2),src.extent(3)); } KOKKOS_INLINE_FUNCTION void operator()(const typename Kokkos::Impl::enable_if::type& i) { dest(i) = source(index(i)); } void operator()(const typename Kokkos::Impl::enable_if::type& i) { - for(int j=0;j::type& i) { - for(int j=0;j::type& i) { - for(int j=0;j()),_list(list.view()),_iswap(iswap), _xprd(xprd),_yprd(yprd),_zprd(zprd), _xy(xy),_xz(xz),_yz(yz) { - const size_t maxsend = (buf.view().dimension_0() - *buf.view().dimension_1())/3; + const size_t maxsend = (buf.view().extent(0) + *buf.view().extent(1))/3; const size_t elements = 3; buffer_view(_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().dimension_0()* - buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)* + buf.template view().extent(1))/elements; buffer_view(_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().dimension_0()* - k_buf.view().dimension_1())/elements) { - int newsize = nsend*elements/k_buf.view().dimension_1()+1; - k_buf.resize(newsize,k_buf.view().dimension_1()); + if(nsend > (int) (k_buf.view().extent(0)* + k_buf.view().extent(1))/elements) { + int newsize = nsend*elements/k_buf.view().extent(1)+1; + k_buf.resize(newsize,k_buf.view().extent(1)); } if(space == Host) { AtomVecAngleKokkos_PackExchangeFunctor @@ -1335,8 +1335,8 @@ struct AtomVecAngleKokkos_UnpackExchangeFunctor { _nlocal(nlocal.template view()),_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().dimension_0()* - buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)* + buf.template view().extent(1))/elements; buffer_view(_buf,buf,maxsendlist,elements); } diff --git a/src/KOKKOS/atom_vec_atomic_kokkos.cpp b/src/KOKKOS/atom_vec_atomic_kokkos.cpp index f021c45db6..0b96d8dbad 100644 --- a/src/KOKKOS/atom_vec_atomic_kokkos.cpp +++ b/src/KOKKOS/atom_vec_atomic_kokkos.cpp @@ -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().dimension_0()*buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)*buf.template view().extent(1))/elements; buffer_view(_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().dimension_0()*k_buf.view().dimension_1())/11) { - int newsize = nsend*11/k_buf.view().dimension_1()+1; - k_buf.resize(newsize,k_buf.view().dimension_1()); + if(nsend > (int) (k_buf.view().extent(0)*k_buf.view().extent(1))/11) { + int newsize = nsend*11/k_buf.view().extent(1)+1; + k_buf.resize(newsize,k_buf.view().extent(1)); } if(space == Host) { AtomVecAtomicKokkos_PackExchangeFunctor f(atomKK,k_buf,k_sendlist,k_copylist,atom->nlocal,dim,lo,hi); @@ -617,7 +617,7 @@ struct AtomVecAtomicKokkos_UnpackExchangeFunctor { _nlocal(nlocal.template view()),_dim(dim), _lo(lo),_hi(hi){ const size_t elements = 11; - const int maxsendlist = (buf.template view().dimension_0()*buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)*buf.template view().extent(1))/elements; buffer_view(_buf,buf,maxsendlist,elements); } diff --git a/src/KOKKOS/atom_vec_bond_kokkos.cpp b/src/KOKKOS/atom_vec_bond_kokkos.cpp index 9441373aa5..adad4f767b 100644 --- a/src/KOKKOS/atom_vec_bond_kokkos.cpp +++ b/src/KOKKOS/atom_vec_bond_kokkos.cpp @@ -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().dimension_0()* - buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)* + buf.template view().extent(1))/elements; buffer_view(_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().dimension_0()* - k_buf.view().dimension_1())/elements) { - int newsize = nsend*elements/k_buf.view().dimension_1()+1; - k_buf.resize(newsize,k_buf.view().dimension_1()); + if(nsend > (int) (k_buf.view().extent(0)* + k_buf.view().extent(1))/elements) { + int newsize = nsend*elements/k_buf.view().extent(1)+1; + k_buf.resize(newsize,k_buf.view().extent(1)); } if(space == Host) { AtomVecBondKokkos_PackExchangeFunctor @@ -794,8 +794,8 @@ struct AtomVecBondKokkos_UnpackExchangeFunctor { _nlocal(nlocal.template view()),_dim(dim), _lo(lo),_hi(hi){ elements = 16+atom->maxspecial+atom->bond_per_atom+atom->bond_per_atom; - const int maxsendlist = (buf.template view().dimension_0()* - buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)* + buf.template view().extent(1))/elements; buffer_view(_buf,buf,maxsendlist,elements); } diff --git a/src/KOKKOS/atom_vec_charge_kokkos.cpp b/src/KOKKOS/atom_vec_charge_kokkos.cpp index a9ae5cc2d1..3e50870c8f 100644 --- a/src/KOKKOS/atom_vec_charge_kokkos.cpp +++ b/src/KOKKOS/atom_vec_charge_kokkos.cpp @@ -169,7 +169,7 @@ struct AtomVecChargeKokkos_PackComm { _x(x.view()),_list(list.view()),_iswap(iswap), _xprd(xprd),_yprd(yprd),_zprd(zprd), _xy(xy),_xz(xz),_yz(yz) { - const size_t maxsend = (buf.view().dimension_0()*buf.view().dimension_1())/3; + const size_t maxsend = (buf.view().extent(0)*buf.view().extent(1))/3; const size_t elements = 3; buffer_view(_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().dimension_0()* - buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)* + buf.template view().extent(1))/elements; buffer_view(_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().dimension_0()*k_buf.view().dimension_1())/12) { - int newsize = nsend*12/k_buf.view().dimension_1()+1; - k_buf.resize(newsize,k_buf.view().dimension_1()); + if(nsend > (int) (k_buf.view().extent(0)*k_buf.view().extent(1))/12) { + int newsize = nsend*12/k_buf.view().extent(1)+1; + k_buf.resize(newsize,k_buf.view().extent(1)); } if(space == Host) { AtomVecChargeKokkos_PackExchangeFunctor @@ -742,7 +742,7 @@ struct AtomVecChargeKokkos_UnpackExchangeFunctor { _nlocal(nlocal.template view()),_dim(dim), _lo(lo),_hi(hi){ const size_t elements = 12; - const int maxsendlist = (buf.template view().dimension_0()*buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)*buf.template view().extent(1))/elements; buffer_view(_buf,buf,maxsendlist,elements); } diff --git a/src/KOKKOS/atom_vec_dpd_kokkos.cpp b/src/KOKKOS/atom_vec_dpd_kokkos.cpp index ca1779d0ee..06ca2083dc 100644 --- a/src/KOKKOS/atom_vec_dpd_kokkos.cpp +++ b/src/KOKKOS/atom_vec_dpd_kokkos.cpp @@ -219,7 +219,7 @@ struct AtomVecDPDKokkos_PackComm { _list(list.view()),_iswap(iswap), _xprd(xprd),_yprd(yprd),_zprd(zprd), _xy(xy),_xz(xz),_yz(yz) { - const size_t maxsend = (buf.view().dimension_0()*buf.view().dimension_1())/3; + const size_t maxsend = (buf.view().extent(0)*buf.view().extent(1))/3; const size_t elements = 3; buffer_view(_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().dimension_0()*buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)*buf.template view().extent(1))/elements; buffer_view(_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().dimension_0()*k_buf.view().dimension_1())/17) { - int newsize = nsend*17/k_buf.view().dimension_1()+1; - k_buf.resize(newsize,k_buf.view().dimension_1()); + if(nsend > (int) (k_buf.view().extent(0)*k_buf.view().extent(1))/17) { + int newsize = nsend*17/k_buf.view().extent(1)+1; + k_buf.resize(newsize,k_buf.view().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()),_dim(dim), _lo(lo),_hi(hi){ const size_t elements = 17; - const int maxsendlist = (buf.template view().dimension_0()*buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)*buf.template view().extent(1))/elements; buffer_view(_buf,buf,maxsendlist,elements); } diff --git a/src/KOKKOS/atom_vec_full_kokkos.cpp b/src/KOKKOS/atom_vec_full_kokkos.cpp index 0560c99037..6492a264ac 100644 --- a/src/KOKKOS/atom_vec_full_kokkos.cpp +++ b/src/KOKKOS/atom_vec_full_kokkos.cpp @@ -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().dimension_0()* - buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)* + buf.template view().extent(1))/elements; buffer_view(_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().dimension_0()* - k_buf.view().dimension_1())/elements) { - int newsize = nsend*elements/k_buf.view().dimension_1()+1; - k_buf.resize(newsize,k_buf.view().dimension_1()); + if(nsend > (int) (k_buf.view().extent(0)* + k_buf.view().extent(1))/elements) { + int newsize = nsend*elements/k_buf.view().extent(1)+1; + k_buf.resize(newsize,k_buf.view().extent(1)); } if(space == Host) { AtomVecFullKokkos_PackExchangeFunctor @@ -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().dimension_0()* - buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)* + buf.template view().extent(1))/elements; buffer_view(_buf,buf,maxsendlist,elements); } diff --git a/src/KOKKOS/atom_vec_kokkos.cpp b/src/KOKKOS/atom_vec_kokkos.cpp index 885f190f84..f2c04bec1b 100644 --- a/src/KOKKOS/atom_vec_kokkos.cpp +++ b/src/KOKKOS/atom_vec_kokkos.cpp @@ -51,7 +51,7 @@ struct AtomVecKokkos_PackComm { _x(x.view()),_list(list.view()),_iswap(iswap), _xprd(xprd),_yprd(yprd),_zprd(zprd), _xy(xy),_xz(xz),_yz(yz) { - const size_t maxsend = (buf.view().dimension_0()*buf.view().dimension_1())/3; + const size_t maxsend = (buf.view().extent(0)*buf.view().extent(1))/3; const size_t elements = 3; buffer_view(_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()),_list(list.view()),_iswap(iswap) { - const size_t maxsend = (buf.view().dimension_0()*buf.view().dimension_1())/3; + const size_t maxsend = (buf.view().extent(0)*buf.view().extent(1))/3; const size_t elements = 3; buffer_view(_buf,buf,maxsend,elements); }; diff --git a/src/KOKKOS/atom_vec_kokkos.h b/src/KOKKOS/atom_vec_kokkos.h index 20a07ec443..b68fdc9b20 100644 --- a/src/KOKKOS/atom_vec_kokkos.h +++ b/src/KOKKOS/atom_vec_kokkos.h @@ -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 @@ -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); diff --git a/src/KOKKOS/atom_vec_molecular_kokkos.cpp b/src/KOKKOS/atom_vec_molecular_kokkos.cpp index 380aa0fb7e..1a0c03cca5 100644 --- a/src/KOKKOS/atom_vec_molecular_kokkos.cpp +++ b/src/KOKKOS/atom_vec_molecular_kokkos.cpp @@ -319,8 +319,8 @@ struct AtomVecMolecularKokkos_PackComm { _x(x.view()),_list(list.view()),_iswap(iswap), _xprd(xprd),_yprd(yprd),_zprd(zprd), _xy(xy),_xz(xz),_yz(yz) { - const size_t maxsend = (buf.view().dimension_0() - *buf.view().dimension_1())/3; + const size_t maxsend = (buf.view().extent(0) + *buf.view().extent(1))/3; const size_t elements = 3; buffer_view(_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().dimension_0()* - buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)* + buf.template view().extent(1))/elements; buffer_view(_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().dimension_0()* - k_buf.view().dimension_1())/elements) { - int newsize = nsend*elements/k_buf.view().dimension_1()+1; - k_buf.resize(newsize,k_buf.view().dimension_1()); + if(nsend > (int) (k_buf.view().extent(0)* + k_buf.view().extent(1))/elements) { + int newsize = nsend*elements/k_buf.view().extent(1)+1; + k_buf.resize(newsize,k_buf.view().extent(1)); } if(space == Host) { AtomVecMolecularKokkos_PackExchangeFunctor @@ -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().dimension_0()* - buf.template view().dimension_1())/elements; + const int maxsendlist = (buf.template view().extent(0)* + buf.template view().extent(1))/elements; buffer_view(_buf,buf,maxsendlist,elements); } diff --git a/src/KOKKOS/bond_class2_kokkos.cpp b/src/KOKKOS/bond_class2_kokkos.cpp index df2f2c1e9b..80f2a6dfbb 100644 --- a/src/KOKKOS/bond_class2_kokkos.cpp +++ b/src/KOKKOS/bond_class2_kokkos.cpp @@ -66,14 +66,14 @@ void BondClass2Kokkos::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); memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view(); //} } if (vflag_atom) { - //if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.template view(); diff --git a/src/KOKKOS/bond_harmonic_kokkos.cpp b/src/KOKKOS/bond_harmonic_kokkos.cpp index c4e0c3a817..e277104ddb 100644 --- a/src/KOKKOS/bond_harmonic_kokkos.cpp +++ b/src/KOKKOS/bond_harmonic_kokkos.cpp @@ -67,14 +67,14 @@ 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); memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view(); //} } if (vflag_atom) { - //if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.template view(); diff --git a/src/KOKKOS/comm_kokkos.cpp b/src/KOKKOS/comm_kokkos.cpp index b41849197b..d24fa8a536 100644 --- a/src/KOKKOS/comm_kokkos.cpp +++ b/src/KOKKOS/comm_kokkos.cpp @@ -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().ptr_on_device() + - firstrecv[iswap]*atomKK->k_x.view().dimension_1(); + buf = atomKK->k_x.view().data() + + firstrecv[iswap]*atomKK->k_x.view().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().ptr_on_device(), + MPI_Send(k_buf_send.view().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().ptr_on_device(), + MPI_Irecv(k_buf_recv.view().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().ptr_on_device(), + MPI_Irecv(k_buf_recv.view().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().ptr_on_device(),n, + MPI_Send(k_buf_send.view().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().ptr_on_device(),size_reverse_recv[iswap],MPI_DOUBLE, + MPI_Irecv(k_buf_recv.view().data(),size_reverse_recv[iswap],MPI_DOUBLE, sendproc[iswap],0,world,&request); if (size_reverse_send[iswap]) { - buf = atomKK->k_f.view().ptr_on_device() + - firstrecv[iswap]*atomKK->k_f.view().dimension_1(); + buf = atomKK->k_f.view().data() + + firstrecv[iswap]*atomKK->k_f.view().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().ptr_on_device(), + MPI_Irecv(k_buf_recv.view().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().ptr_on_device(),n, + MPI_Send(k_buf_send.view().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().ptr_on_device(),nsize*recvnum[iswap],MPI_DOUBLE, + MPI_Irecv(k_buf_recv_pair.view().data(),nsize*recvnum[iswap],MPI_DOUBLE, recvproc[iswap],0,world,&request); if (sendnum[iswap]) - MPI_Send(k_buf_send_pair.view().ptr_on_device(),n,MPI_DOUBLE,sendproc[iswap],0,world); + MPI_Send(k_buf_send_pair.view().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()(); - 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(); k_count.sync(); @@ -587,10 +587,10 @@ void CommKokkos::exchange_device() k_count.modify(); k_count.sync(); - 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(); @@ -655,18 +655,18 @@ void CommKokkos::exchange_device() } if (nrecv > maxrecv) grow_recv_kokkos(nrecv); - MPI_Irecv(k_buf_recv.view().ptr_on_device(),nrecv1, + MPI_Irecv(k_buf_recv.view().data(),nrecv1, MPI_DOUBLE,procneigh[dim][1],0, world,&request); - MPI_Send(k_buf_send.view().ptr_on_device(),nsend, + MPI_Send(k_buf_send.view().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().ptr_on_device()+nrecv1, + MPI_Irecv(k_buf_recv.view().data()+nrecv1, nrecv2,MPI_DOUBLE,procneigh[dim][0],0, world,&request); - MPI_Send(k_buf_send.view().ptr_on_device(),nsend, + MPI_Send(k_buf_send.view().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().ptr_on_device(), + if (nrecv) MPI_Irecv(k_buf_recv.view().data(), nrecv*size_border,MPI_DOUBLE, recvproc[iswap],0,world,&request); - if (n) MPI_Send(k_buf_send.view().ptr_on_device(),n, + if (n) MPI_Send(k_buf_send.view().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(); k_buf_send.resize(maxsend_border,atom->avec->size_border); - buf_send = k_buf_send.view().ptr_on_device(); + buf_send = k_buf_send.view().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().ptr_on_device(); + buf_send = k_buf_send.view().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().ptr_on_device(); + buf_recv = k_buf_recv.view().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(); diff --git a/src/KOKKOS/dihedral_charmm_kokkos.cpp b/src/KOKKOS/dihedral_charmm_kokkos.cpp index 71635ec76c..13e3a6d2d1 100644 --- a/src/KOKKOS/dihedral_charmm_kokkos.cpp +++ b/src/KOKKOS/dihedral_charmm_kokkos.cpp @@ -80,7 +80,7 @@ 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); memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); d_eatom = k_eatom.template view(); @@ -89,7 +89,7 @@ void DihedralCharmmKokkos::compute(int eflag_in, int vflag_in) //} } if (vflag_atom) { - //if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom"); d_vatom = k_vatom.template view(); diff --git a/src/KOKKOS/fix_eos_table_rx_kokkos.cpp b/src/KOKKOS/fix_eos_table_rx_kokkos.cpp index 5c106c19f3..fb4ca443a9 100644 --- a/src/KOKKOS/fix_eos_table_rx_kokkos.cpp +++ b/src/KOKKOS/fix_eos_table_rx_kokkos.cpp @@ -534,11 +534,11 @@ void FixEOStableRXKokkos::create_kokkos_tables() h_table->hi[i] = tb->hi; h_table->invdelta[i] = tb->invdelta; - for(int j = 0; jr.dimension_1(); j++) + for(int j = 0; jr.extent(1); j++) h_table->r(i,j) = tb->r[j]; - for(int j = 0; je.dimension_1(); j++) + for(int j = 0; je.extent(1); j++) h_table->e(i,j) = tb->e[j]; - for(int j = 0; jde.dimension_1(); j++) + for(int j = 0; jde.extent(1); j++) h_table->de(i,j) = tb->de[j]; } diff --git a/src/KOKKOS/fix_property_atom_kokkos.cpp b/src/KOKKOS/fix_property_atom_kokkos.cpp index 9e29b6c35d..fb3f5a3f34 100644 --- a/src/KOKKOS/fix_property_atom_kokkos.cpp +++ b/src/KOKKOS/fix_property_atom_kokkos.cpp @@ -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); diff --git a/src/KOKKOS/fix_rx_kokkos.cpp b/src/KOKKOS/fix_rx_kokkos.cpp index 0319c79cb3..22450e19d9 100644 --- a/src/KOKKOS/fix_rx_kokkos.cpp +++ b/src/KOKKOS/fix_rx_kokkos.cpp @@ -1352,7 +1352,7 @@ void FixRxKokkos::operator()(Tag_FixRxKokkos_solveSystems y( d_scratchSpace.ptr_on_device() + scratchSpaceSize * i ); + StridedArrayType y( d_scratchSpace.data() + scratchSpaceSize * i ); StridedArrayType rwork( &y[nspecies] ); UserRHSDataKokkos<1> userData; @@ -1447,7 +1447,7 @@ void FixRxKokkos::solve_reactions(const int vflag, const bool isPreF { const int count = nlocal + (newton_pair ? nghost : 0); - if (count > k_dpdThetaLocal.template view().dimension_0()) { + if (count > k_dpdThetaLocal.template view().extent(0)) { memoryKK->destroy_kokkos (k_dpdThetaLocal, dpdThetaLocal); memoryKK->create_kokkos (k_dpdThetaLocal, dpdThetaLocal, count, "FixRxKokkos::dpdThetaLocal"); this->d_dpdThetaLocal = k_dpdThetaLocal.template view(); @@ -1544,7 +1544,7 @@ void FixRxKokkos::solve_reactions(const int vflag, const bool isPreF //double *scratchSpace = new double[ scratchSpaceSize * nlocal ]; //typename ArrayTypes::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::solve_reactions(const int vflag, const bool isPreF //StridedArrayType _y( y ); //StridedArrayType _rwork( rwork ); - StridedArrayType y( d_scratchSpace.ptr_on_device() + scratchSpaceSize * i ); + StridedArrayType y( d_scratchSpace.data() + scratchSpaceSize * i ); StridedArrayType rwork( &y[nspecies] ); //UserRHSData userData; @@ -2026,7 +2026,7 @@ void FixRxKokkos::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(); @@ -2047,7 +2047,7 @@ void FixRxKokkos::computeLocalTemperature() int sumWeightsCt = nlocal + (NEWTON_PAIR ? nghost : 0); //memoryKK->create_kokkos (k_sumWeights, sumWeights, sumWeightsCt, "FixRxKokkos::sumWeights"); - if (sumWeightsCt > k_sumWeights.template view().dimension_0()) { + if (sumWeightsCt > k_sumWeights.template view().extent(0)) { memoryKK->destroy_kokkos(k_sumWeights, sumWeights); memoryKK->create_kokkos (k_sumWeights, sumWeightsCt, "FixRxKokkos::sumWeights"); d_sumWeights = k_sumWeights.template view(); diff --git a/src/KOKKOS/fix_shardlow_kokkos.cpp b/src/KOKKOS/fix_shardlow_kokkos.cpp index 571f488023..e973e5c99f 100644 --- a/src/KOKKOS/fix_shardlow_kokkos.cpp +++ b/src/KOKKOS/fix_shardlow_kokkos.cpp @@ -602,9 +602,9 @@ void FixShardlowKokkos::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; diff --git a/src/KOKKOS/gridcomm_kokkos.cpp b/src/KOKKOS/gridcomm_kokkos.cpp index f107370514..847fa5907a 100644 --- a/src/KOKKOS/gridcomm_kokkos.cpp +++ b/src/KOKKOS/gridcomm_kokkos.cpp @@ -241,8 +241,8 @@ void GridCommKokkos::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::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::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::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::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::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::forward_comm(KSpace *kspace, int which) DeviceType::fence(); if (swap[m].sendproc != me) { - MPI_Irecv(k_buf2.view().ptr_on_device(),nforward*swap[m].nunpack,MPI_FFT_SCALAR, + MPI_Irecv(k_buf2.view().data(),nforward*swap[m].nunpack,MPI_FFT_SCALAR, swap[m].recvproc,0,gridcomm,&request); - MPI_Send(k_buf1.view().ptr_on_device(),nforward*swap[m].npack,MPI_FFT_SCALAR, + MPI_Send(k_buf1.view().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::reverse_comm(KSpace *kspace, int which) DeviceType::fence(); if (swap[m].recvproc != me) { - MPI_Irecv(k_buf2.view().ptr_on_device(),nreverse*swap[m].npack,MPI_FFT_SCALAR, + MPI_Irecv(k_buf2.view().data(),nreverse*swap[m].npack,MPI_FFT_SCALAR, swap[m].sendproc,0,gridcomm,&request); - MPI_Send(k_buf1.view().ptr_on_device(),nreverse*swap[m].nunpack,MPI_FFT_SCALAR, + MPI_Send(k_buf1.view().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::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); diff --git a/src/KOKKOS/improper_class2_kokkos.cpp b/src/KOKKOS/improper_class2_kokkos.cpp index 2171305a9c..28f57608bb 100644 --- a/src/KOKKOS/improper_class2_kokkos.cpp +++ b/src/KOKKOS/improper_class2_kokkos.cpp @@ -77,14 +77,14 @@ void ImproperClass2Kokkos::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); memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view(); //} } if (vflag_atom) { - //if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.template view(); diff --git a/src/KOKKOS/improper_harmonic_kokkos.cpp b/src/KOKKOS/improper_harmonic_kokkos.cpp index 49dd36ed19..f1a0ac3631 100644 --- a/src/KOKKOS/improper_harmonic_kokkos.cpp +++ b/src/KOKKOS/improper_harmonic_kokkos.cpp @@ -77,14 +77,14 @@ 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); memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view(); //} } if (vflag_atom) { - //if(k_vatom.dimension_0()destroy_kokkos(k_vatom,vatom); memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.template view(); diff --git a/src/KOKKOS/kokkos_type.h b/src/KOKKOS/kokkos_type.h index da06be98b6..ddea35ca88 100644 --- a/src/KOKKOS/kokkos_type.h +++ b/src/KOKKOS/kokkos_type.h @@ -956,7 +956,7 @@ void buffer_view(BufferView &buf, DualView &view, const size_t n7 = 0) { buf = BufferView( - view.template view().ptr_on_device(), + view.template view().data(), n0,n1,n2,n3,n4,n5,n6,n7); } @@ -973,7 +973,7 @@ struct MemsetZeroFunctor { template void memset_kokkos (ViewType &view) { static MemsetZeroFunctor 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 diff --git a/src/KOKKOS/memory_kokkos.h b/src/KOKKOS/memory_kokkos.h index 9f930faae2..39d30ac785 100644 --- a/src/KOKKOS/memory_kokkos.h +++ b/src/KOKKOS/memory_kokkos.h @@ -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 #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 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); diff --git a/src/KOKKOS/nbin_kokkos.cpp b/src/KOKKOS/nbin_kokkos.cpp index 95ea105ad9..2ca2a83b5a 100644 --- a/src/KOKKOS/nbin_kokkos.cpp +++ b/src/KOKKOS/nbin_kokkos.cpp @@ -68,14 +68,14 @@ NBinKokkos::NBinKokkos(LAMMPS *lmp) : NBinStandard(lmp) { template void NBinKokkos::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(); k_bincount = DAT::tdual_int_1d("Neighbor::d_bincount",mbins); bincount = k_bincount.view(); } - 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(); } @@ -101,7 +101,7 @@ void NBinKokkos::bin_atoms() deep_copy(d_resize, h_resize); MemsetZeroFunctor f_zero; - f_zero.ptr = (void*) k_bincount.view().ptr_on_device(); + f_zero.ptr = (void*) k_bincount.view().data(); Kokkos::parallel_for(mbins, f_zero); atomKK->sync(ExecutionSpaceFromDevice::space,X_MASK); @@ -139,7 +139,7 @@ void NBinKokkos::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; diff --git a/src/KOKKOS/nbin_ssa_kokkos.cpp b/src/KOKKOS/nbin_ssa_kokkos.cpp index 1a88165856..d040ee2c83 100644 --- a/src/KOKKOS/nbin_ssa_kokkos.cpp +++ b/src/KOKKOS/nbin_ssa_kokkos.cpp @@ -72,7 +72,7 @@ NBinSSAKokkos::NBinSSAKokkos(LAMMPS *lmp) : NBinStandard(lmp) template void NBinSSAKokkos::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(); @@ -82,7 +82,7 @@ void NBinSSAKokkos::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(); } @@ -159,7 +159,7 @@ void NBinSSAKokkos::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(); } @@ -188,13 +188,13 @@ void NBinSSAKokkos::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(); } MemsetZeroFunctor f_zero; - f_zero.ptr = (void*) k_bincount.view().ptr_on_device(); + f_zero.ptr = (void*) k_bincount.view().data(); Kokkos::parallel_for(mbins, f_zero); auto bincount_ = bincount; diff --git a/src/KOKKOS/neighbor_kokkos.cpp b/src/KOKKOS/neighbor_kokkos.cpp index ea90a18270..72092aef0b 100644 --- a/src/KOKKOS/neighbor_kokkos.cpp +++ b/src/KOKKOS/neighbor_kokkos.cpp @@ -265,7 +265,7 @@ void NeighborKokkos::build_kokkos(int topoflag) atomKK->sync(ExecutionSpaceFromDevice::space,X_MASK); x = atomKK->k_x; if (includegroup) nlocal = atom->nfirst; - int maxhold_kokkos = xhold.view().dimension_0(); + int maxhold_kokkos = xhold.view().extent(0); if (atom->nmax > maxhold || maxhold_kokkos < maxhold) { maxhold = atom->nmax; xhold = DAT::tdual_x_array("neigh:xhold",maxhold); diff --git a/src/KOKKOS/npair_kokkos.cpp b/src/KOKKOS/npair_kokkos.cpp index d3cdcb0680..4411627a78 100644 --- a/src/KOKKOS/npair_kokkos.cpp +++ b/src/KOKKOS/npair_kokkos.cpp @@ -89,14 +89,14 @@ void NPairKokkos::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(); k_stencil.sync(); 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::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::t_neighbors_2d("neighbors", list->d_neighbors.dimension_0(), list->maxneighs); + list->d_neighbors = typename ArrayTypes::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::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; diff --git a/src/KOKKOS/npair_ssa_kokkos.cpp b/src/KOKKOS/npair_ssa_kokkos.cpp index 9f447bda1a..ba941b8e54 100644 --- a/src/KOKKOS/npair_ssa_kokkos.cpp +++ b/src/KOKKOS/npair_ssa_kokkos.cpp @@ -144,7 +144,7 @@ void NPairSSAKokkos::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(); k_ssa_phaseOff = DAT::tdual_int_1d_3("NPairSSAKokkos:ssa_phaseOff",ssa_phaseCt); @@ -251,8 +251,8 @@ void NPairSSAKokkos::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(); 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::t_neighbors_2d("neighbors", list->d_neighbors.dimension_0(), list->maxneighs); + list->d_neighbors = typename ArrayTypes::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; } } diff --git a/src/KOKKOS/pair_exp6_rx_kokkos.cpp b/src/KOKKOS/pair_exp6_rx_kokkos.cpp index 5f8e724a60..65b36635a3 100644 --- a/src/KOKKOS/pair_exp6_rx_kokkos.cpp +++ b/src/KOKKOS/pair_exp6_rx_kokkos.cpp @@ -189,7 +189,7 @@ void PairExp6rxKokkos::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::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); diff --git a/src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp b/src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp index 0961cf44eb..7cb869a8a5 100644 --- a/src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp +++ b/src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp @@ -181,7 +181,7 @@ void PairMultiLucyRXKokkos::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::create_kokkos_tables() h_table->innersq[i] = tb->innersq; h_table->invdelta[i] = tb->invdelta; - for(int j = 0; jrsq.dimension_1(); j++) + for(int j = 0; jrsq.extent(1); j++) h_table->rsq(i,j) = tb->rsq[j]; - for(int j = 0; je.dimension_1(); j++) + for(int j = 0; je.extent(1); j++) h_table->e(i,j) = tb->e[j]; - for(int j = 0; jde.dimension_1(); j++) + for(int j = 0; jde.extent(1); j++) h_table->de(i,j) = tb->de[j]; - for(int j = 0; jf.dimension_1(); j++) + for(int j = 0; jf.extent(1); j++) h_table->f(i,j) = tb->f[j]; - for(int j = 0; jdf.dimension_1(); j++) + for(int j = 0; jdf.extent(1); j++) h_table->df(i,j) = tb->df[j]; } diff --git a/src/KOKKOS/pair_sw_kokkos.cpp b/src/KOKKOS/pair_sw_kokkos.cpp index cc9307c42e..a19e2fc271 100644 --- a/src/KOKKOS/pair_sw_kokkos.cpp +++ b/src/KOKKOS/pair_sw_kokkos.cpp @@ -122,13 +122,13 @@ void PairSWKokkos::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("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("SW::numneighs_short",ignum); Kokkos::parallel_for(Kokkos::RangePolicy(0,neighflag==FULL?ignum:inum), *this); diff --git a/src/KOKKOS/pair_table_kokkos.cpp b/src/KOKKOS/pair_table_kokkos.cpp index 957214915b..e3e4ee01f3 100644 --- a/src/KOKKOS/pair_table_kokkos.cpp +++ b/src/KOKKOS/pair_table_kokkos.cpp @@ -313,21 +313,21 @@ void PairTableKokkos::create_kokkos_tables() h_table->invdelta[i] = tb->invdelta; h_table->deltasq6[i] = tb->deltasq6; - for(int j = 0; jrsq.dimension_1(); j++) + for(int j = 0; jrsq.extent(1); j++) h_table->rsq(i,j) = tb->rsq[j]; - for(int j = 0; jdrsq.dimension_1(); j++) + for(int j = 0; jdrsq.extent(1); j++) h_table->drsq(i,j) = tb->drsq[j]; - for(int j = 0; je.dimension_1(); j++) + for(int j = 0; je.extent(1); j++) h_table->e(i,j) = tb->e[j]; - for(int j = 0; jde.dimension_1(); j++) + for(int j = 0; jde.extent(1); j++) h_table->de(i,j) = tb->de[j]; - for(int j = 0; jf.dimension_1(); j++) + for(int j = 0; jf.extent(1); j++) h_table->f(i,j) = tb->f[j]; - for(int j = 0; jdf.dimension_1(); j++) + for(int j = 0; jdf.extent(1); j++) h_table->df(i,j) = tb->df[j]; - for(int j = 0; je2.dimension_1(); j++) + for(int j = 0; je2.extent(1); j++) h_table->e2(i,j) = tb->e2[j]; - for(int j = 0; jf2.dimension_1(); j++) + for(int j = 0; jf2.extent(1); j++) h_table->f2(i,j) = tb->f2[j]; } diff --git a/src/KOKKOS/pair_table_rx_kokkos.cpp b/src/KOKKOS/pair_table_rx_kokkos.cpp index 7bc5198d8c..0b1e0ee9e6 100644 --- a/src/KOKKOS/pair_table_rx_kokkos.cpp +++ b/src/KOKKOS/pair_table_rx_kokkos.cpp @@ -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::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("PairTableRXKokkos::mixWtSite1old", ntotal); mixWtSite2old = Kokkos::View("PairTableRXKokkos::mixWtSite2old", ntotal); mixWtSite1 = Kokkos::View("PairTableRXKokkos::mixWtSite1", ntotal); @@ -848,21 +848,21 @@ void PairTableRXKokkos::create_kokkos_tables() h_table->invdelta[i] = tb->invdelta; h_table->deltasq6[i] = tb->deltasq6; - for(int j = 0; jrsq.dimension_1(); j++) + for(int j = 0; jrsq.extent(1); j++) h_table->rsq(i,j) = tb->rsq[j]; - for(int j = 0; jdrsq.dimension_1(); j++) + for(int j = 0; jdrsq.extent(1); j++) h_table->drsq(i,j) = tb->drsq[j]; - for(int j = 0; je.dimension_1(); j++) + for(int j = 0; je.extent(1); j++) h_table->e(i,j) = tb->e[j]; - for(int j = 0; jde.dimension_1(); j++) + for(int j = 0; jde.extent(1); j++) h_table->de(i,j) = tb->de[j]; - for(int j = 0; jf.dimension_1(); j++) + for(int j = 0; jf.extent(1); j++) h_table->f(i,j) = tb->f[j]; - for(int j = 0; jdf.dimension_1(); j++) + for(int j = 0; jdf.extent(1); j++) h_table->df(i,j) = tb->df[j]; - for(int j = 0; je2.dimension_1(); j++) + for(int j = 0; je2.extent(1); j++) h_table->e2(i,j) = tb->e2[j]; - for(int j = 0; jf2.dimension_1(); j++) + for(int j = 0; jf2.extent(1); j++) h_table->f2(i,j) = tb->f2[j]; } diff --git a/src/KOKKOS/pair_tersoff_kokkos.cpp b/src/KOKKOS/pair_tersoff_kokkos.cpp index b64b2f6d43..1fb97f913f 100644 --- a/src/KOKKOS/pair_tersoff_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_kokkos.cpp @@ -207,13 +207,13 @@ void PairTersoffKokkos::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("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("Tersoff::numneighs_short",ignum); Kokkos::parallel_for(Kokkos::RangePolicy(0,neighflag==FULL?ignum:inum), *this); diff --git a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp index 21c7fbb727..a30e1229f2 100644 --- a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp @@ -207,13 +207,13 @@ void PairTersoffMODKokkos::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("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("Tersoff::numneighs_short",ignum); Kokkos::parallel_for(Kokkos::RangePolicy(0,neighflag==FULL?ignum:inum), *this); diff --git a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp index 8a34dc34ab..aeefbf67f6 100644 --- a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp @@ -221,13 +221,13 @@ void PairTersoffZBLKokkos::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("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("Tersoff::numneighs_short",ignum); Kokkos::parallel_for(Kokkos::RangePolicy(0,neighflag==FULL?ignum:inum), *this); diff --git a/src/KOKKOS/pair_vashishta_kokkos.cpp b/src/KOKKOS/pair_vashishta_kokkos.cpp index 5611aa132b..2e4989356e 100644 --- a/src/KOKKOS/pair_vashishta_kokkos.cpp +++ b/src/KOKKOS/pair_vashishta_kokkos.cpp @@ -119,22 +119,22 @@ void PairVashishtaKokkos::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("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("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("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("Vashishta::numneighs_short_3body",ignum); } diff --git a/src/KOKKOS/verlet_kokkos.cpp b/src/KOKKOS/verlet_kokkos.cpp index b2b107284e..fe11e6b3fa 100644 --- a/src/KOKKOS/verlet_kokkos.cpp +++ b/src/KOKKOS/verlet_kokkos.cpp @@ -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::value) { - if(f_merge_copy.dimension_0()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)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(atomKK->k_f.d_view,f_merge_copy)); atomKK->k_f.modified_host() = 0; // special case atomKK->k_f.modify();