Support binning on GPU

Author: mphoward

src/BinningOperation.h

@@ -7,7 +7,7 @@
 
 #include "hoomd/HOOMDMath.h"
 
-#ifndef __HIPCC__
+#ifdef __HIPCC__
 #define HOSTDEVICE __host__ __device__
 #else
 #define HOSTDEVICE

src/CMakeLists.txt

@@ -4,14 +4,14 @@ set(COMPONENT_NAME azplugins)
 # TODO: List all host C++ source code files in _${COMPONENT_NAME}_sources.
 set(_${COMPONENT_NAME}_sources
     ConstantFlow.cc
-    CylindricalVelocityFieldCompute.cc
     GroupVelocityCompute.cc
     module.cc
     ParabolicFlow.cc
     )
 
 # 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.
@@ -123,6 +123,28 @@ foreach(_evaluator ${_aniso_pair_evaluators})
     endif()
 endforeach()
 
+# process velocity field geometries
+set(_binning_geometries
+    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/CylindricalBinningOperation.h

@@ -7,7 +7,7 @@
 
 #include "BinningOperation.h"
 
-#ifndef __HIPCC__
+#ifdef __HIPCC__
 #define HOSTDEVICE __host__ __device__
 #else
 #define HOSTDEVICE

src/CylindricalVelocityFieldCompute.cc

@@ -0,0 +1,86 @@
+// 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_PARTICLE_DATA_LOADER_H_
+#define AZPLUGINS_PARTICLE_DATA_LOADER_H_
+
+#include "hoomd/HOOMDMath.h"
+
+#ifdef __HIPCC__
+#define HOSTDEVICE __host__ __device__
+#else
+#define HOSTDEVICE
+#endif // __HIPCC__
+
+namespace hoomd
+    {
+namespace azplugins
+    {
+namespace detail
+    {
+
+//! Load HOOMD particle in a group from an index
+class LoadHOOMDGroupPositionVelocityMass
+    {
+    public:
+    HOSTDEVICE
+    LoadHOOMDGroupPositionVelocityMass(Scalar4* positions,
+                                       Scalar4* velocities,
+                                       const unsigned int* indexes)
+        : m_positions(positions), m_velocities(velocities), m_indexes(indexes)
+        {
+        }
+
+    HOSTDEVICE void
+    operator()(Scalar3& position, Scalar3& velocity, Scalar& mass, unsigned int idx) const
+        {
+        const unsigned int pidx = m_indexes[idx];
+        const Scalar4 postype = m_positions[pidx];
+        position = make_scalar3(postype.x, postype.y, postype.z);
+
+        const Scalar4 velmass = m_velocities[pidx];
+        velocity = make_scalar3(velmass.x, velmass.y, velmass.z);
+        mass = velmass.w;
+        }
+
+    private:
+    Scalar4* const m_positions;
+    Scalar4* const m_velocities;
+    const unsigned int* const m_indexes;
+    };
+
+//! Load MPCD particle from an index
+class LoadMPCDPositionVelocityMass
+    {
+    public:
+    HOSTDEVICE
+    LoadMPCDPositionVelocityMass(Scalar4* positions, Scalar4* velocities, const Scalar mass)
+        : m_positions(positions), m_velocities(velocities), m_mass(mass)
+        {
+        }
+
+    HOSTDEVICE void
+    operator()(Scalar3& position, Scalar3& velocity, Scalar& mass, unsigned int idx) const
+        {
+        const Scalar4 postype = m_positions[idx];
+        position = make_scalar3(postype.x, postype.y, postype.z);
+
+        const Scalar4 velcell = m_velocities[idx];
+        velocity = make_scalar3(velcell.x, velcell.y, velcell.z);
+        mass = m_mass;
+        }
+
+    private:
+    Scalar4* const m_positions;
+    Scalar4* const m_velocities;
+    const Scalar m_mass;
+    };
+
+    } // end namespace detail
+    } // end namespace azplugins
+    } // end namespace hoomd
+
+#undef HOSTDEVICE
+
+#endif // AZPLUGINS_PARTICLE_DATA_LOADER_H_