tboudreaux/MeanField
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat(newton): first newton solver implementation

Browse Source
This commit is contained in:
Emily Boudreaux
2026-09-08 06:36:39 -04:00
parent 76818f2f82
commit b3c04d507a
98 changed files with 20397 additions and 11040 deletions
Download Patch File Download Diff File Expand all files Collapse all files

225
libmeanfield/impl/operators/kernels/gravity_displacement_force_kernels.cpp
View File

@@ -2,9 +2,12 @@ module;
#include <array>
#include <cmath>
#include <expected>
#include <optional>
#include <stdexcept>
#include <mfem.hpp>
#include <mpi.h>
module mean_field;
@@ -19,6 +22,70 @@ namespace {
enum class GravityDisplacementForceAction { residual, density, gravityGradient, displacement, complete };
using Rejection = mean_field::operators::kernels::GravityDisplacementForceRejection;
using Reason = mean_field::operators::kernels::GravityDisplacementForceRejectionReason;
using Result = mean_field::operators::kernels::GravityDisplacementForceResult;
[[nodiscard]] Rejection mapping_rejection(const mean_field::mapping::MappingStatus status) {
MFEM_VERIFY(
status != mean_field::mapping::MappingStatus::invalid_dimension,
"The gravity-displacement-force mapping reported an invariant dimension mismatch."
);
return {.reason = Reason::invalid_mapping, .mappingStatus = status};
}
[[nodiscard]] Rejection non_finite_rejection() noexcept {
return {.reason = Reason::non_finite_arithmetic};
}
[[nodiscard]] bool vector_is_finite(const mfem::Vector &vector) noexcept {
for (int index = 0; index < vector.Size(); ++index) {
if (!std::isfinite(vector(index))) {
return false;
}
}
return true;
}
[[nodiscard]] int encode_rejection(const std::optional<Rejection> &rejection) noexcept {
if (!rejection.has_value()) {
return 0;
}
if (rejection->reason == Reason::non_finite_arithmetic) {
return 256;
}
return static_cast<int>(rejection->mappingStatus) + 1;
}
[[nodiscard]] Rejection decode_rejection(const int encoded) noexcept {
if (encoded >= 256) {
return non_finite_rejection();
}
return mapping_rejection(static_cast<mean_field::mapping::MappingStatus>(encoded - 1));
}
[[nodiscard]] Result synchronize_rejection(
const std::optional<Rejection> &localRejection,
const MPI_Comm communicator
) {
const int localEncoded = encode_rejection(localRejection);
int globalEncoded = 0;
if (MPI_Allreduce(&localEncoded, &globalEncoded, 1, MPI_INT, MPI_MAX, communicator) != MPI_SUCCESS) {
throw std::runtime_error("Could not synchronize gravity-displacement-force candidate validity.");
}
if (globalEncoded != 0) {
return std::unexpected(decode_rejection(globalEncoded));
}
return {};
}
[[noreturn]] void throw_rejection(const Rejection &rejection) {
if (rejection.reason == Reason::non_finite_arithmetic) {
throw std::domain_error("The gravity-displacement force produced non-finite arithmetic.");
}
throw std::domain_error("The gravity-displacement force encountered an invalid mapped domain.");
}
void true_to_local(
const mfem::ParFiniteElementSpace &finiteElementSpace,
const mfem::Vector &trueVector,
@@ -187,11 +254,6 @@ namespace {
"The gravity-displacement-force displacement dimension does not "
"match the mesh dimension."
);
validate_finite_vector(
displacementTrue, "The gravity-displacement-force displacement contains a "
"non-finite value."
);
}
void validate_density(
@@ -200,7 +262,6 @@ namespace {
const char *message
) {
MFEM_VERIFY(density.Size() == f.densityFes->GetTrueVSize(), message);
validate_finite_vector(density, message);
}
void validate_gravity_gradient(
@@ -209,11 +270,9 @@ namespace {
const char *message
) {
MFEM_VERIFY(gravityGradient.Size() == f.gravityFluxFes->GetTrueVSize(), message);
validate_finite_vector(gravityGradient, message);
}
void apply_gravity_displacement_force_action(
[[nodiscard]] Result apply_gravity_displacement_force_action(
const mean_field::fem::FEM &f,
const mean_field::mapping::DomainMapper &domainMapper,
const GravityDisplacementForceAction requestedAction,
@@ -223,7 +282,8 @@ namespace {
const mfem::Vector *gravityGradientVariationTrue,
const mfem::Vector *displacementVariationTrue,
const mfem::Vector &displacementTrue,
mfem::Vector &actionTrue
mfem::Vector &actionTrue,
const bool reportCandidateRejection
) {
validate_common_inputs(f, domainMapper, displacementTrue);
@@ -308,6 +368,34 @@ namespace {
);
}
bool inputsAreFinite = vector_is_finite(displacementTrue);
if (needsBaseDensity) {
inputsAreFinite = inputsAreFinite && vector_is_finite(*baseDensityTrue);
}
if (needsDensityVariation) {
inputsAreFinite = inputsAreFinite && vector_is_finite(*densityVariationTrue);
}
if (needsBaseGravityGradient) {
inputsAreFinite = inputsAreFinite && vector_is_finite(*baseGravityGradientTrue);
}
if (needsGravityGradientVariation) {
inputsAreFinite = inputsAreFinite && vector_is_finite(*gravityGradientVariationTrue);
}
if (needsDisplacementVariation) {
inputsAreFinite = inputsAreFinite && vector_is_finite(*displacementVariationTrue);
}
if (!reportCandidateRejection) {
MFEM_VERIFY(inputsAreFinite, "The gravity-displacement-force action contains non-finite input data.");
} else {
const std::optional<Rejection> inputRejection =
inputsAreFinite ? std::optional<Rejection>{} : std::optional<Rejection>{non_finite_rejection()};
auto synchronized = synchronize_rejection(inputRejection, f.mesh->GetComm());
if (!synchronized.has_value()) {
return synchronized;
}
}
mfem::Vector baseDensityLocal;
mfem::Vector densityVariationLocal;
mfem::Vector baseGravityGradientLocal;
@@ -374,6 +462,8 @@ namespace {
const mfem::Ordering::Type displacementOrdering = f.displacementFes->GetOrdering();
std::optional<Rejection> candidateRejection;
for (int elementId = 0; elementId < f.mesh->GetNE(); ++elementId) {
mfem::ElementTransformation *transformation = f.mesh->GetElementTransformation(elementId);
@@ -516,13 +606,19 @@ namespace {
mappingData, *transformation, integrationPoint, workspace, mappingContext
);
MFEM_VERIFY(
mappingStatus == mean_field::mapping::MappingStatus::valid,
"Stateless mapping failed in the gravity-displacement-"
"force kernel. Element: "
<< elementId << ", attribute: " << transformation->Attribute
<< ", quadrature point: " << quadratureIndex << ", status: " << static_cast<int>(mappingStatus)
);
if (mappingStatus != mean_field::mapping::MappingStatus::valid) {
if (!reportCandidateRejection) {
MFEM_VERIFY(
false, "Stateless mapping failed in the gravity-displacement-"
"force kernel. Element: "
<< elementId << ", attribute: " << transformation->Attribute
<< ", quadrature point: " << quadratureIndex
<< ", status: " << static_cast<int>(mappingStatus)
);
}
candidateRejection = mapping_rejection(mappingStatus);
continue;
}
if (needsDisplacementVariation) {
const mean_field::mapping::MappingStatus variationStatus = domainMapper.EvaluateVolumeVariation(
@@ -530,13 +626,19 @@ namespace {
workspace, mappingVariation
);
MFEM_VERIFY(
variationStatus == mean_field::mapping::MappingStatus::valid,
"Stateless mapping variation failed in the gravity-"
"displacement-force kernel. Element: "
<< elementId << ", attribute: " << transformation->Attribute << ", quadrature point: "
<< quadratureIndex << ", status: " << static_cast<int>(variationStatus)
);
if (variationStatus != mean_field::mapping::MappingStatus::valid) {
if (!reportCandidateRejection) {
MFEM_VERIFY(
false, "Stateless mapping variation failed in the gravity-"
"displacement-force kernel. Element: "
<< elementId << ", attribute: " << transformation->Attribute
<< ", quadrature point: " << quadratureIndex
<< ", status: " << static_cast<int>(variationStatus)
);
}
candidateRejection = mapping_rejection(variationStatus);
continue;
}
}
densityElement.CalcShape(integrationPoint, densityShape);
@@ -624,10 +726,16 @@ namespace {
const double contribution = displacementShape(scalarDof) * forceValue(component);
MFEM_VERIFY(
std::isfinite(contribution), "The gravity-displacement-force kernel "
"encountered a non-finite contribution."
);
if (!std::isfinite(contribution)) {
if (!reportCandidateRejection) {
MFEM_VERIFY(
false, "The gravity-displacement-force kernel "
"encountered a non-finite contribution."
);
}
candidateRejection = non_finite_rejection();
continue;
}
elementAction(vectorDof) += contribution;
}
@@ -641,11 +749,48 @@ namespace {
localAction.AddElementVector(displacementDofs, elementAction);
}
if (reportCandidateRejection && !vector_is_finite(localAction)) {
candidateRejection = non_finite_rejection();
}
if (reportCandidateRejection) {
auto synchronized = synchronize_rejection(candidateRejection, f.mesh->GetComm());
if (!synchronized.has_value()) {
return synchronized;
}
}
local_to_true(*f.displacementFes, localAction, actionTrue);
if (reportCandidateRejection) {
const std::optional<Rejection> outputRejection = vector_is_finite(actionTrue)
? std::optional<Rejection>{}
: std::optional<Rejection>{non_finite_rejection()};
auto synchronized = synchronize_rejection(outputRejection, f.mesh->GetComm());
if (!synchronized.has_value()) {
return synchronized;
}
}
return {};
}
} // namespace
namespace mean_field::operators::kernels {
GravityDisplacementForceResult try_apply_gravity_displacement_force_residual(
const fem::FEM &f,
const mapping::DomainMapper &domainMapper,
const mfem::Vector &densityTrue,
const mfem::Vector &gravityGradientTrue,
const mfem::Vector &displacementTrue,
mfem::Vector &residualTrue
) {
return apply_gravity_displacement_force_action(
f, domainMapper, GravityDisplacementForceAction::residual, &densityTrue, nullptr, &gravityGradientTrue,
nullptr, nullptr, displacementTrue, residualTrue, true
);
}
void apply_gravity_displacement_force_residual(
const fem::FEM &f,
const mapping::DomainMapper &domainMapper,
@@ -654,10 +799,12 @@ namespace mean_field::operators::kernels {
const mfem::Vector &displacementTrue,
mfem::Vector &residualTrue
) {
apply_gravity_displacement_force_action(
f, domainMapper, GravityDisplacementForceAction::residual, &densityTrue, nullptr, &gravityGradientTrue,
nullptr, nullptr, displacementTrue, residualTrue
auto result = try_apply_gravity_displacement_force_residual(
f, domainMapper, densityTrue, gravityGradientTrue, displacementTrue, residualTrue
);
if (!result.has_value()) {
throw_rejection(result.error());
}
}
void apply_gravity_displacement_force_density_action(
@@ -668,9 +815,9 @@ namespace mean_field::operators::kernels {
const mfem::Vector &displacementTrue,
mfem::Vector &actionTrue
) {
apply_gravity_displacement_force_action(
(void)apply_gravity_displacement_force_action(
f, domainMapper, GravityDisplacementForceAction::density, nullptr, &densityVariationTrue,
&baseGravityGradientTrue, nullptr, nullptr, displacementTrue, actionTrue
&baseGravityGradientTrue, nullptr, nullptr, displacementTrue, actionTrue, false
);
}
@@ -682,9 +829,9 @@ namespace mean_field::operators::kernels {
const mfem::Vector &displacementTrue,
mfem::Vector &actionTrue
) {
apply_gravity_displacement_force_action(
(void)apply_gravity_displacement_force_action(
f, domainMapper, GravityDisplacementForceAction::gravityGradient, &baseDensityTrue, nullptr, nullptr,
&gravityGradientVariationTrue, nullptr, displacementTrue, actionTrue
&gravityGradientVariationTrue, nullptr, displacementTrue, actionTrue, false
);
}
@@ -697,9 +844,9 @@ namespace mean_field::operators::kernels {
const mfem::Vector &displacementTrue,
mfem::Vector &actionTrue
) {
apply_gravity_displacement_force_action(
(void)apply_gravity_displacement_force_action(
f, domainMapper, GravityDisplacementForceAction::displacement, &baseDensityTrue, nullptr,
&baseGravityGradientTrue, nullptr, &displacementVariationTrue, displacementTrue, actionTrue
&baseGravityGradientTrue, nullptr, &displacementVariationTrue, displacementTrue, actionTrue, false
);
}
@@ -714,10 +861,10 @@ namespace mean_field::operators::kernels {
const mfem::Vector &displacementTrue,
mfem::Vector &actionTrue
) {
apply_gravity_displacement_force_action(
(void)apply_gravity_displacement_force_action(
f, domainMapper, GravityDisplacementForceAction::complete, &baseDensityTrue, &densityVariationTrue,
&baseGravityGradientTrue, &gravityGradientVariationTrue, &displacementVariationTrue, displacementTrue,
actionTrue
actionTrue, false
);
}
} // namespace mean_field::operators::kernels

213
libmeanfield/impl/operators/kernels/rotation_displacement_force_kernels.cpp
View File

@@ -2,9 +2,12 @@ module;
#include <array>
#include <cmath>
#include <expected>
#include <optional>
#include <stdexcept>
#include <mfem.hpp>
#include <mpi.h>
module mean_field;
@@ -19,6 +22,70 @@ namespace {
enum class RotationalDisplacementForceAction { residual, density, displacement, complete };
using Rejection = mean_field::operators::kernels::RotationalDisplacementForceRejection;
using Reason = mean_field::operators::kernels::RotationalDisplacementForceRejectionReason;
using Result = mean_field::operators::kernels::RotationalDisplacementForceResult;
[[nodiscard]] Rejection mapping_rejection(const mean_field::mapping::MappingStatus status) {
MFEM_VERIFY(
status != mean_field::mapping::MappingStatus::invalid_dimension,
"The rotational-displacement-force mapping reported an invariant dimension mismatch."
);
return {.reason = Reason::invalid_mapping, .mappingStatus = status};
}
[[nodiscard]] Rejection non_finite_rejection() noexcept {
return {.reason = Reason::non_finite_arithmetic};
}
[[nodiscard]] bool vector_is_finite(const mfem::Vector &vector) noexcept {
for (int index = 0; index < vector.Size(); ++index) {
if (!std::isfinite(vector(index))) {
return false;
}
}
return true;
}
[[nodiscard]] int encode_rejection(const std::optional<Rejection> &rejection) noexcept {
if (!rejection.has_value()) {
return 0;
}
if (rejection->reason == Reason::non_finite_arithmetic) {
return 256;
}
return static_cast<int>(rejection->mappingStatus) + 1;
}
[[nodiscard]] Rejection decode_rejection(const int encoded) noexcept {
if (encoded >= 256) {
return non_finite_rejection();
}
return mapping_rejection(static_cast<mean_field::mapping::MappingStatus>(encoded - 1));
}
[[nodiscard]] Result synchronize_rejection(
const std::optional<Rejection> &localRejection,
const MPI_Comm communicator
) {
const int localEncoded = encode_rejection(localRejection);
int globalEncoded = 0;
if (MPI_Allreduce(&localEncoded, &globalEncoded, 1, MPI_INT, MPI_MAX, communicator) != MPI_SUCCESS) {
throw std::runtime_error("Could not synchronize rotational-displacement-force candidate validity.");
}
if (globalEncoded != 0) {
return std::unexpected(decode_rejection(globalEncoded));
}
return {};
}
[[noreturn]] void throw_rejection(const Rejection &rejection) {
if (rejection.reason == Reason::non_finite_arithmetic) {
throw std::domain_error("The rotational-displacement force produced non-finite arithmetic.");
}
throw std::domain_error("The rotational-displacement force encountered an invalid mapped domain.");
}
void true_to_local(
const mfem::ParFiniteElementSpace &finiteElementSpace,
const mfem::Vector &trueVector,
@@ -186,11 +253,6 @@ namespace {
"The rotational-displacement-force displacement dimension does "
"not match the mesh dimension."
);
validate_finite_vector(
displacementTrue, "The rotational-displacement-force displacement contains a "
"non-finite value."
);
}
void validate_density(
@@ -199,10 +261,9 @@ namespace {
const char *message
) {
MFEM_VERIFY(density.Size() == f.densityFes->GetTrueVSize(), message);
validate_finite_vector(density, message);
}
void apply_rotational_displacement_force_action(
[[nodiscard]] Result apply_rotational_displacement_force_action(
const mean_field::fem::FEM &f,
const mean_field::mapping::DomainMapper &domainMapper,
const mean_field::physics::RigidRotation &rotation,
@@ -211,7 +272,8 @@ namespace {
const mfem::Vector *densityVariationTrue,
const mfem::Vector *displacementVariationTrue,
const mfem::Vector &displacementTrue,
mfem::Vector &actionTrue
mfem::Vector &actionTrue,
const bool reportCandidateRejection
) {
validate_common_inputs(f, domainMapper, displacementTrue);
@@ -261,6 +323,28 @@ namespace {
);
}
bool inputsAreFinite = vector_is_finite(displacementTrue);
if (needsBaseDensity) {
inputsAreFinite = inputsAreFinite && vector_is_finite(*baseDensityTrue);
}
if (needsDensityVariation) {
inputsAreFinite = inputsAreFinite && vector_is_finite(*densityVariationTrue);
}
if (needsDisplacementVariation) {
inputsAreFinite = inputsAreFinite && vector_is_finite(*displacementVariationTrue);
}
if (!reportCandidateRejection) {
MFEM_VERIFY(inputsAreFinite, "The rotational-displacement-force action contains non-finite input data.");
} else {
const std::optional<Rejection> inputRejection =
inputsAreFinite ? std::optional<Rejection>{} : std::optional<Rejection>{non_finite_rejection()};
auto synchronized = synchronize_rejection(inputRejection, f.mesh->GetComm());
if (!synchronized.has_value()) {
return synchronized;
}
}
mfem::Vector baseDensityLocal;
mfem::Vector densityVariationLocal;
mfem::Vector displacementLocal;
@@ -311,6 +395,8 @@ namespace {
const mfem::Ordering::Type displacementOrdering = f.displacementFes->GetOrdering();
std::optional<Rejection> candidateRejection;
for (int elementId = 0; elementId < f.mesh->GetNE(); ++elementId) {
mfem::ElementTransformation *transformation = f.mesh->GetElementTransformation(elementId);
@@ -427,13 +513,19 @@ namespace {
mappingData, *transformation, integrationPoint, workspace, mappingContext
);
MFEM_VERIFY(
mappingStatus == mean_field::mapping::MappingStatus::valid,
"Stateless mapping failed in the rotational-"
"displacement-force kernel. Element: "
<< elementId << ", attribute: " << transformation->Attribute
<< ", quadrature point: " << quadratureIndex << ", status: " << static_cast<int>(mappingStatus)
);
if (mappingStatus != mean_field::mapping::MappingStatus::valid) {
if (!reportCandidateRejection) {
MFEM_VERIFY(
false, "Stateless mapping failed in the rotational-"
"displacement-force kernel. Element: "
<< elementId << ", attribute: " << transformation->Attribute
<< ", quadrature point: " << quadratureIndex
<< ", status: " << static_cast<int>(mappingStatus)
);
}
candidateRejection = mapping_rejection(mappingStatus);
continue;
}
if (needsDisplacementVariation) {
const mean_field::mapping::MappingStatus variationStatus = domainMapper.EvaluateVolumeVariation(
@@ -441,13 +533,19 @@ namespace {
workspace, mappingVariation
);
MFEM_VERIFY(
variationStatus == mean_field::mapping::MappingStatus::valid,
"Stateless mapping variation failed in the "
"rotational-displacement-force kernel. Element: "
<< elementId << ", attribute: " << transformation->Attribute << ", quadrature point: "
<< quadratureIndex << ", status: " << static_cast<int>(variationStatus)
);
if (variationStatus != mean_field::mapping::MappingStatus::valid) {
if (!reportCandidateRejection) {
MFEM_VERIFY(
false, "Stateless mapping variation failed in the "
"rotational-displacement-force kernel. Element: "
<< elementId << ", attribute: " << transformation->Attribute
<< ", quadrature point: " << quadratureIndex
<< ", status: " << static_cast<int>(variationStatus)
);
}
candidateRejection = mapping_rejection(variationStatus);
continue;
}
}
densityElement.CalcShape(integrationPoint, densityShape);
@@ -512,10 +610,16 @@ namespace {
const double contribution = displacementShape(scalarDof) * weightedForce(component);
MFEM_VERIFY(
std::isfinite(contribution), "The rotational-displacement-force kernel "
"encountered a non-finite contribution."
);
if (!std::isfinite(contribution)) {
if (!reportCandidateRejection) {
MFEM_VERIFY(
false, "The rotational-displacement-force kernel "
"encountered a non-finite contribution."
);
}
candidateRejection = non_finite_rejection();
continue;
}
elementAction(vectorDof) += contribution;
}
@@ -529,11 +633,48 @@ namespace {
localAction.AddElementVector(displacementDofs, elementAction);
}
if (reportCandidateRejection && !vector_is_finite(localAction)) {
candidateRejection = non_finite_rejection();
}
if (reportCandidateRejection) {
auto synchronized = synchronize_rejection(candidateRejection, f.mesh->GetComm());
if (!synchronized.has_value()) {
return synchronized;
}
}
local_to_true(*f.displacementFes, localAction, actionTrue);
if (reportCandidateRejection) {
const std::optional<Rejection> outputRejection = vector_is_finite(actionTrue)
? std::optional<Rejection>{}
: std::optional<Rejection>{non_finite_rejection()};
auto synchronized = synchronize_rejection(outputRejection, f.mesh->GetComm());
if (!synchronized.has_value()) {
return synchronized;
}
}
return {};
}
} // namespace
namespace mean_field::operators::kernels {
RotationalDisplacementForceResult try_apply_rotational_displacement_force_residual(
const fem::FEM &f,
const mapping::DomainMapper &domainMapper,
const physics::RigidRotation &rotation,
const mfem::Vector &densityTrue,
const mfem::Vector &displacementTrue,
mfem::Vector &residualTrue
) {
return apply_rotational_displacement_force_action(
f, domainMapper, rotation, RotationalDisplacementForceAction::residual, &densityTrue, nullptr, nullptr,
displacementTrue, residualTrue, true
);
}
void apply_rotational_displacement_force_residual(
const fem::FEM &f,
const mapping::DomainMapper &domainMapper,
@@ -542,10 +683,12 @@ namespace mean_field::operators::kernels {
const mfem::Vector &displacementTrue,
mfem::Vector &residualTrue
) {
apply_rotational_displacement_force_action(
f, domainMapper, rotation, RotationalDisplacementForceAction::residual, &densityTrue, nullptr, nullptr,
displacementTrue, residualTrue
auto result = try_apply_rotational_displacement_force_residual(
f, domainMapper, rotation, densityTrue, displacementTrue, residualTrue
);
if (!result.has_value()) {
throw_rejection(result.error());
}
}
void apply_rotational_displacement_force_density_action(
@@ -556,9 +699,9 @@ namespace mean_field::operators::kernels {
const mfem::Vector &displacementTrue,
mfem::Vector &actionTrue
) {
apply_rotational_displacement_force_action(
(void)apply_rotational_displacement_force_action(
f, domainMapper, rotation, RotationalDisplacementForceAction::density, nullptr, &densityVariationTrue,
nullptr, displacementTrue, actionTrue
nullptr, displacementTrue, actionTrue, false
);
}
@@ -571,9 +714,9 @@ namespace mean_field::operators::kernels {
const mfem::Vector &displacementTrue,
mfem::Vector &actionTrue
) {
apply_rotational_displacement_force_action(
(void)apply_rotational_displacement_force_action(
f, domainMapper, rotation, RotationalDisplacementForceAction::displacement, &baseDensityTrue, nullptr,
&displacementVariationTrue, displacementTrue, actionTrue
&displacementVariationTrue, displacementTrue, actionTrue, false
);
}
@@ -587,9 +730,9 @@ namespace mean_field::operators::kernels {
const mfem::Vector &displacementTrue,
mfem::Vector &actionTrue
) {
apply_rotational_displacement_force_action(
(void)apply_rotational_displacement_force_action(
f, domainMapper, rotation, RotationalDisplacementForceAction::complete, &baseDensityTrue,
&densityVariationTrue, &displacementVariationTrue, displacementTrue, actionTrue
&densityVariationTrue, &displacementVariationTrue, displacementTrue, actionTrue, false
);
}
} // namespace mean_field::operators::kernels