ComputeLonepairsCUDA Class Reference

#include <ComputeLonepairsCUDA.h>

Classes
struct	LonepairBisector
struct	LonepairColinear
struct	LonepairRelative

Public Member Functions
	ComputeLonepairsCUDA ()
	~ComputeLonepairsCUDA ()
void	updateAtoms (std::vector< HomePatch *> patchList, std::vector< AtomMap *> &atomMapsList, const std::vector< CudaLocalRecord > &localRecords, const int *h_globalToLocalID, cudaStream_t stream)
	Prepare the device lone pair lists from lphosts in the Molecule class. More...
void	reposition (double *d_pos_x, double *d_pos_y, double *d_pos_z, cudaStream_t stream) const
	Determine the positions of lone pairs. Should be called before force evaluations. More...
void	redistributeForce (double *d_f_normal_x, double *d_f_normal_y, double *d_f_normal_z, double *d_f_nbond_x, double *d_f_nbond_y, double *d_f_nbond_z, double *d_f_slow_x, double *d_f_slow_y, double *d_f_slow_z, cudaTensor *d_virial_normal, cudaTensor *d_virial_nbond, cudaTensor *d_virial_slow, const double *d_pos_x, const double *d_pos_y, const double *d_pos_z, const int maxForceNumber, const int doVirial, cudaStream_t stream) const
	Project the forces on lone pairs to host atoms. Should be called after force evaluations and before the integration. More...

Detailed Description

Definition at line 13 of file ComputeLonepairsCUDA.h.

Constructor & Destructor Documentation

ComputeLonepairsCUDA()

ComputeLonepairsCUDA::ComputeLonepairsCUDA

(

)

Definition at line 23 of file ComputeLonepairsCUDA.C.

                                          :
  d_lpcolinear_list(nullptr), lpcolinear_list_size(0),
  d_lprelative_list(nullptr), lprelative_list_size(0),
  d_lpbisector_list(nullptr), lpbisector_list_size(0) {}

~ComputeLonepairsCUDA()

ComputeLonepairsCUDA::~ComputeLonepairsCUDA

(

)

Definition at line 28 of file ComputeLonepairsCUDA.C.

References deallocate_device().

                                            {
  deallocate_device(&d_lpcolinear_list);
  lpcolinear_list_size = 0;
  deallocate_device(&d_lprelative_list);
  lprelative_list_size = 0;
  deallocate_device(&d_lpbisector_list);
  lpbisector_list_size = 0;
}

Member Function Documentation

redistributeForce()

void ComputeLonepairsCUDA::redistributeForce	(	double *	d_f_normal_x,
		double *	d_f_normal_y,
		double *	d_f_normal_z,
		double *	d_f_nbond_x,
		double *	d_f_nbond_y,
		double *	d_f_nbond_z,
		double *	d_f_slow_x,
		double *	d_f_slow_y,
		double *	d_f_slow_z,
		cudaTensor *	d_virial_normal,
		cudaTensor *	d_virial_nbond,
		cudaTensor *	d_virial_slow,
		const double *	d_pos_x,
		const double *	d_pos_y,
		const double *	d_pos_z,
		const int	maxForceNumber,
		const int	doVirial,
		cudaStream_t	stream
	)		const

Project the forces on lone pairs to host atoms. Should be called after force evaluations and before the integration.

Definition at line 154 of file ComputeLonepairsCUDA.C.

References redistributeForceBisector(), redistributeForceColinear(), and redistributeForceRelative().

                             {
  if (lpcolinear_list_size > 0) {
    redistributeForceColinear(
      d_f_normal_x, d_f_normal_y, d_f_normal_z,
      d_f_nbond_x, d_f_nbond_y, d_f_nbond_z,
      d_f_slow_x, d_f_slow_y, d_f_slow_z,
      d_virial_normal, d_virial_nbond, d_virial_slow,
      d_pos_x, d_pos_y, d_pos_z, maxForceNumber,
      doVirial, d_lpcolinear_list, lpcolinear_list_size,
      stream);
  }
  if (lpbisector_list_size > 0) {
    redistributeForceBisector(
      d_f_normal_x, d_f_normal_y, d_f_normal_z,
      d_f_nbond_x, d_f_nbond_y, d_f_nbond_z,
      d_f_slow_x, d_f_slow_y, d_f_slow_z,
      d_virial_normal, d_virial_nbond, d_virial_slow,
      d_pos_x, d_pos_y, d_pos_z, maxForceNumber,
      doVirial, d_lpbisector_list, lpbisector_list_size,
      stream);
  }
  if (lprelative_list_size > 0) {
    redistributeForceRelative(
      d_f_normal_x, d_f_normal_y, d_f_normal_z,
      d_f_nbond_x, d_f_nbond_y, d_f_nbond_z,
      d_f_slow_x, d_f_slow_y, d_f_slow_z,
      d_virial_normal, d_virial_nbond, d_virial_slow,
      d_pos_x, d_pos_y, d_pos_z, maxForceNumber,
      doVirial, d_lprelative_list, lprelative_list_size,
      stream);
  }
}

reposition()

void ComputeLonepairsCUDA::reposition	(	double *	d_pos_x,
		double *	d_pos_y,
		double *	d_pos_z,
		cudaStream_t	stream
	)		const

Determine the positions of lone pairs. Should be called before force evaluations.

Definition at line 138 of file ComputeLonepairsCUDA.C.

References repositionBisector(), repositionColinear(), and repositionRelative().

                             {
  if (lpcolinear_list_size > 0) {
    repositionColinear(d_pos_x, d_pos_y, d_pos_z, d_lpcolinear_list, lpcolinear_list_size, stream);
  }
  if (lpbisector_list_size > 0) {
    repositionBisector(d_pos_x, d_pos_y, d_pos_z, d_lpbisector_list, lpbisector_list_size, stream);
  }
  if (lprelative_list_size > 0) {
    repositionRelative(d_pos_x, d_pos_y, d_pos_z, d_lprelative_list, lprelative_list_size, stream);
  }
}

updateAtoms()

void ComputeLonepairsCUDA::updateAtoms	(	std::vector< HomePatch *>	patchList,
		std::vector< AtomMap *> &	atomMapsList,
		const std::vector< CudaLocalRecord > &	localRecords,
		const int *	h_globalToLocalID,
		cudaStream_t	stream
	)

Prepare the device lone pair lists from lphosts in the Molecule class.

Update the SOA indexes of atoms. Should be called on atom migration steps.

Definition at line 57 of file ComputeLonepairsCUDA.C.

References allocate_device_async(), lphost::angle, lphost::atom1, lphost::atom2, lphost::atom3, lphost::atom4, copy_HtoD(), cudaCheck, deallocate_device_async(), lphost::dihedral, lphost::distance, Molecule::get_lphost(), HomePatch::getAtomList(), Patch::getNumAtoms(), globalAtomIDToSOAID(), Node::molecule, NAMD_die(), lphost::numhosts, and Node::Object().

                                                     {
  std::vector<LonepairColinear> lpcolinear_list;
  std::vector<LonepairRelative> lprelative_list;
  std::vector<LonepairBisector> lpbisector_list;
  for (size_t i = 0; i < patchList.size(); ++i) {
    HomePatch *p = patchList[i];
    const int numAtoms = p->getNumAtoms();
    const auto& atomList = p->getAtomList();
    for (int j = 0; j < numAtoms; ++j) {
      const auto mass = atomList[j].mass;
      // Iterate over all atoms with mass < 0.01 to find lone pair IDs
      if (mass < 0.01) {
        const AtomID aid = atomList[j].id;
        // Find the LP host
        Lphost *lph_ptr = Node::Object()->molecule->get_lphost(aid);
        if (lph_ptr == NULL) {
          const std::string error = "rebuildLonePairHostList: no Lphost exists for LP " + std::to_string(aid);
          NAMD_die(error.c_str());
        }
        Lphost lph = *(lph_ptr);
        if (lph.numhosts == 2) { // colinear
          lpcolinear_list.push_back(LonepairColinear{
            globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom1, true),
            globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom2, true),
            globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom3, true),
            lph.distance, lph.angle});
        } else {
          const double distance = lph.distance;
          // Pre-calculate the sine and cosine
          if (distance < 0) { // bisector
            const double dcosa = -distance * std::cos(lph.angle);
            const double dsina = -distance * std::sin(lph.angle);
            const double dsinasint = dsina * std::sin(lph.dihedral);
            const double dsinacost = dsina * std::cos(lph.dihedral);
            lpbisector_list.push_back(LonepairBisector{
              globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom1, true),
              globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom2, true),
              globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom3, true),
              globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom4, true),
              dcosa, dsinacost, dsinasint});
          } else { // relative
            const double dcosa = distance * std::cos(lph.angle);
            const double dsina = distance * std::sin(lph.angle);
            const double dsinasint = dsina * std::sin(lph.dihedral);
            const double dsinacost = dsina * std::cos(lph.dihedral);
            lprelative_list.push_back(LonepairRelative{
              globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom1, true),
              globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom2, true),
              globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom3, true),
              globalAtomIDToSOAID(atomMapsList, localRecords, h_globalToLocalID, lph.atom4, true),
              dcosa, dsinacost, dsinasint});
          }
        }
      }
    }
  }
  lpcolinear_list_size = lpcolinear_list.size();
  if (!lpcolinear_list.empty()) {
    deallocate_device_async(&d_lpcolinear_list, stream);
    allocate_device_async(&d_lpcolinear_list, lpcolinear_list_size, stream);
    copy_HtoD(lpcolinear_list.data(), d_lpcolinear_list, lpcolinear_list_size, stream);
  }
  lprelative_list_size = lprelative_list.size();
  if (!lprelative_list.empty()) {
    deallocate_device_async(&d_lprelative_list, stream);
    allocate_device_async(&d_lprelative_list, lprelative_list_size, stream);
    copy_HtoD(lprelative_list.data(), d_lprelative_list, lprelative_list_size, stream);
  }
  lpbisector_list_size = lpbisector_list.size();
  if (!lpbisector_list.empty()) {
    deallocate_device_async(&d_lpbisector_list, stream);
    allocate_device_async(&d_lpbisector_list, lpbisector_list_size, stream);
    copy_HtoD(lpbisector_list.data(), d_lpbisector_list, lpbisector_list_size, stream);
  }
  cudaCheck(cudaStreamSynchronize(stream));
}

---

The documentation for this class was generated from the following files:

• ComputeLonepairsCUDA.h
• ComputeLonepairsCUDA.C