From 8c98ff4c03ecfdf57a5ade091bc1160dadb6cdcd Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 09:36:29 -0400 Subject: [PATCH 01/13] Fix typo. --- hoomd/md/mesh/bending.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hoomd/md/mesh/bending.py b/hoomd/md/mesh/bending.py index 5ba7798edd..313684790b 100644 --- a/hoomd/md/mesh/bending.py +++ b/hoomd/md/mesh/bending.py @@ -105,7 +105,7 @@ class Helfrich(MeshPotential): mesh (:py:mod:`hoomd.mesh.Mesh`): Mesh data structure constraint. Attributes: - parameter (TypeParameter[dict]): + params (TypeParameter[dict]): The parameter of the Helfrich energy for the defined mesh. As the mesh can only have one type a type name does not have to be stated. The dictionary has the following keys: From fb0eda18858199b7fa31cb3cec2c8a7f000b7f26 Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 09:50:49 -0400 Subject: [PATCH 02/13] Implement volume conservation param_t struct. --- hoomd/md/CMakeLists.txt | 23 +++++----- .../md/VolumeConservationMeshForceCompute.cc | 32 +++++++------- hoomd/md/VolumeConservationMeshForceCompute.h | 32 ++------------ hoomd/md/VolumeConservationMeshParameters.h | 42 +++++++++++++++++++ 4 files changed, 71 insertions(+), 58 deletions(-) create mode 100644 hoomd/md/VolumeConservationMeshParameters.h diff --git a/hoomd/md/CMakeLists.txt b/hoomd/md/CMakeLists.txt index 6490e74dc8..2b35995074 100644 --- a/hoomd/md/CMakeLists.txt +++ b/hoomd/md/CMakeLists.txt @@ -56,7 +56,7 @@ set(_md_sources module-md.cc Thermostat.cc TwoStepNVTAlchemy.cc WallData.cc - VolumeConservationMeshForceCompute.cc + VolumeConservationMeshForceCompute.cc ZeroMomentumUpdater.cc ) @@ -98,7 +98,7 @@ set(_md_headers ActiveForceComputeGPU.h EvaluatorSpecialPairLJ.h EvaluatorSpecialPairCoulomb.h EvaluatorExternalElectricField.h - EvaluatorExternalMagneticField.h + EvaluatorExternalMagneticField.h EvaluatorExternalPeriodic.h EvaluatorPairALJ.h EvaluatorPairBuckingham.h @@ -141,8 +141,8 @@ set(_md_headers ActiveForceComputeGPU.h HarmonicDihedralForceCompute.h HarmonicImproperForceComputeGPU.h HarmonicImproperForceCompute.h - HelfrichMeshForceComputeGPU.h - HelfrichMeshForceCompute.h + HelfrichMeshForceComputeGPU.h + HelfrichMeshForceCompute.h IntegrationMethodTwoStep.h IntegratorTwoStep.h ManifoldZCylinder.h @@ -217,8 +217,9 @@ set(_md_headers ActiveForceComputeGPU.h TwoStepConstantPressure.h AlchemostatTwoStep.h TwoStepNVTAlchemy.h - VolumeConservationMeshForceCompute.h - VolumeConservationMeshForceComputeGPU.h + VolumeConservationMeshParameters.h + VolumeConservationMeshForceCompute.h + VolumeConservationMeshForceComputeGPU.h AlchemostatTwoStep.h WallData.h ZeroMomentumUpdater.h @@ -243,7 +244,7 @@ list(APPEND _md_sources ActiveForceComputeGPU.cc HarmonicAngleForceComputeGPU.cc HarmonicDihedralForceComputeGPU.cc HarmonicImproperForceComputeGPU.cc - HelfrichMeshForceComputeGPU.cc + HelfrichMeshForceComputeGPU.cc MolecularForceCompute.cu NeighborListGPU.cc NeighborListGPUBinned.cc @@ -257,7 +258,7 @@ list(APPEND _md_sources ActiveForceComputeGPU.cc TriangleAreaConservationMeshForceComputeGPU.cc TwoStepBDGPU.cc TwoStepLangevinGPU.cc - VolumeConservationMeshForceComputeGPU.cc + VolumeConservationMeshForceComputeGPU.cc TwoStepConstantVolumeGPU.cc TwoStepConstantPressureGPU.cc MuellerPlatheFlowGPU.cc @@ -271,7 +272,7 @@ set(_md_cu_sources ActiveForceComputeGPU.cu AnisoPotentialPairALJ3GPUKernel.cu AnisoPotentialPairDipoleGPUKernel.cu AnisoPotentialPairGBGPUKernel.cu - BendingRigidityMeshForceComputeGPU.cu + BendingRigidityMeshForceComputeGPU.cu ComputeThermoGPU.cu ComputeThermoHMAGPU.cu ConstantForceComputeGPU.cu @@ -283,7 +284,7 @@ set(_md_cu_sources ActiveForceComputeGPU.cu HarmonicAngleForceGPU.cu HarmonicDihedralForceGPU.cu HarmonicImproperForceGPU.cu - HelfrichMeshForceComputeGPU.cu + HelfrichMeshForceComputeGPU.cu MolecularForceCompute.cu NeighborListGPUBinned.cu NeighborListGPU.cu @@ -302,7 +303,7 @@ set(_md_cu_sources ActiveForceComputeGPU.cu TwoStepConstantVolumeGPU.cu TwoStepNVEGPU.cu TwoStepRATTLENVEGPU.cu - VolumeConservationMeshForceComputeGPU.cu + VolumeConservationMeshForceComputeGPU.cu MuellerPlatheFlowGPU.cu CosineSqAngleForceGPU.cu ) diff --git a/hoomd/md/VolumeConservationMeshForceCompute.cc b/hoomd/md/VolumeConservationMeshForceCompute.cc index 7b9f2594da..5350203b3b 100644 --- a/hoomd/md/VolumeConservationMeshForceCompute.cc +++ b/hoomd/md/VolumeConservationMeshForceCompute.cc @@ -34,7 +34,7 @@ VolumeConservationMeshForceCompute::VolumeConservationMeshForceCompute( if (m_ignore_type) n_types = 1; - GPUArray params(n_types, m_exec_conf); + GPUArray params(n_types, m_exec_conf); m_params.swap(params); GPUArray volume(n_types, m_exec_conf); @@ -52,16 +52,16 @@ VolumeConservationMeshForceCompute::~VolumeConservationMeshForceCompute() Sets parameters for the potential of a particular angle type */ -void VolumeConservationMeshForceCompute::setParams(unsigned int type, Scalar K, Scalar V0) +void VolumeConservationMeshForceCompute::setParams(unsigned int type, const volume_conservation_param_t& params) { if (!m_ignore_type || type == 0) { - ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); - h_params.data[type] = make_scalar2(K, V0); + ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); + h_params.data[type] = params; - if (K <= 0) + if (params.k <= 0) m_exec_conf->msg->warning() << "volume: specified K <= 0" << endl; - if (V0 <= 0) + if (params.V0 <= 0) m_exec_conf->msg->warning() << "volume: specified V0 <= 0" << endl; } } @@ -69,8 +69,7 @@ void VolumeConservationMeshForceCompute::setParams(unsigned int type, Scalar K, void VolumeConservationMeshForceCompute::setParamsPython(std::string type, pybind11::dict params) { auto typ = m_mesh_data->getMeshBondData()->getTypeByName(type); - auto _params = volume_conservation_params(params); - setParams(typ, _params.k, _params.V0); + setParams(typ, volume_conservation_param_t(params)); } pybind11::dict VolumeConservationMeshForceCompute::getParams(std::string type) @@ -81,11 +80,8 @@ pybind11::dict VolumeConservationMeshForceCompute::getParams(std::string type) m_exec_conf->msg->error() << "mesh.volume: Invalid mesh type specified" << endl; throw runtime_error("Error setting parameters in VolumeConservationMeshForceCompute"); } - ArrayHandle h_params(m_params, access_location::host, access_mode::read); - pybind11::dict params; - params["k"] = h_params.data[typ].x; - params["V0"] = h_params.data[typ].y; - return params; + ArrayHandle h_params(m_params, access_location::host, access_mode::read); + return h_params.data[typ].asDict(); } /*! Actually perform the force computation @@ -104,7 +100,7 @@ void VolumeConservationMeshForceCompute::computeForces(uint64_t timestep) ArrayHandle h_force(m_force, access_location::host, access_mode::overwrite); ArrayHandle h_virial(m_virial, access_location::host, access_mode::overwrite); size_t virial_pitch = m_virial.getPitch(); - ArrayHandle h_params(m_params, access_location::host, access_mode::read); + ArrayHandle h_params(m_params, access_location::host, access_mode::read); ArrayHandle h_volume(m_volume, access_location::host, access_mode::read); ArrayHandle h_triangles( @@ -179,12 +175,12 @@ void VolumeConservationMeshForceCompute::computeForces(uint64_t timestep) else triN = h_pts.data[triangle_type]; - Scalar VolDiff = h_volume.data[triangle_type] - h_params.data[triangle_type].y; + Scalar VolDiff = h_volume.data[triangle_type] - h_params.data[triangle_type].V0; - Scalar energy = h_params.data[triangle_type].x * VolDiff * VolDiff - / (6 * h_params.data[triangle_type].y * triN); + Scalar energy = h_params.data[triangle_type].k * VolDiff * VolDiff + / (6 * h_params.data[triangle_type].V0 * triN); - VolDiff = -h_params.data[triangle_type].x / h_params.data[triangle_type].y * VolDiff / 6.0; + VolDiff = -h_params.data[triangle_type].k / h_params.data[triangle_type].V0 * VolDiff / 6.0; Fa.x = VolDiff * dVol_a.x; Fa.y = VolDiff * dVol_a.y; diff --git a/hoomd/md/VolumeConservationMeshForceCompute.h b/hoomd/md/VolumeConservationMeshForceCompute.h index 64e0b7f090..b6d42a17a2 100644 --- a/hoomd/md/VolumeConservationMeshForceCompute.h +++ b/hoomd/md/VolumeConservationMeshForceCompute.h @@ -3,6 +3,7 @@ #include "hoomd/ForceCompute.h" #include "hoomd/MeshDefinition.h" +#include "VolumeConservationMeshParameters.h" #include @@ -31,33 +32,6 @@ namespace md */ class PYBIND11_EXPORT VolumeConservationMeshForceCompute : public ForceCompute { - struct volume_conservation_params - { - Scalar k; - Scalar V0; - -#ifndef __HIPCC__ - volume_conservation_params() : k(0), V0(0) { } - - volume_conservation_params(pybind11::dict params) - : k(params["k"].cast()), V0(params["V0"].cast()) - { - } - - pybind11::dict asDict() - { - pybind11::dict v; - v["k"] = k; - v["V0"] = V0; - return v; - } -#endif - } -#if HOOMD_LONGREAL_SIZE == 32 - __attribute__((aligned(4))); -#else - __attribute__((aligned(8))); -#endif public: //! Constructs the compute VolumeConservationMeshForceCompute(std::shared_ptr sysdef, @@ -68,7 +42,7 @@ class PYBIND11_EXPORT VolumeConservationMeshForceCompute : public ForceCompute virtual ~VolumeConservationMeshForceCompute(); //! Set the parameters - virtual void setParams(unsigned int type, Scalar K, Scalar V0); + virtual void setParams(unsigned int type, const volume_conservation_param_t& params); virtual void setParamsPython(std::string type, pybind11::dict params); @@ -95,7 +69,7 @@ class PYBIND11_EXPORT VolumeConservationMeshForceCompute : public ForceCompute #endif protected: - GPUArray m_params; //!< Parameters + GPUArray m_params; //!< Parameters std::shared_ptr m_mesh_data; //!< Mesh data to use in computing volume energy diff --git a/hoomd/md/VolumeConservationMeshParameters.h b/hoomd/md/VolumeConservationMeshParameters.h new file mode 100644 index 0000000000..60da432854 --- /dev/null +++ b/hoomd/md/VolumeConservationMeshParameters.h @@ -0,0 +1,42 @@ +// Copyright (c) 2009-2024 The Regents of the University of Michigan. +// Part of HOOMD-blue, released under the BSD 3-Clause License. + +#include "hoomd/HOOMDMath.h" + +#ifndef __HIPCC__ +#include +#endif + + +#pragma once + +namespace hoomd::md { + struct volume_conservation_param_t + { + Scalar k; + Scalar V0; + +#ifndef __HIPCC__ + volume_conservation_param_t() : k(0), V0(0) { } + + volume_conservation_param_t(pybind11::dict params) + : k(params["k"].cast()), V0(params["V0"].cast()) + { + } + + pybind11::dict asDict() + { + pybind11::dict v; + v["k"] = k; + v["V0"] = V0; + return v; + } +#endif + } +#if HOOMD_LONGREAL_SIZE == 32 + __attribute__((aligned(4))); +#else + __attribute__((aligned(8))); +#endif + +} From a808c258423e20c4e21353f9198cbd739fd850ce Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 10:01:38 -0400 Subject: [PATCH 03/13] Use param_t struct in TriangleAreaConservationMeshForceCompute. --- hoomd/md/CMakeLists.txt | 1 + ...riangleAreaConservationMeshForceCompute.cc | 35 +++++++--------- ...TriangleAreaConservationMeshForceCompute.h | 5 ++- .../TriangleAreaConservationMeshParameters.h | 42 +++++++++++++++++++ 4 files changed, 61 insertions(+), 22 deletions(-) create mode 100644 hoomd/md/TriangleAreaConservationMeshParameters.h diff --git a/hoomd/md/CMakeLists.txt b/hoomd/md/CMakeLists.txt index 2b35995074..4f463123e3 100644 --- a/hoomd/md/CMakeLists.txt +++ b/hoomd/md/CMakeLists.txt @@ -195,6 +195,7 @@ set(_md_headers ActiveForceComputeGPU.h TableAngleForceCompute.h TableDihedralForceComputeGPU.h TableDihedralForceCompute.h + TriangleAreaConservationMeshParameters.h TriangleAreaConservationMeshForceCompute.h TriangleAreaConservationMeshForceComputeGPU.h TriangleAreaConservationMeshForceComputeGPU.cuh diff --git a/hoomd/md/TriangleAreaConservationMeshForceCompute.cc b/hoomd/md/TriangleAreaConservationMeshForceCompute.cc index 3358ed94d9..753ef33458 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceCompute.cc +++ b/hoomd/md/TriangleAreaConservationMeshForceCompute.cc @@ -30,7 +30,7 @@ TriangleAreaConservationMeshForceCompute::TriangleAreaConservationMeshForceCompu unsigned int n_types = m_mesh_data->getMeshTriangleData()->getNTypes(); - GPUArray params(n_types, m_exec_conf); + GPUArray params(n_types, m_exec_conf); m_params.swap(params); GPUArray area(n_types, m_exec_conf); @@ -48,17 +48,16 @@ TriangleAreaConservationMeshForceCompute::~TriangleAreaConservationMeshForceComp Sets parameters for the potential of a particular mesh type */ -void TriangleAreaConservationMeshForceCompute::setParams(unsigned int type, Scalar K, Scalar A0) +void TriangleAreaConservationMeshForceCompute::setParams(unsigned int type, const triangle_area_conservation_param_t& params) { - ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); - // update the local copy of the memory - h_params.data[type] = make_scalar2(K, A0); + ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); + h_params.data[type] = params; // check for some silly errors a user could make - if (K <= 0) + if (params.k <= 0) m_exec_conf->msg->warning() << "TriangleAreaConservation: specified K <= 0" << endl; - if (A0 <= 0) + if (params.A0 <= 0) m_exec_conf->msg->warning() << "TriangleAreaConservation: specified A0 <= 0" << endl; } @@ -66,8 +65,7 @@ void TriangleAreaConservationMeshForceCompute::setParamsPython(std::string type, pybind11::dict params) { auto typ = m_mesh_data->getMeshTriangleData()->getTypeByName(type); - auto _params = area_conservation_params(params); - setParams(typ, _params.k, _params.A0); + setParams(typ, triangle_area_conservation_param_t(params)); } pybind11::dict TriangleAreaConservationMeshForceCompute::getParams(std::string type) @@ -79,11 +77,8 @@ pybind11::dict TriangleAreaConservationMeshForceCompute::getParams(std::string t << endl; throw runtime_error("Error setting parameters in TriangleAreaConservationMeshForceCompute"); } - ArrayHandle h_params(m_params, access_location::host, access_mode::read); - pybind11::dict params; - params["k"] = h_params.data[typ].x; - params["A0"] = h_params.data[typ].y; - return params; + ArrayHandle h_params(m_params, access_location::host, access_mode::read); + return h_params.data[typ].asDict(); } /*! Actually perform the force computation @@ -100,7 +95,7 @@ void TriangleAreaConservationMeshForceCompute::computeForces(uint64_t timestep) ArrayHandle h_force(m_force, access_location::host, access_mode::overwrite); ArrayHandle h_virial(m_virial, access_location::host, access_mode::overwrite); size_t virial_pitch = m_virial.getPitch(); - ArrayHandle h_params(m_params, access_location::host, access_mode::read); + ArrayHandle h_params(m_params, access_location::host, access_mode::read); ArrayHandle h_triangles( m_mesh_data->getMeshTriangleData()->getMembersArray(), @@ -191,12 +186,12 @@ void TriangleAreaConservationMeshForceCompute::computeForces(uint64_t timestep) unsigned int triangle_type = m_mesh_data->getMeshTriangleData()->getTypeByIndex(i); - Scalar At = h_params.data[triangle_type].y; + Scalar At = h_params.data[triangle_type].A0; Scalar tri_area = rab * rac * s_baac / 6; // triangle area/3 Scalar Ut = 3 * tri_area - At; - Scalar prefactor = h_params.data[triangle_type].x / (2 * At) * Ut; + Scalar prefactor = h_params.data[triangle_type].k / (2 * At) * Ut; Scalar3 Fab, Fac; Fab = prefactor * (-nab * rac * s_baac + ds_drab * rab * rac); @@ -225,7 +220,7 @@ void TriangleAreaConservationMeshForceCompute::computeForces(uint64_t timestep) h_force.data[idx_a].x += (Fab.x + Fac.x); h_force.data[idx_a].y += (Fab.y + Fac.y); h_force.data[idx_a].z += (Fab.z + Fac.z); - h_force.data[idx_a].w += h_params.data[triangle_type].x / (6.0 * At) * Ut + h_force.data[idx_a].w += h_params.data[triangle_type].k / (6.0 * At) * Ut * Ut; // divided by 3 because of three // particles sharing the energy @@ -250,7 +245,7 @@ void TriangleAreaConservationMeshForceCompute::computeForces(uint64_t timestep) h_force.data[idx_b].x -= Fab.x; h_force.data[idx_b].y -= Fab.y; h_force.data[idx_b].z -= Fab.z; - h_force.data[idx_b].w += h_params.data[triangle_type].x / (6.0 * At) * Ut * Ut; + h_force.data[idx_b].w += h_params.data[triangle_type].k / (6.0 * At) * Ut * Ut; for (int j = 0; j < 6; j++) h_virial.data[j * virial_pitch + idx_b] += triangle_area_conservation_virial[j]; } @@ -272,7 +267,7 @@ void TriangleAreaConservationMeshForceCompute::computeForces(uint64_t timestep) h_force.data[idx_c].x -= Fac.x; h_force.data[idx_c].y -= Fac.y; h_force.data[idx_c].z -= Fac.z; - h_force.data[idx_c].w += h_params.data[triangle_type].x / (6.0 * At) * Ut * Ut; + h_force.data[idx_c].w += h_params.data[triangle_type].k / (6.0 * At) * Ut * Ut; for (int j = 0; j < 6; j++) h_virial.data[j * virial_pitch + idx_c] += triangle_area_conservation_virial[j]; } diff --git a/hoomd/md/TriangleAreaConservationMeshForceCompute.h b/hoomd/md/TriangleAreaConservationMeshForceCompute.h index b74e64874b..9fd68f6b59 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceCompute.h +++ b/hoomd/md/TriangleAreaConservationMeshForceCompute.h @@ -3,6 +3,7 @@ #include "hoomd/ForceCompute.h" #include "hoomd/MeshDefinition.h" +#include "TriangleAreaConservationMeshParameters.h" #include @@ -39,7 +40,7 @@ class PYBIND11_EXPORT TriangleAreaConservationMeshForceCompute : public ForceCom virtual ~TriangleAreaConservationMeshForceCompute(); //! Set the parameters - virtual void setParams(unsigned int type, Scalar K, Scalar A0); + virtual void setParams(unsigned int type, const triangle_area_conservation_param_t& params); virtual void setParamsPython(std::string type, pybind11::dict params); @@ -66,7 +67,7 @@ class PYBIND11_EXPORT TriangleAreaConservationMeshForceCompute : public ForceCom #endif protected: - GPUArray m_params; //!< Parameters + GPUArray m_params; //!< Parameters GPUArray m_area; //!< memory space for area std::shared_ptr diff --git a/hoomd/md/TriangleAreaConservationMeshParameters.h b/hoomd/md/TriangleAreaConservationMeshParameters.h new file mode 100644 index 0000000000..60f50c7bad --- /dev/null +++ b/hoomd/md/TriangleAreaConservationMeshParameters.h @@ -0,0 +1,42 @@ +// Copyright (c) 2009-2024 The Regents of the University of Michigan. +// Part of HOOMD-blue, released under the BSD 3-Clause License. + +#include "hoomd/HOOMDMath.h" + +#ifndef __HIPCC__ +#include +#endif + + +#pragma once + +namespace hoomd::md { + struct triangle_area_conservation_param_t + { + Scalar k; + Scalar A0; + +#ifndef __HIPCC__ + triangle_area_conservation_param_t() : k(0), A0(0) { } + + triangle_area_conservation_param_t(pybind11::dict params) + : k(params["k"].cast()), A0(params["A0"].cast()) + { + } + + pybind11::dict asDict() + { + pybind11::dict v; + v["k"] = k; + v["A0"] = A0; + return v; + } +#endif + } +#if HOOMD_LONGREAL_SIZE == 32 + __attribute__((aligned(4))); +#else + __attribute__((aligned(8))); +#endif + +} From e1297e7809d1d73c9c957d2577037cab4ef11307 Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 10:10:10 -0400 Subject: [PATCH 04/13] Use param_t struct in volume conservation. --- hoomd/md/AreaConservationMeshForceCompute.cc | 37 +++++++--------- hoomd/md/AreaConservationMeshForceCompute.h | 33 ++------------- hoomd/md/AreaConservationMeshParameters.h | 42 +++++++++++++++++++ hoomd/md/BendingRigidityMeshForceCompute.cc | 1 - hoomd/md/CMakeLists.txt | 1 + ...riangleAreaConservationMeshForceCompute.cc | 4 +- .../md/VolumeConservationMeshForceCompute.cc | 3 +- 7 files changed, 63 insertions(+), 58 deletions(-) create mode 100644 hoomd/md/AreaConservationMeshParameters.h diff --git a/hoomd/md/AreaConservationMeshForceCompute.cc b/hoomd/md/AreaConservationMeshForceCompute.cc index 89dd22c1ae..51a4bdf1be 100644 --- a/hoomd/md/AreaConservationMeshForceCompute.cc +++ b/hoomd/md/AreaConservationMeshForceCompute.cc @@ -34,7 +34,7 @@ AreaConservationMeshForceCompute::AreaConservationMeshForceCompute( if (m_ignore_type) n_types = 1; - GPUArray params(n_types, m_exec_conf); + GPUArray params(n_types, m_exec_conf); m_params.swap(params); GPUArray area(n_types, m_exec_conf); @@ -47,23 +47,20 @@ AreaConservationMeshForceCompute::~AreaConservationMeshForceCompute() } /*! \param type Type of the angle to set parameters for - \param K Stiffness parameter for the force computation - \param A0 desired surface area to maintain for the force computation + \param params Parameters to set Sets parameters for the potential of a particular mesh type */ -void AreaConservationMeshForceCompute::setParams(unsigned int type, Scalar K, Scalar A0) +void AreaConservationMeshForceCompute::setParams(unsigned int type, const area_conservation_param_t ¶ms) { if (!m_ignore_type || type == 0) { - ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); - // update the local copy of the memory - h_params.data[type] = make_scalar2(K, A0); + ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); + h_params.data[type] = params; - // check for some silly errors a user could make - if (K <= 0) + if (params.k <= 0) m_exec_conf->msg->warning() << "area: specified K <= 0" << endl; - if (A0 <= 0) + if (params.A0 <= 0) m_exec_conf->msg->warning() << "area: specified A0 <= 0" << endl; } } @@ -71,8 +68,7 @@ void AreaConservationMeshForceCompute::setParams(unsigned int type, Scalar K, Sc void AreaConservationMeshForceCompute::setParamsPython(std::string type, pybind11::dict params) { auto typ = m_mesh_data->getMeshBondData()->getTypeByName(type); - auto _params = area_conservation_params(params); - setParams(typ, _params.k, _params.A0); + setParams(typ, area_conservation_param_t(params)); } pybind11::dict AreaConservationMeshForceCompute::getParams(std::string type) @@ -85,11 +81,8 @@ pybind11::dict AreaConservationMeshForceCompute::getParams(std::string type) } if (m_ignore_type) typ = 0; - ArrayHandle h_params(m_params, access_location::host, access_mode::read); - pybind11::dict params; - params["k"] = h_params.data[typ].x; - params["A0"] = h_params.data[typ].y; - return params; + ArrayHandle h_params(m_params, access_location::host, access_mode::read); + return h_params.data[typ].asDict(); } /*! Actually perform the force computation @@ -107,7 +100,7 @@ void AreaConservationMeshForceCompute::computeForces(uint64_t timestep) ArrayHandle h_force(m_force, access_location::host, access_mode::overwrite); ArrayHandle h_virial(m_virial, access_location::host, access_mode::overwrite); size_t virial_pitch = m_virial.getPitch(); - ArrayHandle h_params(m_params, access_location::host, access_mode::read); + ArrayHandle h_params(m_params, access_location::host, access_mode::read); ArrayHandle h_area(m_area, access_location::host, access_mode::read); ArrayHandle h_triangles( @@ -205,12 +198,12 @@ void AreaConservationMeshForceCompute::computeForces(uint64_t timestep) else triN = h_pts.data[triangle_type]; - Scalar AreaDiff = h_area.data[triangle_type] - h_params.data[triangle_type].y; + Scalar AreaDiff = h_area.data[triangle_type] - h_params.data[triangle_type].A0; - Scalar energy = h_params.data[triangle_type].x * AreaDiff * AreaDiff - / (6 * h_params.data[triangle_type].y * triN); + Scalar energy = h_params.data[triangle_type].k * AreaDiff * AreaDiff + / (6 * h_params.data[triangle_type].A0 * triN); - AreaDiff = h_params.data[triangle_type].x / h_params.data[triangle_type].y * AreaDiff / 2.0; + AreaDiff = h_params.data[triangle_type].k / h_params.data[triangle_type].A0 * AreaDiff / 2.0; Fab = AreaDiff * (-nab * rac * s_baac + ds_drab * rab * rac); Fac = AreaDiff * (-nac * rab * s_baac + ds_drac * rab * rac); diff --git a/hoomd/md/AreaConservationMeshForceCompute.h b/hoomd/md/AreaConservationMeshForceCompute.h index ded14009fd..7cdecb0795 100644 --- a/hoomd/md/AreaConservationMeshForceCompute.h +++ b/hoomd/md/AreaConservationMeshForceCompute.h @@ -3,6 +3,7 @@ #include "hoomd/ForceCompute.h" #include "hoomd/MeshDefinition.h" +#include "AreaConservationMeshParameters.h" #include @@ -23,34 +24,6 @@ namespace hoomd { namespace md { -struct area_conservation_params - { - Scalar k; - Scalar A0; - -#ifndef __HIPCC__ - area_conservation_params() : k(0), A0(0) { } - - area_conservation_params(pybind11::dict params) - : k(params["k"].cast()), A0(params["A0"].cast()) - { - } - - pybind11::dict asDict() - { - pybind11::dict v; - v["k"] = k; - v["A0"] = A0; - return v; - } -#endif - } -#if HOOMD_LONGREAL_SIZE == 32 - __attribute__((aligned(4))); -#else - __attribute__((aligned(8))); -#endif - //! Computes area constraint forces on the mesh /*! Area constraint forces are computed on every particle in a mesh. @@ -68,7 +41,7 @@ class PYBIND11_EXPORT AreaConservationMeshForceCompute : public ForceCompute virtual ~AreaConservationMeshForceCompute(); //! Set the parameters - virtual void setParams(unsigned int type, Scalar K, Scalar A0); + virtual void setParams(unsigned int type, const area_conservation_param_t& params); virtual void setParamsPython(std::string type, pybind11::dict params); @@ -95,7 +68,7 @@ class PYBIND11_EXPORT AreaConservationMeshForceCompute : public ForceCompute #endif protected: - GPUArray m_params; //!< Parameters + GPUArray m_params; //!< Parameters GPUArray m_area; //!< memory space for area // std::shared_ptr m_mesh_data; //!< Mesh data to use in computing energy diff --git a/hoomd/md/AreaConservationMeshParameters.h b/hoomd/md/AreaConservationMeshParameters.h new file mode 100644 index 0000000000..53ab62ca3f --- /dev/null +++ b/hoomd/md/AreaConservationMeshParameters.h @@ -0,0 +1,42 @@ +// Copyright (c) 2009-2024 The Regents of the University of Michigan. +// Part of HOOMD-blue, released under the BSD 3-Clause License. + +#include "hoomd/HOOMDMath.h" + +#ifndef __HIPCC__ +#include +#endif + + +#pragma once + +namespace hoomd::md { + struct area_conservation_param_t + { + Scalar k; + Scalar A0; + +#ifndef __HIPCC__ + area_conservation_param_t() : k(0), A0(0) { } + + area_conservation_param_t(pybind11::dict params) + : k(params["k"].cast()), A0(params["A0"].cast()) + { + } + + pybind11::dict asDict() + { + pybind11::dict v; + v["k"] = k; + v["A0"] = A0; + return v; + } +#endif + } +#if HOOMD_LONGREAL_SIZE == 32 + __attribute__((aligned(4))); +#else + __attribute__((aligned(8))); +#endif + +} diff --git a/hoomd/md/BendingRigidityMeshForceCompute.cc b/hoomd/md/BendingRigidityMeshForceCompute.cc index 1cfec414f0..a639f33af5 100644 --- a/hoomd/md/BendingRigidityMeshForceCompute.cc +++ b/hoomd/md/BendingRigidityMeshForceCompute.cc @@ -47,7 +47,6 @@ void BendingRigidityMeshForceCompute::setParams(unsigned int type, Scalar K) ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); h_params.data[type] = K; - // check for some silly errors a user could make if (K <= 0) m_exec_conf->msg->warning() << "rigidity: specified K <= 0" << endl; } diff --git a/hoomd/md/CMakeLists.txt b/hoomd/md/CMakeLists.txt index 4f463123e3..e44ec847f5 100644 --- a/hoomd/md/CMakeLists.txt +++ b/hoomd/md/CMakeLists.txt @@ -70,6 +70,7 @@ set(_md_headers ActiveForceComputeGPU.h AnisoPotentialPairGPU.cuh AnisoPotentialPairGPU.h AnisoPotentialPair.h + AreaConservationMeshParameters.h AreaConservationMeshForceCompute.h AreaConservationMeshForceComputeGPU.h AreaConservationMeshForceComputeGPU.cuh diff --git a/hoomd/md/TriangleAreaConservationMeshForceCompute.cc b/hoomd/md/TriangleAreaConservationMeshForceCompute.cc index 753ef33458..36977647e8 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceCompute.cc +++ b/hoomd/md/TriangleAreaConservationMeshForceCompute.cc @@ -43,8 +43,7 @@ TriangleAreaConservationMeshForceCompute::~TriangleAreaConservationMeshForceComp } /*! \param type Type of the angle to set parameters for - \param K Stiffness parameter for the force computation - \param A0 desired surface area to maintain for the force computation + \param params Parameters to set. Sets parameters for the potential of a particular mesh type */ @@ -53,7 +52,6 @@ void TriangleAreaConservationMeshForceCompute::setParams(unsigned int type, cons ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); h_params.data[type] = params; - // check for some silly errors a user could make if (params.k <= 0) m_exec_conf->msg->warning() << "TriangleAreaConservation: specified K <= 0" << endl; diff --git a/hoomd/md/VolumeConservationMeshForceCompute.cc b/hoomd/md/VolumeConservationMeshForceCompute.cc index 5350203b3b..857b9c2a69 100644 --- a/hoomd/md/VolumeConservationMeshForceCompute.cc +++ b/hoomd/md/VolumeConservationMeshForceCompute.cc @@ -47,8 +47,7 @@ VolumeConservationMeshForceCompute::~VolumeConservationMeshForceCompute() } /*! \param type Type of the angle to set parameters for - \param K Stiffness parameter for the force computation - \param V0 desired volume to maintain for the force computation + \param params Parameters to set. Sets parameters for the potential of a particular angle type */ From 70ecd90609210dbc8b9b3911f3eb47b58fce8914 Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 10:38:41 -0400 Subject: [PATCH 05/13] Use param_t on the GPU. --- hoomd/md/AreaConservationMeshForceComputeGPU.cc | 2 +- hoomd/md/AreaConservationMeshForceComputeGPU.cu | 11 +++++------ hoomd/md/AreaConservationMeshForceComputeGPU.cuh | 3 ++- .../md/TriangleAreaConservationMeshForceComputeGPU.cc | 7 +------ .../md/TriangleAreaConservationMeshForceComputeGPU.cu | 10 +++++----- .../TriangleAreaConservationMeshForceComputeGPU.cuh | 3 ++- hoomd/md/VolumeConservationMeshForceComputeGPU.cc | 6 +----- hoomd/md/VolumeConservationMeshForceComputeGPU.cu | 11 +++++------ hoomd/md/VolumeConservationMeshForceComputeGPU.cuh | 3 ++- 9 files changed, 24 insertions(+), 32 deletions(-) diff --git a/hoomd/md/AreaConservationMeshForceComputeGPU.cc b/hoomd/md/AreaConservationMeshForceComputeGPU.cc index 87d360cf84..5ca15c7611 100644 --- a/hoomd/md/AreaConservationMeshForceComputeGPU.cc +++ b/hoomd/md/AreaConservationMeshForceComputeGPU.cc @@ -89,7 +89,7 @@ void AreaConservationMeshForceComputeGPU::computeForces(uint64_t timestep) ArrayHandle d_force(m_force, access_location::device, access_mode::overwrite); ArrayHandle d_virial(m_virial, access_location::device, access_mode::overwrite); - ArrayHandle d_params(m_params, access_location::device, access_mode::read); + ArrayHandle d_params(m_params, access_location::device, access_mode::read); ArrayHandle d_area(m_area, access_location::device, access_mode::read); diff --git a/hoomd/md/AreaConservationMeshForceComputeGPU.cu b/hoomd/md/AreaConservationMeshForceComputeGPU.cu index 45234f76d8..62e5d4626d 100644 --- a/hoomd/md/AreaConservationMeshForceComputeGPU.cu +++ b/hoomd/md/AreaConservationMeshForceComputeGPU.cu @@ -283,7 +283,7 @@ __global__ void gpu_compute_area_constraint_force_kernel(Scalar4* d_force, const unsigned int* tpos_list, const Index2D tlist_idx, const unsigned int* n_triangles_list, - Scalar2* d_params, + area_conservation_param_t* d_params, const bool ignore_type) { // start by identifying which particle we are to handle @@ -322,10 +322,9 @@ __global__ void gpu_compute_area_constraint_force_kernel(Scalar4* d_force, else triN = gN[cur_triangle_type]; - // get the angle parameters (MEM TRANSFER: 8 bytes) - Scalar2 params = __ldg(d_params + cur_triangle_type); - Scalar K = params.x; - Scalar A_mesh = params.y; + area_conservation_param_t params = d_params[cur_triangle_type]; + Scalar K = params.k; + Scalar A_mesh = params.A0; Scalar AreaDiff = area[cur_triangle_type] - A_mesh; @@ -455,7 +454,7 @@ hipError_t gpu_compute_area_constraint_force(Scalar4* d_force, const unsigned int* tpos_list, const Index2D tlist_idx, const unsigned int* n_triangles_list, - Scalar2* d_params, + area_conservation_param_t* d_params, const bool ignore_type, int block_size) { diff --git a/hoomd/md/AreaConservationMeshForceComputeGPU.cuh b/hoomd/md/AreaConservationMeshForceComputeGPU.cuh index bc46d6923f..b66824d5cd 100644 --- a/hoomd/md/AreaConservationMeshForceComputeGPU.cuh +++ b/hoomd/md/AreaConservationMeshForceComputeGPU.cuh @@ -5,6 +5,7 @@ #include "hoomd/HOOMDMath.h" #include "hoomd/Index1D.h" #include "hoomd/ParticleData.cuh" +#include "AreaConservationMeshParameters.h" #include /*! \file MeshAreaConservationGPU.cuh @@ -50,7 +51,7 @@ hipError_t gpu_compute_area_constraint_force(Scalar4* d_force, const unsigned int* tpos_list, const Index2D tlist_idx, const unsigned int* n_triangles_list, - Scalar2* d_params, + area_conservation_param_t* d_params, const bool ignore_type, int block_size); } // end namespace kernel diff --git a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cc b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cc index 10f15213c4..5335c20605 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cc +++ b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cc @@ -31,11 +31,6 @@ TriangleAreaConservationMeshForceComputeGPU::TriangleAreaConservationMeshForceCo throw std::runtime_error("Error initializing TriangleAreaConservationMeshForceComputeGPU"); } - // allocate and zero device memory - GPUArray params(this->m_mesh_data->getMeshTriangleData()->getNTypes(), - this->m_exec_conf); - m_params.swap(params); - GPUArray sum(this->m_mesh_data->getMeshTriangleData()->getNTypes(), m_exec_conf); m_sum.swap(sum); @@ -81,7 +76,7 @@ void TriangleAreaConservationMeshForceComputeGPU::computeForces(uint64_t timeste ArrayHandle d_force(m_force, access_location::device, access_mode::overwrite); ArrayHandle d_virial(m_virial, access_location::device, access_mode::overwrite); - ArrayHandle d_params(m_params, access_location::device, access_mode::read); + ArrayHandle d_params(m_params, access_location::device, access_mode::read); m_tuner->begin(); kernel::gpu_compute_TriangleAreaConservation_force( diff --git a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cu b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cu index 1cbf5d6bb7..1f93e2cecb 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cu +++ b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cu @@ -42,7 +42,7 @@ gpu_compute_TriangleAreaConservation_force_kernel(Scalar4* d_force, const unsigned int* tpos_list, const Index2D tlist_idx, const unsigned int* n_triangles_list, - Scalar2* d_params, + triangle_area_conservation_param_t* d_params, const unsigned int n_triangle_type) { // start by identifying which particle we are to handle @@ -121,9 +121,9 @@ gpu_compute_TriangleAreaConservation_force_kernel(Scalar4* d_force, Scalar s_baac = sqrt(1.0 - c_baac * c_baac); Scalar inv_s_baac = 1.0 / s_baac; - Scalar2 params = __ldg(d_params + cur_triangle_type); - Scalar K = params.x; - Scalar At = params.y; + triangle_area_conservation_param_t params = d_params[cur_triangle_type]; + Scalar K = params.k; + Scalar At = params.A0; Scalar3 dc_drab = -nac / rab + c_baac / rab * nab; Scalar3 ds_drab = -c_baac * inv_s_baac * dc_drab; @@ -202,7 +202,7 @@ hipError_t gpu_compute_TriangleAreaConservation_force(Scalar4* d_force, const unsigned int* tpos_list, const Index2D tlist_idx, const unsigned int* n_triangles_list, - Scalar2* d_params, + triangle_area_conservation_param_t* d_params, const unsigned int n_triangle_type, int block_size) { diff --git a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cuh b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cuh index 8fac487214..6e287b782f 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cuh +++ b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cuh @@ -4,6 +4,7 @@ #include "hoomd/BondedGroupData.cuh" #include "hoomd/HOOMDMath.h" #include "hoomd/Index1D.h" +#include "TriangleAreaConservationMeshParameters.h" #include "hoomd/ParticleData.cuh" /*! \file TriangleAreaConservationMeshForceComputeGPU.cuh @@ -31,7 +32,7 @@ hipError_t gpu_compute_TriangleAreaConservation_force(Scalar4* d_force, const unsigned int* tpos_list, const Index2D tlist_idx, const unsigned int* n_triangles_list, - Scalar2* d_params, + triangle_area_conservation_param_t* d_params, const unsigned int n_triangle_type, int block_size); } // end namespace kernel diff --git a/hoomd/md/VolumeConservationMeshForceComputeGPU.cc b/hoomd/md/VolumeConservationMeshForceComputeGPU.cc index 8c95336fd7..9fbf8e3726 100644 --- a/hoomd/md/VolumeConservationMeshForceComputeGPU.cc +++ b/hoomd/md/VolumeConservationMeshForceComputeGPU.cc @@ -37,10 +37,6 @@ VolumeConservationMeshForceComputeGPU::VolumeConservationMeshForceComputeGPU( if (this->m_ignore_type) NTypes = 1; - // allocate and zero device memory - GPUArray params(NTypes, m_exec_conf); - m_params.swap(params); - GPUArray sum(NTypes, m_exec_conf); m_sum.swap(sum); @@ -94,7 +90,7 @@ void VolumeConservationMeshForceComputeGPU::computeForces(uint64_t timestep) ArrayHandle d_force(m_force, access_location::device, access_mode::overwrite); ArrayHandle d_virial(m_virial, access_location::device, access_mode::overwrite); - ArrayHandle d_params(m_params, access_location::device, access_mode::read); + ArrayHandle d_params(m_params, access_location::device, access_mode::read); ArrayHandle d_volume(m_volume, access_location::device, access_mode::read); diff --git a/hoomd/md/VolumeConservationMeshForceComputeGPU.cu b/hoomd/md/VolumeConservationMeshForceComputeGPU.cu index 08e7467212..3c9b95ab26 100644 --- a/hoomd/md/VolumeConservationMeshForceComputeGPU.cu +++ b/hoomd/md/VolumeConservationMeshForceComputeGPU.cu @@ -277,7 +277,7 @@ __global__ void gpu_compute_volume_constraint_force_kernel(Scalar4* d_force, const unsigned int* tpos_list, const Index2D tlist_idx, const unsigned int* n_triangles_list, - Scalar2* d_params, + volume_conservation_param_t* d_params, const bool ignore_type) { // start by identifying which particle we are to handle @@ -318,10 +318,9 @@ __global__ void gpu_compute_volume_constraint_force_kernel(Scalar4* d_force, else triN = gN[cur_triangle_type]; - // get the angle parameters (MEM TRANSFER: 8 bytes) - Scalar2 params = __ldg(d_params + cur_triangle_type); - Scalar K = params.x; - Scalar V0 = params.y; + volume_conservation_param_t params = d_params[cur_triangle_type]; + Scalar K = params.k; + Scalar V0 = params.V0; Scalar VolDiff = volume[cur_triangle_type] - V0; @@ -417,7 +416,7 @@ hipError_t gpu_compute_volume_constraint_force(Scalar4* d_force, const unsigned int* tpos_list, const Index2D tlist_idx, const unsigned int* n_triangles_list, - Scalar2* d_params, + volume_conservation_param_t* d_params, const bool ignore_type, int block_size) { diff --git a/hoomd/md/VolumeConservationMeshForceComputeGPU.cuh b/hoomd/md/VolumeConservationMeshForceComputeGPU.cuh index 36088c1f7d..2f12b62a4d 100644 --- a/hoomd/md/VolumeConservationMeshForceComputeGPU.cuh +++ b/hoomd/md/VolumeConservationMeshForceComputeGPU.cuh @@ -5,6 +5,7 @@ #include "hoomd/HOOMDMath.h" #include "hoomd/Index1D.h" #include "hoomd/ParticleData.cuh" +#include "VolumeConservationMeshParameters.h" #include /*! \file MeshVolumeConservationGPU.cuh @@ -52,7 +53,7 @@ hipError_t gpu_compute_volume_constraint_force(Scalar4* d_force, const unsigned int* tpos_list, const Index2D tlist_idx, const unsigned int* n_triangles_list, - Scalar2* d_params, + volume_conservation_param_t* d_params, const bool ignore_type, int block_size); } // end namespace kernel From c7c707648f3dffd39599d84508f5bd8ae252ca4b Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 10:40:01 -0400 Subject: [PATCH 06/13] pre-commit --- hoomd/md/AreaConservationMeshForceCompute.cc | 18 +++++-- hoomd/md/AreaConservationMeshForceCompute.h | 12 ++--- .../md/AreaConservationMeshForceComputeGPU.cc | 4 +- .../AreaConservationMeshForceComputeGPU.cuh | 2 +- hoomd/md/AreaConservationMeshParameters.h | 48 +++++++++---------- ...riangleAreaConservationMeshForceCompute.cc | 16 +++++-- ...TriangleAreaConservationMeshForceCompute.h | 4 +- ...ngleAreaConservationMeshForceComputeGPU.cc | 4 +- ...gleAreaConservationMeshForceComputeGPU.cuh | 2 +- .../TriangleAreaConservationMeshParameters.h | 48 +++++++++---------- .../md/VolumeConservationMeshForceCompute.cc | 17 +++++-- hoomd/md/VolumeConservationMeshForceCompute.h | 2 +- .../VolumeConservationMeshForceComputeGPU.cc | 4 +- .../VolumeConservationMeshForceComputeGPU.cuh | 2 +- hoomd/md/VolumeConservationMeshParameters.h | 48 +++++++++---------- 15 files changed, 130 insertions(+), 101 deletions(-) diff --git a/hoomd/md/AreaConservationMeshForceCompute.cc b/hoomd/md/AreaConservationMeshForceCompute.cc index 51a4bdf1be..03fab98ddb 100644 --- a/hoomd/md/AreaConservationMeshForceCompute.cc +++ b/hoomd/md/AreaConservationMeshForceCompute.cc @@ -51,11 +51,14 @@ AreaConservationMeshForceCompute::~AreaConservationMeshForceCompute() Sets parameters for the potential of a particular mesh type */ -void AreaConservationMeshForceCompute::setParams(unsigned int type, const area_conservation_param_t ¶ms) +void AreaConservationMeshForceCompute::setParams(unsigned int type, + const area_conservation_param_t& params) { if (!m_ignore_type || type == 0) { - ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); + ArrayHandle h_params(m_params, + access_location::host, + access_mode::readwrite); h_params.data[type] = params; if (params.k <= 0) @@ -81,7 +84,9 @@ pybind11::dict AreaConservationMeshForceCompute::getParams(std::string type) } if (m_ignore_type) typ = 0; - ArrayHandle h_params(m_params, access_location::host, access_mode::read); + ArrayHandle h_params(m_params, + access_location::host, + access_mode::read); return h_params.data[typ].asDict(); } @@ -100,7 +105,9 @@ void AreaConservationMeshForceCompute::computeForces(uint64_t timestep) ArrayHandle h_force(m_force, access_location::host, access_mode::overwrite); ArrayHandle h_virial(m_virial, access_location::host, access_mode::overwrite); size_t virial_pitch = m_virial.getPitch(); - ArrayHandle h_params(m_params, access_location::host, access_mode::read); + ArrayHandle h_params(m_params, + access_location::host, + access_mode::read); ArrayHandle h_area(m_area, access_location::host, access_mode::read); ArrayHandle h_triangles( @@ -203,7 +210,8 @@ void AreaConservationMeshForceCompute::computeForces(uint64_t timestep) Scalar energy = h_params.data[triangle_type].k * AreaDiff * AreaDiff / (6 * h_params.data[triangle_type].A0 * triN); - AreaDiff = h_params.data[triangle_type].k / h_params.data[triangle_type].A0 * AreaDiff / 2.0; + AreaDiff + = h_params.data[triangle_type].k / h_params.data[triangle_type].A0 * AreaDiff / 2.0; Fab = AreaDiff * (-nab * rac * s_baac + ds_drab * rab * rac); Fac = AreaDiff * (-nac * rab * s_baac + ds_drac * rab * rac); diff --git a/hoomd/md/AreaConservationMeshForceCompute.h b/hoomd/md/AreaConservationMeshForceCompute.h index 7cdecb0795..f3d72ff3c7 100644 --- a/hoomd/md/AreaConservationMeshForceCompute.h +++ b/hoomd/md/AreaConservationMeshForceCompute.h @@ -1,9 +1,9 @@ // Copyright (c) 2009-2024 The Regents of the University of Michigan. // Part of HOOMD-blue, released under the BSD 3-Clause License. +#include "AreaConservationMeshParameters.h" #include "hoomd/ForceCompute.h" #include "hoomd/MeshDefinition.h" -#include "AreaConservationMeshParameters.h" #include @@ -68,11 +68,11 @@ class PYBIND11_EXPORT AreaConservationMeshForceCompute : public ForceCompute #endif protected: - GPUArray m_params; //!< Parameters - GPUArray m_area; //!< memory space for area - // - std::shared_ptr m_mesh_data; //!< Mesh data to use in computing energy - bool m_ignore_type; //! ignore type to calculate global area if true + GPUArray m_params; //!< Parameters + GPUArray m_area; //!< memory space for area + // + std::shared_ptr m_mesh_data; //!< Mesh data to use in computing energy + bool m_ignore_type; //! ignore type to calculate global area if true //! Actually compute the forces virtual void computeForces(uint64_t timestep); diff --git a/hoomd/md/AreaConservationMeshForceComputeGPU.cc b/hoomd/md/AreaConservationMeshForceComputeGPU.cc index 5ca15c7611..b2a9af55ef 100644 --- a/hoomd/md/AreaConservationMeshForceComputeGPU.cc +++ b/hoomd/md/AreaConservationMeshForceComputeGPU.cc @@ -89,7 +89,9 @@ void AreaConservationMeshForceComputeGPU::computeForces(uint64_t timestep) ArrayHandle d_force(m_force, access_location::device, access_mode::overwrite); ArrayHandle d_virial(m_virial, access_location::device, access_mode::overwrite); - ArrayHandle d_params(m_params, access_location::device, access_mode::read); + ArrayHandle d_params(m_params, + access_location::device, + access_mode::read); ArrayHandle d_area(m_area, access_location::device, access_mode::read); diff --git a/hoomd/md/AreaConservationMeshForceComputeGPU.cuh b/hoomd/md/AreaConservationMeshForceComputeGPU.cuh index b66824d5cd..1853267e8d 100644 --- a/hoomd/md/AreaConservationMeshForceComputeGPU.cuh +++ b/hoomd/md/AreaConservationMeshForceComputeGPU.cuh @@ -1,11 +1,11 @@ // Copyright (c) 2009-2024 The Regents of the University of Michigan. // Part of HOOMD-blue, released under the BSD 3-Clause License. +#include "AreaConservationMeshParameters.h" #include "hoomd/BondedGroupData.cuh" #include "hoomd/HOOMDMath.h" #include "hoomd/Index1D.h" #include "hoomd/ParticleData.cuh" -#include "AreaConservationMeshParameters.h" #include /*! \file MeshAreaConservationGPU.cuh diff --git a/hoomd/md/AreaConservationMeshParameters.h b/hoomd/md/AreaConservationMeshParameters.h index 53ab62ca3f..ce334f97ca 100644 --- a/hoomd/md/AreaConservationMeshParameters.h +++ b/hoomd/md/AreaConservationMeshParameters.h @@ -7,36 +7,36 @@ #include #endif - #pragma once -namespace hoomd::md { - struct area_conservation_param_t - { - Scalar k; - Scalar A0; +namespace hoomd::md + { +struct area_conservation_param_t + { + Scalar k; + Scalar A0; #ifndef __HIPCC__ - area_conservation_param_t() : k(0), A0(0) { } - - area_conservation_param_t(pybind11::dict params) - : k(params["k"].cast()), A0(params["A0"].cast()) - { - } - - pybind11::dict asDict() - { - pybind11::dict v; - v["k"] = k; - v["A0"] = A0; - return v; - } -#endif + area_conservation_param_t() : k(0), A0(0) { } + + area_conservation_param_t(pybind11::dict params) + : k(params["k"].cast()), A0(params["A0"].cast()) + { } + + pybind11::dict asDict() + { + pybind11::dict v; + v["k"] = k; + v["A0"] = A0; + return v; + } +#endif + } #if HOOMD_LONGREAL_SIZE == 32 - __attribute__((aligned(4))); + __attribute__((aligned(4))); #else - __attribute__((aligned(8))); + __attribute__((aligned(8))); #endif -} + } // namespace hoomd::md diff --git a/hoomd/md/TriangleAreaConservationMeshForceCompute.cc b/hoomd/md/TriangleAreaConservationMeshForceCompute.cc index 36977647e8..437bccb800 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceCompute.cc +++ b/hoomd/md/TriangleAreaConservationMeshForceCompute.cc @@ -47,9 +47,13 @@ TriangleAreaConservationMeshForceCompute::~TriangleAreaConservationMeshForceComp Sets parameters for the potential of a particular mesh type */ -void TriangleAreaConservationMeshForceCompute::setParams(unsigned int type, const triangle_area_conservation_param_t& params) +void TriangleAreaConservationMeshForceCompute::setParams( + unsigned int type, + const triangle_area_conservation_param_t& params) { - ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); + ArrayHandle h_params(m_params, + access_location::host, + access_mode::readwrite); h_params.data[type] = params; if (params.k <= 0) @@ -75,7 +79,9 @@ pybind11::dict TriangleAreaConservationMeshForceCompute::getParams(std::string t << endl; throw runtime_error("Error setting parameters in TriangleAreaConservationMeshForceCompute"); } - ArrayHandle h_params(m_params, access_location::host, access_mode::read); + ArrayHandle h_params(m_params, + access_location::host, + access_mode::read); return h_params.data[typ].asDict(); } @@ -93,7 +99,9 @@ void TriangleAreaConservationMeshForceCompute::computeForces(uint64_t timestep) ArrayHandle h_force(m_force, access_location::host, access_mode::overwrite); ArrayHandle h_virial(m_virial, access_location::host, access_mode::overwrite); size_t virial_pitch = m_virial.getPitch(); - ArrayHandle h_params(m_params, access_location::host, access_mode::read); + ArrayHandle h_params(m_params, + access_location::host, + access_mode::read); ArrayHandle h_triangles( m_mesh_data->getMeshTriangleData()->getMembersArray(), diff --git a/hoomd/md/TriangleAreaConservationMeshForceCompute.h b/hoomd/md/TriangleAreaConservationMeshForceCompute.h index 9fd68f6b59..6cab075787 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceCompute.h +++ b/hoomd/md/TriangleAreaConservationMeshForceCompute.h @@ -1,9 +1,9 @@ // Copyright (c) 2009-2024 The Regents of the University of Michigan. // Part of HOOMD-blue, released under the BSD 3-Clause License. +#include "TriangleAreaConservationMeshParameters.h" #include "hoomd/ForceCompute.h" #include "hoomd/MeshDefinition.h" -#include "TriangleAreaConservationMeshParameters.h" #include @@ -68,7 +68,7 @@ class PYBIND11_EXPORT TriangleAreaConservationMeshForceCompute : public ForceCom protected: GPUArray m_params; //!< Parameters - GPUArray m_area; //!< memory space for area + GPUArray m_area; //!< memory space for area std::shared_ptr m_mesh_data; //!< Mesh data to use in computing area conservation energy diff --git a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cc b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cc index 5335c20605..876f53e2d9 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cc +++ b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cc @@ -76,7 +76,9 @@ void TriangleAreaConservationMeshForceComputeGPU::computeForces(uint64_t timeste ArrayHandle d_force(m_force, access_location::device, access_mode::overwrite); ArrayHandle d_virial(m_virial, access_location::device, access_mode::overwrite); - ArrayHandle d_params(m_params, access_location::device, access_mode::read); + ArrayHandle d_params(m_params, + access_location::device, + access_mode::read); m_tuner->begin(); kernel::gpu_compute_TriangleAreaConservation_force( diff --git a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cuh b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cuh index 6e287b782f..a295fbff8a 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cuh +++ b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cuh @@ -1,10 +1,10 @@ // Copyright (c) 2009-2024 The Regents of the University of Michigan. // Part of HOOMD-blue, released under the BSD 3-Clause License. +#include "TriangleAreaConservationMeshParameters.h" #include "hoomd/BondedGroupData.cuh" #include "hoomd/HOOMDMath.h" #include "hoomd/Index1D.h" -#include "TriangleAreaConservationMeshParameters.h" #include "hoomd/ParticleData.cuh" /*! \file TriangleAreaConservationMeshForceComputeGPU.cuh diff --git a/hoomd/md/TriangleAreaConservationMeshParameters.h b/hoomd/md/TriangleAreaConservationMeshParameters.h index 60f50c7bad..1ddddf0200 100644 --- a/hoomd/md/TriangleAreaConservationMeshParameters.h +++ b/hoomd/md/TriangleAreaConservationMeshParameters.h @@ -7,36 +7,36 @@ #include #endif - #pragma once -namespace hoomd::md { - struct triangle_area_conservation_param_t - { - Scalar k; - Scalar A0; +namespace hoomd::md + { +struct triangle_area_conservation_param_t + { + Scalar k; + Scalar A0; #ifndef __HIPCC__ - triangle_area_conservation_param_t() : k(0), A0(0) { } - - triangle_area_conservation_param_t(pybind11::dict params) - : k(params["k"].cast()), A0(params["A0"].cast()) - { - } - - pybind11::dict asDict() - { - pybind11::dict v; - v["k"] = k; - v["A0"] = A0; - return v; - } -#endif + triangle_area_conservation_param_t() : k(0), A0(0) { } + + triangle_area_conservation_param_t(pybind11::dict params) + : k(params["k"].cast()), A0(params["A0"].cast()) + { } + + pybind11::dict asDict() + { + pybind11::dict v; + v["k"] = k; + v["A0"] = A0; + return v; + } +#endif + } #if HOOMD_LONGREAL_SIZE == 32 - __attribute__((aligned(4))); + __attribute__((aligned(4))); #else - __attribute__((aligned(8))); + __attribute__((aligned(8))); #endif -} + } // namespace hoomd::md diff --git a/hoomd/md/VolumeConservationMeshForceCompute.cc b/hoomd/md/VolumeConservationMeshForceCompute.cc index 857b9c2a69..62bb3ecaa7 100644 --- a/hoomd/md/VolumeConservationMeshForceCompute.cc +++ b/hoomd/md/VolumeConservationMeshForceCompute.cc @@ -51,11 +51,14 @@ VolumeConservationMeshForceCompute::~VolumeConservationMeshForceCompute() Sets parameters for the potential of a particular angle type */ -void VolumeConservationMeshForceCompute::setParams(unsigned int type, const volume_conservation_param_t& params) +void VolumeConservationMeshForceCompute::setParams(unsigned int type, + const volume_conservation_param_t& params) { if (!m_ignore_type || type == 0) { - ArrayHandle h_params(m_params, access_location::host, access_mode::readwrite); + ArrayHandle h_params(m_params, + access_location::host, + access_mode::readwrite); h_params.data[type] = params; if (params.k <= 0) @@ -79,8 +82,10 @@ pybind11::dict VolumeConservationMeshForceCompute::getParams(std::string type) m_exec_conf->msg->error() << "mesh.volume: Invalid mesh type specified" << endl; throw runtime_error("Error setting parameters in VolumeConservationMeshForceCompute"); } - ArrayHandle h_params(m_params, access_location::host, access_mode::read); - return h_params.data[typ].asDict(); + ArrayHandle h_params(m_params, + access_location::host, + access_mode::read); + return h_params.data[typ].asDict(); } /*! Actually perform the force computation @@ -99,7 +104,9 @@ void VolumeConservationMeshForceCompute::computeForces(uint64_t timestep) ArrayHandle h_force(m_force, access_location::host, access_mode::overwrite); ArrayHandle h_virial(m_virial, access_location::host, access_mode::overwrite); size_t virial_pitch = m_virial.getPitch(); - ArrayHandle h_params(m_params, access_location::host, access_mode::read); + ArrayHandle h_params(m_params, + access_location::host, + access_mode::read); ArrayHandle h_volume(m_volume, access_location::host, access_mode::read); ArrayHandle h_triangles( diff --git a/hoomd/md/VolumeConservationMeshForceCompute.h b/hoomd/md/VolumeConservationMeshForceCompute.h index b6d42a17a2..370b35901e 100644 --- a/hoomd/md/VolumeConservationMeshForceCompute.h +++ b/hoomd/md/VolumeConservationMeshForceCompute.h @@ -1,9 +1,9 @@ // Copyright (c) 2009-2024 The Regents of the University of Michigan. // Part of HOOMD-blue, released under the BSD 3-Clause License. +#include "VolumeConservationMeshParameters.h" #include "hoomd/ForceCompute.h" #include "hoomd/MeshDefinition.h" -#include "VolumeConservationMeshParameters.h" #include diff --git a/hoomd/md/VolumeConservationMeshForceComputeGPU.cc b/hoomd/md/VolumeConservationMeshForceComputeGPU.cc index 9fbf8e3726..a50897d543 100644 --- a/hoomd/md/VolumeConservationMeshForceComputeGPU.cc +++ b/hoomd/md/VolumeConservationMeshForceComputeGPU.cc @@ -90,7 +90,9 @@ void VolumeConservationMeshForceComputeGPU::computeForces(uint64_t timestep) ArrayHandle d_force(m_force, access_location::device, access_mode::overwrite); ArrayHandle d_virial(m_virial, access_location::device, access_mode::overwrite); - ArrayHandle d_params(m_params, access_location::device, access_mode::read); + ArrayHandle d_params(m_params, + access_location::device, + access_mode::read); ArrayHandle d_volume(m_volume, access_location::device, access_mode::read); diff --git a/hoomd/md/VolumeConservationMeshForceComputeGPU.cuh b/hoomd/md/VolumeConservationMeshForceComputeGPU.cuh index 2f12b62a4d..b1a5d7c148 100644 --- a/hoomd/md/VolumeConservationMeshForceComputeGPU.cuh +++ b/hoomd/md/VolumeConservationMeshForceComputeGPU.cuh @@ -1,11 +1,11 @@ // Copyright (c) 2009-2024 The Regents of the University of Michigan. // Part of HOOMD-blue, released under the BSD 3-Clause License. +#include "VolumeConservationMeshParameters.h" #include "hoomd/BondedGroupData.cuh" #include "hoomd/HOOMDMath.h" #include "hoomd/Index1D.h" #include "hoomd/ParticleData.cuh" -#include "VolumeConservationMeshParameters.h" #include /*! \file MeshVolumeConservationGPU.cuh diff --git a/hoomd/md/VolumeConservationMeshParameters.h b/hoomd/md/VolumeConservationMeshParameters.h index 60da432854..45516ea495 100644 --- a/hoomd/md/VolumeConservationMeshParameters.h +++ b/hoomd/md/VolumeConservationMeshParameters.h @@ -7,36 +7,36 @@ #include #endif - #pragma once -namespace hoomd::md { - struct volume_conservation_param_t - { - Scalar k; - Scalar V0; +namespace hoomd::md + { +struct volume_conservation_param_t + { + Scalar k; + Scalar V0; #ifndef __HIPCC__ - volume_conservation_param_t() : k(0), V0(0) { } - - volume_conservation_param_t(pybind11::dict params) - : k(params["k"].cast()), V0(params["V0"].cast()) - { - } - - pybind11::dict asDict() - { - pybind11::dict v; - v["k"] = k; - v["V0"] = V0; - return v; - } -#endif + volume_conservation_param_t() : k(0), V0(0) { } + + volume_conservation_param_t(pybind11::dict params) + : k(params["k"].cast()), V0(params["V0"].cast()) + { } + + pybind11::dict asDict() + { + pybind11::dict v; + v["k"] = k; + v["V0"] = V0; + return v; + } +#endif + } #if HOOMD_LONGREAL_SIZE == 32 - __attribute__((aligned(4))); + __attribute__((aligned(4))); #else - __attribute__((aligned(8))); + __attribute__((aligned(8))); #endif -} + } // namespace hoomd::md From 9a2de8d4f544fe79dd807454a1aed28604abc9b1 Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 10:48:23 -0400 Subject: [PATCH 07/13] Remove lychee from CI. There are way too many false positives when websites glitch or block what appears to be a bot. We will keep the lychee.toml configuration so that developers can manually check links occaisionally, but we do not need a failing webserver to delay the review of a PR. --- .github/workflows/test.yaml | 28 ---------------------------- 1 file changed, 28 deletions(-) diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml index a4de6870d8..5801365cbb 100644 --- a/.github/workflows/test.yaml +++ b/.github/workflows/test.yaml @@ -97,34 +97,6 @@ jobs: - config: [gcc, 11, -py, 310, -mpi, -tbb] - config: [gcc, 10, -py, 310, -mpi, -tbb] - check-links: - name: "Check links" - runs-on: ubuntu-latest - steps: - - name: Check out repo - uses: actions/checkout@a5ac7e51b41094c92402da3b24376905380afc29 # v4.1.6 - with: - submodules: true - - name: Set up Python - uses: actions/setup-python@82c7e631bb3cdc910f68e0081d67478d79c6982d # v5.1.0 - with: - python-version: "3.12" - - name: Setup uv - uses: glotzerlab/workflows/setup-uv@5cfac9da9cb78e16ae97a9119b6fd13c1c2d6f5e # 0.1.0 - - name: Install dependencies - run: uv pip install -r sphinx-doc/requirements.txt --system - - name: Install tools - run: sudo apt-get install pandoc - - name: Build Sphinx Docs - run: sphinx-build -b html sphinx-doc doc_build - - name: Link Checker - id: lychee - uses: lycheeverse/lychee-action@2b973e86fc7b1f6b36a93795fe2c9c6ae1118621 # v1.10.0 - with: - args: -n './**/*.md' './**/*.html' './**/*.rst' - --timeout 60 - fail: true - tests_complete: name: Unit test if: always() From 891c48578e812796c91a47276cbb5cd3438b8a52 Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 10:55:31 -0400 Subject: [PATCH 08/13] Update change log. --- CHANGELOG.rst | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/CHANGELOG.rst b/CHANGELOG.rst index 93b1d0ec18..189d4daec0 100644 --- a/CHANGELOG.rst +++ b/CHANGELOG.rst @@ -7,6 +7,14 @@ Change Log 4.x --- +4.9.1 (2024-10-31) +^^^^^^^^^^^^^^^^^^ + +*Fixed* + +* Prevent compile errors with ``-DENABLE_GPU=on -DHOOMD_GPU_PLATFORM=HIP`` + (`#1920 `__) + 4.9.0 (2024-10-29) ^^^^^^^^^^^^^^^^^^ From 469a53c380d4b7ffd46e5090b7ce0f49ce603fcc Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 11:11:46 -0400 Subject: [PATCH 09/13] Reduce compiler warnings when building with HIP. --- CMakeLists.txt | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index fb75027279..e66f6859ae 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -68,7 +68,12 @@ set(CMAKE_HIP_STANDARD 14) # Enable compiler warnings on gcc and clang (common compilers used by developers) if(CMAKE_COMPILER_IS_GNUCXX OR CMAKE_CXX_COMPILER_ID MATCHES "Clang") - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -Wconversion -Wno-sign-conversion -Wno-unknown-pragmas -Wno-deprecated-declarations -Wno-unused-result") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -Wno-sign-conversion -Wno-unknown-pragmas -Wno-deprecated-declarations -Wno-unused-result") + + if (NOT (ENABLE_GPU AND HOOMD_GPU_PLATFORM STREQUAL "HIP")) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wconversion") + endif() + set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -Wall") endif() From 09cae4332f0146d694c08a5e9c5ab8e091726537 Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 11:12:55 -0400 Subject: [PATCH 10/13] Fix yet another compile error caused by HIP. --- hoomd/md/BendingRigidityMeshForceComputeGPU.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hoomd/md/BendingRigidityMeshForceComputeGPU.cu b/hoomd/md/BendingRigidityMeshForceComputeGPU.cu index eb3f423fb4..f12277863b 100644 --- a/hoomd/md/BendingRigidityMeshForceComputeGPU.cu +++ b/hoomd/md/BendingRigidityMeshForceComputeGPU.cu @@ -182,7 +182,7 @@ __global__ void gpu_compute_bending_rigidity_force_kernel(Scalar4* d_force, } else { - Fac *= -1; + Fac *= Scalar(-1); } force.x += Fac.x; From 7541d3b5b7fac9c52b2bd1ad41be56c66e8d9ba0 Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 11:40:59 -0400 Subject: [PATCH 11/13] Reduce compiler warnings. --- CMakeLists.txt | 4 ++-- hoomd/md/AreaConservationMeshParameters.h | 7 +++++-- hoomd/md/TriangleAreaConservationMeshParameters.h | 8 ++++++-- hoomd/md/VolumeConservationMeshParameters.h | 7 +++++-- 4 files changed, 18 insertions(+), 8 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index e66f6859ae..22a8271854 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -68,12 +68,12 @@ set(CMAKE_HIP_STANDARD 14) # Enable compiler warnings on gcc and clang (common compilers used by developers) if(CMAKE_COMPILER_IS_GNUCXX OR CMAKE_CXX_COMPILER_ID MATCHES "Clang") - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -Wno-sign-conversion -Wno-unknown-pragmas -Wno-deprecated-declarations -Wno-unused-result") - if (NOT (ENABLE_GPU AND HOOMD_GPU_PLATFORM STREQUAL "HIP")) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wconversion") endif() + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -Wno-sign-conversion -Wno-unknown-pragmas -Wno-deprecated-declarations -Wno-unused-result") + set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -Wall") endif() diff --git a/hoomd/md/AreaConservationMeshParameters.h b/hoomd/md/AreaConservationMeshParameters.h index ce334f97ca..ed1b2bc49b 100644 --- a/hoomd/md/AreaConservationMeshParameters.h +++ b/hoomd/md/AreaConservationMeshParameters.h @@ -9,7 +9,9 @@ #pragma once -namespace hoomd::md +namespace hoomd + { +namespace md { struct area_conservation_param_t { @@ -39,4 +41,5 @@ struct area_conservation_param_t __attribute__((aligned(8))); #endif - } // namespace hoomd::md + } // namespace md + } // namespace hoomd diff --git a/hoomd/md/TriangleAreaConservationMeshParameters.h b/hoomd/md/TriangleAreaConservationMeshParameters.h index 1ddddf0200..6d38e93ad1 100644 --- a/hoomd/md/TriangleAreaConservationMeshParameters.h +++ b/hoomd/md/TriangleAreaConservationMeshParameters.h @@ -9,8 +9,11 @@ #pragma once -namespace hoomd::md +namespace hoomd { +namespace md + { + struct triangle_area_conservation_param_t { Scalar k; @@ -39,4 +42,5 @@ struct triangle_area_conservation_param_t __attribute__((aligned(8))); #endif - } // namespace hoomd::md + } // namespace md + } // namespace hoomd diff --git a/hoomd/md/VolumeConservationMeshParameters.h b/hoomd/md/VolumeConservationMeshParameters.h index 45516ea495..ae32992b48 100644 --- a/hoomd/md/VolumeConservationMeshParameters.h +++ b/hoomd/md/VolumeConservationMeshParameters.h @@ -9,7 +9,9 @@ #pragma once -namespace hoomd::md +namespace hoomd + { +namespace md { struct volume_conservation_param_t { @@ -39,4 +41,5 @@ struct volume_conservation_param_t __attribute__((aligned(8))); #endif - } // namespace hoomd::md + } // namespace md + } // namespace hoomd From f7b39c8e293495eef77aaa4297cd79c5bb0a13b7 Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 11:44:08 -0400 Subject: [PATCH 12/13] Bump version to 4.9.1 --- .bumpversion.cfg | 2 +- .github/ISSUE_TEMPLATE/bug_report.yml | 2 +- .github/ISSUE_TEMPLATE/release.md | 2 +- CMakeLists.txt | 2 +- INSTALLING.rst | 6 +++--- hoomd/hpmc/external/user.py | 4 ++-- hoomd/hpmc/pair/user.py | 2 +- sphinx-doc/conf.py | 4 ++-- 8 files changed, 12 insertions(+), 12 deletions(-) diff --git a/.bumpversion.cfg b/.bumpversion.cfg index 46483a60ea..678262ec30 100644 --- a/.bumpversion.cfg +++ b/.bumpversion.cfg @@ -1,5 +1,5 @@ [bumpversion] -current_version = 4.9.0 +current_version = 4.9.1 commit = False tag = False parse = ^(?P0|[1-9]\d*)\.(?P0|[1-9]\d*)\.(?P0|[1-9]\d*)(?:-(?P(?:0|[1-9]\d*|\d*[a-zA-Z-][0-9a-zA-Z-]*))(?:\.(?P0|[1-9]\d*|\d*[a-zA-Z-][0-9a-zA-Z-]*)))?$ diff --git a/.github/ISSUE_TEMPLATE/bug_report.yml b/.github/ISSUE_TEMPLATE/bug_report.yml index 27a42ef383..6b2cc515b4 100644 --- a/.github/ISSUE_TEMPLATE/bug_report.yml +++ b/.github/ISSUE_TEMPLATE/bug_report.yml @@ -60,7 +60,7 @@ body: attributes: label: HOOMD-blue version description: What version of HOOMD-blue are you using? - placeholder: 4.9.0 + placeholder: 4.9.1 validations: required: true - type: markdown diff --git a/.github/ISSUE_TEMPLATE/release.md b/.github/ISSUE_TEMPLATE/release.md index 78473b5889..04220fd648 100644 --- a/.github/ISSUE_TEMPLATE/release.md +++ b/.github/ISSUE_TEMPLATE/release.md @@ -1,7 +1,7 @@ --- name: Release checklist about: '[for maintainer use]' -title: 'Release 4.9.0' +title: 'Release 4.9.1' labels: '' assignees: 'joaander' diff --git a/CMakeLists.txt b/CMakeLists.txt index 22a8271854..caef1d26ce 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -16,7 +16,7 @@ add_subdirectory (CMake) ################################ ## Version information -set(HOOMD_VERSION_RAW "4.9.0") +set(HOOMD_VERSION_RAW "4.9.1") string(REGEX MATCH "(.*)\\.(.*)\\.(.*)$" _hoomd_version_match ${HOOMD_VERSION_RAW}) set(HOOMD_VERSION_MAJOR ${CMAKE_MATCH_1}) set(HOOMD_VERSION_MINOR ${CMAKE_MATCH_2}) diff --git a/INSTALLING.rst b/INSTALLING.rst index 073105de17..4609ceb9cd 100644 --- a/INSTALLING.rst +++ b/INSTALLING.rst @@ -22,7 +22,7 @@ channel: .. code-block:: bash - micromamba install hoomd=4.9.0 + micromamba install hoomd=4.9.1 .. _conda-forge: https://conda-forge.org/docs/user/introduction.html @@ -32,13 +32,13 @@ appropriate package. Override this and force the GPU enabled package installatio .. code-block:: bash export CONDA_OVERRIDE_CUDA="12.0" - micromamba install "hoomd=4.9.0=*gpu*" "cuda-version=12.0" + micromamba install "hoomd=4.9.1=*gpu*" "cuda-version=12.0" Similarly, you can force CPU-only package installation with: .. code-block:: bash - micromamba install "hoomd=4.9.0=*cpu*" + micromamba install "hoomd=4.9.1=*cpu*" .. note:: diff --git a/hoomd/hpmc/external/user.py b/hoomd/hpmc/external/user.py index 899fbaba02..5900289e0a 100644 --- a/hoomd/hpmc/external/user.py +++ b/hoomd/hpmc/external/user.py @@ -85,9 +85,9 @@ class CPPExternalPotential(ExternalField): Your code *must* return a value. .. _VectorMath.h: https://github.com/glotzerlab/hoomd-blue/blob/\ - v4.9.0/hoomd/VectorMath.h + v4.9.1/hoomd/VectorMath.h .. _BoxDim.h: https://github.com/glotzerlab/hoomd-blue/blob/\ - v4.9.0/hoomd/BoxDim.h + v4.9.1/hoomd/BoxDim.h .. rubric:: Example: diff --git a/hoomd/hpmc/pair/user.py b/hoomd/hpmc/pair/user.py index 76e0a824ac..d042797500 100644 --- a/hoomd/hpmc/pair/user.py +++ b/hoomd/hpmc/pair/user.py @@ -76,7 +76,7 @@ class CPPPotentialBase(AutotunedObject): HOOMD-blue source code. .. _VectorMath.h: https://github.com/glotzerlab/hoomd-blue/blob/\ - v4.9.0/hoomd/VectorMath.h + v4.9.1/hoomd/VectorMath.h Note: Your code *must* return a value. diff --git a/sphinx-doc/conf.py b/sphinx-doc/conf.py index a1641cab44..0c5cf338cc 100644 --- a/sphinx-doc/conf.py +++ b/sphinx-doc/conf.py @@ -68,8 +68,8 @@ copyright = f'2009-{year} The Regents of the University of Michigan' author = 'The Regents of the University of Michigan' -version = '4.9.0' -release = '4.9.0' +version = '4.9.1' +release = '4.9.1' language = 'en' From c50c6971a2ec1bc7354846bbb9e104d9f3956380 Mon Sep 17 00:00:00 2001 From: "Joshua A. Anderson" Date: Thu, 31 Oct 2024 11:45:04 -0400 Subject: [PATCH 13/13] Reword change log entry. --- CHANGELOG.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CHANGELOG.rst b/CHANGELOG.rst index 189d4daec0..c17591f3cb 100644 --- a/CHANGELOG.rst +++ b/CHANGELOG.rst @@ -12,7 +12,7 @@ Change Log *Fixed* -* Prevent compile errors with ``-DENABLE_GPU=on -DHOOMD_GPU_PLATFORM=HIP`` +* Correct compile errors with ``-DENABLE_GPU=on -DHOOMD_GPU_PLATFORM=HIP`` (`#1920 `__) 4.9.0 (2024-10-29)