Particle Loop#
This section contains the API documentation for Particle Loops.
Construct |
Access Modes |
|---|---|
CellDat |
Read, Write, Add |
CellDatConst |
Read, Write, Add, Min, Max, Reduction |
CellInfoNPart |
Read |
DescendantProducts |
Write |
GlobalArray |
Read, Add |
LocalArray |
Read, Write, Add |
LocalMemoryBlock |
Write |
LocalMemoryInterlaced |
Write |
MaskArray |
Read, Write |
NDLocalArray |
Read, Write, Add, Max, Min |
ParticleMask |
Read, Write |
ParticleSetDevice |
Read, Write, Add |
ProductMatrix |
Read, Write |
SymVector |
Read, Write |
ParticleDat |
Read, Write |
ParticleLoopIndex |
Read |
KernelRNG |
Read |
TupleRNG |
Read |
Particle Loop Access Descriptors#
This section contains the API documentation for Particle Loop access interface. For example the main functions and access descriptors. These functions are all in the Access namespace.
-
template<typename T>
inline Read<T> read(T t)# Helper function that allows a loop to be constructed with a read-only parameter passed like:
Access::read(object)
- Parameters:
t – Object to pass with read-only access.
- Returns:
Access::Read object that wraps passed object.
-
template<typename T>
inline Write<T> write(T t)# Helper function that allows a loop to be constructed with a write parameter passed like:
Access::write(object)
- Parameters:
t – Object to pass with write access.
- Returns:
Access::Write object that wraps passed object.
-
template<typename T>
inline Add<T> add(T t)# Helper function that allows a loop to be constructed with a write parameter which is atomic addition passed like:
Access::add(object)
- Parameters:
t – Object to pass with atomic add access.
- Returns:
Access::Add object that wraps passed object.
-
template<typename T>
inline Min<T> min(T t)# Helper function that allows a loop to be constructed with a write parameter which is atomic minimum passed like:
Access::min(object)
- Parameters:
t – Object to pass with atomic min access.
- Returns:
Access::Min object that wraps passed object.
-
template<typename T>
inline Max<T> max(T t)# Helper function that allows a loop to be constructed with a write parameter which is atomic maximum passed like:
Access::max(object)
- Parameters:
t – Object to pass with atomic max access.
- Returns:
Access::Max object that wraps passed object.
-
template<typename T, typename OP>
inline Reduction<T, OP> reduce(
)# Helper function that allows a loop to be constructed with a reduction argument like
Access::reduce(object, binop)
- Parameters:
t – Object to reduce values into.
binop – Binary operation to use for reduction, e.g. Kernel::plus<REAL>.
- Returns:
Access::Max object that wraps passed object.
Particle Loop Functions#
This section contains the API documentation for Particle Loop functions. i.e. functions which create Particle Loop objects.
-
)#
Create a ParticleLoop that executes a kernel for all particles in the ParticleGroup.
- Parameters:
name – Identifier for particle loop.
particle_group – ParticleGroup to execute kernel for all particles.
kernel – Kernel to execute for all particles in the ParticleGroup.
args – The remaining arguments are arguments to be passed to the kernel. All arguments must be wrapped in an access descriptor type.
-
)#
Create a ParticleLoop that executes a kernel for all particles in the ParticleGroup.
- Parameters:
particle_group – ParticleGroup to execute kernel for all particles.
kernel – Kernel to execute for all particles in the ParticleGroup.
args – The remaining arguments are arguments to be passed to the kernel. All arguments must be wrapped in an access descriptor type.
-
)#
Create a ParticleLoop that executes a kernel for all particles in the ParticleDat.
- Parameters:
name – Identifier for particle loop.
particle_dat – ParticleDat to define the iteration set.
kernel – Kernel to execute for all particles in the ParticleGroup.
args – The remaining arguments are arguments to be passed to the kernel. All arguments must be wrapped in an access descriptor type.
-
)#
Create a ParticleLoop that executes a kernel for all particles in the ParticleDat.
- Parameters:
particle_dat – ParticleDat to define the iteration set.
kernel – Kernel to execute for all particles in the ParticleGroup.
args – The remaining arguments are arguments to be passed to the kernel. All arguments must be wrapped in an access descriptor type.
-
)#
Create a ParticleLoop that executes a kernel for all particles in the ParticleSubGroup.
- Parameters:
name – Identifier for particle loop.
particle_sub_group – ParticleSubGroup to execute kernel for all particles.
kernel – Kernel to execute for all particles in the ParticleSubGroup.
args – The remaining arguments are arguments to be passed to the kernel. All arguments must be wrapped in an access descriptor type.
-
)#
Create a ParticleLoop that executes a kernel for all particles in the ParticleSubGroup.
- Parameters:
particle_sub_group – ParticleSubGroup to execute kernel for all particles.
kernel – Kernel to execute for all particles in the ParticleSubGroup.
args – The remaining arguments are arguments to be passed to the kernel. All arguments must be wrapped in an access descriptor type.
-
class ParticleLoopBase#
- #include <particle_loop_base.hpp>
Abstract base class for ParticleLoop such that the templated ParticleLoop can be cast to a base type for storage.
Subclassed by NESO::Particles::ParticleLoop< KERNEL, ARGS… >, NESO::Particles::ParticleLoop< KERNEL, ARGS >, NESO::Particles::ReductionLoopCellwiseBins< KERNEL, ARGS >
Public Functions
- inline virtual void execute(
- const std::optional<int> cell_start = std::nullopt,
- const std::optional<int> cell_end = std::nullopt
Execute the ParticleLoop and block until execution is complete. Must be called collectively on the MPI communicator associated with the SYCLTarget this loop is over.
execute() Launches the ParticleLoop over all cells. execute(i) Launches the ParticleLoop over cell i. execute(i, i+4) Launches the ParticleLoop over cells i, i+1, i+2, i+3. Note cell_end itself is not visited.
- Parameters:
cell_start – Optional starting cell to launch the ParticleLoop over.
cell_end – Optional ending cell to launch the ParticleLoop over.
- inline virtual void submit(
- const std::optional<int> cell_start = std::nullopt,
- const std::optional<int> cell_end = std::nullopt
Launch the ParticleLoop and return. Must be called collectively over the MPI communicator of the ParticleGroup. Loop execution is complete when the corresponding call to wait returns.
submit() Launches the ParticleLoop over all cells. submit(i) Launches the ParticleLoop over cell i. submit(i, i+4) Launches the ParticleLoop over cells i, i+1, i+2, i+3. Note cell_end itself is not visited.
- Parameters:
cell_start – Optional starting cell to launch the ParticleLoop over.
cell_end – Optional ending cell to launch the ParticleLoop over.
-
inline virtual void wait() = 0#
Wait for loop execution to complete. On completion perform post-loop actions. Must be called collectively on communicator.
Public Members
-
std::shared_ptr<void> particle_dat_init = {nullptr}#
This stores the particle dat the loop was created with to prevent use after free errors in the case when the ParticleLoop is created with a ParticleDat.
Kernel Functions#
This section contains the API documentation for device callable functions. These functions can be called within Particle Loops. These functions are all in the Kernel namespace. See the SYCL specification for what these maths functions actually do.
-
template<typename T, typename U>
inline auto min(
)# Returns the minimum of the two arguments.
- Parameters:
x – First argument.
y – second argument.
- Returns:
Minimum of x and y.
-
inline auto min(const REAL &x, const REAL &y)#
Returns the minimum of the two arguments.
- Parameters:
x – First argument.
y – second argument.
- Returns:
Minimum of x and y.
-
template<typename T, typename U>
inline auto max(
)# Returns the maximum of the two arguments.
- Parameters:
x – First argument.
y – second argument.
- Returns:
Maximum of x and y.
-
inline auto max(const REAL &x, const REAL &y)#
Returns the maximum of the two arguments.
- Parameters:
x – First argument.
y – Second argument.
- Returns:
Maximum of x and y.
-
template<typename T>
inline auto abs(const T &x)# Returns the absolute value of x.
- Parameters:
x – Argument.
- Returns:
Absolute value of x.
-
inline auto abs(const REAL &x)#
Returns the absolute value of x.
- Parameters:
x – Argument.
- Returns:
Absolute value of x.
-
inline auto sqrt(const REAL &x)#
Returns the square root of x.
- Parameters:
x – Argument.
- Returns:
Square root of x.
-
template<typename T>
inline auto rsqrt(const T x)# Returns the reciprocal square root of x.
- Parameters:
x – Argument.
- Returns:
Reciprocal square root of x.
-
inline auto exp(const REAL &x)#
Returns exp(x).
- Parameters:
x – Argument.
- Returns:
exp(x).
-
inline auto fmod(const REAL &x, const REAL &y)#
Returns fmod(x, y).
- Parameters:
x – First argument.
y – Second argument.
- Returns:
fmod(x,y).
-
inline auto fma(const REAL &x, const REAL &y, const REAL &z)#
Returns x*y+z.
- Parameters:
x – First argument.
y – Second argument.
z – Third argument.
- Returns:
x * y + z.
-
template<typename T>
inline auto pow(const T x, const T y)# Returns x to the power y.
- Parameters:
x – First argument.
y – Second argument.
- Returns:
x^y.
-
template<typename T>
inline auto atan2(const T y, const T x)# Returns atan2(x, y).
- Parameters:
x – First argument.
y – Second argument.
- Returns:
atan2(x,y).
-
template<typename T, typename U>
inline auto copysign(
)# Returns copysign(x, y).
- Parameters:
x – First argument.
y – Second argument.
- Returns:
copysign(x,y).
-
template<typename T>
inline auto asin(const T x)# Returns asin(x).
- Parameters:
x – Argument.
- Returns:
asin(x).
-
template<typename T>
inline auto sin(const T x)# Returns sin(x).
- Parameters:
x – Argument.
- Returns:
sin(x).
-
template<typename T>
inline auto acos(const T x)# Returns acos(x).
- Parameters:
x – Argument.
- Returns:
acos(x).
-
template<typename T>
inline auto cos(const T x)# Returns cos(x).
- Parameters:
x – Argument.
- Returns:
cos(x).
-
template<typename T>
inline auto atan(const T x)# Returns atan(x).
- Parameters:
x – Argument.
- Returns:
atan(x).
-
template<typename T>
inline auto tan(const T x)# Returns tan(x).
- Parameters:
x – Argument.
- Returns:
tan(x).
-
template<typename T>
inline auto log(const T x)# Returns log(x).
- Parameters:
x – Argument.
- Returns:
log(x).
-
template<typename T>
inline auto log2(const T x)# Returns log2(x).
- Parameters:
x – Argument.
- Returns:
log2(x).
-
template<typename T>
inline auto log10(const T x)# Returns log10(x).
- Parameters:
x – Argument.
- Returns:
log10(x).
-
template<typename T>
inline auto round(const T x)# Returns round(x).
- Parameters:
x – Argument.
- Returns:
round(x).
-
template<typename T>
inline auto tgamma(const T x)# Returns tgamma(x).
- Parameters:
x – Argument.
- Returns:
tgamma(x).
-
template<typename T>
inline auto trunc(const T x)# Returns trunc(x).
- Parameters:
x – Argument.
- Returns:
trunc(x).
-
template<typename T>
inline auto ceil(const T x)# Returns ceil(x).
- Parameters:
x – Argument.
- Returns:
ceil(x).
-
template<typename T>
inline auto floor(const T x)# Returns floor(x).
- Parameters:
x – Argument.
- Returns:
floor(x).
-
template<typename T>
inline auto clamp(
)# Returns x clamped into [minval, maxval].
- Parameters:
x – Argument.
minval – Minimum returned value.
maxval – Maximum returned value.
- Returns:
x clamped into [minval, maxval].
-
template<typename T>
inline auto popcount(const T x)# Returns popcount(x).
- Parameters:
x – Argument.
- Returns:
popcount(x).
- inline void consistent_line_orientation_2d(
- const REAL ax,
- const REAL ay,
- const REAL bx,
- const REAL by,
- REAL *ix,
- REAL *iy,
- REAL *jx,
- REAL *jy
For a line segment [ax, ay] - [bx, by] return ix,iy,jx,jy such that [ix, iy]
[jx, jy] is the same line segment as [ax, ay] - [bx, by] but the outputs are independent of the order in which a and b are specified. i.e. the direction of the output line segment is always in the same direction.
- Parameters:
ax – [in] First point of line segment, x coordinate.
ay – [in] First point of line segment, y coordinate.
bx – [in] Second point of line segment, x coordinate.
by – [in] Second point of line segment, y coordinate.
ix – [inout] First point of output line segment, x coordinate.
iy – [inout] First point of output line segment, y coordinate.
jx – [inout] Second point of output line segment, x coordinate.
jy – [inout] Second point of output line segment, y coordinate.
- inline void line_segment_intersection_2d_lambda(
- const REAL &xa,
- const REAL &ya,
- const REAL &xb,
- const REAL &yb,
- const REAL &x0,
- const REAL &y0,
- const REAL &x1,
- const REAL &y1,
- REAL &l0,
- REAL &l1
Compute the intersection point parameter (lambda0, lambda1) for the lines [(xa, ya), (xb, yb)] and [(x0, y0), (x1, y1)]. Assuming that x1 - x0 != 0 and that the lines are not parallel.
param[in] xa Input coordinate for point a. param[in] ya Input coordinate for point a. param[in] xb Input coordinate for point b. param[in] yb Input coordinate for point b. param[in] x0 Input coordinate for point 0. param[in] y0 Input coordinate for point 0. param[in] x1 Input coordinate for point 1. param[in] y1 Input coordinate for point 1. param[in, out] l0 Output intersection lambda0 if it exists. param[in, out] l1 Output intersection lambda1 if it exists.
- inline bool line_segment_intersection_2d(
- const REAL &xa,
- const REAL &ya,
- const REAL &xb,
- const REAL &yb,
- const REAL &x0,
- const REAL &y0,
- const REAL &x1,
- const REAL &y1,
- REAL &xi,
- REAL &yi,
- REAL &l0_out,
- const REAL tol = 0.0
Compute the intersection point of two line segments [(xa, ya), (xb, yb)] and [(x0, y0), (x1, y1)].
param[in] xa Input coordinate for point a. param[in] ya Input coordinate for point a. param[in] xb Input coordinate for point b. param[in] yb Input coordinate for point b. param[in] x0 Input coordinate for point 0. param[in] y0 Input coordinate for point 0. param[in] x1 Input coordinate for point 1. param[in] y1 Input coordinate for point 1. param[in, out] xi Output intersection point if it exists. param[in, out] yi Output intersection point if it exists. param[in, out] l0_out Proportion of the distance between a and b the intersection point exists at. param[in] tol Tolerance for intersection, e.g. how closely do the lines pass at the ends default 0.0.
- Returns:
True if the line segments intersect otherwise false.
- inline bool line_segment_intersection_2d_x_axis_aligned(
- const REAL &xa,
- const REAL &ya,
- const REAL &xb,
- const REAL &yb,
- const REAL &x0,
- const REAL &y0,
- const REAL &x1,
- REAL &xi,
- REAL &yi,
- const REAL tol = 0.0
Compute the intersection point of two line segments [(xa, ya), (xb, yb)] and [(x0, y0), (x1, y0)].
Note that the second line segment, (x,y) is aligned with the x-axis and we assume that !(ya == y0 && yb == y0).
param[in] xa Input coordinate for point a. param[in] ya Input coordinate for point a. param[in] xb Input coordinate for point b. param[in] yb Input coordinate for point b. param[in] x0 Input coordinate for point x0. param[in] y0 Input coordinate for point y0. param[in] x1 Input coordinate for point x0. param[in, out] xi Output intersection point if it exists. param[in, out] yi Output intersection point if it exists. param[in] tol Tolerance for intersection, e.g. how closely do the lines pass at the ends default 0.0.
- Returns:
True if the line segments intersect otherwise false.
- inline bool line_segment_intersection_2d_y_axis_aligned(
- const REAL &xa,
- const REAL &ya,
- const REAL &xb,
- const REAL &yb,
- const REAL &x0,
- const REAL &y0,
- const REAL &y1,
- REAL &xi,
- REAL &yi,
- const REAL tol = 0.0
Compute the intersection point of two line segments [(xa, ya), (xb, yb)] and [(x0, y0), (x0, y1)].
Note that the second line segment, (x,y) is aligned with the y-axis and we assume that !(xa == x0 && xb == x0).
param[in] xa Input coordinate for point a. param[in] ya Input coordinate for point a. param[in] xb Input coordinate for point b. param[in] yb Input coordinate for point b. param[in] x0 Input coordinate for point x0. param[in] y0 Input coordinate for point y0. param[in] y1 Input coordinate for point y0. param[in, out] xi Output intersection point if it exists. param[in, out] yi Output intersection point if it exists. param[in] tol Tolerance for intersection, e.g. how closely do the lines pass at the ends default 0.0.
- Returns:
True if the line segments intersect otherwise false.
- inline bool plane_intersection_3d_xy_plane_aligned(
- const REAL &ax,
- const REAL &ay,
- const REAL &az,
- const REAL &bx,
- const REAL &by,
- const REAL &bz,
- const REAL &p0x,
- const REAL &p0y,
- const REAL &p0z,
- const REAL &p1x,
- const REAL &p2y,
- REAL &xi,
- REAL &yi,
- REAL &zi,
- const REAL tol = 0.0
Intersection of line segment
with the plane segment[(ax, ay, az), (bx, by, bz)]
Assumes that the line does not lie in the plane.(p0x, p2y, p0z) (p1x, p2y, p0z) ------ y | | ^ | | | ------ ---> x (p0x, p0y, p0z) (p1x, p0y, p0z)
param[in] ax Input coordinate for point a. param[in] ay Input coordinate for point a. param[in] az Input coordinate for point a. param[in] bx Input coordinate for point b. param[in] by Input coordinate for point b. param[in] bz Input coordinate for point b. param[in] p0x Input coordinate for plane. param[in] p0y Input coordinate for plane. param[in] p0z Input coordinate for plane. param[in] p1x Input coordinate for plane. param[in] p2y Input coordinate for plane. param[in, out] xi Output intersection point if it exists. param[in, out] yi Output intersection point if it exists. param[in, out] zi Output intersection point if it exists. param[in] tol Tolerance for intersection, e.g. how closely do the lines pass at the ends default 0.0.
- Returns:
True if the line segments intersect otherwise false.
- inline void evaluate_barycentric_coordinates(
- const sycl::marray<REAL, 3> &bary_coords,
- const sycl::marray<REAL, 3> &triangle_vertex_0,
- const sycl::marray<REAL, 3> &triangle_vertex_1,
- const sycl::marray<REAL, 3> &triangle_vertex_2,
- sycl::marray<REAL, 3> &coords
Helper function to evaluate Barycentric coordinates using a set of vertices for a triangle.
- Parameters:
bary_coords – [in] Barycentric coordinates.
triangle_vertex_0 – [in] First vertex of triangle.
triangle_vertex_1 – [in] Second vertex of triangle.
triangle_vertex_2 – [in] Third vertex of triangle.
coords – [inout] Output evaluation.
- inline bool line_triangle_intersection_moller_trumbore(
- const sycl::marray<REAL, 3> &line_origin,
- const sycl::marray<REAL, 3> &line_direction,
- const sycl::marray<REAL, 3> &triangle_vertex_0,
- const sycl::marray<REAL, 3> &triangle_vertex_1,
- const sycl::marray<REAL, 3> &triangle_vertex_2,
- sycl::marray<REAL, 3> &bary_coords,
- REAL ¶meterised_distance,
- const REAL tol_plane = 0.0,
- const REAL tol_contained = 0.0
Line segment - Triangle intersection test, in 3D, using the Möller–Trumbore intersection algorithm.
- Parameters:
line_origin – [in] Origin point of line.
line_direction – [in] Direction of line.
triangle_vertex_0 – [in] First vertex of triangle.
triangle_vertex_1 – [in] Second vertex of triangle.
triangle_vertex_2 – [in] Third vertex of triangle.
bary_coords – [inout] Point of intersection in Barycentric coordinates.
parameterised_distance – [inout] Point of intersection in terms of direction vector.
tol_plane – [in] Tolerance for the line segment embedded in the plane of the triangle, default 0.0.
tol_contained – [in] Tolerance for intersection, default 0.0.
-
template<std::size_t N>
inline void naive_matrix_inverse( - const REAL *M,
- REAL *L
Naively invert a matrix. The error bars on this call may be quite large. This function uses row-major format.
- Parameters:
M – [in] Matrix to invert.
L – [inout] Output space for M^-1.
-
template<>
inline void naive_matrix_inverse<3>( - const REAL *M,
- REAL *L
Naively invert a matrix. The error bars on this call may be quite large. This function uses row-major format.
- Parameters:
M – [in] Matrix to invert.
L – [inout] Output space for M^-1.
-
template<>
inline void naive_matrix_inverse<4>( - const REAL *M,
- REAL *L
Naively invert a matrix. The error bars on this call may be quite large. This function uses row-major format.
- Parameters:
M – [in] Matrix to invert.
L – [inout] Output space for M^-1.
-
inline REAL relative_error(const REAL correct, const REAL to_test)#
- Parameters:
correct – Correct reference value.
to_test – Value to test against reference value.
- Returns:
|correct - to_test| / |correct| if |correct| > 0.0 else |correct - to_test|.
CellDat#
The CellDat object is a 2D matrix per mesh cell with a variable number of rows across mesh cells and a fixed number of columns across all mesh cells.
-
template<typename T>
struct Read# - #include <cell_dat.hpp>
Access:CellDat::Read<T> read access cellwise.
Public Functions
-
Read() = default#
Pointer to underlying data for the array.
-
Read() = default#
-
template<typename T>
struct Add# - #include <cell_dat.hpp>
Access:CellDat::Add<T> add access cellwise.
Public Functions
-
Add() = default#
Pointer to underlying data for the array.
-
Add() = default#
-
template<typename T>
struct Write# - #include <cell_dat.hpp>
Access:CellDat::Add<T> write access cellwise.
Public Functions
-
Write() = default#
Pointer to underlying data for the array.
-
Write() = default#
-
template<typename T>
class CellDat# - #include <cell_dat.hpp>
Store data on each cell where the number of columns required per cell is constant but the number of rows is variable. Data is stored in a column major manner with a new device pointer per column.
Public Functions
-
inline T *col_device_ptr(const int cell, const int col)#
Get the device pointer for a column in a cell. Do NOT use this method to access particle data - it will break in a hard to detect manner. Instead use the methods for accessing particle data directly which are designed for accessing particle data directly.
- Parameters:
cell – Cell index to get pointer for.
col – Column in cell to get pointer for.
- Returns:
Device pointer to data for the specified column.
- SYCLTargetSharedPtr sycl_target,
- const int ncells,
- const int ncol
Create new CellDat on a specified compute target with a specified number of cells and number of columns per cell.
- Parameters:
sycl_target – SYCLTargetSharedPtr to use as compute device.
ncells – Number of cells (fixed).
ncol – Number of columns in each cell (fixed).
-
template<typename U>
inline void reduce_nrow( - const U *h_nrow_required
Pass a number of nrows for each cell where the user promises the new number of rows are equal to or less than the current number of rows.
-
inline void set_nrow(const INT cell, const INT nrow_required)#
Set the number of rows required in a provided cell. This will realloc if needed and copy the existing data into the new space. May not shrink the array if the requested size is smaller than the existing size. wait_set_nrow should be called before using the dat.
-
inline void wait_set_nrow()#
Wait for set_nrow to complete
-
inline int compute_nrow_max()#
Recompute nrow_max from current row counts.
- Returns:
The maximum number of rows across all cells.
-
inline int get_nrow_max()#
Get the maximum number of rows across all cells.
- Returns:
The maximum number of rows across all cells.
-
inline int get_nrow_min()#
Get the minimum number of rows across all cells.
- Returns:
The minimum number of rows across all cells.
-
inline CellData<T> get_cell(const int cell)#
Get the contents of a provided cell on the host as a CellData instance.
- Parameters:
cell – Cell to get data from.
- Returns:
Cell contents of specified cell as CellData instance.
-
inline T get_value(const int cell, const int row, const int col)#
Get a value from a cell, row and column.
- Parameters:
cell – Cell index.
row – Row index.
col – Column index.
- Returns:
Value at location.
- inline void set_value(
- const int cell,
- const int row,
- const int col,
- const T value
Set a value directly using cell, row, column and value.
- Parameters:
cell – Cell index.
row – Row index.
col – Column index.
value – Value to set at location.
- inline void get_cell_async(
- const int cell,
- CellData<T> cell_data,
- EventStack &event_stack
Get the contents of a provided cell on the host as a CellData instance.
- Parameters:
cell – Cell to get data from.
cell_data – CellData instance to populate, must be sufficiently sized.
event_stack – EventStack instance to call wait on for copy.
-
inline void set_cell(const int cell, CellData<T> cell_data)#
Set the contents of a cell on the device using a CellData instance.
- Parameters:
cell – Cell index to set data in.
cell_data – New cell data to set.
- inline void set_cell_async(
- const int cell,
- CellData<T> cell_data,
- EventStack &event_stack
Set the contents of a cell on the device using a CellData instance.
- Parameters:
cell – Cell index to set data in.
cell_data – New cell data to set.
event_stack – EventStack instance to wait on.
-
inline T ***device_ptr()#
Get the root device pointer for the data storage. Data can be accessed on the device in SYCL kernels with access like: d[cell_index][column_index][row_index]
- Returns:
Device pointer that can be used to access the underlying data.
-
inline void print(int start = -1, int end = -1)#
Helper function to print the contents of all cells or a specified range of cells.
- Parameters:
start – (optional) First cell to print.
end – (option) Last cell minus one to print.
-
inline T *col_device_ptr(const int cell, const int col)#
CellDatConst#
A CellDatConst is a 2D matrix per mesh cell where all mesh cells have a matrix with the same number of rows and columns.
-
template<typename T>
struct Read# - #include <cell_dat_const.hpp>
Access:CellDatConst::Read<T> is a kernel argument type for accessing CellDatConst data in a kernel.
Public Functions
-
Read() = default#
Pointer to underlying data for the array.
-
Read() = default#
-
template<typename T>
struct Write# - #include <cell_dat_const.hpp>
Access:CellDatConst::Write<T> is a kernel argument type for accessing CellDatConst data in a kernel.
Public Functions
-
Write() = default#
Pointer to underlying data for the array.
-
Write() = default#
-
template<typename T>
struct Add# - #include <cell_dat_const.hpp>
Access:CellDatConst::Add<T> is a kernel argument type for accessing CellDatConst data in a kernel.
Public Functions
-
Add() = default#
Pointer to underlying data for the array.
-
Add() = default#
-
template<typename T>
struct Min# - #include <cell_dat_const.hpp>
Access:CellDatConst::Min<T> is a kernel argument type for accessing CellDatConst data in a kernel.
Public Functions
-
Min() = default#
Pointer to underlying data for the array.
-
Min() = default#
-
template<typename T>
struct Max# - #include <cell_dat_const.hpp>
Access:CellDatConst::Max<T> is a kernel argument type for accessing CellDatConst data in a kernel.
Public Functions
-
Max() = default#
Pointer to underlying data for the array.
-
Max() = default#
-
template<typename T, typename OP>
struct Reduction# - #include <cell_dat_const.hpp>
Kernel type for Reduction for CellDatConst.
-
template<typename T>
class CellDatConst# - #include <cell_dat_const.hpp>
Container that allocates on the device a matrix of fixed size nrow X ncol for N cells. Data stored in column major format. i.e. Data order from slowest to fastest is: cell, column, row.
Public Functions
-
CellDatConst(const CellDatConst &st) = delete#
Disable (implicit) copies.
-
CellDatConst &operator=(CellDatConst const &a) = delete#
Disable (implicit) copies.
-
inline void fill(const T value)#
Fill all the entries with a given value.
- Parameters:
value – Value to place in all entries.
- SYCLTargetSharedPtr sycl_target,
- const int ncells,
- const int nrow,
- const int ncol
Create new CellDatConst on the specified compute target with a fixed cell count, fixed number of rows per cell and fixed number of columns per cell.
- Parameters:
sycl_target – SYCLTargetSharedPtr to use as compute device.
ncells – Number of cells.
nrow – Number of rows.
ncol – Number of columns.
-
inline int idx(const int cell, const int row, const int col)#
Helper function to index into the stored data. Note column major format.
- Parameters:
cell – Cell index to index into.
row – Row to access.
col – Column to access.
- Returns:
Linear index into data structure for the specified row and column.
-
inline T *device_ptr()#
Get the device pointer for the underlying data. Only accessible on the device.
- Returns:
Returns a device pointer on the compute target.
-
inline CellData<T> get_cell(const int cell)#
Get the data stored in a provided cell on the host as a CellData instance.
- Parameters:
cell – Cell to access the underlying data for.
- Returns:
Cell data for cell.
-
inline void set_cell(const int cell, CellData<T> cell_data)#
Set the data in a cell using a CellData instance.
- Parameters:
cell – Cell to set data for.
cell_data – Source data.
-
inline std::vector<CellData<T>> get_all_cells()#
Get the cell data for all cells.
- Returns:
CellData instances for all cells.
- inline void set_all_cells(
- const std::vector<CellData<T>> &cell_data
Set the cell data for all cells.
- Parameters:
cell_data – Vector of CellData instances each of size nrow, ncol.
-
CellDatConst(const CellDatConst &st) = delete#
DescendantProducts#
Particle Loop construct for creating particles from particles.
-
struct Write#
- #include <descendant_products.hpp>
ParticleLoop access type for DescendantProducts Write access.
Public Functions
-
template<typename INDEX_TYPE>
inline REAL &at_real( - const INDEX_TYPE &particle_index,
- const int product,
- const int property,
- const int component
Write access to REAL particle property.
- Parameters:
particle_index – The index of the particle creating this product.
product – The index of the product to access, i.e. the row in the product matrix.
property – The particle property to access using the ordering REAL ParticleProp instances were passed to the ProductMatrixSpec
component – The component of the property to access.
- Returns:
Modifiable reference to property value.
-
template<typename INDEX_TYPE>
inline INT &at_int( - const INDEX_TYPE &particle_index,
- const int product,
- const int property,
- const int component
Write access to INT particle property.
- Parameters:
particle_index – The index of the particle creating this product.
product – The index of the product to access, i.e. the row in the product matrix.
property – The particle property to access using the ordering INT ParticleProp instances were passed to the ProductMatrixSpec
component – The component of the property to access.
- Returns:
Modifiable reference to property value.
-
template<typename INDEX_TYPE>
inline void set_parent( - const INDEX_TYPE &particle_index,
- const int product
Set the parent of a product.
- Parameters:
particle_index – Index of parent particle.
product – Index of child property to set as descendant of parent.
-
template<typename INDEX_TYPE>
-
class DescendantProducts : public NESO::Particles::ProductMatrix#
- #include <descendant_products.hpp>
Class to create M products from each parent particle. Products may inherit property values from the parent.
Public Functions
- SYCLTargetSharedPtr sycl_target,
- std::shared_ptr<ProductMatrixSpec> spec,
- const int num_products_per_parent
Create an instance from a ProductMatrixSpec that defines which particle properties should be available for writing in the particle loop kernel.
- Parameters:
sycl_target – Compute device on which products will be made.
spec – Specification of properties which require specialisation in the creating kernel.
num_products_per_parent – The maximum number of products required per parent particle.
-
virtual void reset(const int num_parent_particles) override#
Resize the container to hold space for the products of a number of parent particles.
- Parameters:
num_parent_particles – Number of parent particles space is required for.
GlobalArray#
An array type with globally fixed size where the result from each MPI rank is reduced across all MPI ranks on loop completetion.
-
template<typename T>
struct Read# - #include <global_array.hpp>
Access:GlobalArray::Read<T> and Access:GlobalArray::Add<T> are the kernel argument types for accessing GlobalArray data in a kernel.
-
template<typename T>
struct Add# - #include <global_array.hpp>
Access:GlobalArray::Read<T> and Access:GlobalArray::Add<T> are the kernel argument types for accessing GlobalArray data in a kernel.
Public Functions
-
template<typename T>
class GlobalArray# - #include <global_array.hpp>
GlobalArray is an array type which can be accessed from kernels in read or atomic add mode. Post loop execution, with add access mode, the global array values are automatically reduced across the MPI communicator.
Public Functions
-
GlobalArray<T> &operator=(const GlobalArray<T>&) = default#
Note that the copy operator creates shallow copies of the array.
- SYCLTargetSharedPtr sycl_target,
- const std::size_t size,
- const std::optional<T> init_value = std::nullopt
Create a new GlobalArray on a compute target and given size.
- Parameters:
sycl_target – Device to create GlobalArray on.
size – Number of elements in array.
init_value – Default value to initialise values to.
-
inline void fill(const T value)#
Fill the array with a value.
- Parameters:
value – Value to fill the array with - not reduced across MPI ranks.
-
inline sycl::event get_async(std::vector<T> &data)#
Asynchronously get the values in the local array into a std::vector.
- Parameters:
data – [inout] Input vector to copy values from GlobalArray into.
- Returns:
Event to wait on before using new values in the std::vector.
-
inline void get(std::vector<T> &data)#
Get the values in the local array into a std::vector. Blocks until copy is complete.
- Parameters:
data – [inout] Input vector to copy values from GlobalArray into.
-
inline std::vector<T> get()#
Get the values in the local array into a std::vector.
- Returns:
std::vector of values in the GlobalArray.
Public Members
-
SYCLTargetSharedPtr sycl_target#
The SYCLTarget the GlobalArray is created on.
-
MPI_Comm comm#
The MPI communicator reductions are made over.
-
std::size_t size = {0}#
The number of elements in the array.
-
GlobalArray<T> &operator=(const GlobalArray<T>&) = default#
LocalArray#
Creates an array local to each MPI rank that can be accessed from the Particle Loop.
-
template<typename T>
struct Read# - #include <local_array.hpp>
Access:LocalArray::Read<T>, Access::LocalArray::Write<T> and Access:LocalArray::Add<T> are the kernel argument types for accessing LocalArray data in a kernel. ParticleLoop access type for LocalArray Read access.
Public Functions
-
template<typename T>
struct Add# - #include <local_array.hpp>
ParticleLoop access type for LocalArray Add access.
Public Functions
-
template<typename T>
struct Write# - #include <local_array.hpp>
ParticleLoop access type for LocalArray Write access.
Public Functions
-
template<typename T>
class LocalArray# - #include <local_array.hpp>
Container to hold an array of values on each MPI rank.
Public Functions
-
LocalArray<T> &operator=(const LocalArray<T>&) = default#
Note that the copy operator creates shallow copies of the array.
- SYCLTargetSharedPtr sycl_target,
- const std::size_t size,
- const std::optional<T> init_value = std::nullopt
Create a new LocalArray on a compute target and given size.
- Parameters:
sycl_target – Device to create LocalArray on.
size – Number of elements in array.
init_value – Default value to initialise values to.
-
inline void fill(const T value)#
Fill the array with a value.
- Parameters:
value – Value to fill the array with.
- SYCLTargetSharedPtr sycl_target,
- const std::vector<T> &data
Create a new LocalArray on a compute target and given size.
- Parameters:
sycl_target – Device to create LocalArray on.
data – Vector to initialise array values to.
-
inline sycl::event set_async(const std::vector<T> &data)#
Asynchronously set the values in the local array to those in a std::vector.
- Parameters:
data – Input vector to copy values from.
- Returns:
Event to wait on before using new values in LocalArray.
-
inline void set(const std::vector<T> &data)#
Set the values in the local array to those in a std::vector. Blocks until the copy is complete.
- Parameters:
data – Input vector to copy values from.
-
inline sycl::event get_async(std::vector<T> &data)#
Asynchronously get the values in the local array into a std::vector.
- Parameters:
data – [inout] Input vector to copy values from LocalArray into.
- Returns:
Event to wait on before using new values in the std::vector.
-
inline void get(std::vector<T> &data)#
Get the values in the local array into a std::vector. Blocks until copy is complete.
- Parameters:
data – [inout] Input vector to copy values from LocalArray into.
-
inline std::vector<T> get()#
Get the values in the local array into a std::vector.
- Returns:
std::vector of values in the LocalArray.
-
inline void realloc_no_copy(const size_t size)#
Reallocate the buffer to hold the requested number of elements. Current contents is not copied to the new buffer.
- Parameters:
size – Number of elements this buffer should hold.
Public Members
-
SYCLTargetSharedPtr sycl_target#
The SYCLTarget the LocalArray is created on.
-
std::size_t size#
The number of elements in the array.
-
LocalArray<T> &operator=(const LocalArray<T>&) = default#
LocalMemoryBlock#
Construct for creating a contiguous block of local memory per particle in the Particle Loop.
-
template<typename T>
class LocalMemoryBlock# - #include <local_memory_block.hpp>
Type to provide kernel local memory on a per particle basis.
Public Functions
-
inline LocalMemoryBlock(const std::size_t size)#
Constructor with specified number of elements per particle.
- Parameters:
size – Number of elements required per particle.
-
inline void set_size(const std::size_t size)#
Set the local memory required to a new number of elements.
- Parameters:
size – Number of required elements per particle.
Public Members
-
std::size_t size#
Number of elements of type T required per particle.
-
inline LocalMemoryBlock(const std::size_t size)#
LocalMemoryInterlaced#
Construct for creating a fixed number of elements per particle in the Particle Loop where elements are stored with a non-unit stride.
-
template<typename T>
class LocalMemoryInterlaced# - #include <local_memory_interlaced.hpp>
Type to provide kernel local memory on a per particle basis which is interlaced between particles in the workgroup.
Public Functions
-
inline LocalMemoryInterlaced(const std::size_t size)#
Constructor with specified number of elements per particle. With this container values are interlaced.
- Parameters:
size – Number of elements required per particle.
-
inline void set_size(const std::size_t size)#
Set the local memory required to a new number of elements.
- Parameters:
size – Number of required elements per particle.
Public Members
-
std::size_t size#
Number of elements of type T required per particle.
-
inline LocalMemoryInterlaced(const std::size_t size)#
MaskArray#
Container for per particle masks that can be accessed in Particle Loops.
-
using Read = MaskArrayDeviceBase<MaskArrayBaseType const*>#
Read access descriptor for MaskArray.
-
template<typename T>
struct MaskArrayDeviceBase# - #include <mask_array.hpp>
Read access descriptor for MaskArray.
Public Functions
- inline std::size_t get_base_index(
- const std::size_t index_entry,
- const std::size_t index_bit
- Parameters:
index_entry – Index of the entry in [0, size)·
index_bit – Index of the mask (bit) in the entry in [0, num_masks_per_entry).
- Returns:
Offset to the MaskArrayBaseType containing the mask.
- inline bool get(
- const std::size_t index_entry,
- const std::size_t index_bit
Get mask at location.
- Parameters:
index_entry – Index of the entry in [0, size)·
index_bit – Index of the mask (bit) in the entry in [0, num_masks_per_entry).
- Returns:
Value of mask (bit) at location.
Public Static Functions
- static inline MaskArrayBaseType get_inner_index(
- const std::size_t index_entry,
- const std::size_t index_bit
- Parameters:
index_entry – Index of the entry in [0, size)·
index_bit – Index of the mask (bit) in the entry in [0, num_masks_per_entry).
- Returns:
The entry within a single MaskArrayBaseType that corresponds to the provided entry and bit.
- static inline std::size_t get_outer_index(
- const std::size_t index_entry,
- const std::size_t index_bit
- Parameters:
index_entry – Index of the entry in [0, size)·
index_bit – Index of the mask (bit) in the entry in [0, num_masks_per_entry).
- Returns:
The index to a MaskArrayBaseType.
- static inline bool get_inner(
- MaskArrayBaseType const *const base,
- const std::size_t index_bit
Get mask at location in base type.
- Parameters:
base – Pointer to a base type instance·
index_bit – Index of the mask (bit) in the entry in [0, num_masks_per_entry).
- Returns:
Value of mask (bit) at location.
-
struct Write : public NESO::Particles::MaskArrayDeviceBase<MaskArrayBaseType*>#
- #include <mask_array.hpp>
Write access descriptor for MaskArray.
Public Functions
- inline void set(
- const std::size_t index_entry,
- const std::size_t index_bit,
- const bool value
Set mask at location.
- Parameters:
index_entry – Index of the entry in [0, size)·
index_bit – Index of the mask (bit) in the entry in [0, num_masks_per_entry).
value – Value of mask (bit) to set.
Public Static Functions
- static inline void set_inner(
- MaskArrayBaseType *base,
- const std::size_t index_bit,
- const bool value
Set mask at location in base type.
- Parameters:
base – Pointer to a base type instance·
index_bit – Index of the mask (bit) in the entry in [0, num_masks_per_entry).
value – Value of mask (bit) to set.
-
class MaskArray#
- #include <mask_array.hpp>
Type that stores N bits per entry (particle).
Subclassed by NESO::Particles::PairMask, NESO::Particles::ParticleMask
Public Functions
- SYCLTargetSharedPtr sycl_target,
- const std::size_t num_masks_per_entry
Create a mask array on a given compute device with a set number of bits per entry.
- Parameters:
sycl_target – Compute device.
num_masks_per_entry – Number of bits stored per entry.
- void reset(
- const bool value,
- std::optional<std::size_t> new_size = std::nullopt
Zero the masks optionally provide a new array size.
- Parameters:
value – Provide the mask value to reset all entries to.
new_size – Optional new array size.
-
MaskArrayDevice get_device()#
- Returns:
Access to the masks via the device type.
-
std::size_t get_num_masks_true(const std::size_t mask_index = 0)#
- Parameters:
mask_index – Index of mask in entry to count, default 0.
- Returns:
Number of masks set to true.
Public Static Functions
-
static MaskArrayBaseType get_reset_mask(const bool value)#
- Parameters:
value – Mask value to set all masks to.
- Returns:
A base element that is either all 1 or 0.
NDLocalArray#
Creates an N-D array local to each MPI rank that can be accessed from the Particle Loop.
-
template<typename T, std::size_t N>
struct Read# - #include <nd_local_array.hpp>
ParticleLoop access type for NDLocalArray Read access.
-
template<typename T, std::size_t N>
struct Write# - #include <nd_local_array.hpp>
ParticleLoop access type for NDLocalArray Write access.
Public Functions
-
template<typename T, std::size_t N>
struct Add# - #include <nd_local_array.hpp>
ParticleLoop access type for NDLocalArray Add access.
Public Functions
-
template<typename T, std::size_t N>
struct Max# - #include <nd_local_array.hpp>
ParticleLoop access type for NDLocalArray Max access.
Public Functions
-
template<typename T, std::size_t N>
struct Min# - #include <nd_local_array.hpp>
ParticleLoop access type for NDLocalArray Min access.
Public Functions
-
template<typename T, std::size_t N>
class NDLocalArray# - #include <nd_local_array.hpp>
Generic N-Dimensional array type which is accessible on the host and in a ParticleLoop kernel.
Public Functions
- SYCLTargetSharedPtr sycl_target,
- NDIndex<N> index
Create a NDLocalArray on a compute device with a given shape.
- Parameters:
sycl_target – Compute device to create local array on.
index – Specification of the extent of each dimension.
-
)#
Create a NDLocalArray on a compute device from a NDHostArray.
- Parameters:
sycl_target – Compute device to create local array on.
nd_host_array – NDHostArray to copy to device.
- SYCLTargetSharedPtr sycl_target,
- SHAPE... shape
Create a NDLocalArray on a compute device with a given shape.
- Parameters:
sycl_target – Compute device to create local array on.
shape – Parameter pack of size N which defines the extent of the array in each of the N dimensions.
-
inline void fill(const T value)#
Fill the array with a value.
- Parameters:
value – Value to fill the array with.
- std::shared_ptr<RNGGenerationFunction<T>> rng_generation_function
Fill the array from samples generated by a RNG generation function that conforms to the RNGGenerationFunction interface.
- Parameters:
rng_generation_function – RNG sampler to use.
-
inline sycl::event set_async(const std::vector<T> &data)#
Asynchronously set the values in the local array to those in a std::vector.
- Parameters:
data – Input vector to copy values from.
- Returns:
Event to wait on before using new values in NDLocalArray.
-
inline void set(const std::vector<T> &data)#
Set the values in the local array to those in a std::vector. Blocks until the copy is complete.
- Parameters:
data – Input vector to copy values from.
Copy the values from the NDHostArray into an NDLocalArray.
- Parameters:
nd_host_array – Source array.
-
inline sycl::event get_async(std::vector<T> &data)#
Asynchronously get the values in the local array into a std::vector.
- Parameters:
data – [inout] Input vector to copy values from NDLocalArray into.
- Returns:
Event to wait on before using new values in the std::vector.
-
inline void get(std::vector<T> &data)#
Get the values in the local array into a std::vector. Blocks until copy is complete.
- Parameters:
data – [inout] Input vector to copy values from NDLocalArray into.
-
inline std::vector<T> get()#
Get the values in the local array into a std::vector.
- Returns:
std::vector of values in the NDLocalArray.
Copy the values from the NDLocalArray into an NDHostArray.
- Parameters:
nd_host_array – [inout] Destination array, will be allocated if nullptr.
- std::shared_ptr<NDLocalArray<U, N>> second_array,
- BINOP binop
Update each held entry a as
a <- binop(a, b)
where the entries b are supplied by another NDLocalArray and binop is a provided binary operator.
- Parameters:
second_array – Second array, i.e. b, for the binary combination.
binop – Binary operation to use to combine elements, must be a device copyable object with a cell method that takes two arguments of type T and U.
Public Members
-
SYCLTargetSharedPtr sycl_target#
Compute device for the array.
-
NDIndex<N> index#
Indexing function for the NDLocalArray.
ParticleMask#
A per particle mask.
-
struct Read#
- #include <particle_mask.hpp>
Read kernel type for ParticleMask.
Public Functions
-
inline bool get(const Access::LoopIndex::Read &index) const#
- Parameters:
index – ParticleLoopIndex of particle to get mask for.
- Returns:
The current mask value for the particle.
-
inline bool get(const Access::LoopIndex::Read &index) const#
-
struct Write#
- #include <particle_mask.hpp>
Write kernel type for ParticleMask.
Public Functions
-
inline bool get(const Access::LoopIndex::Read &index) const#
- Parameters:
index – ParticleLoopIndex of particle to get mask for.
- Returns:
The current mask value for the particle.
-
inline void set_on(const Access::LoopIndex::Read &index)#
Set the mask to on.
- Parameters:
index – ParticleLoopIndex of particle to set mask for.
-
inline void set_off(const Access::LoopIndex::Read &index)#
Set the mask to off.
- Parameters:
index – ParticleLoopIndex of particle to set mask for.
- inline void set(
- const Access::LoopIndex::Read &index,
- const bool value
Set the mask.
- Parameters:
index – ParticleLoopIndex of particle to set mask for.
value – Mask value to set..
-
inline bool get(const Access::LoopIndex::Read &index) const#
-
class ParticleMask : public NESO::Particles::MaskArray#
- #include <particle_mask.hpp>
Container to store a single mask per particle which is accessible from ParticleLoop.
Public Functions
Create a new ParticleMask on a compute device.
- Parameters:
sycl_target – Compute device to create mask on.
Reset the container to be able to store masks for all particles within the passed ParticleGroup.
- Parameters:
particle_group – ParticleGroup to reset size for.
- ParticleGroupSharedPtr particle_group,
- Sym<INT> sym,
- const int component
Set the masks from a ParticleDat Sym and component. Calls reset using the ParticleGroup. Calls reset before setting.
- Parameters:
particle_group – ParticleGroup to set from.
sym – Sym for ParticleDat to set from.
component – ParticleDat component to set from.
Set all the masks to a specified value. Calls reset before setting.
- Parameters:
particle_group – ParticleGroup to set from.
value – Value to set.
- ParticleGroupSharedPtr particle_group,
- Sym<INT> sym,
- const int component
Store the masks in a ParticleDat Sym and component.
- Parameters:
particle_group – ParticleGroup to set into.
sym – Sym for ParticleDat to set int.
component – ParticleDat component to set into.
ParticleSetDevice#
A storage container for multiple particles data on device accessible in a Particle Loop.
-
struct Read#
- #include <particle_set_device.hpp>
Access:ParticleSetDevice::Read<T>, Access::ParticleSetDevice::Write<T> and Access:ParticleSetDevice::Add<T> are the kernel argument types for accessing ParticleSetDevice data in a kernel. ParticleLoop access type for ParticleSetDevice Read access.
Public Functions
- inline const REAL &at_real(
- const int particle,
- const int property,
- const int component
Access a REAL particle property.
- Parameters:
particle – The index of the particle to access, i.e. the row in the device particle set.
property – The particle property to access using the ordering REAL ParticleProp instances were passed to the ParticleSetDeviceSpec
component – The component of the property to access.
- Returns:
Constant reference to particle value.
- inline const INT &at_int(
- const int particle,
- const int property,
- const int component
Access a INT particle property.
- Parameters:
particle – The index of the particle to access, i.e. the row in the device particle set.
property – The particle property to access using the ordering INT ParticleProp instances were passed to the ParticleSetDeviceSpec
component – The component of the property to access.
- Returns:
Constant reference to particle value.
-
struct Add#
- #include <particle_set_device.hpp>
ParticleLoop access type for ParticleSetDevice Add access.
Public Functions
- inline REAL fetch_add_real(
- const int particle,
- const int property,
- const int component,
- const REAL value
Atomically increment a REAL particle property.
- Parameters:
particle – The index of the particle to access, i.e. the row in the device particle set.
property – The particle property to access using the ordering REAL ParticleProp instances were passed to the ParticleSetDeviceSpec
component – The component of the property to access.
value – Value to increment property value by.
- Returns:
Value prior to increment.
- inline INT fetch_add_int(
- const int particle,
- const int property,
- const int component,
- const INT value
Atomically increment a INT particle property.
- Parameters:
particle – The index of the particle to access, i.e. the row in the device particle set.
property – The particle property to access using the ordering INT ParticleProp instances were passed to the ParticleSetDeviceSpec
component – The component of the property to access.
value – Value to increment property value by.
- Returns:
Value prior to increment.
-
struct Write#
- #include <particle_set_device.hpp>
ParticleLoop access type for ParticleSetDevice Write access.
Public Functions
-
Write() = default#
Pointer to underlying data for the array.
- inline REAL &at_real(
- const int particle,
- const int property,
- const int component
Write access to REAL particle property.
- Parameters:
particle – The index of the particle to access, i.e. the row in the device particle set.
property – The particle property to access using the ordering REAL ParticleProp instances were passed to the ParticleSetDeviceSpec
component – The component of the property to access.
- Returns:
Modifiable reference to property value.
- inline INT &at_int(
- const int particle,
- const int property,
- const int component
Write access to INT particle property.
- Parameters:
particle – The index of the particle to access, i.e. the row in the device particle set.
property – The particle property to access using the ordering INT ParticleProp instances were passed to the ParticleSetDeviceSpec
component – The component of the property to access.
- Returns:
Modifiable reference to property value.
-
Write() = default#
-
struct ParticleSetDeviceSpec#
- #include <particle_set_device.hpp>
Type to describe the particle properties of particles.
Public Functions
-
ParticleSetDeviceSpec(ParticleSpec &particle_spec)#
Create a specification for particles based on a ParticleSpec.
- Parameters:
particle_spec – Specification for particle particle properties.
- void set_default_value(
- Sym<INT> sym,
- const int component,
- const INT value
Set the default value for a particle property.
- Parameters:
sym – Sym of particle property.
component – Component of particle property.
value – Default value to set.
- void set_default_value(
- Sym<REAL> sym,
- const int component,
- const REAL value
Set the default value for a particle property.
- Parameters:
sym – Sym of particle property.
component – Component of particle property.
value – Default value to set.
Public Members
-
int num_components_real#
The total number of components of type REAL stored.
-
int num_components_int#
The total number of components of type INT stored.
-
int num_properties_real#
The number of REAL particle properties.
-
int num_properties_int#
The number of INT particle properties.
-
std::vector<int> components_real#
The number of components per output property for REAL.
-
std::vector<int> components_int#
The number of components per output property for INT.
-
std::map<std::pair<Sym<REAL>, int>, REAL> default_values_real#
Default values applied on call to reset for REAL.
-
std::map<std::pair<Sym<INT>, int>, INT> default_values_int#
Default values applied on call to reset for INT.
-
ParticleSpec particle_spec#
The ParticleSpec the instance was created from.
-
ParticleSetDeviceSpec(ParticleSpec &particle_spec)#
-
class ParticleSetDevice#
- #include <particle_set_device.hpp>
Type to store N particles which can be passed to a ParticleLoop. Fundamentally this class allocates two matrices, one for REAL valued properties and one for INT value properties. These matrices are allocated column major. Each output particle populates a row in these two matrices. The column ordering is based on the ordering of properties and there components in the input particle specification.
Subclassed by NESO::Particles::ProductMatrix
Public Functions
-
ParticleSetDevice &operator=(const ParticleSetDevice&) = default#
Note that the copy operator creates shallow copies of the array.
- SYCLTargetSharedPtr sycl_target,
- std::shared_ptr<ParticleSetDeviceSpec> spec
Create a new device particle set on the SYCLTarget. The reset method should be called with the desired number of output particles before a loop is executed which requires space for those output particles.
- Parameters:
sycl_target – Device on which particle loops will be executed using the device particle set.
spec – A specification for the output particle properties.
- SYCLTargetSharedPtr sycl_target,
- const int num_particles,
- ParticleSpec &particle_spec
Create a new device particle set from a ParticleSec.
- Parameters:
sycl_target – Device to store particles on.
num_particles – Number of particles.
particle_spec – ParticleSpec for particles.
-
virtual void reset(const int num_particles)#
Allocate space for a number of particle properties and fill the matrix with the default values.
- Parameters:
num_particles – Number of output particles to set in matrix.
-
ParticleSetSharedPtr get()#
- Returns:
a ParticleSet of the contained data.
Set the data in the ParticleSetDevice from a ParticleSet.
- Parameters:
particle_set – ParticleSet to set data from. Properties defined in the ParticleSet which do not exist in the ParticleSetDevice are ignored.
Public Members
-
SYCLTargetSharedPtr sycl_target#
The SYCLTarget particles are created on.
-
int num_particles#
The number of particles stored.
-
std::shared_ptr<ParticleSetDeviceSpec> spec#
The specification of the particles.
-
ParticleSetDevice &operator=(const ParticleSetDevice&) = default#
ProductMatrix#
A construct to store products of particles. i.e. New particles created in a Particle Loop.
-
typedef ParticleSetDeviceSpec ProductMatrixSpec#
Type to describe the particle properties of products.
- inline std::shared_ptr<ProductMatrixSpec> product_matrix_spec(
- ParticleSpec particle_spec
Helper function to create ProductMatrixSpec instances.
- Parameters:
particle_spec – Specification for product particle properties.
- SYCLTargetSharedPtr sycl_target,
- std::shared_ptr<ProductMatrixSpec> spec
Helper function to create ProductMatrix shared pointer.
- Parameters:
sycl_target – Device on which particle loops will be executed using the product matrix.
spec – A specification for the output particle properties.
-
struct Read#
- #include <product_matrix.hpp>
Access:ProductMatrix::Read<T>, Access::ProductMatrix::Write<T> and Access:ProductMatrix::Add<T> are the kernel argument types for accessing ProductMatrix data in a kernel. ParticleLoop access type for ProductMatrix Read access.
Public Functions
- inline const REAL &at_real(
- const int product,
- const int property,
- const int component
Access a REAL product property.
- Parameters:
product – The index of the product to access, i.e. the row in the product matrix.
property – The particle property to access using the ordering REAL ParticleProp instances were passed to the ProductMatrixSpec
component – The component of the property to access.
- Returns:
Constant reference to product value.
- inline const INT &at_int(
- const int product,
- const int property,
- const int component
Access a INT product property.
- Parameters:
product – The index of the product to access, i.e. the row in the product matrix.
property – The particle property to access using the ordering INT ParticleProp instances were passed to the ProductMatrixSpec
component – The component of the property to access.
- Returns:
Constant reference to product value.
-
struct Add#
- #include <product_matrix.hpp>
ParticleLoop access type for ProductMatrix Add access.
Public Functions
- inline REAL fetch_add_real(
- const int product,
- const int property,
- const int component,
- const REAL value
Atomically increment a REAL product property.
- Parameters:
product – The index of the product to access, i.e. the row in the product matrix.
property – The particle property to access using the ordering REAL ParticleProp instances were passed to the ProductMatrixSpec
component – The component of the property to access.
value – Value to increment property value by.
- Returns:
Value prior to increment.
- inline INT fetch_add_int(
- const int product,
- const int property,
- const int component,
- const INT value
Atomically increment a INT product property.
- Parameters:
product – The index of the product to access, i.e. the row in the product matrix.
property – The particle property to access using the ordering INT ParticleProp instances were passed to the ProductMatrixSpec
component – The component of the property to access.
value – Value to increment property value by.
- Returns:
Value prior to increment.
-
struct Write#
- #include <product_matrix.hpp>
ParticleLoop access type for ProductMatrix Write access.
Public Functions
-
Write() = default#
Pointer to underlying data for the array.
- inline REAL &at_real(
- const int product,
- const int property,
- const int component
Write access to REAL particle property.
- Parameters:
product – The index of the product to access, i.e. the row in the product matrix.
property – The particle property to access using the ordering REAL ParticleProp instances were passed to the ProductMatrixSpec
component – The component of the property to access.
- Returns:
Modifiable reference to property value.
- inline INT &at_int(
- const int product,
- const int property,
- const int component
Write access to INT particle property.
- Parameters:
product – The index of the product to access, i.e. the row in the product matrix.
property – The particle property to access using the ordering INT ParticleProp instances were passed to the ProductMatrixSpec
component – The component of the property to access.
- Returns:
Modifiable reference to property value.
-
Write() = default#
-
class ProductMatrix : public NESO::Particles::ParticleSetDevice#
- #include <product_matrix.hpp>
Type to store N products which can be passed to a ParticleLoop. Fundamentally this class allocates two matrices, one for REAL valued properties and one for INT value properties. These matrices are allocated column major. Each output particle populates a row in these two matrices. The column ordering is based on the ordering of properties and there components in the input particle specification.
Subclassed by NESO::Particles::DescendantProducts
Public Functions
-
ProductMatrix &operator=(const ProductMatrix&) = default#
Note that the copy operator creates shallow copies of the array.
- SYCLTargetSharedPtr sycl_target,
- std::shared_ptr<ProductMatrixSpec> spec
Create a new product matrix on the SYCLTarget. The reset method should be called with the desired number of output particles before a loop is executed which requires space for those output particles.
- Parameters:
sycl_target – Device on which particle loops will be executed using the product matrix.
spec – A specification for the output particle properties.
-
inline virtual void reset(const int num_products) override#
Allocate space for a number of particle properties and fill the matrix with the default values.
- Parameters:
num_products – Number of output particles to set in matrix.
Public Members
-
int num_products#
The number of products stored.
-
ProductMatrix &operator=(const ProductMatrix&) = default#
KernelRNG#
Construct for making available RNG samples inside Particle Loop kernels.
Helper function to create a NullKernelRNG.
- Returns:
NullKernelRNG that can be passed as a ParticleLoop argument.
-
template<typename T>
struct Read# - #include <kernel_rng.hpp>
This is the kernel type for KernelRNG which is used for all implementations which present RNG values to the kernel via an allocated device buffer.
Public Functions
- inline auto at(
- const Access::LoopIndex::Read &particle_index,
- const int component,
- bool *valid_sample
Access the RNG data for this particle.
- Parameters:
particle_index – [in] Particle index to access.
component – [in] RNG component to access.
valid_sample – [inout] On return this bool is set to true if the returned sample is good (i.e. RNG is in a valid state).
- Returns:
Constant reference to RNG data.
- inline auto at(
- const Access::LoopIndex::Read &particle_index,
- const int component
Access the RNG data for this particle.
- Parameters:
particle_index – Particle index to access.
component – RNG component to access.
- Returns:
Constant reference to RNG data.
- inline auto at(
- const Access::PairLoopIndex::Read &pair_index,
- const int component,
- bool *valid_sample
Access the RNG data for this particle pair.
- Parameters:
pair_index – [in] Particle pair index to access.
component – [in] RNG component to access.
valid_sample – [inout] On return this bool is set to true if the returned sample is good (i.e. RNG is in a valid state).
- Returns:
Constant reference to RNG data.
-
template<typename T>
struct NullKernelRNGDevice# - #include <null_kernel_rng.hpp>
Null Kernel RNG that performs no operations. This is the kernel type for Access::read(std::shared_ptr<NullKernelRNG>).
Public Functions
-
inline const T at(const int &row, const int &col) const#
- Parameters:
row – Unused, exists to match API.
col – Unused, exists to match API.
- Returns:
0.
-
inline const T at(const int &row, const int &col) const#
-
template<typename T>
class NullKernelRNG : public NESO::Particles::KernelRNG<NullKernelRNGDevice<T>># - #include <null_kernel_rng.hpp>
KernelRNG implementation for RNG implementations which provide no RNG samples. This type exists as an RNG implementation for downstream interfaces which must provide a KernelRNG but the specific specialisation of the downstream interface has no RNG requirements.
Public Functions
- inline virtual Access::KernelRNG::Read<NullKernelRNGDevice<T>> impl_get_const(
- ParticleLoopImplementation::ParticleLoopGlobalInfo *global_info
No-op impl_get_const.
- Parameters:
global_info – Unused, exists to match API.
- Returns:
Null kernel type.
- inline virtual void impl_pre_loop_read(
- ParticleLoopImplementation::ParticleLoopGlobalInfo *global_info
No-op impl_pre_loop_read.
- Parameters:
global_info – Unused, exists to match API.
- inline virtual void impl_post_loop_read(
- ParticleLoopImplementation::ParticleLoopGlobalInfo *global_info
No-op impl_post_loop_read.
- Parameters:
global_info – Unused, exists to match API.
-
inline virtual bool valid_internal_state() override#
- Returns:
True.
TupleRNG#
Construct for passing multiple KernelRNG instances as one Particle Loop argument.
-
template<std::size_t INDEX, typename ...KERNELRNGS>
auto &get( - Read<KERNELRNGS...> &tuple_rng
Helper function that calls Access::TupleRNG::get in a way that avoids using “.template”.
- Parameters:
tuple_rng – Access::TupleRNG::Read instance to access.
- Returns:
The reference returned by Access::TupleRNG::Read::get.
-
template<typename ...RNGPTRS>
auto tuple_rng(RNGPTRS... rng_ptrs)# Helper function to create a tuple of KernelRNG/TupleRNG instances.
-
template<typename ...KERNELRNGS>
struct Read# - #include <tuple_rng.hpp>
This is the kernel type for TupleRNG.
Public Functions
-
template<std::size_t INDEX>
inline auto &get()# Access the N-th KernelRNG instance.
-
template<std::size_t INDEX>
-
template<typename ...RNGPTRS>
class TupleRNG# - #include <tuple_rng.hpp>
Container type which holds multiple RNG types which inherit from KernelRNG or are also TupleRNG instances.
Public Functions
-
inline TupleRNG(RNGPTRS... args)#
Constructor to create a tuple of KernelRNG/TupleRNG instances.
- Parameters:
args – Shared pointers to instances of KernelRNG descendent types and TupleRNG instances.
- inline KernelType impl_get_const(
- ParticleLoopImplementation::ParticleLoopGlobalInfo *global_info
Called by ParticleLoop to create the loop arguments.
- Parameters:
global_info – Global information for the loop which is to be executed.
- inline void impl_pre_loop_read(
- ParticleLoopImplementation::ParticleLoopGlobalInfo *global_info
Executed by the loop pre execution.
- Parameters:
global_info – Global information for the loop which is to be executed.
- inline void impl_post_loop_read(
- ParticleLoopImplementation::ParticleLoopGlobalInfo *global_info
Executed by the loop post execution.
- Parameters:
global_info – Global information for the loop which completed.
-
inline TupleRNG(RNGPTRS... args)#
SymVector#
Allows passing a runtime number of Sym arguments to a ParticleLoop.
- ParticleGroupSharedPtr particle_group,
- std::vector<Sym<T>> syms
Helper function to create a SymVector.
- Parameters:
particle_group – ParticleGroup to use.
syms – Vector of Syms to use from particle_group.
- ParticleGroupSharedPtr particle_group,
- std::initializer_list<Sym<T>> syms
Helper function to create a SymVector.
- Parameters:
particle_group – ParticleGroup to use.
syms – Syms to use from particle_group.
- ParticleSubGroupSharedPtr particle_sub_group,
- std::vector<Sym<T>> syms
Helper function to create a SymVector.
- Parameters:
particle_sub_group – ParticleSubGroup to use.
syms – Vector of Syms to use from particle_group.
- ParticleSubGroupSharedPtr particle_sub_group,
- std::initializer_list<Sym<T>> syms
Helper function to create a SymVector.
- Parameters:
particle_sub_group – ParticleSubGroup to use.
syms – Syms to use from particle_group.
-
template<typename T>
struct Read# - #include <sym_vector.hpp>
Kernel type for SymVector accessed with read mode.
-
template<typename T>
struct Write# - #include <sym_vector.hpp>
Kernel type for SymVector accessed with write mode.
-
template<typename T>
class SymVector# - #include <sym_vector.hpp>
Enables ParticleDats to be accessed in ParticleLoops with a number of ParticleDats determined at runtime.
Public Functions
-
)#
Create a SymVector using a ParticleGroup and a std::vector of Syms.
- Parameters:
particle_group – ParticleGroup to use.
syms – Vector of Syms to use from particle_group.
-
)#
Create a SymVector using a ParticleGroup and an initialiser list of syms, e.g.
SymVector(particle_group, {Sym<INT>(“a”), Sym<INT>(“b”)});
- Parameters:
particle_group – ParticleGroup to use.
syms – Syms to use from particle_group.
CellInfoNPart#
Particle Loop construct for accessing the index of particles within a Particle Loop. For example the particle cell and layer along with loop indices.
-
struct CellInfoNPart#
- #include <cell_info_npart.hpp>
The type to pass to a ParticleLoop with read access to access the number of particles in the cell for the iteration set passed to the ParticleLoop.
ParticleLoopIndex#
Particle Loop construct for accessing the index of particles within a Particle Loop. For example the particle cell and layer along with loop indices.
-
struct ParticleLoopIndex#
- #include <particle_loop_index.hpp>
The type to pass to a ParticleLoop to read the ParticleLoop loop index in a kernel.
-
struct Read#
- #include <particle_loop_index.hpp>
ParticleLoop index containing the cell and layer.
Public Functions
-
inline INT get_local_linear_index() const#
- Returns:
The local linear index of the particle on this MPI rank in the ParticleGroup
-
inline INT get_loop_linear_index() const#
- Returns:
The linear index of the particle within the current ParticleLoop.
-
inline INT get_sub_linear_index() const#
- Returns:
The local linear index of the particle in the ParticleSubGroup. This call is identical to get_local_linear_index when the iteration set is a ParticleGroup.
Public Members
-
INT cell#
The cell containing the particle.
-
INT layer#
The layer of the particle.
-
INT loop_layer#
The looping layer of the particle.
-
int loop_type_int#
The type of the ParticleLoop (intended for internal use).
-
INT const *npart_cell_es#
pointer to the exclusive sum of particle counts in each cell (intended for internal use).
-
INT const *npart_cell_es_lb#
pointer to the exclusive sum of particle counts in each cell for loop bounds (intended for internal use).
-
int starting_cell#
Loop iteration index (intended for internal use - see get_loop_linear_index) Starting cell for ParticleLoop called cell wise (intended for internal use).
-
inline INT get_local_linear_index() const#