diff --git a/src/KOKKOS/angle_charmm_kokkos.cpp b/src/KOKKOS/angle_charmm_kokkos.cpp
index 59a20c25df36184d4a23a9bf9f6071c625f5eb38..f048738489ea8efd7ce749c5477fdf1a391b3f6f 100644
--- a/src/KOKKOS/angle_charmm_kokkos.cpp
+++ b/src/KOKKOS/angle_charmm_kokkos.cpp
@@ -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>();
diff --git a/src/KOKKOS/atom_kokkos.cpp b/src/KOKKOS/atom_kokkos.cpp
index 4ecead5b1d052be975c40069d169a7fad43cbf63..b54719e852cb939ffc313df6a81235054781bccc 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<LMPHostType>();
diff --git a/src/KOKKOS/atom_kokkos.h b/src/KOKKOS/atom_kokkos.h
index 224502318996aa12da4139928b5e323569d8cbbe..a83b299ebdb8c16ae1a464369714797429d4c3b7 100644
--- a/src/KOKKOS/atom_kokkos.h
+++ b/src/KOKKOS/atom_kokkos.h
@@ -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);
}
};
diff --git a/src/KOKKOS/atom_vec_angle_kokkos.cpp b/src/KOKKOS/atom_vec_angle_kokkos.cpp
index 644ed5280ae529b46875615b823554b7d221a7ff..8ff6d721a8a7b1fd4a8b06fdd7f534d0720530eb 100644
--- a/src/KOKKOS/atom_vec_angle_kokkos.cpp
+++ b/src/KOKKOS/atom_vec_angle_kokkos.cpp
@@ -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);
}
diff --git a/src/KOKKOS/atom_vec_atomic_kokkos.cpp b/src/KOKKOS/atom_vec_atomic_kokkos.cpp
index f021c45db6feb9c1407115dd1cbeadebf9ddf1d6..0b96d8dbadc98f8fcf3559e21464d6102b8af088 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<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);
}
diff --git a/src/KOKKOS/atom_vec_bond_kokkos.cpp b/src/KOKKOS/atom_vec_bond_kokkos.cpp
index 9441373aa5521cd354c770c1086808fbbb66bc57..adad4f767bb1d2e9a55021c32a71c3bb0b450d24 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<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);
}
diff --git a/src/KOKKOS/atom_vec_charge_kokkos.cpp b/src/KOKKOS/atom_vec_charge_kokkos.cpp
index a9ae5cc2d1e999d8dc94964e43a3f0d8be4a2a84..3e50870c8f9db9a399355ef60724d6ac85f41ee2 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<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);
}
diff --git a/src/KOKKOS/atom_vec_dpd_kokkos.cpp b/src/KOKKOS/atom_vec_dpd_kokkos.cpp
index ca1779d0eeb7807e84d17461631b9bafeee7d0a5..06ca2083dce28372a6d9a05235ef224766420208 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<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);
}
diff --git a/src/KOKKOS/atom_vec_full_kokkos.cpp b/src/KOKKOS/atom_vec_full_kokkos.cpp
index 0560c99037761f38a39523826893422cde95cf64..6492a264ac153d330c2d1670b67c790c7b163aa6 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<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);
}
diff --git a/src/KOKKOS/atom_vec_kokkos.cpp b/src/KOKKOS/atom_vec_kokkos.cpp
index 885f190f84bc6542dd222ac6c8f8adbe1d46e044..f2c04bec1b6281b3922513628b2a655e44c9baf7 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<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);
};
diff --git a/src/KOKKOS/atom_vec_kokkos.h b/src/KOKKOS/atom_vec_kokkos.h
index 20a07ec443690451a963ce14c51ed0f71873ae93..b68fdc9b209dcefa6e8ee36c21053a396bc73396 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<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);
diff --git a/src/KOKKOS/atom_vec_molecular_kokkos.cpp b/src/KOKKOS/atom_vec_molecular_kokkos.cpp
index 380aa0fb7ec31162633be6028b1c09b3ecbb4447..1a0c03cca5631ce58a86321757d5610323319336 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<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);
}
diff --git a/src/KOKKOS/bond_class2_kokkos.cpp b/src/KOKKOS/bond_class2_kokkos.cpp
index df2f2c1e9be24a6d807ecf12e67c239735dde7f6..80f2a6dfbbe8b6e423d0b5fd7174f53f7eeb2da1 100644
--- a/src/KOKKOS/bond_class2_kokkos.cpp
+++ b/src/KOKKOS/bond_class2_kokkos.cpp
@@ -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>();
diff --git a/src/KOKKOS/bond_harmonic_kokkos.cpp b/src/KOKKOS/bond_harmonic_kokkos.cpp
index c4e0c3a81763e765ee593ba13b84c4b0fcec3616..e277104ddb317217dfd5dcddeef9bac584b4dd6f 100644
--- a/src/KOKKOS/bond_harmonic_kokkos.cpp
+++ b/src/KOKKOS/bond_harmonic_kokkos.cpp
@@ -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>();
diff --git a/src/KOKKOS/comm_kokkos.cpp b/src/KOKKOS/comm_kokkos.cpp
index b41849197be85d30188cbce317fadea4608008f4..d24fa8a536d56cbb550d5b7e547b32aef57bef0b 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<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>();
diff --git a/src/KOKKOS/dihedral_charmm_kokkos.cpp b/src/KOKKOS/dihedral_charmm_kokkos.cpp
index 71635ec76c48d5642e7eccccef8d360c2077fc60..13e3a6d2d1c600d62fd6cbeb56efa0f88b5553c1 100644
--- a/src/KOKKOS/dihedral_charmm_kokkos.cpp
+++ b/src/KOKKOS/dihedral_charmm_kokkos.cpp
@@ -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>();
diff --git a/src/KOKKOS/fix_eos_table_rx_kokkos.cpp b/src/KOKKOS/fix_eos_table_rx_kokkos.cpp
index 5c106c19f3432edcc5a432d96031b12270bb39bb..fb4ca443a9dac49a9566733a0db474e2022b2fc7 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<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];
}
diff --git a/src/KOKKOS/fix_property_atom_kokkos.cpp b/src/KOKKOS/fix_property_atom_kokkos.cpp
index 9e29b6c35dca9d1ea5d27c00056b48220e497eff..fb3f5a3f341fbcf4f51ab555c1686a0f7de2d147 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 0319c79cb35b8a1ecfddae39ecb7d7fec68f3258..22450e19d948ac53a64837c021bf6c32c879968a 100644
--- a/src/KOKKOS/fix_rx_kokkos.cpp
+++ b/src/KOKKOS/fix_rx_kokkos.cpp
@@ -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>();
diff --git a/src/KOKKOS/fix_shardlow_kokkos.cpp b/src/KOKKOS/fix_shardlow_kokkos.cpp
index 571f4880230caefaaed8c2561f16f0681807ac54..e973e5c99f5d4035207689c297c0ba1bb4f92cd3 100644
--- a/src/KOKKOS/fix_shardlow_kokkos.cpp
+++ b/src/KOKKOS/fix_shardlow_kokkos.cpp
@@ -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;
diff --git a/src/KOKKOS/gridcomm_kokkos.cpp b/src/KOKKOS/gridcomm_kokkos.cpp
index f107370514001d22c2f054f52b2c98ddd312ded6..847fa5907a9f05cae3f8e4777097bc92029ebd09 100644
--- a/src/KOKKOS/gridcomm_kokkos.cpp
+++ b/src/KOKKOS/gridcomm_kokkos.cpp
@@ -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);
diff --git a/src/KOKKOS/improper_class2_kokkos.cpp b/src/KOKKOS/improper_class2_kokkos.cpp
index 2171305a9ce0599858dfcbc502de9700061aa6bb..28f57608bbce990793b47b5f66f95d28dbec6ec5 100644
--- a/src/KOKKOS/improper_class2_kokkos.cpp
+++ b/src/KOKKOS/improper_class2_kokkos.cpp
@@ -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>();
diff --git a/src/KOKKOS/improper_harmonic_kokkos.cpp b/src/KOKKOS/improper_harmonic_kokkos.cpp
index 49dd36ed194da6ea94df99144f75602ed11af5ea..f1a0ac3631c1a862dab0172b4f58dadcfeff8038 100644
--- a/src/KOKKOS/improper_harmonic_kokkos.cpp
+++ b/src/KOKKOS/improper_harmonic_kokkos.cpp
@@ -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>();
diff --git a/src/KOKKOS/kokkos_type.h b/src/KOKKOS/kokkos_type.h
index da06be98b6775056fe3e8e744ee933e75e2d0c2e..ddea35ca8892244c053f232af7fda8f6c73ca82c 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<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
diff --git a/src/KOKKOS/memory_kokkos.h b/src/KOKKOS/memory_kokkos.h
index 9f930faae2e4f973a0849064a585888e29a8483f..39d30ac78540c66724d8a00f0a1a9eff9687416a 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 <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);
diff --git a/src/KOKKOS/nbin_kokkos.cpp b/src/KOKKOS/nbin_kokkos.cpp
index 95ea105ad9df6d95f0d2ddba64154f1af64d0579..2ca2a83b5ae5a7284ec5c65af956fdcf9fc131c3 100644
--- a/src/KOKKOS/nbin_kokkos.cpp
+++ b/src/KOKKOS/nbin_kokkos.cpp
@@ -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;
diff --git a/src/KOKKOS/nbin_ssa_kokkos.cpp b/src/KOKKOS/nbin_ssa_kokkos.cpp
index 1a881658567524b181bbeb52af3b97a17246fc12..d040ee2c83f55194402830e76173d600474b275d 100644
--- a/src/KOKKOS/nbin_ssa_kokkos.cpp
+++ b/src/KOKKOS/nbin_ssa_kokkos.cpp
@@ -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;
diff --git a/src/KOKKOS/neighbor_kokkos.cpp b/src/KOKKOS/neighbor_kokkos.cpp
index ea90a182709160fc80e50a53b364463eb6f4783f..72092aef0b17bfffaea8865c9156f010f9f829c0 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<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);
diff --git a/src/KOKKOS/npair_kokkos.cpp b/src/KOKKOS/npair_kokkos.cpp
index d3cdcb0680efef3420d0e1a44f966a265759defc..4411627a784f01329cb35da5afbb9457597d45e1 100644
--- a/src/KOKKOS/npair_kokkos.cpp
+++ b/src/KOKKOS/npair_kokkos.cpp
@@ -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;
diff --git a/src/KOKKOS/npair_ssa_kokkos.cpp b/src/KOKKOS/npair_ssa_kokkos.cpp
index 9f447bda1a8038874064a54035e0c9de569c0a0c..ba941b8e547d7f524de64025ad34803cf9024a05 100644
--- a/src/KOKKOS/npair_ssa_kokkos.cpp
+++ b/src/KOKKOS/npair_ssa_kokkos.cpp
@@ -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;
}
}
diff --git a/src/KOKKOS/pair_exp6_rx_kokkos.cpp b/src/KOKKOS/pair_exp6_rx_kokkos.cpp
index 5f8e724a60a9a6828a0bdfbab9e7cfcaedbff4be..65b36635a34f88f1eb1ca6eaf6abfae2c1681086 100644
--- a/src/KOKKOS/pair_exp6_rx_kokkos.cpp
+++ b/src/KOKKOS/pair_exp6_rx_kokkos.cpp
@@ -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);
diff --git a/src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp b/src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp
index 0961cf44eb09ea1211ab7898c9c6532cbd518a14..7cb869a8a558bc524d172bf31b529554cfd0866d 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<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];
}
diff --git a/src/KOKKOS/pair_sw_kokkos.cpp b/src/KOKKOS/pair_sw_kokkos.cpp
index cc9307c42edbef3451b4bc8ad7cc0b46966879df..a19e2fc271d0292e6e6c2d9d487c29a82c886135 100644
--- a/src/KOKKOS/pair_sw_kokkos.cpp
+++ b/src/KOKKOS/pair_sw_kokkos.cpp
@@ -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);
diff --git a/src/KOKKOS/pair_table_kokkos.cpp b/src/KOKKOS/pair_table_kokkos.cpp
index 957214915bf403ac790b032d9179bc55e10161c3..e3e4ee01f397c750636fde67868fc7bf17f9555d 100644
--- a/src/KOKKOS/pair_table_kokkos.cpp
+++ b/src/KOKKOS/pair_table_kokkos.cpp
@@ -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];
}
diff --git a/src/KOKKOS/pair_table_rx_kokkos.cpp b/src/KOKKOS/pair_table_rx_kokkos.cpp
index 7bc5198d8cef6e8bbfcf40ba24e4d0bb575b9cec..0b1e0ee9e6f6adbaaef11d5011287269d8f4ed64 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<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];
}
diff --git a/src/KOKKOS/pair_tersoff_kokkos.cpp b/src/KOKKOS/pair_tersoff_kokkos.cpp
index b64b2f6d439a68fa7a76162efa281334ec8ab4bd..1fb97f913f8d560d7c9c3605fe577d08efffe711 100644
--- a/src/KOKKOS/pair_tersoff_kokkos.cpp
+++ b/src/KOKKOS/pair_tersoff_kokkos.cpp
@@ -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);
diff --git a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp
index 21c7fbb7278ba0422167ba866925e3b27ebec4ba..a30e1229f2a6959d35d462ff006e8ceabb07747b 100644
--- a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp
+++ b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp
@@ -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);
diff --git a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp
index 8a34dc34abfef481e9ccbb4c07e84fe7f21ad226..aeefbf67f688caf403ca4e6a10d9a49ffc3e08cf 100644
--- a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp
+++ b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp
@@ -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);
diff --git a/src/KOKKOS/pair_vashishta_kokkos.cpp b/src/KOKKOS/pair_vashishta_kokkos.cpp
index 5611aa132bb825e970304a323e4798ab093fc095..2e4989356eab8b296f8ff87e882907101e82ae91 100644
--- a/src/KOKKOS/pair_vashishta_kokkos.cpp
+++ b/src/KOKKOS/pair_vashishta_kokkos.cpp
@@ -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);
}
diff --git a/src/KOKKOS/verlet_kokkos.cpp b/src/KOKKOS/verlet_kokkos.cpp
index b2b107284ec26109843db5092bf1dc9594acee8e..fe11e6b3fa4986459ad9847ec3f41b6d72ccd72e 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<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>();