Merge pull request #1920 from glotzerlab/4.9.1-hot-fixes

Release 4.9.1

.bumpversion.cfg

@@ -1,5 +1,5 @@
 [bumpversion]
-current_version = 4.9.0
+current_version = 4.9.1
 commit = False
 tag = False
 parse = ^(?P<major>0|[1-9]\d*)\.(?P<minor>0|[1-9]\d*)\.(?P<patch>0|[1-9]\d*)(?:-(?P<prerelease>(?:0|[1-9]\d*|\d*[a-zA-Z-][0-9a-zA-Z-]*))(?:\.(?P<number>0|[1-9]\d*|\d*[a-zA-Z-][0-9a-zA-Z-]*)))?$

.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

.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'
 

.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()

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 <https://github.com/glotzerlab/hoomd-blue/pull/1915>`__)
+
 4.9.0 (2024-10-29)
 ^^^^^^^^^^^^^^^^^^
 

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()
 

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::
 

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:
 

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.

hoomd/md/AreaConservationMeshForceCompute.cc

@@ -34,7 +34,7 @@ AreaConservationMeshForceCompute::AreaConservationMeshForceCompute(
     if (m_ignore_type)
         n_types = 1;
 
-    GPUArray<Scalar2> params(n_types, m_exec_conf);
+    GPUArray<area_conservation_param_t> params(n_types, m_exec_conf);
     m_params.swap(params);
 
     GPUArray<Scalar> area(n_types, m_exec_conf);
@@ -47,32 +47,31 @@ 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<Scalar2> 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<area_conservation_param_t> 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;
         }
     }
 
 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<Scalar2> 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<area_conservation_param_t> 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<Scalar4> h_force(m_force, access_location::host, access_mode::overwrite);
     ArrayHandle<Scalar> h_virial(m_virial, access_location::host, access_mode::overwrite);
     size_t virial_pitch = m_virial.getPitch();
-    ArrayHandle<Scalar2> h_params(m_params, access_location::host, access_mode::read);
+    ArrayHandle<area_conservation_param_t> h_params(m_params,
+                                                    access_location::host,
+                                                    access_mode::read);
     ArrayHandle<Scalar> h_area(m_area, access_location::host, access_mode::read);
 
     ArrayHandle<typename Angle::members_t> 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);

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<Scalar>()), A0(params["A0"].cast<Scalar>())
-        {
-        }
-
-    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<Scalar2> m_params;                  //!< Parameters
-    GPUArray<Scalar> m_area;                     //!< memory space for area
-                                                 //
-    std::shared_ptr<MeshDefinition> m_mesh_data; //!< Mesh data to use in computing energy
-    bool m_ignore_type;                          //! ignore type to calculate global area if true
+    GPUArray<area_conservation_param_t> m_params; //!< Parameters
+    GPUArray<Scalar> m_area;                      //!< memory space for area
+                                                  //
+    std::shared_ptr<MeshDefinition> 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);

hoomd/md/AreaConservationMeshForceComputeGPU.cc

@@ -89,7 +89,9 @@ void AreaConservationMeshForceComputeGPU::computeForces(uint64_t timestep)
 
     ArrayHandle<Scalar4> d_force(m_force, access_location::device, access_mode::overwrite);
     ArrayHandle<Scalar> d_virial(m_virial, access_location::device, access_mode::overwrite);
-    ArrayHandle<Scalar2> d_params(m_params, access_location::device, access_mode::read);
+    ArrayHandle<area_conservation_param_t> d_params(m_params,
+                                                    access_location::device,
+                                                    access_mode::read);
 
     ArrayHandle<Scalar> d_area(m_area, access_location::device, access_mode::read);
 

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)
     {

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