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/.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() diff --git a/CHANGELOG.rst b/CHANGELOG.rst index 93b1d0ec18..c17591f3cb 100644 --- a/CHANGELOG.rst +++ b/CHANGELOG.rst @@ -7,6 +7,14 @@ Change Log 4.x --- +4.9.1 (2024-10-31) +^^^^^^^^^^^^^^^^^^ + +*Fixed* + +* Correct compile errors with ``-DENABLE_GPU=on -DHOOMD_GPU_PLATFORM=HIP`` + (`#1920 `__) + 4.9.0 (2024-10-29) ^^^^^^^^^^^^^^^^^^ diff --git a/CMakeLists.txt b/CMakeLists.txt index fb75027279..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}) @@ -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") + 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/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/hoomd/md/AreaConservationMeshForceCompute.cc b/hoomd/md/AreaConservationMeshForceCompute.cc index 89dd22c1ae..03fab98ddb 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,23 @@ 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& params) { 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 +71,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 +84,10 @@ 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 +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( @@ -205,12 +205,13 @@ 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..f3d72ff3c7 100644 --- a/hoomd/md/AreaConservationMeshForceCompute.h +++ b/hoomd/md/AreaConservationMeshForceCompute.h @@ -1,6 +1,7 @@ // 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" @@ -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,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 87d360cf84..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.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..1853267e8d 100644 --- a/hoomd/md/AreaConservationMeshForceComputeGPU.cuh +++ b/hoomd/md/AreaConservationMeshForceComputeGPU.cuh @@ -1,6 +1,7 @@ // 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" @@ -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/AreaConservationMeshParameters.h b/hoomd/md/AreaConservationMeshParameters.h new file mode 100644 index 0000000000..ed1b2bc49b --- /dev/null +++ b/hoomd/md/AreaConservationMeshParameters.h @@ -0,0 +1,45 @@ +// 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 + { +namespace 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 + + } // namespace md + } // namespace hoomd 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/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; diff --git a/hoomd/md/CMakeLists.txt b/hoomd/md/CMakeLists.txt index 6490e74dc8..e44ec847f5 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 ) @@ -70,6 +70,7 @@ set(_md_headers ActiveForceComputeGPU.h AnisoPotentialPairGPU.cuh AnisoPotentialPairGPU.h AnisoPotentialPair.h + AreaConservationMeshParameters.h AreaConservationMeshForceCompute.h AreaConservationMeshForceComputeGPU.h AreaConservationMeshForceComputeGPU.cuh @@ -98,7 +99,7 @@ set(_md_headers ActiveForceComputeGPU.h EvaluatorSpecialPairLJ.h EvaluatorSpecialPairCoulomb.h EvaluatorExternalElectricField.h - EvaluatorExternalMagneticField.h + EvaluatorExternalMagneticField.h EvaluatorExternalPeriodic.h EvaluatorPairALJ.h EvaluatorPairBuckingham.h @@ -141,8 +142,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 @@ -195,6 +196,7 @@ set(_md_headers ActiveForceComputeGPU.h TableAngleForceCompute.h TableDihedralForceComputeGPU.h TableDihedralForceCompute.h + TriangleAreaConservationMeshParameters.h TriangleAreaConservationMeshForceCompute.h TriangleAreaConservationMeshForceComputeGPU.h TriangleAreaConservationMeshForceComputeGPU.cuh @@ -217,8 +219,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 +246,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 +260,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 +274,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 +286,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 +305,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/TriangleAreaConservationMeshForceCompute.cc b/hoomd/md/TriangleAreaConservationMeshForceCompute.cc index 3358ed94d9..437bccb800 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); @@ -43,22 +43,23 @@ 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 */ -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 +67,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 +79,10 @@ 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 +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(), @@ -191,12 +192,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 +226,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 +251,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 +273,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..6cab075787 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceCompute.h +++ b/hoomd/md/TriangleAreaConservationMeshForceCompute.h @@ -1,6 +1,7 @@ // 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" @@ -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,8 +67,8 @@ class PYBIND11_EXPORT TriangleAreaConservationMeshForceCompute : public ForceCom #endif protected: - GPUArray m_params; //!< Parameters - GPUArray m_area; //!< memory space for area + GPUArray m_params; //!< Parameters + 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 10f15213c4..876f53e2d9 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,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.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..a295fbff8a 100644 --- a/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cuh +++ b/hoomd/md/TriangleAreaConservationMeshForceComputeGPU.cuh @@ -1,6 +1,7 @@ // 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" @@ -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/TriangleAreaConservationMeshParameters.h b/hoomd/md/TriangleAreaConservationMeshParameters.h new file mode 100644 index 0000000000..6d38e93ad1 --- /dev/null +++ b/hoomd/md/TriangleAreaConservationMeshParameters.h @@ -0,0 +1,46 @@ +// 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 + { +namespace 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 + + } // namespace md + } // namespace hoomd diff --git a/hoomd/md/VolumeConservationMeshForceCompute.cc b/hoomd/md/VolumeConservationMeshForceCompute.cc index 7b9f2594da..62bb3ecaa7 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); @@ -47,21 +47,23 @@ 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 */ -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 +71,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 +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); - 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 +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( @@ -179,12 +181,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..370b35901e 100644 --- a/hoomd/md/VolumeConservationMeshForceCompute.h +++ b/hoomd/md/VolumeConservationMeshForceCompute.h @@ -1,6 +1,7 @@ // 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" @@ -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/VolumeConservationMeshForceComputeGPU.cc b/hoomd/md/VolumeConservationMeshForceComputeGPU.cc index 8c95336fd7..a50897d543 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,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.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..b1a5d7c148 100644 --- a/hoomd/md/VolumeConservationMeshForceComputeGPU.cuh +++ b/hoomd/md/VolumeConservationMeshForceComputeGPU.cuh @@ -1,6 +1,7 @@ // 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" @@ -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 diff --git a/hoomd/md/VolumeConservationMeshParameters.h b/hoomd/md/VolumeConservationMeshParameters.h new file mode 100644 index 0000000000..ae32992b48 --- /dev/null +++ b/hoomd/md/VolumeConservationMeshParameters.h @@ -0,0 +1,45 @@ +// 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 + { +namespace 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 + + } // namespace md + } // namespace hoomd 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: 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'