Skip to content

Commit

Permalink
Merge pull request #223 from ExCALIBUR-NEPTUNE/feature/composite-inte…
Browse files Browse the repository at this point in the history
…raction-base

Feature/composite interaction base
  • Loading branch information
JamesEdgeley authored Oct 9, 2024
2 parents d48ed0b + c07e051 commit c13c49e
Show file tree
Hide file tree
Showing 68 changed files with 25,264 additions and 792 deletions.
21 changes: 21 additions & 0 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -110,18 +110,24 @@ set(LIB_SRC_FILES
${SRC_DIR}/linear_interpolator_1D.cpp
${SRC_DIR}/mesh.cpp
${SRC_DIR}/nektar_interface/basis_reference.cpp
${SRC_DIR}/nektar_interface/composite_interaction/composite_collections.cpp
${SRC_DIR}/nektar_interface/composite_interaction/composite_intersection.cpp
${SRC_DIR}/nektar_interface/composite_interaction/composite_utility.cpp
${SRC_DIR}/nektar_interface/composite_interaction/composite_transport.cpp
${SRC_DIR}/nektar_interface/expansion_looping/jacobi_coeff_mod_basis.cpp
${SRC_DIR}/nektar_interface/function_evaluation.cpp
${SRC_DIR}/nektar_interface/geometry_transport/geometry_transport_2d.cpp
${SRC_DIR}/nektar_interface/geometry_transport/geometry_transport_3d.cpp
${SRC_DIR}/nektar_interface/geometry_transport/halo_extension.cpp
${SRC_DIR}/nektar_interface/particle_boundary_conditions.cpp
${SRC_DIR}/nektar_interface/particle_cell_mapping/map_particles_2d.cpp
${SRC_DIR}/nektar_interface/particle_cell_mapping/map_particles_2d_regular.cpp
${SRC_DIR}/nektar_interface/particle_cell_mapping/map_particles_3d.cpp
${SRC_DIR}/nektar_interface/particle_cell_mapping/map_particles_3d_regular.cpp
${SRC_DIR}/nektar_interface/particle_cell_mapping/map_particles_common.cpp
${SRC_DIR}/nektar_interface/particle_cell_mapping/map_particles_host.cpp
${SRC_DIR}/nektar_interface/particle_cell_mapping/nektar_graph_local_mapper.cpp
${SRC_DIR}/nektar_interface/utilities.cpp
${SRC_DIR}/nektar_interface/solver_base/partsys_base.cpp
${SRC_DIR}/plasma.cpp
${SRC_DIR}/run_info.cpp
Expand All @@ -148,6 +154,14 @@ set(HEADER_FILES
${INC_DIR}/nektar_interface/bounding_box_intersection.hpp
${INC_DIR}/nektar_interface/cell_id_translation.hpp
${INC_DIR}/nektar_interface/coordinate_mapping.hpp
${INC_DIR}/nektar_interface/composite_interaction/composite_collection.hpp
${INC_DIR}/nektar_interface/composite_interaction/composite_collections.hpp
${INC_DIR}/nektar_interface/composite_interaction/composite_interaction.hpp
${INC_DIR}/nektar_interface/composite_interaction/composite_intersection.hpp
${INC_DIR}/nektar_interface/composite_interaction/composite_utility.hpp
${INC_DIR}/nektar_interface/composite_interaction/composite_transport.hpp
${INC_DIR}/nektar_interface/composite_interaction/line_plane_intersection.hpp
${INC_DIR}/nektar_interface/composite_interaction/line_line_intersection.hpp
${INC_DIR}/nektar_interface/solver_base/eqnsys_base.hpp
${INC_DIR}/nektar_interface/expansion_looping/basis_evaluate_base.hpp
${INC_DIR}/nektar_interface/expansion_looping/expansion_looping.hpp
Expand All @@ -165,6 +179,7 @@ set(HEADER_FILES
${INC_DIR}/nektar_interface/function_basis_projection.hpp
${INC_DIR}/nektar_interface/function_evaluation.hpp
${INC_DIR}/nektar_interface/function_projection.hpp
${INC_DIR}/nektar_interface/geometry_transport/geometry_packing_utility.hpp
${INC_DIR}/nektar_interface/geometry_transport/geometry_container_3d.hpp
${INC_DIR}/nektar_interface/geometry_transport/geometry_local_remote_3d.hpp
${INC_DIR}/nektar_interface/geometry_transport/geometry_transport.hpp
Expand All @@ -174,6 +189,7 @@ set(HEADER_FILES
${INC_DIR}/nektar_interface/geometry_transport/halo_extension.hpp
${INC_DIR}/nektar_interface/geometry_transport/packed_geom_2d.hpp
${INC_DIR}/nektar_interface/geometry_transport/packed_geoms_2d.hpp
${INC_DIR}/nektar_interface/geometry_transport/remote_geom.hpp
${INC_DIR}/nektar_interface/geometry_transport/remote_geom_2d.hpp
${INC_DIR}/nektar_interface/geometry_transport/remote_geom_3d.hpp
${INC_DIR}/nektar_interface/geometry_transport/shape_mapping.hpp
Expand All @@ -187,6 +203,7 @@ set(HEADER_FILES
${INC_DIR}/nektar_interface/particle_cell_mapping/generated_linear/prism.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/generated_linear/pyramid.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/generated_linear/quadrilateral.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/generated_linear/quadrilateralembed3d.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/generated_linear/tetrahedron.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/map_particles_2d.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/map_particles_2d_regular.hpp
Expand All @@ -202,17 +219,21 @@ set(HEADER_FILES
${INC_DIR}/nektar_interface/particle_cell_mapping/newton_prism.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/newton_pyr.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/newton_quad.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/newton_quad_embed_3d.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/newton_triangle_embed_3d.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/newton_tet.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/particle_cell_mapping.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/particle_cell_mapping_2d.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/particle_cell_mapping_3d.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/particle_cell_mapping_common.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/particle_cell_mapping_newton.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/x_map_newton.hpp
${INC_DIR}/nektar_interface/particle_cell_mapping/x_map_newton_kernel.hpp
${INC_DIR}/nektar_interface/particle_interface.hpp
${INC_DIR}/nektar_interface/particle_mesh_interface.hpp
${INC_DIR}/nektar_interface/special_functions.hpp
${INC_DIR}/nektar_interface/solver_base/time_evolved_eqnsys_base.hpp
${INC_DIR}/nektar_interface/typedefs.hpp
${INC_DIR}/nektar_interface/utilities.hpp
${INC_DIR}/nektar_interface/utility_mesh_cartesian.hpp
${INC_DIR}/nektar_interface/utility_mesh_plotting.hpp
Expand Down
34 changes: 10 additions & 24 deletions include/nektar_interface/cell_id_translation.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -128,32 +128,18 @@ class CellIDTranslation {
*/
inline void execute() {
auto t0 = profile_timestamp();

auto pl_iter_range = this->cell_id_dat->get_particle_loop_iter_range();
auto pl_stride = this->cell_id_dat->get_particle_loop_cell_stride();
auto pl_npart_cell = this->cell_id_dat->get_particle_loop_npart_cell();

auto k_cell_id_dat = this->cell_id_dat->cell_dat.device_ptr();
const auto k_lookup_map = this->id_map.d_buffer.ptr;
const INT k_shift = this->shift;

this->sycl_target->queue
.submit([&](sycl::handler &cgh) {
cgh.parallel_for<>(
sycl::range<1>(pl_iter_range), [=](sycl::id<1> idx) {
NESO_PARTICLES_KERNEL_START
const INT cellx = NESO_PARTICLES_KERNEL_CELL;
const INT layerx = NESO_PARTICLES_KERNEL_LAYER;

const INT nektar_cell = k_cell_id_dat[cellx][0][layerx];
const INT shifted_nektar_cell = nektar_cell - k_shift;
const INT neso_cell = k_lookup_map[shifted_nektar_cell];
k_cell_id_dat[cellx][0][layerx] = neso_cell;

NESO_PARTICLES_KERNEL_END
});
})
.wait_and_throw();
particle_loop(
"CellIDTranslation::execute", this->cell_id_dat,
[=](auto k_cell_id_dat) {
const INT nektar_cell = k_cell_id_dat.at(0);
const INT shifted_nektar_cell = nektar_cell - k_shift;
const INT neso_cell = k_lookup_map[shifted_nektar_cell];
k_cell_id_dat.at(0) = neso_cell;
},
Access::write(this->cell_id_dat))
->execute();
sycl_target->profile_map.inc("CellIDTranslation", "execute", 1,
profile_elapsed(t0, profile_timestamp()));
};
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,39 @@
#ifndef __NEKTAR_INTERFACE_COMPOSITE_INTERACTION_COMPOSITE_COLLECTION_H_
#define __NEKTAR_INTERFACE_COMPOSITE_INTERACTION_COMPOSITE_COLLECTION_H_

#include "line_line_intersection.hpp"
#include "line_plane_intersection.hpp"

namespace NESO::CompositeInteraction {

/**
* Struct pointed to by a BlockedBinaryTree for each MeshHierarchy cell.
*/
struct CompositeCollection {
// Face members
int num_quads;
int num_tris;
LinePlaneIntersection *lpi_quads;
LinePlaneIntersection *lpi_tris;
size_t stride_quads;
size_t stride_tris;
unsigned char *buf_quads;
unsigned char *buf_tris;
int *composite_ids_quads;
int *composite_ids_tris;
int *geom_ids_quads;
int *geom_ids_tris;
// Segment members
int num_segments;
LineLineIntersection *lli_segments;
int *composite_ids_segments;
int *geom_ids_segments;
// group ids
int *group_ids_quads;
int *group_ids_tris;
int *group_ids_segments;
};

} // namespace NESO::CompositeInteraction

#endif
Original file line number Diff line number Diff line change
@@ -0,0 +1,161 @@
#ifndef __NEKTAR_INTERFACE_COMPOSITE_INTERACTION_COMPOSITE_COLLECTIONS_H_
#define __NEKTAR_INTERFACE_COMPOSITE_INTERACTION_COMPOSITE_COLLECTIONS_H_

#include <SpatialDomains/MeshGraph.h>
using namespace Nektar;

#include "composite_collection.hpp"
#include "composite_transport.hpp"
#include "composite_utility.hpp"
#include <nektar_interface/geometry_transport/packed_geom_2d.hpp>
#include <nektar_interface/particle_cell_mapping/newton_geom_interfaces.hpp>
#include <nektar_interface/particle_mesh_interface.hpp>
#include <nektar_interface/typedefs.hpp>

#include <algorithm>
#include <map>
#include <memory>
#include <set>
#include <utility>
#include <vector>

namespace NESO::CompositeInteraction {

/**
* Type to point to the normal vector for a boundary element.
*/
struct NormalData {
REAL *d_normal;
};

/**
* Device type to help finding the normal vector for a particle-boundary
* interaction.
*/
struct NormalMapper {
// Root of the tree containing normal data.
BlockedBinaryNode<INT, NormalData, 8> *root;

/**
* Get a pointer to the normal data for an edge element.
*
* @param[in] global_id Global index of boundary element to retrive normal
* for.
* @param[in, out] normal Pointer to populate with address of normal vector.
* @returns True if the boundary element is found otherwise false.
*/
inline bool get(const INT global_id, REAL **normal) const {
NormalData *node;
bool *exists;
const bool e = root->get_location(global_id, &exists, &node);
*normal = node->d_normal;
return e && (*exists);
}
};

/**
* Distributed data structure to hold the geometry information that intersects
* each Mesh Hierarchy cell.
*/
class CompositeCollections {
protected:
ParticleMeshInterfaceSharedPtr particle_mesh_interface;

// Stack for device buffers (Newton X mapping data)
std::stack<std::shared_ptr<BufferDevice<unsigned char>>> stack_geometry_data;
// Stack for device buffers (LinePlaneIntersections)
std::stack<std::shared_ptr<BufferDevice<LinePlaneIntersection>>>
stack_lpi_data;
// Stack for device buffers (LineLineIntersections)
std::stack<std::shared_ptr<BufferDevice<LineLineIntersection>>>
stack_lli_data;
// Stack for the device buffers (CompositeCollection)
std::stack<std::shared_ptr<BufferDevice<CompositeCollection>>>
stack_collection_data;
// Stack for the device buffers (composite ids)
std::stack<std::shared_ptr<BufferDevice<int>>> stack_composite_ids;
// stack for device buffers (REAL data)
std::stack<std::shared_ptr<BufferDevice<REAL>>> stack_real;

// Cells already collected.
std::set<INT> collected_cells;

/**
* This method takes the packed geoms and creates the node the in blocked
* binary tree such that the kernel can access the geometry info.
*/
void collect_cell(const INT cell);

/// The normal data for each geometry object collected.
std::shared_ptr<BlockedBinaryTree<INT, NormalData, 8>> map_normals;

/// Inverse map from composite to boundary group.
std::map<int, int> map_composite_to_group;

public:
/// Disable (implicit) copies.
CompositeCollections(const CompositeCollections &st) = delete;
/// Disable (implicit) copies.
CompositeCollections &operator=(CompositeCollections const &a) = delete;

/// SYCLTarget to use for computation.
SYCLTargetSharedPtr sycl_target;

/// Map from composites to geometry objects held.
std::map<int, std::map<int, std::shared_ptr<Geometry>>>
map_composites_to_geoms;

/// The composite transport instance.
std::shared_ptr<CompositeTransport> composite_transport;

/// The container that holds the map from MeshHierarchy cells to the geometry
/// objects for the composites in those cells (faces).
std::shared_ptr<BlockedBinaryTree<INT, CompositeCollection *, 4>>
map_cells_collections;

/// Map from boundary group id to composites in the group.
std::map<int, std::vector<int>> boundary_groups;

/// Maximum number of modes for X maps.
int max_num_modes;

/**
* Free the data structure. Must be called collectively on the communicator.
*/
void free();

/**
* Create new distributed container for geometry objects which are members
* of the passed composite indices.
*
* @param sycl_target Compute device on which intersection computation will
* take place.
* @param particle_mesh_interface ParticleMeshInterface for the domain.
* @param boundary_groups Map from boundary group id to composite ids which
* form the group.
*/
CompositeCollections(SYCLTargetSharedPtr sycl_target,
ParticleMeshInterfaceSharedPtr particle_mesh_interface,
std::map<int, std::vector<int>> boundary_groups);

/**
* Must be called collectively on the communicator. Collect on this MPI rank
* all the geometry objects which intersect certain MeshHierarchy cells.
*
* @param cells MeshHierarchy cells to collect all geometry objects for.
*/
void collect_geometry(std::set<INT> &cells);

/**
* Get a device callable mapper to map from global boundary indices to normal
* vectors. This method should be called after @ref collect_geometry as the
* map is populated with potentially relvant normal data on the fly.
*
* @returns Device copyable and callable mapper.
*/
NormalMapper get_device_normal_mapper();
};

} // namespace NESO::CompositeInteraction

#endif
Original file line number Diff line number Diff line change
@@ -0,0 +1,8 @@
#ifndef __COMPOSITE_INTERACTION_H_
#define __COMPOSITE_INTERACTION_H_

#include "composite_intersection.hpp"
#include "line_line_intersection.hpp"
#include "line_plane_intersection.hpp"

#endif
Loading

0 comments on commit c13c49e

Please sign in to comment.