Skip to content

Commit a63bac6

Browse files
mphowardmphoward
authored
Velocity field compute (#98)
* Draft CPU implementation of velocity field compute * Fix MPI errors * Improve unit tests * Improve documentation * Rename C++ class to match Python * Fix typos in docs * Support binning on GPU * Add Cartesian flow field * Ensure particles lie in global box * Fix minor Python issues * Fix miscellaneous issues
1 parent 69099a5 commit a63bac6

20 files changed

+1918
-3
lines changed

‎README.md‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22

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

88
## Compiling azplugins

‎doc/index.rst‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@ azplugins
88

99
azplugins is a component for `HOOMD-blue`_ which expands its functionality for
1010
tackling a variety of problems in soft matter physics. Currently, azplugins is
11-
tested against v4.8.2 of HOOMD-blue.
11+
tested against v5.0.0 of HOOMD-blue.
1212

1313
Compiling
1414
=========
@@ -52,6 +52,7 @@ Contents
5252
:caption: API
5353

5454
module-azplugins-bond
55+
module-azplugins-compute
5556
module-azplugins-flow
5657
module-azplugins-pair
5758

‎doc/module-azplugins-compute.rst‎

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
.. Copyright (c) 2018-2020, Michael P. Howard
2+
.. Copyright (c) 2021-2024, Auburn University
3+
.. Part of azplugins, released under the BSD 3-Clause License.
4+
5+
azplugins.compute
6+
-----------------
7+
8+
.. rubric:: Overview
9+
10+
.. py:currentmodule:: hoomd.azplugins.compute
11+
12+
.. autosummary::
13+
:nosignatures:
14+
15+
CartesianVelocityFieldCompute
16+
CylindricalVelocityFieldCompute
17+
VelocityFieldCompute
18+
19+
.. rubric:: Details
20+
21+
.. automodule:: hoomd.azplugins.compute
22+
:synopsis: Computes.
23+
:members: CartesianVelocityFieldCompute,
24+
CylindricalVelocityFieldCompute,
25+
VelocityFieldCompute
26+
:no-inherited-members:
27+
:show-inheritance:

‎src/BinningOperation.h‎

Lines changed: 95 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,95 @@
1+
// Copyright (c) 2018-2020, Michael P. Howard
2+
// Copyright (c) 2021-2024, Auburn University
3+
// Part of azplugins, released under the BSD 3-Clause License.
4+
5+
#ifndef AZPLUGINS_BINNING_OPERATION_H_
6+
#define AZPLUGINS_BINNING_OPERATION_H_
7+
8+
#include "hoomd/HOOMDMath.h"
9+
10+
#ifdef __HIPCC__
11+
#define HOSTDEVICE __host__ __device__
12+
#else
13+
#define HOSTDEVICE
14+
#endif // __HIPCC__
15+
16+
namespace hoomd
17+
{
18+
namespace azplugins
19+
{
20+
21+
//! Interface for binning operations in different 3D coordinate systems.
22+
class BinningOperation
23+
{
24+
public:
25+
HOSTDEVICE BinningOperation(const uint3& num_bins,
26+
const Scalar3& lower_bounds,
27+
const Scalar3& upper_bounds)
28+
: m_lower_bounds {lower_bounds.x, lower_bounds.y, lower_bounds.z},
29+
m_upper_bounds {upper_bounds.x, upper_bounds.y, upper_bounds.z}, m_num_bins {1, 1, 1},
30+
m_should_bin {false, false, false}
31+
{
32+
if (num_bins.x > 0)
33+
{
34+
m_num_bins[0] = num_bins.x;
35+
m_should_bin[0] = true;
36+
}
37+
if (num_bins.y > 0)
38+
{
39+
m_num_bins[1] = num_bins.y;
40+
m_should_bin[1] = true;
41+
}
42+
if (num_bins.z > 0)
43+
{
44+
m_num_bins[2] = num_bins.z;
45+
m_should_bin[2] = true;
46+
}
47+
}
48+
49+
HOSTDEVICE bool bin(uint3& bin,
50+
Scalar3& transformed_vector,
51+
const Scalar3& coordinates,
52+
const Scalar3& vector) const
53+
{
54+
return false;
55+
}
56+
57+
HOSTDEVICE size_t getTotalNumBins() const
58+
{
59+
return static_cast<size_t>(m_num_bins[0]) * m_num_bins[1] * m_num_bins[2];
60+
}
61+
62+
HOSTDEVICE size_t ravelBin(const uint3& bin) const
63+
{
64+
return static_cast<size_t>(bin.z) + m_num_bins[2] * (bin.y + m_num_bins[1] * bin.x);
65+
}
66+
67+
protected:
68+
Scalar m_lower_bounds[3];
69+
Scalar m_upper_bounds[3];
70+
unsigned int m_num_bins[3];
71+
bool m_should_bin[3];
72+
73+
HOSTDEVICE bool bin1D(unsigned int& bin_1d, const Scalar x, unsigned int dim) const
74+
{
75+
int bin_ = static_cast<int>(
76+
slow::floor(((x - m_lower_bounds[dim]) / (m_upper_bounds[dim] - m_lower_bounds[dim]))
77+
* m_num_bins[dim]));
78+
if (bin_ >= 0 && bin_ < static_cast<int>(m_num_bins[dim]))
79+
{
80+
bin_1d = bin_;
81+
return true;
82+
}
83+
else
84+
{
85+
return false;
86+
}
87+
}
88+
};
89+
90+
} // end namespace azplugins
91+
} // end namespace hoomd
92+
93+
#undef HOSTDEVICE
94+
95+
#endif // AZPLUGINS_BINNING_OPERATION_H_

‎src/CMakeLists.txt‎

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,11 +11,13 @@ set(_${COMPONENT_NAME}_sources
1111

1212
# TODO: List all GPU C++ source code files in _${COMPONENT_NAME}_cu_sources.
1313
set(_${COMPONENT_NAME}_cu_sources
14+
VelocityFieldComputeGPU.cu
1415
)
1516

1617
# TODO: List all Python modules in python_files.
1718
set(python_files
1819
__init__.py
20+
compute.py
1921
conftest.py
2022
bond.py
2123
flow.py
@@ -121,6 +123,29 @@ foreach(_evaluator ${_aniso_pair_evaluators})
121123
endif()
122124
endforeach()
123125

126+
# process velocity field geometries
127+
set(_binning_geometries
128+
Cartesian
129+
Cylindrical
130+
)
131+
foreach(_geometry ${_binning_geometries})
132+
configure_file(export_VelocityFieldCompute.cc.inc
133+
export_${_geometry}VelocityFieldCompute.cc
134+
@ONLY)
135+
set(_${COMPONENT_NAME}_sources ${_${COMPONENT_NAME}_sources} export_${_geometry}VelocityFieldCompute.cc)
136+
137+
if (ENABLE_HIP)
138+
configure_file(export_VelocityFieldComputeGPU.cc.inc
139+
export_${_geometry}VelocityFieldComputeGPU.cc
140+
@ONLY)
141+
configure_file(VelocityFieldComputeGPUKernel.cu.inc
142+
${_geometry}VelocityFieldComputeGPUKernel.cu
143+
@ONLY)
144+
set(_${COMPONENT_NAME}_sources ${_${COMPONENT_NAME}_sources} export_${_geometry}VelocityFieldComputeGPU.cc)
145+
set(_${COMPONENT_NAME}_cu_sources ${_${COMPONENT_NAME}_cu_sources} ${_geometry}VelocityFieldComputeGPUKernel.cu)
146+
endif()
147+
endforeach()
148+
124149
if (ENABLE_HIP)
125150
set(_cuda_sources ${_${COMPONENT_NAME}_cu_sources})
126151
endif (ENABLE_HIP)

‎src/CartesianBinningOperation.h‎

Lines changed: 61 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,61 @@
1+
// Copyright (c) 2018-2020, Michael P. Howard
2+
// Copyright (c) 2021-2024, Auburn University
3+
// Part of azplugins, released under the BSD 3-Clause License.
4+
5+
#ifndef AZPLUGINS_CARTESIAN_BINNING_OPERATION_H_
6+
#define AZPLUGINS_CARTESIAN_BINNING_OPERATION_H_
7+
8+
#include "BinningOperation.h"
9+
10+
#ifdef __HIPCC__
11+
#define HOSTDEVICE __host__ __device__
12+
#else
13+
#define HOSTDEVICE
14+
#endif // __HIPCC__
15+
16+
namespace hoomd
17+
{
18+
namespace azplugins
19+
{
20+
21+
//! Binning operation in Cartesian coordinates
22+
class CartesianBinningOperation : public BinningOperation
23+
{
24+
public:
25+
using BinningOperation::BinningOperation;
26+
27+
HOSTDEVICE bool bin(uint3& bin,
28+
Scalar3& transformed_vector,
29+
const Scalar3& coordinates,
30+
const Scalar3& vector) const
31+
{
32+
uint3 bin_ = make_uint3(0, 0, 0);
33+
34+
if (m_should_bin[0] && !bin1D(bin_.x, coordinates.x, 0))
35+
{
36+
return false;
37+
}
38+
39+
if (m_should_bin[1] && !bin1D(bin_.y, coordinates.y, 1))
40+
{
41+
return false;
42+
}
43+
44+
if (m_should_bin[2] && !bin1D(bin_.z, coordinates.z, 2))
45+
{
46+
return false;
47+
}
48+
49+
bin = bin_;
50+
transformed_vector = vector;
51+
52+
return true;
53+
}
54+
};
55+
56+
} // end namespace azplugins
57+
} // end namespace hoomd
58+
59+
#undef HOSTDEVICE
60+
61+
#endif // AZPLUGINS_CARTESIAN_BINNING_OPERATION_H_

‎src/CylindricalBinningOperation.h‎

Lines changed: 87 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,87 @@
1+
// Copyright (c) 2018-2020, Michael P. Howard
2+
// Copyright (c) 2021-2024, Auburn University
3+
// Part of azplugins, released under the BSD 3-Clause License.
4+
5+
#ifndef AZPLUGINS_CYLINDRICAL_BINNING_OPERATION_H_
6+
#define AZPLUGINS_CYLINDRICAL_BINNING_OPERATION_H_
7+
8+
#include "BinningOperation.h"
9+
10+
#ifdef __HIPCC__
11+
#define HOSTDEVICE __host__ __device__
12+
#else
13+
#define HOSTDEVICE
14+
#endif // __HIPCC__
15+
16+
namespace hoomd
17+
{
18+
namespace azplugins
19+
{
20+
21+
//! Binning operation in cylindrical coordinates
22+
class CylindricalBinningOperation : public BinningOperation
23+
{
24+
public:
25+
using BinningOperation::BinningOperation;
26+
27+
HOSTDEVICE bool bin(uint3& bin,
28+
Scalar3& transformed_vector,
29+
const Scalar3& coordinates,
30+
const Scalar3& vector) const
31+
{
32+
uint3 bin_ = make_uint3(0, 0, 0);
33+
34+
// bin with respect to z first because it is cheapest in case of early exit
35+
if (m_should_bin[2] && !bin1D(bin_.z, coordinates.z, 2))
36+
{
37+
return false;
38+
}
39+
40+
// bin with respect to theta second because its internals aren't needed later
41+
if (m_should_bin[1])
42+
{
43+
#if HOOMD_LONGREAL_SIZE == 32
44+
Scalar theta = ::atan2f(coordinates.y, coordinates.x);
45+
#else
46+
Scalar theta = ::atan2(coordinates.y, coordinates.x);
47+
#endif
48+
if (theta < Scalar(0))
49+
{
50+
theta += Scalar(2) * M_PI;
51+
}
52+
if (!bin1D(bin_.y, theta, 1))
53+
{
54+
return false;
55+
}
56+
}
57+
58+
// bin with respect to r last, r will be used for vector transform so always do it
59+
const Scalar r = fast::sqrt(coordinates.x * coordinates.x + coordinates.y * coordinates.y);
60+
if (m_should_bin[0] && !bin1D(bin_.x, r, 0))
61+
{
62+
return false;
63+
}
64+
65+
// transform Cartesian vector to cylindrical coordinates,
66+
// defaulting angle to 0 if r == 0 (point is exactly at center)
67+
Scalar cos_theta(1), sin_theta(0);
68+
if (r > Scalar(0))
69+
{
70+
cos_theta = coordinates.x / r;
71+
sin_theta = coordinates.y / r;
72+
}
73+
transformed_vector = make_scalar3(cos_theta * vector.x + sin_theta * vector.y,
74+
-sin_theta * vector.x + cos_theta * vector.y,
75+
vector.z);
76+
77+
bin = bin_;
78+
return true;
79+
}
80+
};
81+
82+
} // end namespace azplugins
83+
} // end namespace hoomd
84+
85+
#undef HOSTDEVICE
86+
87+
#endif // AZPLUGINS_CYLINDRICAL_BINNING_OPERATION_H_

0 commit comments

Comments
 (0)