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());
324 return policy_stride * runtime_view_stride(o);
338 if constexpr (has_runtime_extents) {
339 return runtime_extent1(o);
350 if constexpr (has_runtime_extents) {
351 return runtime_row_stride(o) * runtime_view_stride(o);
353 return extent1 * stride(o);
362 if constexpr (has_runtime_extents) {
363 return runtime_col_stride(o) * runtime_view_stride(o);
376template<
typename T,
typename =
void>
386template<
typename MatrixType>
389 "tdls adaptors: A must be a dense object exposing data() and an indexing_policy "
390 "type (a gather view holding one pointer per element is not one)");
397template<
typename MatrixType>
404template<
typename ConfigType>
408template<
typename Scalar>
419template<
typename ConfigType>
420concept solver_config = is_solver_config<std::remove_cvref_t<ConfigType>>::value;
427template<
typename MatrixType, auto UserConfig>
429 static_assert(std::is_same_v<std::remove_cvref_t<
decltype(UserConfig)>,
431 "tdls adaptors: the config scalar type must match the matrix element type");
440template<
typename MatrixType, TiledLUppConfig<matrix_scalar<MatrixType>> Config>
445 using scalar =
typename mtraits::value_type;
446 static_assert(std::is_floating_point_v<scalar>,
447 "tdls adaptors: the element type must be float, double or long double");
448 static_assert(mtraits::arity == 2,
"tdls adaptors: A must be matrix-like");
451 static_assert(
runtime_sized || mtraits::extent0 == mtraits::extent1,
452 "tdls adaptors: A must be square");
453 static_assert(Config.layout == MatrixLayout::RowMajor,
454 "tdls adaptors: dense objects are addressed row-major, the TFEL convention; "
455 "a column-major configuration cannot be used through the adaptors");
457 static constexpr int N = mtraits::extent0;
474template<
typename VectorType,
typename Scalar,
int N>
477 static_assert(vtraits::arity == 1,
"tdls adaptors: expected a vector-like object");
478 static_assert(N == 0 || vtraits::has_runtime_extents || vtraits::extent0 == N,
479 "tdls adaptors: vector extent does not match the system dimension");
480 static_assert(std::is_same_v<typename vtraits::value_type, Scalar>,
481 "tdls adaptors: mixed scalar types");
492template<
typename RhsType,
typename SolutionType>
496 static_assert(bt::is_internal == xt::is_internal && bt::policy_stride == xt::policy_stride &&
497 bt::has_runtime_stride == xt::has_runtime_stride,
498 "tdls adaptors: b and x must share the same residency and "
499 "stride (the raw API carries a single rhs_stride for both)");
500 static_assert(bt::has_runtime_extents || xt::has_runtime_extents || bt::extent0 == xt::extent0,
501 "tdls adaptors: b and x extents do not match");
513template<
typename RhsType,
typename SolutionType,
typename Scalar,
int N>
517 static_assert(xt::arity == 2,
"tdls adaptors: a matrix-like right-hand side requires a "
518 "matrix-like solution");
519 static_assert(std::is_same_v<typename bt::value_type, Scalar> &&
520 std::is_same_v<typename xt::value_type, Scalar>,
521 "tdls adaptors: mixed scalar types");
522 static_assert(bt::has_runtime_extents == xt::has_runtime_extents,
523 "tdls adaptors: B and X must both be fixed-size or both runtime-sized");
524 static_assert(N == 0 || !bt::has_runtime_extents,
525 "tdls adaptors: a runtime-sized right-hand-side block requires a "
526 "runtime-sized matrix A");
527 static_assert(bt::has_runtime_extents ||
528 (bt::extent0 == xt::extent0 && bt::extent1 == xt::extent1 &&
529 (N == 0 || bt::extent0 == N)),
530 "tdls adaptors: B and X must have N rows and the same column count");
531 static_assert(bt::policy_stride == xt::policy_stride &&
532 bt::has_runtime_stride == xt::has_runtime_stride,
533 "tdls adaptors: B and X must share the same layout (the raw API "
534 "carries a single stride pair for both)");
543template<
typename PivotType,
int N>
547 std::is_pointer_v<std::decay_t<PivotType>> &&
548 std::is_same_v<std::remove_cv_t<std::remove_pointer_t<std::decay_t<PivotType>>>,
int>;
551 static_assert(
is_raw ||
is_dense,
"tdls adaptors: the pivot must be an int pointer/array or "
552 "a dense int object");
561 return pt::arity == 1 && (N == 0 || pt::has_runtime_extents || pt::extent0 == N);
564 static_assert(
has_matching_shape,
"tdls adaptors: the pivot must be vector-like, of the "
565 "system dimension when both are fixed-size");
571 return !std::is_const_v<std::remove_pointer_t<std::decay_t<PivotType>>>;
573 return !std::is_const_v<PivotType> &&
591 return const_cast<int*
>(
static_cast<const int*
>(p));
594 static_assert(std::is_same_v<typename pt::value_type, int>,
595 "tdls adaptors: the pivot element type must be int");
596 return pt::pointer(p);
621template<
typename Ctx,
int pass_width,
bool inplace,
typename MatrixType,
typename PivotType,
622 typename RhsType,
typename SolutionType>
626 using mt =
typename Ctx::mtraits;
630 static_assert(pass_width >= 0,
"tdls adaptors: pass_width must not be negative");
631 if constexpr (Ctx::runtime_sized) {
632 const int n = mt::runtime_extent0(A);
633 const int m = bt::columns(b);
634 const int rs = xt::row_stride(x);
635 const int xcol = xt::col_stride(x);
636 if constexpr (inplace) {
637 Ctx::dynamic_solver::template substitute_inplace_multirhs<pass_width>(
638 n, m, mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv),
639 xt::pointer(x), rs, xcol);
641 Ctx::dynamic_solver::template substitute_multirhs<pass_width>(
642 n, m, mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv),
643 bt::pointer(b), xt::pointer(x), rs, xcol);
646 constexpr int M = bt::extent1;
647 static_assert(M >= 1,
"tdls adaptors: B must have at least one column");
648 const int rs = xt::row_stride(x);
649 const int xcol = xt::col_stride(x);
650 if constexpr (inplace) {
651 Ctx::solver::template substitute_inplace_multirhs<M,
false, pa::is_internal,
652 mt::is_internal, pass_width>(
653 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), xt::pointer(x),
656 Ctx::solver::template substitute_multirhs<M,
false, pa::is_internal, mt::is_internal,
658 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), bt::pointer(b),
659 xt::pointer(x), rs, xcol);
673template<auto UserConfig,
bool oot_diagnostics,
typename MatrixType,
typename PivotType>
677 using mt =
typename ctx::mtraits;
679 static_assert(mt::is_mutable,
"tdls adaptors: factorize writes into A");
680 static_assert(!std::is_const_v<MatrixType>,
"tdls adaptors: A must not be const here "
681 "(factorize writes it)");
682 static_assert(pa::is_mutable,
"tdls adaptors: the pivot must be mutable here "
683 "(factorize writes it)");
684 if constexpr (ctx::runtime_sized) {
685 return ctx::dynamic_solver::template factorize<oot_diagnostics>(
686 mt::runtime_extent0(A), mt::pointer(A), mt::stride(A), pa::pointer(piv),
687 pa::stride(piv), oot_count);
689 return ctx::solver::template factorize<pa::is_internal, mt::is_internal, oot_diagnostics>(
690 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), oot_count);
705template<
auto UserConfig,
int pass_width,
bool oot_diagnostics,
typename MatrixType,
706 typename PivotType,
typename RhsType,
typename SolutionType>
708solve_dispatch(MatrixType& A, PivotType& piv,
const RhsType& b, SolutionType& x,
int& oot_count) {
710 using mt =
typename ctx::mtraits;
714 static_assert(mt::is_mutable,
"tdls adaptors: solve factors A in place");
715 static_assert(xt::is_mutable,
"tdls adaptors: solve writes into x");
716 static_assert(!std::is_const_v<MatrixType> && !std::is_const_v<SolutionType>,
717 "tdls adaptors: A and x must not be const here (solve writes them)");
718 static_assert(pa::is_mutable,
"tdls adaptors: the pivot must be mutable here "
719 "(solve writes it)");
720 if constexpr (bt::arity == 2) {
721 check_multirhs_pair<RhsType, SolutionType, typename ctx::scalar, ctx::N>();
722 if (!factorize_dispatch<UserConfig, oot_diagnostics>(A, piv, oot_count))
return false;
723 substitute_multirhs_dispatch<ctx, pass_width, false>(A, piv, b, x);
726 static_assert(pass_width == 0,
727 "tdls adaptors: pass_width only applies to matrix-like right-hand sides");
728 check_vector<RhsType, typename ctx::scalar, ctx::N>();
729 check_vector<SolutionType, typename ctx::scalar, ctx::N>();
730 check_rhs_pair<RhsType, SolutionType>();
731 if constexpr (ctx::runtime_sized) {
732 return ctx::dynamic_solver::template solve<oot_diagnostics>(
733 mt::runtime_extent0(A), mt::pointer(A), mt::stride(A), pa::pointer(piv),
734 pa::stride(piv), bt::pointer(b), xt::pointer(x), xt::stride(x), oot_count);
736 return ctx::solver::template
solve<xt::is_internal, pa::is_internal, mt::is_internal,
738 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), bt::pointer(b),
739 xt::pointer(x), xt::stride(x), oot_count);
755template<
auto UserConfig,
int pass_width,
bool oot_diagnostics,
typename MatrixType,
756 typename PivotType,
typename VectorType>
760 using mt =
typename ctx::mtraits;
763 static_assert(mt::is_mutable,
"tdls adaptors: solve_inplace factors A in place");
764 static_assert(yt::is_mutable,
"tdls adaptors: solve_inplace writes into y");
765 static_assert(!std::is_const_v<MatrixType> && !std::is_const_v<VectorType>,
766 "tdls adaptors: A and y must not be const here (solve_inplace writes them)");
767 static_assert(pa::is_mutable,
"tdls adaptors: the pivot must be mutable here "
768 "(solve_inplace writes it)");
769 if constexpr (yt::arity == 2) {
770 check_multirhs_pair<VectorType, VectorType, typename ctx::scalar, ctx::N>();
771 if (!factorize_dispatch<UserConfig, oot_diagnostics>(A, piv, oot_count))
return false;
772 substitute_multirhs_dispatch<ctx, pass_width, true>(A, piv, y, y);
775 static_assert(pass_width == 0,
776 "tdls adaptors: pass_width only applies to matrix-like right-hand sides");
777 check_vector<VectorType, typename ctx::scalar, ctx::N>();
778 if constexpr (ctx::runtime_sized) {
779 return ctx::dynamic_solver::template solve_inplace<oot_diagnostics>(
780 mt::runtime_extent0(A), mt::pointer(A), mt::stride(A), pa::pointer(piv),
781 pa::stride(piv), yt::pointer(y), yt::stride(y), oot_count);
783 return ctx::solver::template
solve_inplace<yt::is_internal, pa::is_internal,
784 mt::is_internal, oot_diagnostics>(
785 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), yt::pointer(y),
786 yt::stride(y), oot_count);
801template<detail::solver_config auto UserConfig,
typename MatrixType,
typename PivotType>
805 return detail::factorize_dispatch<UserConfig, false>(A, piv, unused);
812template<
typename MatrixType,
typename PivotType>
815 return factorize<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}>(A, piv);
827template<detail::solver_config auto UserConfig,
typename MatrixType,
typename PivotType>
829factorize(MatrixType& A, PivotType& piv,
int& oot_count) {
830 return detail::factorize_dispatch<UserConfig, true>(A, piv, oot_count);
839template<
typename MatrixType,
typename PivotType>
841factorize(MatrixType& A, PivotType& piv,
int& oot_count) {
842 return factorize<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}>(A, piv, oot_count);
862template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
863 typename PivotType,
typename RhsType,
typename SolutionType>
865solve(MatrixType& A, PivotType& piv,
const RhsType& b, SolutionType& x) {
867 return detail::solve_dispatch<UserConfig, pass_width, false>(A, piv, b, x, unused);
879template<
int pass_width = 0,
typename MatrixType,
typename PivotType,
typename RhsType,
880 typename SolutionType>
882solve(MatrixType& A, PivotType& piv,
const RhsType& b, SolutionType& x) {
883 return solve<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, b, x);
899template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
900 typename PivotType,
typename RhsType,
typename SolutionType>
902solve(MatrixType& A, PivotType& piv,
const RhsType& b, SolutionType& x,
int& oot_count) {
903 return detail::solve_dispatch<UserConfig, pass_width, true>(A, piv, b, x, oot_count);
916template<
int pass_width = 0,
typename MatrixType,
typename PivotType,
typename RhsType,
917 typename SolutionType>
919solve(MatrixType& A, PivotType& piv,
const RhsType& b, SolutionType& x,
int& oot_count) {
920 return solve<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, b, x,
943template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
944 typename PivotType,
typename VectorType>
948 return detail::solve_inplace_dispatch<UserConfig, pass_width, false>(A, piv, y, unused);
961template<
int pass_w
idth = 0,
typename MatrixType,
typename PivotType,
typename VectorType>
964 return solve_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv,
981template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
982 typename PivotType,
typename VectorType>
984solve_inplace(MatrixType& A, PivotType& piv, VectorType& y,
int& oot_count) {
985 return detail::solve_inplace_dispatch<UserConfig, pass_width, true>(A, piv, y, oot_count);
998template<
int pass_w
idth = 0,
typename MatrixType,
typename PivotType,
typename VectorType>
1001 return solve_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(
1002 A, piv, y, oot_count);
1021template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
1022 typename PivotType,
typename RhsType,
typename SolutionType>
1024substitute(
const MatrixType& A,
const PivotType& piv,
const RhsType& b, SolutionType& x) {
1027 using mt =
typename ctx::mtraits;
1031 static_assert(xt::is_mutable,
"tdls adaptors: substitute writes into x");
1032 static_assert(!std::is_const_v<SolutionType>,
1033 "tdls adaptors: x must not be const here (substitute writes it)");
1034 if constexpr (bt::arity == 2) {
1035 detail::check_multirhs_pair<RhsType, SolutionType, typename ctx::scalar, ctx::N>();
1036 detail::substitute_multirhs_dispatch<ctx, pass_width, false>(A, piv, b, x);
1038 static_assert(pass_width == 0,
1039 "tdls adaptors: pass_width only applies to matrix-like right-hand sides");
1040 detail::check_vector<RhsType, typename ctx::scalar, ctx::N>();
1041 detail::check_vector<SolutionType, typename ctx::scalar, ctx::N>();
1042 detail::check_rhs_pair<RhsType, SolutionType>();
1043 if constexpr (ctx::runtime_sized) {
1044 ctx::dynamic_solver::substitute(mt::runtime_extent0(A), mt::pointer(A), mt::stride(A),
1045 pa::pointer(piv), pa::stride(piv), bt::pointer(b),
1046 xt::pointer(x), xt::stride(x));
1048 ctx::solver::template substitute<xt::is_internal, pa::is_internal, mt::is_internal>(
1049 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), bt::pointer(b),
1050 xt::pointer(x), xt::stride(x));
1063template<
int pass_width = 0,
typename MatrixType,
typename PivotType,
typename RhsType,
1064 typename SolutionType>
1066substitute(
const MatrixType& A,
const PivotType& piv,
const RhsType& b, SolutionType& x) {
1067 substitute<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, b, x);
1084template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
1085 typename PivotType,
typename SolutionType>
1090 using mt =
typename ctx::mtraits;
1093 static_assert(xt::is_mutable,
"tdls adaptors: substitute_inplace writes into x");
1094 static_assert(!std::is_const_v<SolutionType>,
1095 "tdls adaptors: x must not be const here (substitute_inplace writes it)");
1096 if constexpr (xt::arity == 2) {
1097 detail::check_multirhs_pair<SolutionType, SolutionType, typename ctx::scalar, ctx::N>();
1098 detail::substitute_multirhs_dispatch<ctx, pass_width, true>(A, piv, x, x);
1100 static_assert(pass_width == 0,
1101 "tdls adaptors: pass_width only applies to matrix-like right-hand sides");
1102 detail::check_vector<SolutionType, typename ctx::scalar, ctx::N>();
1103 if constexpr (ctx::runtime_sized) {
1104 ctx::dynamic_solver::substitute_inplace(mt::runtime_extent0(A), mt::pointer(A),
1105 mt::stride(A), pa::pointer(piv),
1106 pa::stride(piv), xt::pointer(x), xt::stride(x));
1110 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), xt::pointer(x),
1123template<
int pass_w
idth = 0,
typename MatrixType,
typename PivotType,
typename SolutionType>
1126 substitute_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, x);
1150template<detail::solver_config
auto UserConfig,
int pass_width = 0,
typename MatrixType,
1151 typename PivotType,
typename SolutionType>
1156 using mt =
typename ctx::mtraits;
1159 static_assert(xt::is_mutable,
"tdls adaptors: substitute_canonical writes into x");
1160 static_assert(!std::is_const_v<SolutionType>,
1161 "tdls adaptors: x must not be const here (substitute_canonical writes it)");
1162 if constexpr (xt::arity == 2) {
1163 detail::check_multirhs_pair<SolutionType, SolutionType, typename ctx::scalar, ctx::N>();
1164 static_assert(pass_width >= 0,
"tdls adaptors: pass_width must not be negative");
1165 if constexpr (ctx::runtime_sized) {
1166 ctx::dynamic_solver::template substitute_canonical_multirhs<pass_width>(
1167 mt::runtime_extent0(A), xt::columns(x), mt::pointer(A), mt::stride(A),
1168 pa::pointer(piv), pa::stride(piv), col, xt::pointer(x), xt::row_stride(x),
1171 constexpr int M = xt::extent1;
1172 static_assert(M >= 1,
"tdls adaptors: x must have at least one column");
1173 ctx::solver::template substitute_canonical_multirhs<M,
false, pa::is_internal,
1174 mt::is_internal, pass_width>(
1175 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), col,
1176 xt::pointer(x), xt::row_stride(x), xt::col_stride(x));
1179 static_assert(pass_width == 0,
1180 "tdls adaptors: pass_width only applies to matrix-like right-hand sides");
1181 detail::check_vector<SolutionType, typename ctx::scalar, ctx::N>();
1182 if constexpr (ctx::runtime_sized) {
1183 ctx::dynamic_solver::substitute_canonical(
1184 mt::runtime_extent0(A), mt::pointer(A), mt::stride(A), pa::pointer(piv),
1185 pa::stride(piv), col, xt::pointer(x), xt::stride(x));
1189 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), col,
1190 xt::pointer(x), xt::stride(x));
1204template<
int pass_w
idth = 0,
typename MatrixType,
typename PivotType,
typename SolutionType>
1207 substitute_canonical<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv,
TDLS_HOST_DEVICE TDLS_FORCEINLINE 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:1153
TDLS_HOST_DEVICE TDLS_FORCEINLINE 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:946
TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool factorize(MatrixType &A, PivotType &piv)
Factor a dense matrix object in place, A := P*L*U.
Definition adaptors.hpp:802
TDLS_HOST_DEVICE TDLS_FORCEINLINE 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:758
typename matrix_scalar_of< MatrixType >::type matrix_scalar
Scalar type of a dense matrix argument, see matrix_scalar_of.
Definition adaptors.hpp:398
TDLS_HOST_DEVICE TDLS_FORCEINLINE 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:708
TDLS_HOST_DEVICE TDLS_FORCEINLINE 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:493
TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr void check_vector()
Checks that a vector-like argument matches the system: arity 1, extent N, same scalar type....
Definition adaptors.hpp:475
consteval auto checked_config()
Validates a user configuration value against the matrix scalar type and returns it with its concrete ...
Definition adaptors.hpp:428
TDLS_HOST_DEVICE TDLS_FORCEINLINE 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:1087
TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool solve(MatrixType &A, PivotType &piv, const RhsType &b, SolutionType &x)
Solve A x = b on dense objects: factorize + substitute.
Definition adaptors.hpp:865
TDLS_HOST_DEVICE TDLS_FORCEINLINE 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:1024
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_FORCEINLINE 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:675
TDLS_HOST_DEVICE TDLS_FORCEINLINE 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:514
TDLS_HOST_DEVICE TDLS_FORCEINLINE 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:624
Concept gating the explicit-configuration entry points.
Definition adaptors.hpp:420
Toolchain detection and portability macros.
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:45
#define TDLS_FORCEINLINE
__forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS_HOST_DEVICE for why),...
Definition macros.hpp:67
Runtime-size variant of the TiledLUpp solver.
TiledLUpp solver: LU with partial pivoting on tile grids, one thread / work-item per system.
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:159
Tiled dense LU factorization with logical partial pivoting and out-of-tile pivot recovery,...
Definition solver_static.hpp:200
Common compile-time context of the adaptor entry points: resolves the scalar type,...
Definition adaptors.hpp:441
static constexpr int N
system dimension (fixed-size path; zero on the runtime path)
Definition adaptors.hpp:457
typename mtraits::value_type scalar
scalar type of the system
Definition adaptors.hpp:445
static constexpr bool runtime_sized
true when the matrix is runtime-sized (dynamic solver path)
Definition adaptors.hpp:450
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:377
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:405
Scalar type of a dense matrix argument, read from its storage description; a matrix argument without ...
Definition adaptors.hpp:387
typename storage_traits< std::remove_cv_t< MatrixType > >::value_type type
scalar type of the matrix
Definition adaptors.hpp:392
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:544
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int * pointer(const PivotType &p) noexcept
Definition adaptors.hpp:589
static constexpr bool is_internal
true when the pivot storage is compile-time contiguous
Definition adaptors.hpp:579
static constexpr bool is_raw
true for a raw int pointer or int array
Definition adaptors.hpp:546
static constexpr bool is_dense
true for an object with storage_traits
Definition adaptors.hpp:550
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:569
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:556
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int stride(const PivotType &p) noexcept
Definition adaptors.hpp:602
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int stride(const DenseType &o) noexcept
Definition adaptors.hpp:323
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int runtime_extent0(const DenseType &o) noexcept
Definition adaptors.hpp:257
typename DenseType::indexing_policy indexing_policy
indexing policy of the object
Definition adaptors.hpp:190
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_FORCEINLINE constexpr int runtime_row_stride(const DenseType &o) noexcept
Definition adaptors.hpp:276
static constexpr bool has_uniform_rows()
Definition adaptors.hpp:223
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int runtime_view_stride(const DenseType &o) noexcept
Definition adaptors.hpp:308
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int columns(const DenseType &o) noexcept
Definition adaptors.hpp:337
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int runtime_col_stride(const DenseType &o) noexcept
Definition adaptors.hpp:287
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int col_stride(const DenseType &o) noexcept
Definition adaptors.hpp:361
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_FORCEINLINE constexpr int runtime_extent1(const DenseType &o) noexcept
Definition adaptors.hpp:266
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr value_type * pointer(const DenseType &o) noexcept
Definition adaptors.hpp:296
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int row_stride(const DenseType &o) noexcept
Definition adaptors.hpp:349
Storage description of a dense object: element type, extents, element pointer and element stride.
Definition adaptors.hpp:181