QMCPACK
SoaDistanceTableAAOMPTarget.h
Go to the documentation of this file.
1 //////////////////////////////////////////////////////////////////////////////////////
2 // This file is distributed under the University of Illinois/NCSA Open Source License.
3 // See LICENSE file in top directory for details.
4 //
5 // Copyright (c) 2021 QMCPACK developers.
6 //
7 // File developed by: Jeongnim Kim, jeongnim.kim@intel.com, Intel Corp.
8 // Amrita Mathuriya, amrita.mathuriya@intel.com, Intel Corp.
9 // Ye Luo, yeluo@anl.gov, Argonne National Laboratory
10 //
11 // File created by: Ye Luo, yeluo@anl.gov, Argonne National Laboratory
12 //////////////////////////////////////////////////////////////////////////////////////
13 // -*- C++ -*-
14 #ifndef QMCPLUSPLUS_DTDIMPL_AA_OMPTARGET_H
15 #define QMCPLUSPLUS_DTDIMPL_AA_OMPTARGET_H
16 
18 #include "DistanceTable.h"
22 #include "ResourceCollection.h"
24 
25 namespace qmcplusplus
26 {
27 /**@ingroup nnlist
28  * @brief A derived classe from DistacneTableData, specialized for dense case
29  */
30 template<typename T, unsigned D, int SC>
31 struct SoaDistanceTableAAOMPTarget : public DTD_BConds<T, D, SC>, public DistanceTableAA
32 {
33  /// actual memory for dist and displacements_
35 
36  /// actual memory for temp_r_
38  /// actual memory for temp_dr_
40  /// actual memory for old_r_
42  /// actual memory for old_dr_
44 
45  ///multi walker shared memory buffer
46  struct DTAAMultiWalkerMem : public Resource
47  {
48  ///dist displ for temporary and old pairs
50 
51  /** distances from a range of indics to the source.
52  * for original particle index i (row) and source particle id j (col)
53  * j < i, the element data is dist(r_i - r_j)
54  * j > i, the element data is dist(r_(n - 1 - i) - r_(n - 1 - j))
55  */
57 
58  DTAAMultiWalkerMem() : Resource("DTAAMultiWalkerMem") {}
59 
61 
62  std::unique_ptr<Resource> makeClone() const override { return std::make_unique<DTAAMultiWalkerMem>(*this); }
63  };
64 
66 
68  : DTD_BConds<T, D, SC>(target.getLattice()),
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  }
86 
87  SoaDistanceTableAAOMPTarget() = delete;
89  ~SoaDistanceTableAAOMPTarget(){PRAGMA_OFFLOAD("omp target exit data map(delete : this[:1])")}
90 
91  size_t compute_size(int N) const
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  }
97 
98  void resize()
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  }
116 
117  const RealType* getMultiWalkerTempDataPtr() const override
118  {
119  return mw_mem_handle_.getResource().mw_new_old_dist_displ.data();
120  }
121 
122  void createResource(ResourceCollection& collection) const override
123  {
124  auto resource_index = collection.addResource(std::make_unique<DTAAMultiWalkerMem>());
125  }
126 
127  void acquireResource(ResourceCollection& collection, const RefVectorWithLeader<DistanceTable>& dt_list) const override
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  }
157 
158  void releaseResource(ResourceCollection& collection, const RefVectorWithLeader<DistanceTable>& dt_list) const override
159  {
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  }
171 
172  inline void evaluate(ParticleSet& P) override
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)
179  displacements_[iat], 0, iat, iat);
180  }
181 
182  /** compute distances from particles in [range_begin, range_end) to all the particles.
183  * Although [range_begin, range_end) and be any particle [0, num_sources), it is only necessary to compute
184  * half of the table due to the symmetry of AA table. See note of the output data object mw_distances_subset
185  * To keep resident memory minimal on the device, range_end - range_begin < num_particls_stored is required.
186  */
188  const RefVectorWithLeader<ParticleSet>& p_list,
189  size_t range_begin,
190  size_t range_end) const override
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  }
255 
256  ///evaluate the temporary pair relations
257  inline void move(const ParticleSet& P, const PosType& rnew, const IndexType iat, bool prepare_old) override
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);
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
277  0, num_targets_, iat);
278  old_r_[iat] = std::numeric_limits<T>::max(); //assign a big number
279  }
280  }
281 
282  /** evaluate the temporary pair relations when a move is proposed
283  * this implementation is asynchronous and the synchronization is managed at ParticleSet.
284  * Transferring results to host depends on DTModes::NEED_TEMP_DATA_ON_HOST.
285  * If the temporary pair distance are consumed on the device directly, the device to host data transfer can be
286  * skipped as an optimization.
287  */
289  const RefVectorWithLeader<ParticleSet>& p_list,
290  const std::vector<PosType>& rnew_list,
291  const IndexType iat,
292  bool prepare_old = true) const override
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  }
375 
376  int get_first_neighbor(IndexType iat, RealType& r, PosType& dr, bool newpos) const override
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  }
416 
417  /** After accepting the iat-th particle, update the iat-th row of distances_ and displacements_.
418  * Upper triangle is not needed in the later computation and thus not updated
419  */
420  inline void update(IndexType iat) override
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  }
437 
438  void updatePartial(IndexType jat, bool from_temp) override
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  }
462 
464  IndexType jat,
465  const std::vector<bool>& from_temp) override
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  }
474 
476  const RefVectorWithLeader<ParticleSet>& p_list) const override
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  }
483 
484  size_t get_num_particls_stored() const override { return num_particls_stored; }
485 
486 private:
487  ///number of targets with padding
488  const size_t num_targets_padded_;
489 #if !defined(NDEBUG)
490  /** set to particle id after move() with prepare_old = true. -1 means not prepared.
491  * It is intended only for safety checks, not for codepath selection.
492  */
494 #endif
495  /// timer for offload portion
497  /// timer for evaluate()
499  /// timer for move()
501  /// timer for update()
503  /// the particle count of the internal stored distances.
504  const size_t num_particls_stored = 64;
505 };
506 } // namespace qmcplusplus
507 #endif
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_
const RealType * getMultiWalkerTempDataPtr() const override
return multi walker temporary pair distance table data pointer
size_t addResource(std::unique_ptr< Resource > &&res, bool noprint=false)
size_t getAlignedSize(size_t n)
return size in T&#39;s of allocated aligned memory
size_type capacity() const
return the physical size
std::vector< T, aligned_allocator< T > > aligned_vector
int get_first_neighbor(IndexType iat, RealType &r, PosType &dr, bool newpos) const override
void takebackResource(ResourceHandle< RS > &res_handle)
helper functions for EinsplineSetBuilder
Definition: Configuration.h:43
void acquireResource(ResourceCollection &collection, const RefVectorWithLeader< DistanceTable > &dt_list) const override
acquire a shared resource from a collection
virtual const PosVectorSoa & getAllParticlePos() const =0
all particle position accessor
if(c->rank()==0)
const size_t num_targets_padded_
number of targets with padding
void evaluate(ParticleSet &P) override
evaluate the full Distance Table
ResourceHandle manages the temporary resource referenced from a collection.
void releaseResource(ResourceCollection &collection, const RefVectorWithLeader< DistanceTable > &dt_list) const override
return a shared resource to a collection
SoA adaptor class for Vector<TinyVector<T,D> >
Timer accumulates time and call counts.
Definition: NewTimer.h:135
void free()
free
Definition: OhmmsVector.h:196
ResourceHandle< DTAAMultiWalkerMem > mw_mem_handle_
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...
A derived classe from DistacneTableData, specialized for dense case.
const std::string name_
name of the table
Definition: DistanceTable.h:57
T min(T a, T b)
void createResource(ResourceCollection &collection) const override
initialize a shared resource and hand it to a collection
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
Specialized paritlce class for atomistic simulations.
Definition: ParticleSet.h:55
DistRow temp_r_mem_
actual memory for temp_r_
size_type size() const
return the current size
Definition: OhmmsVector.h:162
Vector< RealType, OMPallocator< RealType, PinnedAlignedAllocator< RealType > > > mw_distances_subset
distances from a range of indics to the source.
const RealType * mw_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.
Introduced to handle virtual moves and ratio computations, e.g.
void mw_updatePartial(const RefVectorWithLeader< DistanceTable > &dt_list, IndexType jat, const std::vector< bool > &from_temp) override
walker batched version of updatePartial.
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 da...
CASTTYPE & getCastedElement(size_t i) const
NewTimer & createGlobalTimer(const std::string &myname, timer_levels mylevel)
void move(const ParticleSet &P, const PosType &rnew, const IndexType iat, bool prepare_old) override
evaluate the temporary pair relations
DisplRow old_dr_mem_
actual memory for old_dr_
ParticlePos R
Position.
Definition: ParticleSet.h:79
AA type of DistanceTable containing storage.
const DynamicCoordinates & getCoordinates() const
Definition: ParticleSet.h:246
Vector< RealType, OMPallocator< RealType, PinnedAlignedAllocator< RealType > > > mw_new_old_dist_displ
dist displ for temporary and old pairs
QMCTraits::IndexType IndexType
Definition: DistanceTable.h:43
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
virtual void mw_evaluate(const RefVectorWithLeader< DistanceTable > &dt_list, const RefVectorWithLeader< ParticleSet > &p_list) const
handle math function mapping inside OpenMP offload regions.
void update(IndexType iat) override
After accepting the iat-th particle, update the iat-th row of distances_ and displacements_.
ResourceHandle< RS > lendResource()
void updatePartial(IndexType jat, bool from_temp) override
fill partially the distance table by the pair relations from the temporary or old particle position...
void resize(size_type n)
resize myData
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...
size_type size() const
return the physical size
void attachReference(size_type n, size_type n_padded, T *ptr)
attach to pre-allocated data
std::vector< DistRow > distances_
distances_[num_targets_][num_sources_], [i][3][j] = |r_A2[j] - r_A1[i]| Note: Derived classes decide ...
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. ...
NewTimer & offload_timer_
timer for offload portion
DTModes modes_
operation modes defined by DTModes
Definition: DistanceTable.h:60
const size_t num_particls_stored
the particle count of the internal stored distances.