Velocity field compute

README.md

@@ -2,7 +2,7 @@
 
 azplugins is a component for [HOOMD-blue][1] which expands its functionality for
 tackling a variety of problems in soft matter physics. Currently, azplugins is
-tested against v4.8.2 of HOOMD-blue. See [CHANGELOG.rst](CHANGELOG.rst) for a
+tested against v5.0.0 of HOOMD-blue. See [CHANGELOG.rst](CHANGELOG.rst) for a
 list of recent development.
 
 ## Compiling azplugins

doc/index.rst

@@ -8,7 +8,7 @@ azplugins
 
 azplugins is a component for `HOOMD-blue`_ which expands its functionality for
 tackling a variety of problems in soft matter physics. Currently, azplugins is
-tested against v4.8.2 of HOOMD-blue.
+tested against v5.0.0 of HOOMD-blue.
 
 Compiling
 =========
@@ -52,6 +52,7 @@ Contents
     :caption: API
 
     module-azplugins-bond
+    module-azplugins-compute
     module-azplugins-flow
     module-azplugins-pair
 

doc/module-azplugins-compute.rst

@@ -0,0 +1,27 @@
+.. Copyright (c) 2018-2020, Michael P. Howard
+.. Copyright (c) 2021-2024, Auburn University
+.. Part of azplugins, released under the BSD 3-Clause License.
+
+azplugins.compute
+-----------------
+
+.. rubric:: Overview
+
+.. py:currentmodule:: hoomd.azplugins.compute
+
+.. autosummary::
+    :nosignatures:
+
+    CartesianVelocityFieldCompute
+    CylindricalVelocityFieldCompute
+    VelocityFieldCompute
+
+.. rubric:: Details
+
+.. automodule:: hoomd.azplugins.compute
+    :synopsis: Computes.
+    :members: CartesianVelocityFieldCompute,
+        CylindricalVelocityFieldCompute,
+        VelocityFieldCompute
+    :no-inherited-members:
+    :show-inheritance:

src/BinningOperation.h

@@ -0,0 +1,95 @@
+// Copyright (c) 2018-2020, Michael P. Howard
+// Copyright (c) 2021-2024, Auburn University
+// Part of azplugins, released under the BSD 3-Clause License.
+
+#ifndef AZPLUGINS_BINNING_OPERATION_H_
+#define AZPLUGINS_BINNING_OPERATION_H_
+
+#include "hoomd/HOOMDMath.h"
+
+#ifdef __HIPCC__
+#define HOSTDEVICE __host__ __device__
+#else
+#define HOSTDEVICE
+#endif // __HIPCC__
+
+namespace hoomd
+    {
+namespace azplugins
+    {
+
+//! Interface for binning operations in different 3D coordinate systems.
+class BinningOperation
+    {
+    public:
+    HOSTDEVICE BinningOperation(const uint3& num_bins,
+                                const Scalar3& lower_bounds,
+                                const Scalar3& upper_bounds)
+        : m_lower_bounds {lower_bounds.x, lower_bounds.y, lower_bounds.z},
+          m_upper_bounds {upper_bounds.x, upper_bounds.y, upper_bounds.z}, m_num_bins {1, 1, 1},
+          m_should_bin {false, false, false}
+        {
+        if (num_bins.x > 0)
+            {
+            m_num_bins[0] = num_bins.x;
+            m_should_bin[0] = true;
+            }
+        if (num_bins.y > 0)
+            {
+            m_num_bins[1] = num_bins.y;
+            m_should_bin[1] = true;
+            }
+        if (num_bins.z > 0)
+            {
+            m_num_bins[2] = num_bins.z;
+            m_should_bin[2] = true;
+            }
+        }
+
+    HOSTDEVICE bool bin(uint3& bin,
+                        Scalar3& transformed_vector,
+                        const Scalar3& coordinates,
+                        const Scalar3& vector) const
+        {
+        return false;
+        }
+
+    HOSTDEVICE size_t getTotalNumBins() const
+        {
+        return static_cast<size_t>(m_num_bins[0]) * m_num_bins[1] * m_num_bins[2];
+        }
+
+    HOSTDEVICE size_t ravelBin(const uint3& bin) const
+        {
+        return static_cast<size_t>(bin.z) + m_num_bins[2] * (bin.y + m_num_bins[1] * bin.x);
+        }
+
+    protected:
+    Scalar m_lower_bounds[3];
+    Scalar m_upper_bounds[3];
+    unsigned int m_num_bins[3];
+    bool m_should_bin[3];
+
+    HOSTDEVICE bool bin1D(unsigned int& bin_1d, const Scalar x, unsigned int dim) const
+        {
+        int bin_ = static_cast<int>(
+            slow::floor(((x - m_lower_bounds[dim]) / (m_upper_bounds[dim] - m_lower_bounds[dim]))
+                        * m_num_bins[dim]));
+        if (bin_ >= 0 && bin_ < static_cast<int>(m_num_bins[dim]))
+            {
+            bin_1d = bin_;
+            return true;
+            }
+        else
+            {
+            return false;
+            }
+        }
+    };
+
+    } // end namespace azplugins
+    } // end namespace hoomd
+
+#undef HOSTDEVICE
+
+#endif // AZPLUGINS_BINNING_OPERATION_H_

src/CMakeLists.txt

@@ -11,11 +11,13 @@ set(_${COMPONENT_NAME}_sources
 
 # TODO: List all GPU C++ source code files in _${COMPONENT_NAME}_cu_sources.
 set(_${COMPONENT_NAME}_cu_sources
+    VelocityFieldComputeGPU.cu
     )
 
 # TODO: List all Python modules in python_files.
 set(python_files
     __init__.py
+    compute.py
     conftest.py
     bond.py
     flow.py
@@ -121,6 +123,29 @@ foreach(_evaluator ${_aniso_pair_evaluators})
     endif()
 endforeach()
 
+# process velocity field geometries
+set(_binning_geometries
+    Cartesian
+    Cylindrical
+    )
+foreach(_geometry ${_binning_geometries})
+    configure_file(export_VelocityFieldCompute.cc.inc
+                   export_${_geometry}VelocityFieldCompute.cc
+                   @ONLY)
+    set(_${COMPONENT_NAME}_sources ${_${COMPONENT_NAME}_sources} export_${_geometry}VelocityFieldCompute.cc)
+
+    if (ENABLE_HIP)
+        configure_file(export_VelocityFieldComputeGPU.cc.inc
+                       export_${_geometry}VelocityFieldComputeGPU.cc
+                       @ONLY)
+        configure_file(VelocityFieldComputeGPUKernel.cu.inc
+                       ${_geometry}VelocityFieldComputeGPUKernel.cu
+                       @ONLY)
+        set(_${COMPONENT_NAME}_sources ${_${COMPONENT_NAME}_sources} export_${_geometry}VelocityFieldComputeGPU.cc)
+        set(_${COMPONENT_NAME}_cu_sources ${_${COMPONENT_NAME}_cu_sources} ${_geometry}VelocityFieldComputeGPUKernel.cu)
+    endif()
+endforeach()
+
 if (ENABLE_HIP)
 set(_cuda_sources ${_${COMPONENT_NAME}_cu_sources})
 endif (ENABLE_HIP)

src/CartesianBinningOperation.h

@@ -0,0 +1,61 @@
+// Copyright (c) 2018-2020, Michael P. Howard
+// Copyright (c) 2021-2024, Auburn University
+// Part of azplugins, released under the BSD 3-Clause License.
+
+#ifndef AZPLUGINS_CARTESIAN_BINNING_OPERATION_H_
+#define AZPLUGINS_CARTESIAN_BINNING_OPERATION_H_
+
+#include "BinningOperation.h"
+
+#ifdef __HIPCC__
+#define HOSTDEVICE __host__ __device__
+#else
+#define HOSTDEVICE
+#endif // __HIPCC__
+
+namespace hoomd
+    {
+namespace azplugins
+    {
+
+//! Binning operation in Cartesian coordinates
+class CartesianBinningOperation : public BinningOperation
+    {
+    public:
+    using BinningOperation::BinningOperation;
+
+    HOSTDEVICE bool bin(uint3& bin,
+                        Scalar3& transformed_vector,
+                        const Scalar3& coordinates,
+                        const Scalar3& vector) const
+        {
+        uint3 bin_ = make_uint3(0, 0, 0);
+
+        if (m_should_bin[0] && !bin1D(bin_.x, coordinates.x, 0))
+            {
+            return false;
+            }
+
+        if (m_should_bin[1] && !bin1D(bin_.y, coordinates.y, 1))
+            {
+            return false;
+            }
+
+        if (m_should_bin[2] && !bin1D(bin_.z, coordinates.z, 2))
+            {
+            return false;
+            }
+
+        bin = bin_;
+        transformed_vector = vector;
+
+        return true;
+        }
+    };
+
+    } // end namespace azplugins
+    } // end namespace hoomd
+
+#undef HOSTDEVICE
+
+#endif // AZPLUGINS_CARTESIAN_BINNING_OPERATION_H_

src/CylindricalBinningOperation.h

@@ -0,0 +1,87 @@
+// Copyright (c) 2018-2020, Michael P. Howard
+// Copyright (c) 2021-2024, Auburn University
+// Part of azplugins, released under the BSD 3-Clause License.
+
+#ifndef AZPLUGINS_CYLINDRICAL_BINNING_OPERATION_H_
+#define AZPLUGINS_CYLINDRICAL_BINNING_OPERATION_H_
+
+#include "BinningOperation.h"
+
+#ifdef __HIPCC__
+#define HOSTDEVICE __host__ __device__
+#else
+#define HOSTDEVICE
+#endif // __HIPCC__
+
+namespace hoomd
+    {
+namespace azplugins
+    {
+
+//! Binning operation in cylindrical coordinates
+class CylindricalBinningOperation : public BinningOperation
+    {
+    public:
+    using BinningOperation::BinningOperation;
+
+    HOSTDEVICE bool bin(uint3& bin,
+                        Scalar3& transformed_vector,
+                        const Scalar3& coordinates,
+                        const Scalar3& vector) const
+        {
+        uint3 bin_ = make_uint3(0, 0, 0);
+
+        // bin with respect to z first because it is cheapest in case of early exit
+        if (m_should_bin[2] && !bin1D(bin_.z, coordinates.z, 2))
+            {
+            return false;
+            }
+
+        // bin with respect to theta second because its internals aren't needed later
+        if (m_should_bin[1])
+            {
+#if HOOMD_LONGREAL_SIZE == 32
+            Scalar theta = ::atan2f(coordinates.y, coordinates.x);
+#else
+            Scalar theta = ::atan2(coordinates.y, coordinates.x);
+#endif
+            if (theta < Scalar(0))
+                {
+                theta += Scalar(2) * M_PI;
+                }
+            if (!bin1D(bin_.y, theta, 1))
+                {
+                return false;
+                }
+            }
+
+        // bin with respect to r last, r will be used for vector transform so always do it
+        const Scalar r = fast::sqrt(coordinates.x * coordinates.x + coordinates.y * coordinates.y);
+        if (m_should_bin[0] && !bin1D(bin_.x, r, 0))
+            {
+            return false;
+            }
+
+        // transform Cartesian vector to cylindrical coordinates,
+        // defaulting angle to 0 if r == 0 (point is exactly at center)
+        Scalar cos_theta(1), sin_theta(0);
+        if (r > Scalar(0))
+            {
+            cos_theta = coordinates.x / r;
+            sin_theta = coordinates.y / r;
+            }
+        transformed_vector = make_scalar3(cos_theta * vector.x + sin_theta * vector.y,
+                                          -sin_theta * vector.x + cos_theta * vector.y,
+                                          vector.z);
+
+        bin = bin_;
+        return true;
+        }
+    };
+
+    } // end namespace azplugins
+    } // end namespace hoomd
+
+#undef HOSTDEVICE
+
+#endif // AZPLUGINS_CYLINDRICAL_BINNING_OPERATION_H_