QMCPACK
SoaDistanceTableAAOMPTarget< T, D, SC > Struct Template Reference

A derived classe from DistacneTableData, specialized for dense case. More...

+ Inheritance diagram for SoaDistanceTableAAOMPTarget< T, D, SC >:
+ Collaboration diagram for SoaDistanceTableAAOMPTarget< T, D, SC >:

Classes

struct  DTAAMultiWalkerMem
 multi walker shared memory buffer More...
 

Public Member Functions

 SoaDistanceTableAAOMPTarget (ParticleSet &target)
 
 SoaDistanceTableAAOMPTarget ()=delete
 
 SoaDistanceTableAAOMPTarget (const SoaDistanceTableAAOMPTarget &)=delete
 
 ~SoaDistanceTableAAOMPTarget ()
 
size_t compute_size (int N) const
 
void resize ()
 
const RealTypegetMultiWalkerTempDataPtr () const override
 return multi walker temporary pair distance table data pointer More...
 
void createResource (ResourceCollection &collection) const override
 initialize a shared resource and hand it to a collection More...
 
void acquireResource (ResourceCollection &collection, const RefVectorWithLeader< DistanceTable > &dt_list) const override
 acquire a shared resource from a collection More...
 
void releaseResource (ResourceCollection &collection, const RefVectorWithLeader< DistanceTable > &dt_list) const override
 return a shared resource to a collection More...
 
void evaluate (ParticleSet &P) override
 evaluate the full Distance Table More...
 
const RealTypemw_evalDistsInRange (const RefVectorWithLeader< DistanceTable > &dt_list, const RefVectorWithLeader< ParticleSet > &p_list, size_t range_begin, size_t range_end) const override
 compute distances from particles in [range_begin, range_end) to all the particles. More...
 
void move (const ParticleSet &P, const PosType &rnew, const IndexType iat, bool prepare_old) override
 evaluate the temporary pair relations More...
 
void mw_move (const RefVectorWithLeader< DistanceTable > &dt_list, const RefVectorWithLeader< ParticleSet > &p_list, const std::vector< PosType > &rnew_list, const IndexType iat, bool prepare_old=true) const override
 evaluate the temporary pair relations when a move is proposed this implementation is asynchronous and the synchronization is managed at ParticleSet. More...
 
int get_first_neighbor (IndexType iat, RealType &r, PosType &dr, bool newpos) const override
 
void update (IndexType iat) override
 After accepting the iat-th particle, update the iat-th row of distances_ and displacements_. More...
 
void updatePartial (IndexType jat, bool from_temp) override
 fill partially the distance table by the pair relations from the temporary or old particle position. More...
 
void mw_updatePartial (const RefVectorWithLeader< DistanceTable > &dt_list, IndexType jat, const std::vector< bool > &from_temp) override
 walker batched version of updatePartial. More...
 
void mw_finalizePbyP (const RefVectorWithLeader< DistanceTable > &dt_list, const RefVectorWithLeader< ParticleSet > &p_list) const override
 walker batched version of finalizePbyP If not DTModes::NEED_TEMP_DATA_ON_HOST, host distance table data is not updated at all during p-by-p Thus, a recompute is necessary to update the whole host distance table for consumers like the Coulomb potential. More...
 
size_t get_num_particls_stored () const override
 
- Public Member Functions inherited from DTD_BConds< T, D, SC >
 DTD_BConds (const CrystalLattice< T, D > &lat)
 constructor: doing nothing More...
 
apply_bc (TinyVector< T, D > &displ) const
 apply BC on displ and return |displ|^2 More...
 
void apply_bc (std::vector< TinyVector< T, D >> &dr, std::vector< T > &r, std::vector< T > &rinv) const
 apply BC on dr and evaluate r and rinv More...
 
void apply_bc (std::vector< TinyVector< T, D >> &dr, std::vector< T > &r) const
 
void evaluate_rsquared (TinyVector< T, D > *restrict dr, T *restrict rr, int n)
 
- Public Member Functions inherited from DistanceTableAA
 DistanceTableAA (const ParticleSet &target, DTModes modes)
 constructor using source and target ParticleSet More...
 
const std::vector< DistRow > & getDistances () const
 return full table distances More...
 
const std::vector< DisplRow > & getDisplacements () const
 return full table displacements More...
 
const DistRowgetDistRow (int iel) const
 return a row of distances for a given target particle More...
 
const DisplRowgetDisplRow (int iel) const
 return a row of displacements for a given target particle More...
 
const DistRowgetTempDists () const
 return the temporary distances when a move is proposed More...
 
const DisplRowgetTempDispls () const
 return the temporary displacements when a move is proposed More...
 
const DistRowgetOldDists () const
 return old distances set up by move() for optimized distance table consumers More...
 
const DisplRowgetOldDispls () const
 return old displacements set up by move() for optimized distance table consumers More...
 
- Public Member Functions inherited from DistanceTable
 DistanceTable (const ParticleSet &source, const ParticleSet &target, DTModes modes)
 constructor using source and target ParticleSet More...
 
 DistanceTable (const DistanceTable &)=delete
 copy constructor. deleted More...
 
virtual ~DistanceTable ()=default
 virutal destructor More...
 
DTModes getModes () const
 get modes More...
 
void setModes (DTModes modes)
 set modes More...
 
const std::string & getName () const
 return the name of table More...
 
const ParticleSetget_origin () const
 returns the reference the origin particleset More...
 
size_t centers () const
 returns the number of centers More...
 
size_t targets () const
 returns the number of centers More...
 
size_t sources () const
 returns the number of source particles More...
 
virtual void mw_evaluate (const RefVectorWithLeader< DistanceTable > &dt_list, const RefVectorWithLeader< ParticleSet > &p_list) const
 
virtual void mw_recompute (const RefVectorWithLeader< DistanceTable > &dt_list, const RefVectorWithLeader< ParticleSet > &p_list, const std::vector< bool > &recompute) const
 recompute multi walker internal data, recompute More...
 
virtual void finalizePbyP (const ParticleSet &P)
 finalize distance table calculation after particle-by-particle moves if update() doesn't make the table up-to-date during p-by-p moves finalizePbyP takes action to bring the table up-to-date More...
 
virtual int get_first_neighbor (IndexType iat, RealType &r, PosType &dr, bool newpos) const =0
 find the first nearest neighbor More...
 
void print (std::ostream &os)
 

Public Attributes

aligned_vector< RealTypememory_pool_
 actual memory for dist and displacements_ More...
 
DistRow temp_r_mem_
 actual memory for temp_r_ More...
 
DisplRow temp_dr_mem_
 actual memory for temp_dr_ More...
 
DistRow old_r_mem_
 actual memory for old_r_ More...
 
DisplRow old_dr_mem_
 actual memory for old_dr_ More...
 
ResourceHandle< DTAAMultiWalkerMemmw_mem_handle_
 

Private Attributes

const size_t num_targets_padded_
 number of targets with padding More...
 
int old_prepared_elec_id_
 set to particle id after move() with prepare_old = true. More...
 
NewTimeroffload_timer_
 timer for offload portion More...
 
NewTimerevaluate_timer_
 timer for evaluate() More...
 
NewTimermove_timer_
 timer for move() More...
 
NewTimerupdate_timer_
 timer for update() More...
 
const size_t num_particls_stored = 64
 the particle count of the internal stored distances. More...
 

Additional Inherited Members

- Public Types inherited from DistanceTable
using IndexType = QMCTraits::IndexType
 
using RealType = QMCTraits::RealType
 
using PosType = QMCTraits::PosType
 
using DistRow = Vector< RealType, aligned_allocator< RealType > >
 
using DisplRow = VectorSoaContainer< RealType, DIM >
 
- Static Public Attributes inherited from DistanceTable
static constexpr unsigned DIM = OHMMS_DIM
 
- Protected Attributes inherited from DistanceTableAA
std::vector< DistRowdistances_
 distances_[num_targets_][num_sources_], [i][3][j] = |r_A2[j] - r_A1[i]| Note: Derived classes decide if it is a memory view or the actual storage For only the lower triangle (j<i) data can be accessed safely. More...
 
std::vector< DisplRowdisplacements_
 displacements_[num_targets_][3][num_sources_], [i][3][j] = r_A2[j] - r_A1[i] Note: Derived classes decide if it is a memory view or the actual storage only the lower triangle (j<i) is defined. More...
 
DistRow temp_r_
 temp_r More...
 
DisplRow temp_dr_
 temp_dr More...
 
DistRow old_r_
 old distances More...
 
DisplRow old_dr_
 old displacements More...
 
- Protected Attributes inherited from DistanceTable
const ParticleSetorigin_
 
const size_t num_sources_
 
const size_t num_targets_
 
const std::string name_
 name of the table More...
 
DTModes modes_
 operation modes defined by DTModes More...
 

Detailed Description

template<typename T, unsigned D, int SC>
struct qmcplusplus::SoaDistanceTableAAOMPTarget< T, D, SC >

A derived classe from DistacneTableData, specialized for dense case.

Definition at line 31 of file SoaDistanceTableAAOMPTarget.h.

Constructor & Destructor Documentation

◆ SoaDistanceTableAAOMPTarget() [1/3]

Definition at line 67 of file SoaDistanceTableAAOMPTarget.h.

References ParticleSet::getCoordinates(), and SoaDistanceTableAAOMPTarget< T, D, SC >::resize().

68  : DTD_BConds<T, D, SC>(target.getLattice()),
70  num_targets_padded_(getAlignedSize<T>(num_targets_)),
71 #if !defined(NDEBUG)
73 #endif
74  offload_timer_(createGlobalTimer(std::string("DTAAOMPTarget::offload_") + name_, timer_level_fine)),
75  evaluate_timer_(createGlobalTimer(std::string("DTAAOMPTarget::evaluate_") + name_, timer_level_fine)),
76  move_timer_(createGlobalTimer(std::string("DTAAOMPTarget::move_") + name_, timer_level_fine)),
77  update_timer_(createGlobalTimer(std::string("DTAAOMPTarget::update_") + name_, timer_level_fine))
78 
79  {
80  auto* coordinates_soa = dynamic_cast<const RealSpacePositionsOMPTarget*>(&target.getCoordinates());
81  if (!coordinates_soa)
82  throw std::runtime_error("Source particle set doesn't have OpenMP offload. Contact developers!");
83  resize();
84  PRAGMA_OFFLOAD("omp target enter data map(to : this[:1])")
85  }
const size_t num_targets_padded_
number of targets with padding
const std::string name_
name of the table
Definition: DistanceTable.h:57
int old_prepared_elec_id_
set to particle id after move() with prepare_old = true.
DistanceTableAA(const ParticleSet &target, DTModes modes)
constructor using source and target ParticleSet
NewTimer & createGlobalTimer(const std::string &myname, timer_levels mylevel)
NewTimer & offload_timer_
timer for offload portion

◆ SoaDistanceTableAAOMPTarget() [2/3]

◆ SoaDistanceTableAAOMPTarget() [3/3]

SoaDistanceTableAAOMPTarget ( const SoaDistanceTableAAOMPTarget< T, D, SC > &  )
delete

◆ ~SoaDistanceTableAAOMPTarget()

Definition at line 89 of file SoaDistanceTableAAOMPTarget.h.

89 {PRAGMA_OFFLOAD("omp target exit data map(delete : this[:1])")}

Member Function Documentation

◆ acquireResource()

void acquireResource ( ResourceCollection collection,
const RefVectorWithLeader< DistanceTable > &  dt_list 
) const
inlineoverridevirtual

acquire a shared resource from a collection

Reimplemented from DistanceTable.

Definition at line 127 of file SoaDistanceTableAAOMPTarget.h.

References Vector< T, Alloc >::attachReference(), Vector< T, Alloc >::free(), RefVectorWithLeader< T >::getCastedElement(), RefVectorWithLeader< T >::getCastedLeader(), RefVectorWithLeader< T >::getLeader(), ResourceCollection::lendResource(), SoaDistanceTableAAOMPTarget< T, D, SC >::mw_mem_handle_, DistanceTable::num_targets_, SoaDistanceTableAAOMPTarget< T, D, SC >::num_targets_padded_, and DistanceTableAA::temp_r_.

128  {
129  assert(this == &dt_list.getLeader());
130  auto& dt_leader = dt_list.getCastedLeader<SoaDistanceTableAAOMPTarget>();
131  dt_leader.mw_mem_handle_ = collection.lendResource<DTAAMultiWalkerMem>();
132  const size_t nw = dt_list.size();
133  const size_t stride_size = num_targets_padded_ * (D + 1);
134 
135  for (int iw = 0; iw < nw; iw++)
136  {
137  auto& dt = dt_list.getCastedElement<SoaDistanceTableAAOMPTarget>(iw);
138  dt.temp_r_.free();
139  dt.temp_dr_.free();
140  dt.old_r_.free();
141  dt.old_dr_.free();
142  }
143 
144  auto& mw_new_old_dist_displ = dt_leader.mw_mem_handle_.getResource().mw_new_old_dist_displ;
145  mw_new_old_dist_displ.resize(nw * 2 * stride_size);
146  for (int iw = 0; iw < nw; iw++)
147  {
148  auto& dt = dt_list.getCastedElement<SoaDistanceTableAAOMPTarget>(iw);
149  dt.temp_r_.attachReference(mw_new_old_dist_displ.data() + stride_size * iw, num_targets_padded_);
150  dt.temp_dr_.attachReference(num_targets_, num_targets_padded_,
151  mw_new_old_dist_displ.data() + stride_size * iw + num_targets_padded_);
152  dt.old_r_.attachReference(mw_new_old_dist_displ.data() + stride_size * (iw + nw), num_targets_padded_);
153  dt.old_dr_.attachReference(num_targets_, num_targets_padded_,
154  mw_new_old_dist_displ.data() + stride_size * (iw + nw) + num_targets_padded_);
155  }
156  }
const size_t num_targets_padded_
number of targets with padding

◆ compute_size()

size_t compute_size ( int  N) const
inline

Definition at line 91 of file SoaDistanceTableAAOMPTarget.h.

References qmcplusplus::Units::force::N.

Referenced by SoaDistanceTableAAOMPTarget< T, D, SC >::resize().

92  {
93  const size_t num_padded = getAlignedSize<T>(N);
94  const size_t Alignment = getAlignment<T>();
95  return (num_padded * (2 * N - num_padded + 1) + (Alignment - 1) * num_padded) / 2;
96  }

◆ createResource()

void createResource ( ResourceCollection collection) const
inlineoverridevirtual

initialize a shared resource and hand it to a collection

Reimplemented from DistanceTable.

Definition at line 122 of file SoaDistanceTableAAOMPTarget.h.

References ResourceCollection::addResource().

123  {
124  auto resource_index = collection.addResource(std::make_unique<DTAAMultiWalkerMem>());
125  }

◆ evaluate()

void evaluate ( ParticleSet P)
inlineoverridevirtual

evaluate the full Distance Table

Parameters
Pthe target particle set

Implements DistanceTable.

Definition at line 172 of file SoaDistanceTableAAOMPTarget.h.

References DistanceTableAA::displacements_, DistanceTableAA::distances_, SoaDistanceTableAAOMPTarget< T, D, SC >::evaluate_timer_, DynamicCoordinates::getAllParticlePos(), ParticleSet::getCoordinates(), DistanceTable::num_targets_, and ParticleSet::R.

173  {
174  ScopedTimer local_timer(evaluate_timer_);
175 
176  constexpr T BigR = std::numeric_limits<T>::max();
177  for (int iat = 1; iat < num_targets_; ++iat)
178  DTD_BConds<T, D, SC>::computeDistances(P.R[iat], P.getCoordinates().getAllParticlePos(), distances_[iat].data(),
179  displacements_[iat], 0, iat, iat);
180  }
ScopeGuard< NewTimer > ScopedTimer
Definition: NewTimer.h:257
std::vector< DisplRow > displacements_
displacements_[num_targets_][3][num_sources_], [i][3][j] = r_A2[j] - r_A1[i] Note: Derived classes de...
std::vector< DistRow > distances_
distances_[num_targets_][num_sources_], [i][3][j] = |r_A2[j] - r_A1[i]| Note: Derived classes decide ...

◆ get_first_neighbor()

int get_first_neighbor ( IndexType  iat,
RealType r,
PosType dr,
bool  newpos 
) const
inlineoverride

Definition at line 376 of file SoaDistanceTableAAOMPTarget.h.

References DistanceTableAA::displacements_, DistanceTableAA::distances_, DistanceTable::num_targets_, DistanceTableAA::temp_dr_, and DistanceTableAA::temp_r_.

377  {
378  //ensure there are neighbors
379  assert(num_targets_ > 1);
380  RealType min_dist = std::numeric_limits<RealType>::max();
381  int index = -1;
382  if (newpos)
383  {
384  for (int jat = 0; jat < num_targets_; ++jat)
385  if (temp_r_[jat] < min_dist && jat != iat)
386  {
387  min_dist = temp_r_[jat];
388  index = jat;
389  }
390  assert(index >= 0);
391  dr = temp_dr_[index];
392  }
393  else
394  {
395  for (int jat = 0; jat < iat; ++jat)
396  if (distances_[iat][jat] < min_dist)
397  {
398  min_dist = distances_[iat][jat];
399  index = jat;
400  }
401  for (int jat = iat + 1; jat < num_targets_; ++jat)
402  if (distances_[jat][iat] < min_dist)
403  {
404  min_dist = distances_[jat][iat];
405  index = jat;
406  }
407  assert(index != iat && index >= 0);
408  if (index < iat)
409  dr = displacements_[iat][index];
410  else
411  dr = displacements_[index][iat];
412  }
413  r = min_dist;
414  return index;
415  }
std::vector< DisplRow > displacements_
displacements_[num_targets_][3][num_sources_], [i][3][j] = r_A2[j] - r_A1[i] Note: Derived classes de...
QMCTraits::RealType RealType
std::vector< DistRow > distances_
distances_[num_targets_][num_sources_], [i][3][j] = |r_A2[j] - r_A1[i]| Note: Derived classes decide ...

◆ get_num_particls_stored()

size_t get_num_particls_stored ( ) const
inlineoverridevirtual

Reimplemented from DistanceTableAA.

Definition at line 484 of file SoaDistanceTableAAOMPTarget.h.

References SoaDistanceTableAAOMPTarget< T, D, SC >::num_particls_stored.

484 { return num_particls_stored; }
const size_t num_particls_stored
the particle count of the internal stored distances.

◆ getMultiWalkerTempDataPtr()

const RealType* getMultiWalkerTempDataPtr ( ) const
inlineoverridevirtual

return multi walker temporary pair distance table data pointer

Reimplemented from DistanceTableAA.

Definition at line 117 of file SoaDistanceTableAAOMPTarget.h.

References SoaDistanceTableAAOMPTarget< T, D, SC >::mw_mem_handle_.

118  {
119  return mw_mem_handle_.getResource().mw_new_old_dist_displ.data();
120  }
ResourceHandle< DTAAMultiWalkerMem > mw_mem_handle_

◆ move()

void move ( const ParticleSet P,
const PosType rnew,
const IndexType  iat,
bool  prepare_old 
)
inlineoverridevirtual

evaluate the temporary pair relations

Implements DistanceTable.

Definition at line 257 of file SoaDistanceTableAAOMPTarget.h.

References Vector< T, Alloc >::attachReference(), VectorSoaContainer< T, D, Alloc >::attachReference(), VectorSoaContainer< T, D, Alloc >::capacity(), Vector< T, Alloc >::data(), VectorSoaContainer< T, D, Alloc >::data(), DynamicCoordinates::getAllParticlePos(), ParticleSet::getCoordinates(), SoaDistanceTableAAOMPTarget< T, D, SC >::move_timer_, DistanceTable::num_targets_, DistanceTableAA::old_dr_, SoaDistanceTableAAOMPTarget< T, D, SC >::old_dr_mem_, SoaDistanceTableAAOMPTarget< T, D, SC >::old_prepared_elec_id_, DistanceTableAA::old_r_, SoaDistanceTableAAOMPTarget< T, D, SC >::old_r_mem_, ParticleSet::R, Vector< T, Alloc >::size(), VectorSoaContainer< T, D, Alloc >::size(), DistanceTableAA::temp_dr_, SoaDistanceTableAAOMPTarget< T, D, SC >::temp_dr_mem_, DistanceTableAA::temp_r_, and SoaDistanceTableAAOMPTarget< T, D, SC >::temp_r_mem_.

258  {
259  ScopedTimer local_timer(move_timer_);
260 
261 #if !defined(NDEBUG)
262  old_prepared_elec_id_ = prepare_old ? iat : -1;
263 #endif
266 
267  assert((prepare_old && iat >= 0 && iat < num_targets_) || !prepare_old);
268  DTD_BConds<T, D, SC>::computeDistances(rnew, P.getCoordinates().getAllParticlePos(), temp_r_.data(), temp_dr_, 0,
269  num_targets_, iat);
270  // set up old_r_ and old_dr_ for moves may get accepted.
271  if (prepare_old)
272  {
275  //recompute from scratch
276  DTD_BConds<T, D, SC>::computeDistances(P.R[iat], P.getCoordinates().getAllParticlePos(), old_r_.data(), old_dr_,
277  0, num_targets_, iat);
278  old_r_[iat] = std::numeric_limits<T>::max(); //assign a big number
279  }
280  }
DisplRow temp_dr_mem_
actual memory for temp_dr_
size_type capacity() const
return the physical size
ScopeGuard< NewTimer > ScopedTimer
Definition: NewTimer.h:257
int old_prepared_elec_id_
set to particle id after move() with prepare_old = true.
DisplRow old_dr_
old displacements
DistRow old_r_
old distances
void attachReference(T *ref, size_type n)
Definition: OhmmsVector.h:131
DistRow temp_r_mem_
actual memory for temp_r_
size_type size() const
return the current size
Definition: OhmmsVector.h:162
DisplRow old_dr_mem_
actual memory for old_dr_
size_type size() const
return the physical size
void attachReference(size_type n, size_type n_padded, T *ptr)
attach to pre-allocated data

◆ mw_evalDistsInRange()

const RealType* mw_evalDistsInRange ( const RefVectorWithLeader< DistanceTable > &  dt_list,
const RefVectorWithLeader< ParticleSet > &  p_list,
size_t  range_begin,
size_t  range_end 
) const
inlineoverridevirtual

compute distances from particles in [range_begin, range_end) to all the particles.

Although [range_begin, range_end) and be any particle [0, num_sources), it is only necessary to compute half of the table due to the symmetry of AA table. See note of the output data object mw_distances_subset To keep resident memory minimal on the device, range_end - range_begin < num_particls_stored is required.

Reimplemented from DistanceTableAA.

Definition at line 187 of file SoaDistanceTableAAOMPTarget.h.

References RefVectorWithLeader< T >::getCastedLeader(), RefVectorWithLeader< T >::getLeader(), RealSpacePositionsOMPTarget::getMultiWalkerRSoADevicePtrs(), omptarget::min(), and SoaDistanceTableAAOMPTarget< T, D, SC >::DTAAMultiWalkerMem::mw_distances_subset.

191  {
192  auto& dt_leader = dt_list.getCastedLeader<SoaDistanceTableAAOMPTarget>();
193  const size_t subset_size = range_end - range_begin;
194  if (subset_size > dt_leader.num_particls_stored)
195  throw std::runtime_error("not enough internal buffer");
196 
197  ScopedTimer local_timer(dt_leader.evaluate_timer_);
198 
199  DTAAMultiWalkerMem& mw_mem = dt_leader.mw_mem_handle_;
200  auto& pset_leader = p_list.getLeader();
201 
202  const size_t nw = dt_list.size();
203  const auto num_sources_local = dt_leader.num_targets_;
204  const auto num_padded = dt_leader.num_targets_padded_;
205  mw_mem.mw_distances_subset.resize(nw * subset_size * num_padded);
206 
207  const int ChunkSizePerTeam = 512;
208  const size_t num_teams = (num_sources_local + ChunkSizePerTeam - 1) / ChunkSizePerTeam;
209 
210  auto& coordinates_leader = static_cast<const RealSpacePositionsOMPTarget&>(pset_leader.getCoordinates());
211 
212  auto* rsoa_dev_list_ptr = coordinates_leader.getMultiWalkerRSoADevicePtrs().data();
213  auto* dist_ranged = mw_mem.mw_distances_subset.data();
214  {
215  ScopedTimer offload(dt_leader.offload_timer_);
216  PRAGMA_OFFLOAD("omp target teams distribute collapse(2) num_teams(nw * num_teams)")
217  for (int iw = 0; iw < nw; ++iw)
218  for (int team_id = 0; team_id < num_teams; team_id++)
219  {
220  auto* source_pos_ptr = rsoa_dev_list_ptr[iw];
221  const size_t first = ChunkSizePerTeam * team_id;
222  const size_t last = omptarget::min(first + ChunkSizePerTeam, num_sources_local);
223 
224  PRAGMA_OFFLOAD("omp parallel for")
225  for (int iel = first; iel < last; iel++)
226  {
227  for (int irow = 0; irow < subset_size; irow++)
228  {
229  T* dist = dist_ranged + (irow + subset_size * iw) * num_padded;
230  size_t id_target = irow + range_begin;
231 
232  T dx, dy, dz;
233  if (id_target < iel)
234  {
235  dx = source_pos_ptr[id_target] - source_pos_ptr[iel];
236  dy = source_pos_ptr[id_target + num_padded] - source_pos_ptr[iel + num_padded];
237  dz = source_pos_ptr[id_target + num_padded * 2] - source_pos_ptr[iel + num_padded * 2];
238  }
239  else
240  {
241  const size_t id_target_reverse = num_sources_local - 1 - id_target;
242  const size_t iel_reverse = num_sources_local - 1 - iel;
243  dx = source_pos_ptr[id_target_reverse] - source_pos_ptr[iel_reverse];
244  dy = source_pos_ptr[id_target_reverse + num_padded] - source_pos_ptr[iel_reverse + num_padded];
245  dz = source_pos_ptr[id_target_reverse + num_padded * 2] - source_pos_ptr[iel_reverse + num_padded * 2];
246  }
247 
248  dist[iel] = DTD_BConds<T, D, SC>::computeDist(dx, dy, dz);
249  }
250  }
251  }
252  }
253  return mw_mem.mw_distances_subset.data();
254  }
ScopeGuard< NewTimer > ScopedTimer
Definition: NewTimer.h:257
T min(T a, T b)
for(int i=0;i< size_test;++i) CHECK(Approx(gauss_random_vals[offset_for_rs+i])

◆ mw_finalizePbyP()

void mw_finalizePbyP ( const RefVectorWithLeader< DistanceTable > &  dt_list,
const RefVectorWithLeader< ParticleSet > &  p_list 
) const
inlineoverridevirtual

walker batched version of finalizePbyP If not DTModes::NEED_TEMP_DATA_ON_HOST, host distance table data is not updated at all during p-by-p Thus, a recompute is necessary to update the whole host distance table for consumers like the Coulomb potential.

Reimplemented from DistanceTable.

Definition at line 475 of file SoaDistanceTableAAOMPTarget.h.

References DistanceTable::modes_, DistanceTable::mw_evaluate(), qmcplusplus::NEED_FULL_TABLE_ON_HOST_AFTER_DONEPBYP, and qmcplusplus::NEED_TEMP_DATA_ON_HOST.

477  {
478  // if the distance table is not updated by mw_move during p-by-p, needs to recompute the whole table
479  // before being used by Hamiltonian if requested
481  mw_evaluate(dt_list, p_list);
482  }
virtual void mw_evaluate(const RefVectorWithLeader< DistanceTable > &dt_list, const RefVectorWithLeader< ParticleSet > &p_list) const
whether full table needs to be ready at anytime or not after donePbyP Optimization can be implemented...
whether temporary data set on the host is updated or not when a move is proposed. ...
DTModes modes_
operation modes defined by DTModes
Definition: DistanceTable.h:60

◆ mw_move()

void mw_move ( const RefVectorWithLeader< DistanceTable > &  dt_list,
const RefVectorWithLeader< ParticleSet > &  p_list,
const std::vector< PosType > &  rnew_list,
const IndexType  iat,
bool  prepare_old = true 
) const
inlineoverridevirtual

evaluate the temporary pair relations when a move is proposed this implementation is asynchronous and the synchronization is managed at ParticleSet.

Transferring results to host depends on DTModes::NEED_TEMP_DATA_ON_HOST. If the temporary pair distance are consumed on the device directly, the device to host data transfer can be skipped as an optimization.

Reimplemented from DistanceTable.

Definition at line 288 of file SoaDistanceTableAAOMPTarget.h.

References RefVectorWithLeader< T >::getCastedElement(), RefVectorWithLeader< T >::getCastedLeader(), RefVectorWithLeader< T >::getLeader(), omptarget::min(), DistanceTable::modes_, SoaDistanceTableAAOMPTarget< T, D, SC >::move_timer_, qmcplusplus::NEED_TEMP_DATA_ON_HOST, DistanceTable::num_targets_, SoaDistanceTableAAOMPTarget< T, D, SC >::num_targets_padded_, SoaDistanceTableAAOMPTarget< T, D, SC >::offload_timer_, and SoaDistanceTableAAOMPTarget< T, D, SC >::old_prepared_elec_id_.

293  {
294  assert(this == &dt_list.getLeader());
295  auto& dt_leader = dt_list.getCastedLeader<SoaDistanceTableAAOMPTarget>();
296  DTAAMultiWalkerMem& mw_mem = dt_leader.mw_mem_handle_;
297  auto& pset_leader = p_list.getLeader();
298 
299  ScopedTimer local_timer(move_timer_);
300  const size_t nw = dt_list.size();
301  const size_t stride_size = num_targets_padded_ * (D + 1);
302 
303  auto& mw_new_old_dist_displ = mw_mem.mw_new_old_dist_displ;
304 
305  for (int iw = 0; iw < nw; iw++)
306  {
307  auto& dt = dt_list.getCastedElement<SoaDistanceTableAAOMPTarget>(iw);
308 #if !defined(NDEBUG)
309  dt.old_prepared_elec_id_ = prepare_old ? iat : -1;
310 #endif
311  auto& coordinates_soa = static_cast<const RealSpacePositionsOMPTarget&>(p_list[iw].getCoordinates());
312  }
313 
314  const int ChunkSizePerTeam = 512;
315  const size_t num_teams = (num_targets_ + ChunkSizePerTeam - 1) / ChunkSizePerTeam;
316 
317  auto& coordinates_leader = static_cast<const RealSpacePositionsOMPTarget&>(pset_leader.getCoordinates());
318 
319  const auto num_sources_local = num_targets_;
320  const auto num_padded = num_targets_padded_;
321  auto* rsoa_dev_list_ptr = coordinates_leader.getMultiWalkerRSoADevicePtrs().data();
322  auto* r_dr_ptr = mw_new_old_dist_displ.data();
323  auto* new_pos_ptr = coordinates_leader.getFusedNewPosBuffer().data();
324  const size_t new_pos_stride = coordinates_leader.getFusedNewPosBuffer().capacity();
325 
326  {
327  ScopedTimer offload(offload_timer_);
328  PRAGMA_OFFLOAD("omp target teams distribute collapse(2) num_teams(nw * num_teams) \
329  depend(out: r_dr_ptr[:mw_new_old_dist_displ.size()])")
330  for (int iw = 0; iw < nw; ++iw)
331  for (int team_id = 0; team_id < num_teams; team_id++)
332  {
333  auto* source_pos_ptr = rsoa_dev_list_ptr[iw];
334  const size_t first = ChunkSizePerTeam * team_id;
335  const size_t last = omptarget::min(first + ChunkSizePerTeam, num_sources_local);
336 
337  { // temp
338  auto* r_iw_ptr = r_dr_ptr + iw * stride_size;
339  auto* dr_iw_ptr = r_dr_ptr + iw * stride_size + num_padded;
340 
341  T pos[D];
342  for (int idim = 0; idim < D; idim++)
343  pos[idim] = new_pos_ptr[idim * new_pos_stride + iw];
344 
345  PRAGMA_OFFLOAD("omp parallel for")
346  for (int iel = first; iel < last; iel++)
347  DTD_BConds<T, D, SC>::computeDistancesOffload(pos, source_pos_ptr, num_padded, r_iw_ptr, dr_iw_ptr,
348  num_padded, iel, iat);
349  }
350 
351  if (prepare_old)
352  { // old
353  auto* r_iw_ptr = r_dr_ptr + (iw + nw) * stride_size;
354  auto* dr_iw_ptr = r_dr_ptr + (iw + nw) * stride_size + num_padded;
355 
356  T pos[D];
357  for (int idim = 0; idim < D; idim++)
358  pos[idim] = source_pos_ptr[idim * num_padded + iat];
359 
360  PRAGMA_OFFLOAD("omp parallel for")
361  for (int iel = first; iel < last; iel++)
362  DTD_BConds<T, D, SC>::computeDistancesOffload(pos, source_pos_ptr, num_padded, r_iw_ptr, dr_iw_ptr,
363  num_padded, iel, iat);
364  r_iw_ptr[iat] = std::numeric_limits<T>::max(); //assign a big number
365  }
366  }
367  }
368 
370  {
371  PRAGMA_OFFLOAD("omp target update nowait depend(inout: r_dr_ptr[:mw_new_old_dist_displ.size()]) \
372  from(r_dr_ptr[:mw_new_old_dist_displ.size()])")
373  }
374  }
DTD_BConds(const CrystalLattice< T, D > &lat)
constructor: doing nothing
if(c->rank()==0)
const size_t num_targets_padded_
number of targets with padding
ScopeGuard< NewTimer > ScopedTimer
Definition: NewTimer.h:257
T min(T a, T b)
for(int i=0;i< size_test;++i) CHECK(Approx(gauss_random_vals[offset_for_rs+i])
whether temporary data set on the host is updated or not when a move is proposed. ...
NewTimer & offload_timer_
timer for offload portion
DTModes modes_
operation modes defined by DTModes
Definition: DistanceTable.h:60

◆ mw_updatePartial()

void mw_updatePartial ( const RefVectorWithLeader< DistanceTable > &  dt_list,
IndexType  jat,
const std::vector< bool > &  from_temp 
)
inlineoverridevirtual

walker batched version of updatePartial.

If not DTModes::NEED_TEMP_DATA_ON_HOST, host data is not up-to-date and host distance table will not be updated.

Reimplemented from DistanceTable.

Definition at line 463 of file SoaDistanceTableAAOMPTarget.h.

References DistanceTable::modes_, qmcplusplus::NEED_TEMP_DATA_ON_HOST, and SoaDistanceTableAAOMPTarget< T, D, SC >::updatePartial().

466  {
467  // if temp data on host is not updated by mw_move during p-by-p moves, there is no need to update distance table
469  return;
470 
471  for (int iw = 0; iw < dt_list.size(); iw++)
472  dt_list[iw].updatePartial(jat, from_temp[iw]);
473  }
void updatePartial(IndexType jat, bool from_temp) override
fill partially the distance table by the pair relations from the temporary or old particle position...
whether temporary data set on the host is updated or not when a move is proposed. ...
DTModes modes_
operation modes defined by DTModes
Definition: DistanceTable.h:60

◆ releaseResource()

void releaseResource ( ResourceCollection collection,
const RefVectorWithLeader< DistanceTable > &  dt_list 
) const
inlineoverridevirtual

return a shared resource to a collection

Reimplemented from DistanceTable.

Definition at line 158 of file SoaDistanceTableAAOMPTarget.h.

References Vector< T, Alloc >::free(), RefVectorWithLeader< T >::getCastedElement(), RefVectorWithLeader< T >::getCastedLeader(), SoaDistanceTableAAOMPTarget< T, D, SC >::mw_mem_handle_, ResourceCollection::takebackResource(), and DistanceTableAA::temp_r_.

159  {
160  collection.takebackResource(dt_list.getCastedLeader<SoaDistanceTableAAOMPTarget>().mw_mem_handle_);
161  const size_t nw = dt_list.size();
162  for (int iw = 0; iw < nw; iw++)
163  {
164  auto& dt = dt_list.getCastedElement<SoaDistanceTableAAOMPTarget>(iw);
165  dt.temp_r_.free();
166  dt.temp_dr_.free();
167  dt.old_r_.free();
168  dt.old_dr_.free();
169  }
170  }

◆ resize()

void resize ( )
inline

Definition at line 98 of file SoaDistanceTableAAOMPTarget.h.

References SoaDistanceTableAAOMPTarget< T, D, SC >::compute_size(), DistanceTableAA::displacements_, DistanceTableAA::distances_, SoaDistanceTableAAOMPTarget< T, D, SC >::memory_pool_, DistanceTable::num_targets_, SoaDistanceTableAAOMPTarget< T, D, SC >::old_dr_mem_, SoaDistanceTableAAOMPTarget< T, D, SC >::old_r_mem_, VectorSoaContainer< T, D, Alloc >::resize(), Vector< T, Alloc >::resize(), SoaDistanceTableAAOMPTarget< T, D, SC >::temp_dr_mem_, and SoaDistanceTableAAOMPTarget< T, D, SC >::temp_r_mem_.

Referenced by SoaDistanceTableAAOMPTarget< T, D, SC >::SoaDistanceTableAAOMPTarget().

99  {
100  // initialize memory containers and views
101  const size_t total_size = compute_size(num_targets_);
102  memory_pool_.resize(total_size * (1 + D));
103  distances_.resize(num_targets_);
105  for (int i = 0; i < num_targets_; ++i)
106  {
107  distances_[i].attachReference(memory_pool_.data() + compute_size(i), i);
108  displacements_[i].attachReference(i, total_size, memory_pool_.data() + total_size + compute_size(i));
109  }
110 
115  }
void resize(size_type n, Type_t val=Type_t())
Resize the container.
Definition: OhmmsVector.h:166
DisplRow temp_dr_mem_
actual memory for temp_dr_
aligned_vector< RealType > memory_pool_
actual memory for dist and displacements_
std::vector< DisplRow > displacements_
displacements_[num_targets_][3][num_sources_], [i][3][j] = r_A2[j] - r_A1[i] Note: Derived classes de...
DistRow temp_r_mem_
actual memory for temp_r_
DisplRow old_dr_mem_
actual memory for old_dr_
void resize(size_type n)
resize myData
std::vector< DistRow > distances_
distances_[num_targets_][num_sources_], [i][3][j] = |r_A2[j] - r_A1[i]| Note: Derived classes decide ...

◆ update()

void update ( IndexType  iat)
inlineoverridevirtual

After accepting the iat-th particle, update the iat-th row of distances_ and displacements_.

Upper triangle is not needed in the later computation and thus not updated

Implements DistanceTable.

Definition at line 420 of file SoaDistanceTableAAOMPTarget.h.

References qmcplusplus::syclBLAS::copy_n(), Vector< T, Alloc >::data(), VectorSoaContainer< T, D, Alloc >::data(), DistanceTableAA::displacements_, DistanceTableAA::distances_, DistanceTable::num_targets_, Vector< T, Alloc >::size(), DistanceTableAA::temp_dr_, DistanceTableAA::temp_r_, and SoaDistanceTableAAOMPTarget< T, D, SC >::update_timer_.

421  {
422  ScopedTimer local_timer(update_timer_);
423  //update [0, iat) columns
424  const int nupdate = iat;
425  //copy row
426  assert(nupdate <= temp_r_.size());
427  std::copy_n(temp_r_.data(), nupdate, distances_[iat].data());
428  for (int idim = 0; idim < D; ++idim)
429  std::copy_n(temp_dr_.data(idim), nupdate, displacements_[iat].data(idim));
430  //copy column
431  for (size_t i = iat + 1; i < num_targets_; ++i)
432  {
433  distances_[i][iat] = temp_r_[i];
434  displacements_[i](iat) = -temp_dr_[i];
435  }
436  }
ScopeGuard< NewTimer > ScopedTimer
Definition: NewTimer.h:257
std::vector< DisplRow > displacements_
displacements_[num_targets_][3][num_sources_], [i][3][j] = r_A2[j] - r_A1[i] Note: Derived classes de...
size_type size() const
return the current size
Definition: OhmmsVector.h:162
sycl::event copy_n(sycl::queue &aq, const T1 *restrict VA, size_t array_size, T2 *restrict VC, const std::vector< sycl::event > &events)
Definition: syclBLAS.cpp:548
std::vector< DistRow > distances_
distances_[num_targets_][num_sources_], [i][3][j] = |r_A2[j] - r_A1[i]| Note: Derived classes decide ...

◆ updatePartial()

void updatePartial ( IndexType  jat,
bool  from_temp 
)
inlineoverridevirtual

fill partially the distance table by the pair relations from the temporary or old particle position.

Used in forward mode when a move is reject

Parameters
iatthe particle with an accepted move
from_tempif true, copy from temp. if false, copy from old

Reimplemented from DistanceTable.

Definition at line 438 of file SoaDistanceTableAAOMPTarget.h.

References qmcplusplus::syclBLAS::copy_n(), Vector< T, Alloc >::data(), VectorSoaContainer< T, D, Alloc >::data(), DistanceTableAA::displacements_, DistanceTableAA::distances_, DistanceTableAA::old_dr_, SoaDistanceTableAAOMPTarget< T, D, SC >::old_prepared_elec_id_, DistanceTableAA::old_r_, Vector< T, Alloc >::size(), DistanceTableAA::temp_dr_, DistanceTableAA::temp_r_, and SoaDistanceTableAAOMPTarget< T, D, SC >::update_timer_.

Referenced by SoaDistanceTableAAOMPTarget< T, D, SC >::mw_updatePartial().

439  {
440  ScopedTimer local_timer(update_timer_);
441 
442  //update [0, jat)
443  const int nupdate = jat;
444  if (from_temp)
445  {
446  //copy row
447  assert(nupdate <= temp_r_.size());
448  std::copy_n(temp_r_.data(), nupdate, distances_[jat].data());
449  for (int idim = 0; idim < D; ++idim)
450  std::copy_n(temp_dr_.data(idim), nupdate, displacements_[jat].data(idim));
451  }
452  else
453  {
454  assert(old_prepared_elec_id_ == jat);
455  //copy row
456  assert(nupdate <= old_r_.size());
457  std::copy_n(old_r_.data(), nupdate, distances_[jat].data());
458  for (int idim = 0; idim < D; ++idim)
459  std::copy_n(old_dr_.data(idim), nupdate, displacements_[jat].data(idim));
460  }
461  }
ScopeGuard< NewTimer > ScopedTimer
Definition: NewTimer.h:257
std::vector< DisplRow > displacements_
displacements_[num_targets_][3][num_sources_], [i][3][j] = r_A2[j] - r_A1[i] Note: Derived classes de...
int old_prepared_elec_id_
set to particle id after move() with prepare_old = true.
DisplRow old_dr_
old displacements
DistRow old_r_
old distances
size_type size() const
return the current size
Definition: OhmmsVector.h:162
sycl::event copy_n(sycl::queue &aq, const T1 *restrict VA, size_t array_size, T2 *restrict VC, const std::vector< sycl::event > &events)
Definition: syclBLAS.cpp:548
std::vector< DistRow > distances_
distances_[num_targets_][num_sources_], [i][3][j] = |r_A2[j] - r_A1[i]| Note: Derived classes decide ...

Member Data Documentation

◆ evaluate_timer_

NewTimer& evaluate_timer_
private

◆ memory_pool_

aligned_vector<RealType> memory_pool_

actual memory for dist and displacements_

Definition at line 34 of file SoaDistanceTableAAOMPTarget.h.

Referenced by SoaDistanceTableAAOMPTarget< T, D, SC >::resize().

◆ move_timer_

◆ mw_mem_handle_

◆ num_particls_stored

const size_t num_particls_stored = 64
private

the particle count of the internal stored distances.

Definition at line 504 of file SoaDistanceTableAAOMPTarget.h.

Referenced by SoaDistanceTableAAOMPTarget< T, D, SC >::get_num_particls_stored().

◆ num_targets_padded_

const size_t num_targets_padded_
private

◆ offload_timer_

NewTimer& offload_timer_
private

timer for offload portion

Definition at line 496 of file SoaDistanceTableAAOMPTarget.h.

Referenced by SoaDistanceTableAAOMPTarget< T, D, SC >::mw_move().

◆ old_dr_mem_

DisplRow old_dr_mem_

◆ old_prepared_elec_id_

int old_prepared_elec_id_
private

set to particle id after move() with prepare_old = true.

-1 means not prepared. It is intended only for safety checks, not for codepath selection.

Definition at line 493 of file SoaDistanceTableAAOMPTarget.h.

Referenced by SoaDistanceTableAAOMPTarget< T, D, SC >::move(), SoaDistanceTableAAOMPTarget< T, D, SC >::mw_move(), and SoaDistanceTableAAOMPTarget< T, D, SC >::updatePartial().

◆ old_r_mem_

DistRow old_r_mem_

◆ temp_dr_mem_

DisplRow temp_dr_mem_

◆ temp_r_mem_

DistRow temp_r_mem_

◆ update_timer_


The documentation for this struct was generated from the following file: