Remove deprecated Kokkos code

This commit is contained in:
Stan Moore 2020-03-25 13:55:31 -06:00
parent 2ac79d4483
commit 299f79c919
8 changed files with 57 additions and 48 deletions

View File

@ -205,7 +205,7 @@ void CommKokkos::forward_comm_device(int dummy)
}
n = avec->pack_comm_kokkos(sendnum[iswap],k_sendlist,
iswap,k_buf_send,pbc_flag[iswap],pbc[iswap]);
DeviceType::fence();
DeviceType().fence();
if (n) {
MPI_Send(k_buf_send.view<DeviceType>().data(),
n,MPI_DOUBLE,sendproc[iswap],0,world);
@ -224,14 +224,14 @@ void CommKokkos::forward_comm_device(int dummy)
}
n = avec->pack_comm_vel_kokkos(sendnum[iswap],k_sendlist,iswap,
k_buf_send,pbc_flag[iswap],pbc[iswap]);
DeviceType::fence();
DeviceType().fence();
if (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_vel_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv);
DeviceType::fence();
DeviceType().fence();
} else {
if (size_forward_recv[iswap])
MPI_Irecv(k_buf_recv.view<DeviceType>().data(),
@ -239,26 +239,26 @@ void CommKokkos::forward_comm_device(int dummy)
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();
DeviceType().fence();
if (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);
DeviceType::fence();
DeviceType().fence();
}
} else {
if (!ghost_velocity) {
if (sendnum[iswap])
n = avec->pack_comm_self(sendnum[iswap],k_sendlist,iswap,
firstrecv[iswap],pbc_flag[iswap],pbc[iswap]);
DeviceType::fence();
DeviceType().fence();
} else {
n = avec->pack_comm_vel_kokkos(sendnum[iswap],k_sendlist,iswap,
k_buf_send,pbc_flag[iswap],pbc[iswap]);
DeviceType::fence();
DeviceType().fence();
avec->unpack_comm_vel_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_send);
DeviceType::fence();
DeviceType().fence();
}
}
}
@ -334,7 +334,7 @@ void CommKokkos::reverse_comm_device()
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();
DeviceType().fence();
if (n)
MPI_Send(k_buf_send.view<DeviceType>().data(),n,
MPI_DOUBLE,recvproc[iswap],0,world);
@ -342,7 +342,7 @@ void CommKokkos::reverse_comm_device()
}
avec->unpack_reverse_kokkos(sendnum[iswap],k_sendlist,iswap,
k_buf_recv);
DeviceType::fence();
DeviceType().fence();
} else {
if (sendnum[iswap])
n = avec->unpack_reverse_self(sendnum[iswap],k_sendlist,iswap,
@ -410,7 +410,7 @@ void CommKokkos::forward_comm_pair_device(Pair *pair)
n = pairKKBase->pack_forward_comm_kokkos(sendnum[iswap],k_sendlist,
iswap,k_buf_send_pair,pbc_flag[iswap],pbc[iswap]);
DeviceType::fence();
DeviceType().fence();
// exchange with another proc
// if self, set recv buffer to send buffer
@ -445,7 +445,7 @@ void CommKokkos::forward_comm_pair_device(Pair *pair)
// unpack buffer
pairKKBase->unpack_forward_comm_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv_pair);
DeviceType::fence();
DeviceType().fence();
}
}
@ -647,7 +647,7 @@ void CommKokkos::exchange_device()
k_exchange_sendlist,k_exchange_copylist,
ExecutionSpaceFromDevice<DeviceType>::space,
dim,lo,hi);
DeviceType::fence();
DeviceType().fence();
} else {
while (i < nlocal) {
if (x[i][dim] < lo || x[i][dim] >= hi) {
@ -671,7 +671,7 @@ void CommKokkos::exchange_device()
atom->nlocal=avec->
unpack_exchange_kokkos(k_buf_send,nrecv,atom->nlocal,dim,lo,hi,
ExecutionSpaceFromDevice<DeviceType>::space);
DeviceType::fence();
DeviceType().fence();
}
} else {
MPI_Sendrecv(&nsend,1,MPI_INT,procneigh[dim][0],0,
@ -704,7 +704,7 @@ void CommKokkos::exchange_device()
atom->nlocal = avec->
unpack_exchange_kokkos(k_buf_recv,nrecv,atom->nlocal,dim,lo,hi,
ExecutionSpaceFromDevice<DeviceType>::space);
DeviceType::fence();
DeviceType().fence();
}
}
@ -964,13 +964,13 @@ void CommKokkos::borders_device() {
n = avec->
pack_border_vel_kokkos(nsend,k_sendlist,k_buf_send,iswap,
pbc_flag[iswap],pbc[iswap],exec_space);
DeviceType::fence();
DeviceType().fence();
}
else {
n = avec->
pack_border_kokkos(nsend,k_sendlist,k_buf_send,iswap,
pbc_flag[iswap],pbc[iswap],exec_space);
DeviceType::fence();
DeviceType().fence();
}
// swap atoms with other proc
@ -1000,21 +1000,21 @@ void CommKokkos::borders_device() {
if (sendproc[iswap] != me) {
avec->unpack_border_vel_kokkos(nrecv,atom->nlocal+atom->nghost,
k_buf_recv,exec_space);
DeviceType::fence();
DeviceType().fence();
} else {
avec->unpack_border_vel_kokkos(nrecv,atom->nlocal+atom->nghost,
k_buf_send,exec_space);
DeviceType::fence();
DeviceType().fence();
}
} else {
if (sendproc[iswap] != me) {
avec->unpack_border_kokkos(nrecv,atom->nlocal+atom->nghost,
k_buf_recv,exec_space);
DeviceType::fence();
DeviceType().fence();
} else {
avec->unpack_border_kokkos(nrecv,atom->nlocal+atom->nghost,
k_buf_send,exec_space);
DeviceType::fence();
DeviceType().fence();
}
}
// set all pointers & counters

View File

@ -230,7 +230,7 @@ void FFT3dKokkos<DeviceType>::fft_3d_kokkos(typename FFT_AT::t_FFT_DATA_1d d_in,
cufftExec(plan->plan_fast,d_data.data(),d_data.data(),flag);
#else
typename FFT_AT::t_FFT_DATA_1d d_tmp =
typename FFT_AT::t_FFT_DATA_1d(Kokkos::view_alloc("fft_3d:tmp",Kokkos::WithoutInitializing),d_in.dimension_0());
typename FFT_AT::t_FFT_DATA_1d(Kokkos::view_alloc("fft_3d:tmp",Kokkos::WithoutInitializing),d_in.extent(0));
kiss_fft_functor<DeviceType> f;
if (flag == -1)
f = kiss_fft_functor<DeviceType>(d_data,d_tmp,plan->cfg_fast_forward,length);
@ -238,7 +238,7 @@ void FFT3dKokkos<DeviceType>::fft_3d_kokkos(typename FFT_AT::t_FFT_DATA_1d d_in,
f = kiss_fft_functor<DeviceType>(d_data,d_tmp,plan->cfg_fast_backward,length);
Kokkos::parallel_for(total/length,f);
d_data = d_tmp;
d_tmp = typename FFT_AT::t_FFT_DATA_1d(Kokkos::view_alloc("fft_3d:tmp",Kokkos::WithoutInitializing),d_in.dimension_0());
d_tmp = typename FFT_AT::t_FFT_DATA_1d(Kokkos::view_alloc("fft_3d:tmp",Kokkos::WithoutInitializing),d_in.extent(0));
#endif
@ -281,7 +281,7 @@ void FFT3dKokkos<DeviceType>::fft_3d_kokkos(typename FFT_AT::t_FFT_DATA_1d d_in,
f = kiss_fft_functor<DeviceType>(d_data,d_tmp,plan->cfg_mid_backward,length);
Kokkos::parallel_for(total/length,f);
d_data = d_tmp;
d_tmp = typename FFT_AT::t_FFT_DATA_1d(Kokkos::view_alloc("fft_3d:tmp",Kokkos::WithoutInitializing),d_in.dimension_0());
d_tmp = typename FFT_AT::t_FFT_DATA_1d(Kokkos::view_alloc("fft_3d:tmp",Kokkos::WithoutInitializing),d_in.extent(0));
#endif
// 2nd mid-remap to prepare for 3rd FFTs
@ -864,7 +864,7 @@ void FFT3dKokkos<DeviceType>::fft_3d_1d_only_kokkos(typename FFT_AT::t_FFT_DATA_
cufftExec(plan->plan_slow,d_data.data(),d_data.data(),flag);
#else
kiss_fft_functor<DeviceType> f;
typename FFT_AT::t_FFT_DATA_1d d_tmp = typename FFT_AT::t_FFT_DATA_1d("fft_3d:tmp",d_data.dimension_0());
typename FFT_AT::t_FFT_DATA_1d d_tmp = typename FFT_AT::t_FFT_DATA_1d("fft_3d:tmp",d_data.extent(0));
if (flag == -1) {
f = kiss_fft_functor<DeviceType>(d_data,d_tmp,plan->cfg_fast_forward,length1);
Kokkos::parallel_for(total1/length1,f);

View File

@ -524,7 +524,7 @@ void GridCommKokkos<DeviceType>::forward_comm(KSpace *kspace, int which)
kspaceKKBase->pack_forward_kspace_kokkos(which,k_buf2,swap[m].npack,k_packlist,m);
else
kspaceKKBase->pack_forward_kspace_kokkos(which,k_buf1,swap[m].npack,k_packlist,m);
DeviceType::fence();
DeviceType().fence();
if (swap[m].sendproc != me) {
FFT_SCALAR* buf1;
@ -552,7 +552,7 @@ void GridCommKokkos<DeviceType>::forward_comm(KSpace *kspace, int which)
}
kspaceKKBase->unpack_forward_kspace_kokkos(which,k_buf2,swap[m].nunpack,k_unpacklist,m);
DeviceType::fence();
DeviceType().fence();
}
}
@ -574,7 +574,7 @@ void GridCommKokkos<DeviceType>::reverse_comm(KSpace *kspace, int which)
kspaceKKBase->pack_reverse_kspace_kokkos(which,k_buf2,swap[m].nunpack,k_unpacklist,m);
else
kspaceKKBase->pack_reverse_kspace_kokkos(which,k_buf1,swap[m].nunpack,k_unpacklist,m);
DeviceType::fence();
DeviceType().fence();
if (swap[m].recvproc != me) {
FFT_SCALAR* buf1;
@ -602,7 +602,7 @@ void GridCommKokkos<DeviceType>::reverse_comm(KSpace *kspace, int which)
}
kspaceKKBase->unpack_reverse_kspace_kokkos(which,k_buf2,swap[m].npack,k_packlist,m);
DeviceType::fence();
DeviceType().fence();
}
}

View File

@ -1012,7 +1012,7 @@ void memset_kokkos (ViewType &view) {
#else
Kokkos::parallel_for(view.span()*sizeof(typename ViewType::value_type)/4, f);
#endif
ViewType::execution_space::fence();
ViewType::execution_space().fence();
}
struct params_lj_coul {

View File

@ -86,7 +86,7 @@ static void pack_3d(typename FFT_AT::t_FFT_SCALAR_1d_um d_data, int data_offset,
const int nfast = plan->nfast;
pack_3d_functor f(d_buf,buf_offset,d_data,data_offset,plan);
Kokkos::parallel_for(nslow*nmid*nfast,f);
DeviceType::fence();
DeviceType().fence();
}
/* ----------------------------------------------------------------------
@ -140,7 +140,7 @@ static void unpack_3d(typename FFT_AT::t_FFT_SCALAR_1d_um d_buf, int buf_offset,
const int nfast = plan->nfast;
unpack_3d_functor f(d_buf,buf_offset,d_data,data_offset,plan);
Kokkos::parallel_for(nslow*nmid*nfast,f);
DeviceType::fence();
DeviceType().fence();
}
/* ----------------------------------------------------------------------
@ -195,7 +195,7 @@ static void unpack_3d_permute1_1(typename FFT_AT::t_FFT_SCALAR_1d_um d_buf, int
const int nfast = plan->nfast;
unpack_3d_permute1_1_functor f(d_buf,buf_offset,d_data,data_offset,plan);
Kokkos::parallel_for(nslow*nmid*nfast,f);
DeviceType::fence();
DeviceType().fence();
}
/* ----------------------------------------------------------------------
unpack from buf -> data, one axis permutation, 2 values/element
@ -249,7 +249,7 @@ static void unpack_3d_permute1_2(typename FFT_AT::t_FFT_SCALAR_1d_um d_buf, int
const int nfast = plan->nfast;
unpack_3d_permute1_2_functor f(d_buf,buf_offset,d_data,data_offset,plan);
Kokkos::parallel_for(nslow*nmid*nfast,f);
DeviceType::fence();
DeviceType().fence();
}
/* ----------------------------------------------------------------------
@ -305,7 +305,7 @@ static void unpack_3d_permute1_n(typename FFT_AT::t_FFT_SCALAR_1d_um d_buf, int
const int nfast = plan->nfast;
unpack_3d_permute1_n_functor f(d_buf,buf_offset,d_data,data_offset,plan);
Kokkos::parallel_for(nslow*nmid*nfast,f);
DeviceType::fence();
DeviceType().fence();
}
/* ----------------------------------------------------------------------
@ -358,7 +358,7 @@ static void unpack_3d_permute2_1(typename FFT_AT::t_FFT_SCALAR_1d_um d_buf, int
const int nfast = plan->nfast;
unpack_3d_permute2_1_functor f(d_buf,buf_offset,d_data,data_offset,plan);
Kokkos::parallel_for(nslow*nmid*nfast,f);
DeviceType::fence();
DeviceType().fence();
}
/* ----------------------------------------------------------------------
@ -412,7 +412,7 @@ static void unpack_3d_permute2_2(typename FFT_AT::t_FFT_SCALAR_1d_um d_buf, int
const int nfast = plan->nfast;
unpack_3d_permute2_2_functor f(d_buf,buf_offset,d_data,data_offset,plan);
Kokkos::parallel_for(nslow*nmid*nfast,f);
DeviceType::fence();
DeviceType().fence();
}
/* ----------------------------------------------------------------------
unpack from buf -> data, two axis permutation, nqty values/element
@ -466,7 +466,7 @@ static void unpack_3d_permute2_n(typename FFT_AT::t_FFT_SCALAR_1d_um d_buf, int
const int nfast = plan->nfast;
unpack_3d_permute2_n_functor f(d_buf,buf_offset,d_data,data_offset,plan);
Kokkos::parallel_for(nslow*nmid*nfast,f);
DeviceType::fence();
DeviceType().fence();
}
};

View File

@ -444,7 +444,7 @@ struct PairComputeFunctor {
ev.evdwl += fev.evdwl;
if (c.eflag_atom)
d_eatom(i,0) += fev.evdwl;
d_eatom(i) += fev.evdwl;
if (c.vflag_global) {
ev.v[0] += fev.v[0];
@ -554,7 +554,7 @@ struct PairComputeFunctor {
}
if (c.eflag_atom)
d_eatom(i,0) += fev.evdwl + fev.ecoul;
d_eatom(i) += fev.evdwl + fev.ecoul;
if (c.vflag_global) {
ev.v[0] += fev.v[0];
@ -850,8 +850,14 @@ EV_FLOAT pair_compute_neighlist (PairStyle* fpair, typename Kokkos::Impl::enable
}
template<class FunctorStyle>
int GetTeamSize(FunctorStyle& functor, int team_size, int vector_length) {
int team_size_max = Kokkos::TeamPolicy<>::team_size_max(functor);
int GetTeamSize(FunctorStyle& functor, int inum, int reduce_flag, int team_size, int vector_length) {
int team_size_max;
if (reduce_flag) {
EV_FLOAT ev;
team_size_max = Kokkos::TeamPolicy<>(inum,Kokkos::AUTO).team_size_max(functor,ev,Kokkos::ParallelReduceTag());
} else {
team_size_max = Kokkos::TeamPolicy<>(inum,Kokkos::AUTO).team_size_max(functor,Kokkos::ParallelForTag());
}
#ifdef KOKKOS_ENABLE_CUDA
if(team_size*vector_length > team_size_max)
@ -877,13 +883,13 @@ EV_FLOAT pair_compute_neighlist (PairStyle* fpair, typename Kokkos::Impl::enable
if(fpair->atom->ntypes > MAX_TYPES_STACKPARAMS) {
PairComputeFunctor<PairStyle,NEIGHFLAG,false,Specialisation > ff(fpair,list);
atoms_per_team = GetTeamSize(ff, atoms_per_team, vector_length);
atoms_per_team = GetTeamSize(ff, list->inum, (fpair->eflag || fpair->vflag), atoms_per_team, vector_length);
Kokkos::TeamPolicy<Kokkos::IndexType<int> > policy(list->inum,atoms_per_team,vector_length);
if (fpair->eflag || fpair->vflag) Kokkos::parallel_reduce(policy,ff,ev);
else Kokkos::parallel_for(policy,ff);
} else {
PairComputeFunctor<PairStyle,NEIGHFLAG,true,Specialisation > ff(fpair,list);
atoms_per_team = GetTeamSize(ff, atoms_per_team, vector_length);
atoms_per_team = GetTeamSize(ff, list->inum, (fpair->eflag || fpair->vflag), atoms_per_team, vector_length);
Kokkos::TeamPolicy<Kokkos::IndexType<int> > policy(list->inum,atoms_per_team,vector_length);
if (fpair->eflag || fpair->vflag) Kokkos::parallel_reduce(policy,ff,ev);
else Kokkos::parallel_for(policy,ff);

View File

@ -182,11 +182,14 @@ void PairSNAPKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
if (max_neighs<num_neighs) max_neighs = num_neighs;
}*/
max_neighs = 0;
Kokkos::parallel_reduce("PairSNAPKokkos::find_max_neighs",inum, FindMaxNumNeighs<DeviceType>(k_list), Kokkos::Experimental::Max<int>(max_neighs));
Kokkos::parallel_reduce("PairSNAPKokkos::find_max_neighs",inum, FindMaxNumNeighs<DeviceType>(k_list), Kokkos::Max<int>(max_neighs));
int chunk_size = MIN(2000,inum);
chunk_offset = 0;
int vector_length = 1;
int team_size = 1;
int team_size_max = Kokkos::TeamPolicy<DeviceType>::team_size_max(*this);
int team_size_max = Kokkos::TeamPolicy<DeviceType>(chunk_size,Kokkos::AUTO).team_size_max(*this,Kokkos::ParallelForTag());
#ifdef KOKKOS_ENABLE_CUDA
team_size = 32;//max_neighs;
if (team_size*vector_length > team_size_max)

View File

@ -120,7 +120,7 @@ void RemapKokkos<DeviceType>::remap_3d_kokkos(typename FFT_AT::t_FFT_SCALAR_1d d
// post all recvs into scratch space
for (irecv = 0; irecv < plan->nrecv; irecv++) {
FFT_SCALAR* scratch = d_scratch.ptr_on_device() + plan->recv_bufloc[irecv];
FFT_SCALAR* scratch = d_scratch.data() + plan->recv_bufloc[irecv];
MPI_Irecv(scratch,plan->recv_size[irecv],
MPI_FFT_SCALAR,plan->recv_proc[irecv],0,
plan->comm,&plan->request[irecv]);
@ -132,7 +132,7 @@ void RemapKokkos<DeviceType>::remap_3d_kokkos(typename FFT_AT::t_FFT_SCALAR_1d d
int in_offset = plan->send_offset[isend];
plan->pack(d_in,in_offset,
plan->d_sendbuf,0,&plan->packplan[isend]);
MPI_Send(plan->d_sendbuf.ptr_on_device(),plan->send_size[isend],MPI_FFT_SCALAR,
MPI_Send(plan->d_sendbuf.data(),plan->send_size[isend],MPI_FFT_SCALAR,
plan->send_proc[isend],0,plan->comm);
}