1#ifndef TDLS_TFEL_ADAPTORS_HPP
2#define TDLS_TFEL_ADAPTORS_HPP
94template<
typename T,
typename =
void>
97struct has_data<T, std::void_t<decltype(std::declval<const T&>().data())>>
98 : std::is_pointer<decltype(std::declval<const T&>().data())> {};
103template<
typename T,
typename =
void>
106struct has_pair_data<T, std::void_t<decltype(std::declval<const T&>().data().first),
107 decltype(std::declval<const T&>().data().second)>>
108 : std::bool_constant<std::is_pointer_v<decltype(std::declval<const T&>().data().first)> &&
109 std::is_integral_v<decltype(std::declval<const T&>().data().second)>> {};
112template<typename T, bool = has_pair_data<T>::value>
115 using type =
decltype(std::declval<const T&>().data());
122 using type =
decltype(std::declval<const T&>().data().first);
126template<typename T, bool = has_pair_data<T>::value>
129 using type =
decltype(std::declval<T&>().data());
136 using type =
decltype(std::declval<T&>().data().first);
140template<
typename T,
typename =
void>
143struct has_stride<T, std::void_t<decltype(std::declval<const T&>().stride())>> : std::true_type {};
146template<
typename T,
typename =
void>
149struct has_get_stride<T, std::void_t<decltype(std::declval<const T&>().getStride())>>
153template<
typename T,
typename =
void>
180template<
typename DenseType,
typename =
void>
187template<
typename DenseType>
188struct storage_traits<DenseType, std::enable_if_t<detail::is_dense_v<DenseType>>> {
195 std::remove_pointer_t<typename detail::const_data_pointer<DenseType>::type>>;
197 static constexpr bool is_mutable = !std::is_const_v<
198 std::remove_pointer_t<typename detail::mutable_data_pointer<DenseType>::type>>;
200 static constexpr int arity =
static_cast<int>(indexing_policy::arity);
201 static_assert(arity == 1 || arity == 2,
"tdls adaptors: only vector-like (arity 1) and "
202 "matrix-like (arity 2) objects are supported");
206 static constexpr int extent1 = (arity == 2) ?
static_cast<int>(
indexing_policy{}.size(1)) : 1;
211 if constexpr (arity == 1) {
214 return static_cast<int>(
219 static constexpr int policy_stride = compute_policy_stride();
224 if constexpr (arity == 1) {
235 static_assert(has_uniform_rows(),
236 "tdls adaptors: row-strided matrix views (sub-matrix views) cannot "
237 "be expressed by the single-stride addressing of the TiledLUpp solvers");
245 static constexpr bool is_internal = (policy_stride == 1) && !has_runtime_stride;
250 static constexpr bool has_runtime_extents = (extent0 == 0);
258 return static_cast<int>(o.getIndexingPolicy().size(0));
267 return static_cast<int>(o.getIndexingPolicy().size(1));
277 return static_cast<int>(
288 return static_cast<int>(
310 return static_cast<int>(o.data().second);
312 return static_cast<int>(o.stride());
314 return static_cast<int>(o.getStride());
325 return policy_stride * runtime_view_stride(o);
340 if constexpr (has_runtime_extents) {
341 return runtime_extent1(o);
353 if constexpr (has_runtime_extents) {
354 return runtime_row_stride(o) * runtime_view_stride(o);
356 return extent1 * stride(o);
366 if constexpr (has_runtime_extents) {
367 return static_cast<std::size_t
>(runtime_col_stride(o)) *
368 static_cast<std::size_t
>(runtime_view_stride(o));
370 return static_cast<std::size_t
>(stride(o));
381template<
typename T,
typename =
void>
391template<
typename MatrixType>
394 "tdls adaptors: A must be a dense object exposing data() and an indexing_policy "
395 "type (a gather view holding one pointer per element is not one)");
402template<
typename MatrixType>
409template<
typename ConfigType>
413template<
typename Scalar>
424template<
typename ConfigType>
425concept solver_config = is_solver_config<std::remove_cvref_t<ConfigType>>::value;
432template<
typename MatrixType, auto UserConfig>
434 static_assert(std::is_same_v<std::remove_cvref_t<
decltype(UserConfig)>,
436 "tdls adaptors: the config scalar type must match the matrix element type");
445template<
typename MatrixType, TiledLUppConfig<matrix_scalar<MatrixType>> Config>
450 using scalar =
typename mtraits::value_type;
451 static_assert(std::is_floating_point_v<scalar>,
452 "tdls adaptors: the element type must be float, double or long double");
453 static_assert(mtraits::arity == 2,
"tdls adaptors: A must be matrix-like");
456 static_assert(
runtime_sized || mtraits::extent0 == mtraits::extent1,
457 "tdls adaptors: A must be square");
458 static_assert(Config.layout == MatrixLayout::RowMajor,
459 "tdls adaptors: dense objects are addressed row-major, the TFEL convention; "
460 "a column-major configuration cannot be used through the adaptors");
462 static constexpr int N = mtraits::extent0;
479template<
typename VectorType,
typename Scalar,
int N>
482 static_assert(vtraits::arity == 1,
"tdls adaptors: expected a vector-like object");
483 static_assert(N == 0 || vtraits::has_runtime_extents || vtraits::extent0 == N,
484 "tdls adaptors: vector extent does not match the system dimension");
485 static_assert(std::is_same_v<typename vtraits::value_type, Scalar>,
486 "tdls adaptors: mixed scalar types");
497template<
typename RhsType,
typename SolutionType>
501 static_assert(bt::is_internal == xt::is_internal && bt::policy_stride == xt::policy_stride &&
502 bt::has_runtime_stride == xt::has_runtime_stride,
503 "tdls adaptors: b and x must share the same residency and "
504 "stride (the raw API carries a single rhs_stride for both)");
505 static_assert(bt::has_runtime_extents || xt::has_runtime_extents || bt::extent0 == xt::extent0,
506 "tdls adaptors: b and x extents do not match");
518template<
typename RhsType,
typename SolutionType,
typename Scalar,
int N>
522 static_assert(xt::arity == 2,
"tdls adaptors: a matrix-like right-hand side requires a "
523 "matrix-like solution");
524 static_assert(std::is_same_v<typename bt::value_type, Scalar> &&
525 std::is_same_v<typename xt::value_type, Scalar>,
526 "tdls adaptors: mixed scalar types");
527 static_assert(bt::has_runtime_extents == xt::has_runtime_extents,
528 "tdls adaptors: B and X must both be fixed-size or both runtime-sized");
529 static_assert(N == 0 || !bt::has_runtime_extents,
530 "tdls adaptors: a runtime-sized right-hand-side block requires a "
531 "runtime-sized matrix A");
532 static_assert(bt::has_runtime_extents ||
533 (bt::extent0 == xt::extent0 && bt::extent1 == xt::extent1 &&
534 (N == 0 || bt::extent0 == N)),
535 "tdls adaptors: B and X must have N rows and the same column count");
536 static_assert(bt::policy_stride == xt::policy_stride &&
537 bt::has_runtime_stride == xt::has_runtime_stride,
538 "tdls adaptors: B and X must share the same layout (the raw API "
539 "carries a single stride pair for both)");
548template<
typename PivotType,
int N>
552 std::is_pointer_v<std::decay_t<PivotType>> &&
553 std::is_same_v<std::remove_cv_t<std::remove_pointer_t<std::decay_t<PivotType>>>,
int>;
556 static_assert(
is_raw ||
is_dense,
"tdls adaptors: the pivot must be an int pointer/array or "
557 "a dense int object");
566 return pt::arity == 1 && (N == 0 || pt::has_runtime_extents || pt::extent0 == N);
569 static_assert(
has_matching_shape,
"tdls adaptors: the pivot must be vector-like, of the "
570 "system dimension when both are fixed-size");
576 return !std::is_const_v<std::remove_pointer_t<std::decay_t<PivotType>>>;
578 return !std::is_const_v<PivotType> &&
597 return const_cast<int*
>(
static_cast<const int*
>(p));
600 static_assert(std::is_same_v<typename pt::value_type, int>,
601 "tdls adaptors: the pivot element type must be int");
602 return pt::pointer(p);
634template<
typename Ctx,
int pass_width,
bool inplace,
typename MatrixType,
typename PivotType,
635 typename RhsType,
typename SolutionType>
639 using mt =
typename Ctx::mtraits;
643 static_assert(pass_width >= 0,
"tdls adaptors: pass_width must not be negative");
644 if constexpr (Ctx::runtime_sized) {
645 const int n = mt::runtime_extent0(A);
646 const int m = bt::columns(b);
647 const int rs = xt::row_stride(x);
648 const std::size_t xcol = xt::col_stride(x);
649 if constexpr (inplace) {
650 Ctx::dynamic_solver::template substitute_inplace_multirhs<pass_width>(
651 n, m, mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv),
652 xt::pointer(x), rs, xcol);
654 Ctx::dynamic_solver::template substitute_multirhs<pass_width>(
655 n, m, mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv),
656 bt::pointer(b), xt::pointer(x), rs, xcol);
659 constexpr int M = bt::extent1;
660 static_assert(M >= 1,
"tdls adaptors: B must have at least one column");
661 const int rs = xt::row_stride(x);
662 const std::size_t xcol = xt::col_stride(x);
663 if constexpr (inplace) {
664 Ctx::solver::template substitute_inplace_multirhs<M,
false, pa::is_internal,
665 mt::is_internal, pass_width>(
666 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), xt::pointer(x),
669 Ctx::solver::template substitute_multirhs<M,
false, pa::is_internal, mt::is_internal,
671 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), bt::pointer(b),
672 xt::pointer(x), rs, xcol);
686template<auto UserConfig,
bool oot_diagnostics,
typename MatrixType,
typename PivotType>
690 using mt =
typename ctx::mtraits;
692 static_assert(mt::is_mutable,
"tdls adaptors: factorize writes into A");
693 static_assert(!std::is_const_v<MatrixType>,
"tdls adaptors: A must not be const here "
694 "(factorize writes it)");
695 static_assert(pa::is_mutable,
"tdls adaptors: the pivot must be mutable here "
696 "(factorize writes it)");
697 if constexpr (ctx::runtime_sized) {
698 return ctx::dynamic_solver::template factorize<oot_diagnostics>(
699 mt::runtime_extent0(A), mt::pointer(A), mt::stride(A), pa::pointer(piv),
700 pa::stride(piv), oot_count);
702 return ctx::solver::template factorize<pa::is_internal, mt::is_internal, oot_diagnostics>(
703 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), oot_count);
718template<
auto UserConfig,
int pass_width,
bool oot_diagnostics,
typename MatrixType,
719 typename PivotType,
typename RhsType,
typename SolutionType>
721solve_dispatch(MatrixType& A, PivotType& piv,
const RhsType& b, SolutionType& x,
int& oot_count) {
723 using mt =
typename ctx::mtraits;
727 static_assert(mt::is_mutable,
"tdls adaptors: solve factors A in place");
728 static_assert(xt::is_mutable,
"tdls adaptors: solve writes into x");
729 static_assert(!std::is_const_v<MatrixType> && !std::is_const_v<SolutionType>,
730 "tdls adaptors: A and x must not be const here (solve writes them)");
731 static_assert(pa::is_mutable,
"tdls adaptors: the pivot must be mutable here "
732 "(solve writes it)");
733 if constexpr (bt::arity == 2) {
734 check_multirhs_pair<RhsType, SolutionType, typename ctx::scalar, ctx::N>();
735 if (!factorize_dispatch<UserConfig, oot_diagnostics>(A, piv, oot_count))
return false;
736 substitute_multirhs_dispatch<ctx, pass_width, false>(A, piv, b, x);
739 static_assert(pass_width == 0,
740 "tdls adaptors: pass_width only applies to matrix-like right-hand sides");
741 check_vector<RhsType, typename ctx::scalar, ctx::N>();
742 check_vector<SolutionType, typename ctx::scalar, ctx::N>();
743 check_rhs_pair<RhsType, SolutionType>();
744 if constexpr (ctx::runtime_sized) {
745 return ctx::dynamic_solver::template solve<oot_diagnostics>(
746 mt::runtime_extent0(A), mt::pointer(A), mt::stride(A), pa::pointer(piv),
747 pa::stride(piv), bt::pointer(b), xt::pointer(x), xt::stride(x), oot_count);
749 return ctx::solver::template
solve<xt::is_internal, pa::is_internal, mt::is_internal,
751 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), bt::pointer(b),
752 xt::pointer(x), xt::stride(x), oot_count);
768template<
auto UserConfig,
int pass_width,
bool oot_diagnostics,
typename MatrixType,
769 typename PivotType,
typename VectorType>
773 using mt =
typename ctx::mtraits;
776 static_assert(mt::is_mutable,
"tdls adaptors: solve_inplace factors A in place");
777 static_assert(yt::is_mutable,
"tdls adaptors: solve_inplace writes into y");
778 static_assert(!std::is_const_v<MatrixType> && !std::is_const_v<VectorType>,
779 "tdls adaptors: A and y must not be const here (solve_inplace writes them)");
780 static_assert(pa::is_mutable,
"tdls adaptors: the pivot must be mutable here "
781 "(solve_inplace writes it)");
782 if constexpr (yt::arity == 2) {
783 check_multirhs_pair<VectorType, VectorType, typename ctx::scalar, ctx::N>();
784 if (!factorize_dispatch<UserConfig, oot_diagnostics>(A, piv, oot_count))
return false;
785 substitute_multirhs_dispatch<ctx, pass_width, true>(A, piv, y, y);
788 static_assert(pass_width == 0,
789 "tdls adaptors: pass_width only applies to matrix-like right-hand sides");
790 check_vector<VectorType, typename ctx::scalar, ctx::N>();
791 if constexpr (ctx::runtime_sized) {
792 return ctx::dynamic_solver::template solve_inplace<oot_diagnostics>(
793 mt::runtime_extent0(A), mt::pointer(A), mt::stride(A), pa::pointer(piv),
794 pa::stride(piv), yt::pointer(y), yt::stride(y), oot_count);
796 return ctx::solver::template
solve_inplace<yt::is_internal, pa::is_internal,
797 mt::is_internal, oot_diagnostics>(
798 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), yt::pointer(y),
799 yt::stride(y), oot_count);
814template<detail::solver_config auto UserConfig,
typename MatrixType,
typename PivotType>
818 return detail::factorize_dispatch<UserConfig, false>(A, piv, unused);
825template<
typename MatrixType,
typename PivotType>
828 return factorize<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}>(A, piv);
840template<detail::solver_config auto UserConfig,
typename MatrixType,
typename PivotType>
842factorize(MatrixType& A, PivotType& piv,
int& oot_count) {
843 return detail::factorize_dispatch<UserConfig, true>(A, piv, oot_count);
852template<
typename MatrixType,
typename PivotType>
854factorize(MatrixType& A, PivotType& piv,
int& oot_count) {
855 return factorize<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}>(A, piv, oot_count);
875template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
876 typename PivotType,
typename RhsType,
typename SolutionType>
878solve(MatrixType& A, PivotType& piv,
const RhsType& b, SolutionType& x) {
880 return detail::solve_dispatch<UserConfig, pass_width, false>(A, piv, b, x, unused);
892template<
int pass_width = 0,
typename MatrixType,
typename PivotType,
typename RhsType,
893 typename SolutionType>
895solve(MatrixType& A, PivotType& piv,
const RhsType& b, SolutionType& x) {
896 return solve<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, b, x);
912template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
913 typename PivotType,
typename RhsType,
typename SolutionType>
915solve(MatrixType& A, PivotType& piv,
const RhsType& b, SolutionType& x,
int& oot_count) {
916 return detail::solve_dispatch<UserConfig, pass_width, true>(A, piv, b, x, oot_count);
929template<
int pass_width = 0,
typename MatrixType,
typename PivotType,
typename RhsType,
930 typename SolutionType>
932solve(MatrixType& A, PivotType& piv,
const RhsType& b, SolutionType& x,
int& oot_count) {
933 return solve<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, b, x,
956template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
957 typename PivotType,
typename VectorType>
961 return detail::solve_inplace_dispatch<UserConfig, pass_width, false>(A, piv, y, unused);
974template<
int pass_w
idth = 0,
typename MatrixType,
typename PivotType,
typename VectorType>
977 return solve_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv,
994template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
995 typename PivotType,
typename VectorType>
997solve_inplace(MatrixType& A, PivotType& piv, VectorType& y,
int& oot_count) {
998 return detail::solve_inplace_dispatch<UserConfig, pass_width, true>(A, piv, y, oot_count);
1011template<
int pass_w
idth = 0,
typename MatrixType,
typename PivotType,
typename VectorType>
1014 return solve_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(
1015 A, piv, y, oot_count);
1034template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
1035 typename PivotType,
typename RhsType,
typename SolutionType>
1037substitute(
const MatrixType& A,
const PivotType& piv,
const RhsType& b, SolutionType& x) {
1040 using mt =
typename ctx::mtraits;
1044 static_assert(xt::is_mutable,
"tdls adaptors: substitute writes into x");
1045 static_assert(!std::is_const_v<SolutionType>,
1046 "tdls adaptors: x must not be const here (substitute writes it)");
1047 if constexpr (bt::arity == 2) {
1048 detail::check_multirhs_pair<RhsType, SolutionType, typename ctx::scalar, ctx::N>();
1049 detail::substitute_multirhs_dispatch<ctx, pass_width, false>(A, piv, b, x);
1051 static_assert(pass_width == 0,
1052 "tdls adaptors: pass_width only applies to matrix-like right-hand sides");
1053 detail::check_vector<RhsType, typename ctx::scalar, ctx::N>();
1054 detail::check_vector<SolutionType, typename ctx::scalar, ctx::N>();
1055 detail::check_rhs_pair<RhsType, SolutionType>();
1056 if constexpr (ctx::runtime_sized) {
1057 ctx::dynamic_solver::substitute(mt::runtime_extent0(A), mt::pointer(A), mt::stride(A),
1058 pa::pointer(piv), pa::stride(piv), bt::pointer(b),
1059 xt::pointer(x), xt::stride(x));
1061 ctx::solver::template substitute<xt::is_internal, pa::is_internal, mt::is_internal>(
1062 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), bt::pointer(b),
1063 xt::pointer(x), xt::stride(x));
1076template<
int pass_width = 0,
typename MatrixType,
typename PivotType,
typename RhsType,
1077 typename SolutionType>
1079substitute(
const MatrixType& A,
const PivotType& piv,
const RhsType& b, SolutionType& x) {
1080 substitute<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, b, x);
1097template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
1098 typename PivotType,
typename SolutionType>
1103 using mt =
typename ctx::mtraits;
1106 static_assert(xt::is_mutable,
"tdls adaptors: substitute_inplace writes into x");
1107 static_assert(!std::is_const_v<SolutionType>,
1108 "tdls adaptors: x must not be const here (substitute_inplace writes it)");
1109 if constexpr (xt::arity == 2) {
1110 detail::check_multirhs_pair<SolutionType, SolutionType, typename ctx::scalar, ctx::N>();
1111 detail::substitute_multirhs_dispatch<ctx, pass_width, true>(A, piv, x, x);
1113 static_assert(pass_width == 0,
1114 "tdls adaptors: pass_width only applies to matrix-like right-hand sides");
1115 detail::check_vector<SolutionType, typename ctx::scalar, ctx::N>();
1116 if constexpr (ctx::runtime_sized) {
1117 ctx::dynamic_solver::substitute_inplace(mt::runtime_extent0(A), mt::pointer(A),
1118 mt::stride(A), pa::pointer(piv),
1119 pa::stride(piv), xt::pointer(x), xt::stride(x));
1123 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), xt::pointer(x),
1136template<
int pass_w
idth = 0,
typename MatrixType,
typename PivotType,
typename SolutionType>
1139 substitute_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, x);
1163template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
1164 typename PivotType,
typename SolutionType>
1169 using mt =
typename ctx::mtraits;
1172 static_assert(xt::is_mutable,
"tdls adaptors: substitute_canonical writes into x");
1173 static_assert(!std::is_const_v<SolutionType>,
1174 "tdls adaptors: x must not be const here (substitute_canonical writes it)");
1175 if constexpr (xt::arity == 2) {
1176 detail::check_multirhs_pair<SolutionType, SolutionType, typename ctx::scalar, ctx::N>();
1177 static_assert(pass_width >= 0,
"tdls adaptors: pass_width must not be negative");
1178 if constexpr (ctx::runtime_sized) {
1179 ctx::dynamic_solver::template substitute_canonical_multirhs<pass_width>(
1180 mt::runtime_extent0(A), xt::columns(x), mt::pointer(A), mt::stride(A),
1181 pa::pointer(piv), pa::stride(piv), col, xt::pointer(x), xt::row_stride(x),
1184 constexpr int M = xt::extent1;
1185 static_assert(M >= 1,
"tdls adaptors: x must have at least one column");
1186 ctx::solver::template substitute_canonical_multirhs<M,
false, pa::is_internal,
1187 mt::is_internal, pass_width>(
1188 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), col,
1189 xt::pointer(x), xt::row_stride(x), xt::col_stride(x));
1192 static_assert(pass_width == 0,
1193 "tdls adaptors: pass_width only applies to matrix-like right-hand sides");
1194 detail::check_vector<SolutionType, typename ctx::scalar, ctx::N>();
1195 if constexpr (ctx::runtime_sized) {
1196 ctx::dynamic_solver::substitute_canonical(
1197 mt::runtime_extent0(A), mt::pointer(A), mt::stride(A), pa::pointer(piv),
1198 pa::stride(piv), col, xt::pointer(x), xt::stride(x));
1202 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), col,
1203 xt::pointer(x), xt::stride(x));
1217template<
int pass_w
idth = 0,
typename MatrixType,
typename PivotType,
typename SolutionType>
1220 substitute_canonical<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv,
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool factorize_dispatch(MatrixType &A, PivotType &piv, int &oot_count)
Shared engine of the factorize entry points: reduces the dense objects to raw arguments and forwards ...
Definition adaptors.hpp:688
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace(MatrixType &A, PivotType &piv, VectorType &y)
Solve A y = y on dense objects with the fused factorization (forward substitution folded into factori...
Definition adaptors.hpp:959
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool solve_dispatch(MatrixType &A, PivotType &piv, const RhsType &b, SolutionType &x, int &oot_count)
Shared engine of the solve entry points: factorize + substitute on dense objects, with or without the...
Definition adaptors.hpp:721
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace_dispatch(MatrixType &A, PivotType &piv, VectorType &y, int &oot_count)
Shared engine of the solve_inplace entry points: fused factorization on dense objects,...
Definition adaptors.hpp:771
typename matrix_scalar_of< MatrixType >::type matrix_scalar
Scalar type of a dense matrix argument, see matrix_scalar_of.
Definition adaptors.hpp:403
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void check_rhs_pair()
Checks that the right-hand side and the solution share one residency and one compile-time stride: the...
Definition adaptors.hpp:498
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void check_vector()
Checks that a vector-like argument matches the system: arity 1, extent N, same scalar type....
Definition adaptors.hpp:480
consteval auto checked_config()
Validates a user configuration value against the matrix scalar type and returns it with its concrete ...
Definition adaptors.hpp:433
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void substitute_inplace(const MatrixType &A, const PivotType &piv, SolutionType &x)
Solve in place on dense objects: x holds the unpermuted right-hand side on entry and the solution on ...
Definition adaptors.hpp:1100
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void check_multirhs_pair()
Checks a matrix-like right-hand-side block pair (B, X): matching scalar, matching extents where compi...
Definition adaptors.hpp:519
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool solve(MatrixType &A, PivotType &piv, const RhsType &b, SolutionType &x)
Solve A x = b on dense objects: factorize + substitute.
Definition adaptors.hpp:878
constexpr bool is_dense_v
A type is "dense" when it exposes both a data pointer and an indexing policy, the structural core of ...
Definition adaptors.hpp:161
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void substitute_multirhs_dispatch(const MatrixType &A, const PivotType &piv, const RhsType &b, SolutionType &x)
Shared translator of the matrix right-hand-side entry points: reduces the dense objects to raw argume...
Definition adaptors.hpp:637
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void substitute(const MatrixType &A, const PivotType &piv, const RhsType &b, SolutionType &x)
Solve x := U^-1 L^-1 P b on dense objects, from a prior factorize. b and x must not alias.
Definition adaptors.hpp:1037
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool factorize(MatrixType &A, PivotType &piv)
Factor a dense matrix object in place, A := P*L*U.
Definition adaptors.hpp:815
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void substitute_canonical(const MatrixType &A, const PivotType &piv, const int col, SolutionType &x)
Solve A x = e_col on dense objects, from a prior factorize: the consistent-tangent-operator path.
Definition adaptors.hpp:1166
Concept gating the explicit-configuration entry points.
Definition adaptors.hpp:425
Toolchain detection and portability macros.
#define TDLS_FORCE_INLINE_GPU
Inlining forced on GPU only: __forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS...
Definition macros.hpp:77
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:50
Compile-time knobs of the TiledLUpp solvers, carrying the tuned defaults.
Definition config.hpp:50
Runtime-size tiled dense LU factorization with logical partial pivoting and out-of-tile pivot recover...
Definition solver_dynamic.hpp:176
Tiled dense LU factorization with logical partial pivoting and out-of-tile pivot recovery,...
Definition solver_static.hpp:255
Common compile-time context of the adaptor entry points: resolves the scalar type,...
Definition adaptors.hpp:446
static constexpr int N
system dimension (fixed-size path; zero on the runtime path)
Definition adaptors.hpp:462
typename mtraits::value_type scalar
scalar type of the system
Definition adaptors.hpp:450
static constexpr bool runtime_sized
true when the matrix is runtime-sized (dynamic solver path)
Definition adaptors.hpp:455
decltype(std::declval< const T & >().data().first) type
pointer type of the first pair member of data()
Definition adaptors.hpp:122
Element pointer type of a dense object, const flavour.
Definition adaptors.hpp:113
decltype(std::declval< const T & >().data()) type
pointer type returned by data()
Definition adaptors.hpp:115
Detects a data() member returning a pointer.
Definition adaptors.hpp:95
Detects a getStride() member.
Definition adaptors.hpp:147
Detects a nested indexing_policy type.
Definition adaptors.hpp:154
Detects a data() member returning a (pointer, stride) pair: the shape of strided views,...
Definition adaptors.hpp:104
Detects a complete storage_traits: the structural contract or a user specialization.
Definition adaptors.hpp:382
Detects a stride() member.
Definition adaptors.hpp:141
Detects the configuration value types accepted by the entry points. Only the TiledLUpp family exists ...
Definition adaptors.hpp:410
Scalar type of a dense matrix argument, read from its storage description; a matrix argument without ...
Definition adaptors.hpp:392
typename storage_traits< std::remove_cv_t< MatrixType > >::value_type type
scalar type of the matrix
Definition adaptors.hpp:397
decltype(std::declval< T & >().data().first) type
pointer type of the first pair member of data()
Definition adaptors.hpp:136
Element pointer type of a dense object, mutable flavour.
Definition adaptors.hpp:127
decltype(std::declval< T & >().data()) type
pointer type returned by data()
Definition adaptors.hpp:129
Pivot argument unwrapping: accepts a raw int pointer/array (treated as contiguous caller-local storag...
Definition adaptors.hpp:549
static constexpr bool is_internal
true when the pivot storage is compile-time contiguous
Definition adaptors.hpp:584
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int * pointer(const PivotType &p) noexcept
Definition adaptors.hpp:595
static constexpr bool is_raw
true for a raw int pointer or int array
Definition adaptors.hpp:551
static constexpr bool is_dense
true for an object with storage_traits
Definition adaptors.hpp:555
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int stride(const PivotType &p) noexcept
Definition adaptors.hpp:609
static constexpr bool is_mutable
true when the pivot argument may be written (non-const object, non-const pointee for raw pointers)
Definition adaptors.hpp:574
static constexpr bool has_matching_shape
true when the pivot is vector-like, of the system dimension when both are fixed-size
Definition adaptors.hpp:561
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int runtime_view_stride(const DenseType &o) noexcept
Definition adaptors.hpp:308
typename DenseType::indexing_policy indexing_policy
indexing policy of the object
Definition adaptors.hpp:190
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int runtime_extent0(const DenseType &o) noexcept
Definition adaptors.hpp:257
static constexpr int compute_policy_stride()
Definition adaptors.hpp:210
typename indexing_policy::size_type policy_size_type
a shorthand for the indexing size type
Definition adaptors.hpp:192
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int columns(const DenseType &o) noexcept
Definition adaptors.hpp:339
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int stride(const DenseType &o) noexcept
Definition adaptors.hpp:324
static constexpr bool has_uniform_rows()
Definition adaptors.hpp:223
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr std::size_t col_stride(const DenseType &o) noexcept
Definition adaptors.hpp:365
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int runtime_col_stride(const DenseType &o) noexcept
Definition adaptors.hpp:287
std::remove_cv_t< std::remove_pointer_t< typename detail::const_data_pointer< DenseType >::type > > value_type
element type (without cv-qualifiers)
Definition adaptors.hpp:194
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int row_stride(const DenseType &o) noexcept
Definition adaptors.hpp:352
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr value_type * pointer(const DenseType &o) noexcept
Definition adaptors.hpp:296
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int runtime_extent1(const DenseType &o) noexcept
Definition adaptors.hpp:266
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int runtime_row_stride(const DenseType &o) noexcept
Definition adaptors.hpp:276
Storage description of a dense object: element type, extents, element pointer and element stride.
Definition adaptors.hpp:181
Runtime-size variant of the TiledLUpp solver.
TiledLUpp solver: LU with partial pivoting on tile grids, one thread / work-item per system.