Line data    Source code

       1             : #ifndef AMREX_PCI_H_
       2             : #define AMREX_PCI_H_
       3             : 
       4             : template <class FAB>
       5             : void
       6      539914 : FabArray<FAB>::PC_local_cpu (const CPC& thecpc, FabArray<FAB> const& src,
       7             :                              int scomp, int dcomp, int ncomp, CpOp op)
       8             : {
       9      539914 :     auto const N_locs = static_cast<int>(thecpc.m_LocTags->size());
      10      539914 :     if (N_locs == 0) { return; }
      11      539914 :     bool is_thread_safe = thecpc.m_threadsafe_loc;
      12             : 
      13      539914 :     if (is_thread_safe)
      14             :     {
      15             : #ifdef AMREX_USE_OMP
      16             : #pragma omp parallel for
      17             : #endif
      18     5460756 :         for (int i = 0; i < N_locs; ++i)
      19             :         {
      20     4920839 :             const CopyComTag& tag = (*thecpc.m_LocTags)[i];
      21     4920839 :             if (this != &src || tag.dstIndex != tag.srcIndex || tag.sbox != tag.dbox) {
      22             :                 // avoid self copy or plus
      23     4920839 :                 const FAB* sfab = &(src[tag.srcIndex]);
      24     4920839 :                       FAB* dfab = &(get(tag.dstIndex));
      25     4920839 :                 if (op == FabArrayBase::COPY)
      26             :                 {
      27     4920839 :                     dfab->template copy<RunOn::Host>(*sfab, tag.sbox, scomp, tag.dbox, dcomp, ncomp);
      28             :                 }
      29             :                 else
      30             :                 {
      31           0 :                     dfab->template plus<RunOn::Host>(*sfab, tag.sbox, tag.dbox, scomp, dcomp, ncomp);
      32             :                 }
      33             :             }
      34             :         }
      35             :     }
      36             :     else
      37             :     {
      38           0 :         LayoutData<Vector<FabCopyTag<FAB> > > loc_copy_tags(boxArray(),DistributionMap());
      39           0 :         for (int i = 0; i < N_locs; ++i)
      40             :         {
      41           0 :             const CopyComTag& tag = (*thecpc.m_LocTags)[i];
      42           0 :             if (this != &src || tag.dstIndex != tag.srcIndex || tag.sbox != tag.dbox) {
      43           0 :                 loc_copy_tags[tag.dstIndex].push_back
      44           0 :                     ({src.fabPtr(tag.srcIndex), tag.dbox, tag.sbox.smallEnd()-tag.dbox.smallEnd()});
      45             :             }
      46             :         }
      47             : 
      48             : #ifdef AMREX_USE_OMP
      49             : #pragma omp parallel
      50             : #endif
      51           0 :         for (MFIter mfi(*this); mfi.isValid(); ++mfi)
      52             :         {
      53           0 :             const auto& tags = loc_copy_tags[mfi];
      54           0 :             auto dfab = this->array(mfi);
      55           0 :             if (op == FabArrayBase::COPY)
      56             :             {
      57           0 :                 for (auto const & tag : tags)
      58             :                 {
      59           0 :                     auto const sfab = tag.sfab->array();
      60           0 :                     Dim3 offset = tag.offset.dim3();
      61           0 :                     amrex::LoopConcurrentOnCpu (tag.dbox, ncomp,
      62           0 :                     [=] (int i, int j, int k, int n) noexcept
      63             :                     {
      64           0 :                         dfab(i,j,k,dcomp+n) = sfab(i+offset.x,j+offset.y,k+offset.z,scomp+n);
      65             :                     });
      66             :                 }
      67             :             }
      68             :             else
      69             :             {
      70           0 :                 for (auto const & tag : tags)
      71             :                 {
      72           0 :                     auto const sfab = tag.sfab->array();
      73           0 :                     Dim3 offset = tag.offset.dim3();
      74           0 :                     amrex::LoopConcurrentOnCpu (tag.dbox, ncomp,
      75           0 :                     [=] (int i, int j, int k, int n) noexcept
      76             :                     {
      77           0 :                         dfab(i,j,k,dcomp+n) += sfab(i+offset.x,j+offset.y,k+offset.z,scomp+n);
      78             :                     });
      79             :                 }
      80             :             }
      81             :         }
      82             :     }
      83             : }
      84             : 
      85             : #ifdef AMREX_USE_GPU
      86             : template <class FAB>
      87             : void
      88             : FabArray<FAB>::PC_local_gpu (const CPC& thecpc, FabArray<FAB> const& src,
      89             :                              int scomp, int dcomp, int ncomp, CpOp op)
      90             : {
      91             :     int N_locs = thecpc.m_LocTags->size();
      92             :     if (N_locs == 0) { return; }
      93             :     bool is_thread_safe = thecpc.m_threadsafe_loc;
      94             : 
      95             :     using TagType = Array4CopyTag<value_type>;
      96             :     Vector<TagType> loc_copy_tags;
      97             :     loc_copy_tags.reserve(N_locs);
      98             : 
      99             :     Vector<BaseFab<int> > maskfabs;
     100             :     Vector<Array4Tag<int> > masks;
     101             :     if (!is_thread_safe)
     102             :     {
     103             :         if ((op == FabArrayBase::COPY && !amrex::IsStoreAtomic<value_type>::value) ||
     104             :             (op == FabArrayBase::ADD  && !amrex::HasAtomicAdd <value_type>::value))
     105             :         {
     106             :             maskfabs.resize(this->local_size());
     107             :             masks.reserve(N_locs);
     108             :         }
     109             :     }
     110             : 
     111             :     for (int i = 0; i < N_locs; ++i)
     112             :     {
     113             :         const CopyComTag& tag = (*thecpc.m_LocTags)[i];
     114             :         if (this != &src || tag.dstIndex != tag.srcIndex || tag.sbox != tag.dbox) {
     115             :             int li = this->localindex(tag.dstIndex);
     116             :             loc_copy_tags.push_back
     117             :                 ({this->atLocalIdx(li).array(),
     118             :                   src.fabPtr(tag.srcIndex)->const_array(),
     119             :                   tag.dbox,
     120             :                   (tag.sbox.smallEnd()-tag.dbox.smallEnd()).dim3()});
     121             : 
     122             :             if (maskfabs.size() > 0) {
     123             :                 if (!maskfabs[li].isAllocated()) {
     124             :                     maskfabs[li].resize(this->atLocalIdx(li).box());
     125             :                 }
     126             :                 masks.emplace_back(Array4Tag<int>{maskfabs[li].array()});
     127             :             }
     128             :         }
     129             :     }
     130             : 
     131             :     if (maskfabs.size() > 0) {
     132             :         amrex::ParallelFor(masks,
     133             :         [=] AMREX_GPU_DEVICE (int i, int j, int k, Array4Tag<int> const& msk) noexcept
     134             :         {
     135             :             msk.dfab(i,j,k) = 0;
     136             :         });
     137             :     }
     138             : 
     139             :     if (op == FabArrayBase::COPY)
     140             :     {
     141             :         if (is_thread_safe) {
     142             :             detail::fab_to_fab<value_type, value_type>(loc_copy_tags, scomp,
     143             :                 dcomp, ncomp, detail::CellStore<value_type, value_type>());
     144             :         } else {
     145             :             detail::fab_to_fab_atomic_cpy<value_type, value_type>(
     146             :                 loc_copy_tags, scomp, dcomp, ncomp, masks);
     147             :         }
     148             :     }
     149             :     else
     150             :     {
     151             :         if (is_thread_safe) {
     152             :             detail::fab_to_fab<value_type, value_type>(loc_copy_tags, scomp,
     153             :                 dcomp, ncomp, detail::CellAdd<value_type, value_type>());
     154             :         } else {
     155             :             detail::fab_to_fab_atomic_add<value_type, value_type>(
     156             :                 loc_copy_tags, scomp, dcomp, ncomp, masks);
     157             :         }
     158             :     }
     159             : }
     160             : #endif
     161             : 
     162             : #endif