TDLS 0.2.0
Tiny Device-callable Linear Solvers
Loading...
Searching...
No Matches
adaptors.hpp
Go to the documentation of this file.
1#ifndef TDLS_TFEL_ADAPTORS_HPP
2#define TDLS_TFEL_ADAPTORS_HPP
3
4
5
75
76
77
78#include <cstddef>
79#include <type_traits>
80
81#include <tdls/core/macros.hpp>
84
85
86
87namespace tdls {
88
89
90
91namespace detail {
92
94template<typename T, typename = void>
95struct has_data : std::false_type {};
96template<typename T>
97struct has_data<T, std::void_t<decltype(std::declval<const T&>().data())>>
98 : std::is_pointer<decltype(std::declval<const T&>().data())> {};
99
103template<typename T, typename = void>
104struct has_pair_data : std::false_type {};
105template<typename T>
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)>> {};
110
112template<typename T, bool = has_pair_data<T>::value>
115 using type = decltype(std::declval<const T&>().data());
116};
119template<typename T>
120struct const_data_pointer<T, true> {
122 using type = decltype(std::declval<const T&>().data().first);
123};
124
126template<typename T, bool = has_pair_data<T>::value>
129 using type = decltype(std::declval<T&>().data());
130};
133template<typename T>
134struct mutable_data_pointer<T, true> {
136 using type = decltype(std::declval<T&>().data().first);
137};
138
140template<typename T, typename = void>
141struct has_stride : std::false_type {};
142template<typename T>
143struct has_stride<T, std::void_t<decltype(std::declval<const T&>().stride())>> : std::true_type {};
144
146template<typename T, typename = void>
147struct has_get_stride : std::false_type {};
148template<typename T>
149struct has_get_stride<T, std::void_t<decltype(std::declval<const T&>().getStride())>>
150 : std::true_type {};
151
153template<typename T, typename = void>
154struct has_indexing_policy : std::false_type {};
155template<typename T>
156struct has_indexing_policy<T, std::void_t<typename T::indexing_policy>> : std::true_type {};
157
160template<typename T>
161inline constexpr bool is_dense_v =
163
164} // namespace detail
165
166
167
180template<typename DenseType, typename = void>
182
187template<typename DenseType>
188struct storage_traits<DenseType, std::enable_if_t<detail::is_dense_v<DenseType>>> {
190 using indexing_policy = typename DenseType::indexing_policy;
192 using policy_size_type = typename indexing_policy::size_type;
194 using value_type = std::remove_cv_t<
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");
204 static constexpr int extent0 = static_cast<int>(indexing_policy{}.size(0));
206 static constexpr int extent1 = (arity == 2) ? static_cast<int>(indexing_policy{}.size(1)) : 1;
207
210 static constexpr int compute_policy_stride() {
211 if constexpr (arity == 1) {
212 return static_cast<int>(indexing_policy{}.getIndex(policy_size_type(1)));
213 } else {
214 return static_cast<int>(
216 }
217 }
219 static constexpr int policy_stride = compute_policy_stride();
220
223 [[nodiscard]] static constexpr bool has_uniform_rows() {
224 if constexpr (arity == 1) {
225 return true;
226 } else {
227 return static_cast<int>(indexing_policy{}.getIndex(
228 policy_size_type(1), policy_size_type(0))) == extent1 * policy_stride;
229 }
230 }
231 // Single-stride addressing requires uniform rows. Row-strided matrix
232 // views (sub-matrix views) violate this and cannot be passed to the
233 // solvers; solve the full matrix or use an element-strided
234 // (coalesced) view instead.
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");
238
240 static constexpr bool has_runtime_stride = detail::has_pair_data<DenseType>::value ||
245 static constexpr bool is_internal = (policy_stride == 1) && !has_runtime_stride;
246
250 static constexpr bool has_runtime_extents = (extent0 == 0);
251
256 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int
257 runtime_extent0(const DenseType& o) noexcept {
258 return static_cast<int>(o.getIndexingPolicy().size(0));
259 }
260
265 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int
266 runtime_extent1(const DenseType& o) noexcept {
267 return static_cast<int>(o.getIndexingPolicy().size(1));
268 }
269
275 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int
276 runtime_row_stride(const DenseType& o) noexcept {
277 return static_cast<int>(
278 o.getIndexingPolicy().getIndex(policy_size_type(1), policy_size_type(0)));
279 }
280
286 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int
287 runtime_col_stride(const DenseType& o) noexcept {
288 return static_cast<int>(
289 o.getIndexingPolicy().getIndex(policy_size_type(0), policy_size_type(1)));
290 }
291
295 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr value_type*
296 pointer(const DenseType& o) noexcept {
298 return const_cast<value_type*>(static_cast<const value_type*>(o.data().first));
299 } else {
300 return const_cast<value_type*>(static_cast<const value_type*>(o.data()));
301 }
302 }
303
307 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int
308 runtime_view_stride(const DenseType& o) noexcept {
310 return static_cast<int>(o.data().second);
311 } else if constexpr (detail::has_stride<DenseType>::value) {
312 return static_cast<int>(o.stride());
313 } else if constexpr (detail::has_get_stride<DenseType>::value) {
314 return static_cast<int>(o.getStride());
315 } else {
316 return 1;
317 }
318 }
319
323 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int stride(const DenseType& o) noexcept {
324 return policy_stride * runtime_view_stride(o);
325 }
326
327 // The three accessors below describe a matrix-like block as the
328 // _multirhs entry points of the raw API see it: a column count, a
329 // total distance between row-consecutive elements and a total
330 // distance between column-consecutive elements. Each reads the
331 // policy type when the extents are compile-time and the policy
332 // instance of the object otherwise, so a fixed-size block needs no
333 // getIndexingPolicy() member, as the structural contract states.
334
337 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int columns(const DenseType& o) noexcept {
338 if constexpr (has_runtime_extents) {
339 return runtime_extent1(o);
340 } else {
341 return extent1;
342 }
343 }
344
349 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int row_stride(const DenseType& o) noexcept {
350 if constexpr (has_runtime_extents) {
351 return runtime_row_stride(o) * runtime_view_stride(o);
352 } else {
353 return extent1 * stride(o);
354 }
355 }
356
361 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int col_stride(const DenseType& o) noexcept {
362 if constexpr (has_runtime_extents) {
363 return runtime_col_stride(o) * runtime_view_stride(o);
364 } else {
365 return stride(o);
366 }
367 }
368};
369
370
371
372namespace detail {
373
376template<typename T, typename = void>
377struct has_storage_traits : std::false_type {};
378template<typename T>
379struct has_storage_traits<T, std::void_t<typename storage_traits<T>::value_type>> : std::true_type {
380};
381
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)");
393};
394
397template<typename MatrixType>
399
404template<typename ConfigType>
405struct is_solver_config : std::false_type {};
408template<typename Scalar>
409struct is_solver_config<TiledLUppConfig<Scalar>> : std::true_type {};
410
419template<typename ConfigType>
420concept solver_config = is_solver_config<std::remove_cvref_t<ConfigType>>::value;
421
427template<typename MatrixType, auto UserConfig>
428consteval auto checked_config() {
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");
432 return UserConfig;
433}
434
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");
450 static constexpr bool runtime_sized = mtraits::has_runtime_extents;
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;
464};
465
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");
482}
483
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");
502}
503
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)");
535}
536
543template<typename PivotType, int N>
546 static constexpr bool is_raw =
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");
553
556 static constexpr bool has_matching_shape = [] {
557 if constexpr (is_raw) {
558 return true;
559 } else {
561 return pt::arity == 1 && (N == 0 || pt::has_runtime_extents || pt::extent0 == N);
562 }
563 }();
564 static_assert(has_matching_shape, "tdls adaptors: the pivot must be vector-like, of the "
565 "system dimension when both are fixed-size");
566
569 static constexpr bool is_mutable = [] {
570 if constexpr (is_raw) {
571 return !std::is_const_v<std::remove_pointer_t<std::decay_t<PivotType>>>;
572 } else {
573 return !std::is_const_v<PivotType> &&
575 }
576 }();
577
579 static constexpr bool is_internal = [] {
580 if constexpr (is_raw) {
581 return true;
582 } else {
584 }
585 }();
586
589 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int* pointer(const PivotType& p) noexcept {
590 if constexpr (is_raw) {
591 return const_cast<int*>(static_cast<const int*>(p));
592 } else {
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);
597 }
598 }
599
602 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int stride(const PivotType& p) noexcept {
603 if constexpr (is_raw) {
604 return 1;
605 } else {
607 }
608 }
609};
610
621template<typename Ctx, int pass_width, bool inplace, typename MatrixType, typename PivotType,
622 typename RhsType, typename SolutionType>
624substitute_multirhs_dispatch(const MatrixType& A, const PivotType& piv, const RhsType& b,
625 SolutionType& x) {
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);
640 } else {
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);
644 }
645 } else {
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),
654 rs, xcol);
655 } else {
656 Ctx::solver::template substitute_multirhs<M, false, pa::is_internal, mt::is_internal,
657 pass_width>(
658 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), bt::pointer(b),
659 xt::pointer(x), rs, xcol);
660 }
661 }
662}
663
673template<auto UserConfig, bool oot_diagnostics, typename MatrixType, typename PivotType>
674[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
675factorize_dispatch(MatrixType& A, PivotType& piv, int& oot_count) {
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);
688 } else {
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);
691 }
692}
693
705template<auto UserConfig, int pass_width, bool oot_diagnostics, typename MatrixType,
706 typename PivotType, typename RhsType, typename SolutionType>
707[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
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);
724 return true;
725 } else {
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);
735 } else {
736 return ctx::solver::template solve<xt::is_internal, pa::is_internal, mt::is_internal,
737 oot_diagnostics>(
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);
740 }
741 }
742}
743
755template<auto UserConfig, int pass_width, bool oot_diagnostics, typename MatrixType,
756 typename PivotType, typename VectorType>
757[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
758solve_inplace_dispatch(MatrixType& A, PivotType& piv, VectorType& y, int& oot_count) {
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);
773 return true;
774 } else {
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);
782 } else {
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);
787 }
788 }
789}
790
791} // namespace detail
792
793
794
801template<detail::solver_config auto UserConfig, typename MatrixType, typename PivotType>
802[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool factorize(MatrixType& A,
803 PivotType& piv) {
804 int unused = 0;
805 return detail::factorize_dispatch<UserConfig, false>(A, piv, unused);
806}
807
812template<typename MatrixType, typename PivotType>
813[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool factorize(MatrixType& A,
814 PivotType& piv) {
815 return factorize<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}>(A, piv);
816}
817
827template<detail::solver_config auto UserConfig, typename MatrixType, typename PivotType>
828[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
829factorize(MatrixType& A, PivotType& piv, int& oot_count) {
830 return detail::factorize_dispatch<UserConfig, true>(A, piv, oot_count);
831}
832
839template<typename MatrixType, typename PivotType>
840[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
841factorize(MatrixType& A, PivotType& piv, int& oot_count) {
842 return factorize<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}>(A, piv, oot_count);
843}
844
862template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
863 typename PivotType, typename RhsType, typename SolutionType>
864[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
865solve(MatrixType& A, PivotType& piv, const RhsType& b, SolutionType& x) {
866 int unused = 0;
867 return detail::solve_dispatch<UserConfig, pass_width, false>(A, piv, b, x, unused);
868}
869
879template<int pass_width = 0, typename MatrixType, typename PivotType, typename RhsType,
880 typename SolutionType>
881[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
882solve(MatrixType& A, PivotType& piv, const RhsType& b, SolutionType& x) {
883 return solve<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, b, x);
884}
885
899template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
900 typename PivotType, typename RhsType, typename SolutionType>
901[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
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);
904}
905
916template<int pass_width = 0, typename MatrixType, typename PivotType, typename RhsType,
917 typename SolutionType>
918[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
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,
921 oot_count);
922}
923
943template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
944 typename PivotType, typename VectorType>
945[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
946solve_inplace(MatrixType& A, PivotType& piv, VectorType& y) {
947 int unused = 0;
948 return detail::solve_inplace_dispatch<UserConfig, pass_width, false>(A, piv, y, unused);
949}
950
961template<int pass_width = 0, typename MatrixType, typename PivotType, typename VectorType>
962[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
963solve_inplace(MatrixType& A, PivotType& piv, VectorType& y) {
964 return solve_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv,
965 y);
966}
967
981template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
982 typename PivotType, typename VectorType>
983[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
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);
986}
987
998template<int pass_width = 0, typename MatrixType, typename PivotType, typename VectorType>
999[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE constexpr bool
1000solve_inplace(MatrixType& A, PivotType& piv, VectorType& y, int& oot_count) {
1001 return solve_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(
1002 A, piv, y, oot_count);
1003}
1004
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) {
1025 using ctx =
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);
1037 } else {
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));
1047 } else {
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));
1051 }
1052 }
1053}
1054
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);
1068}
1069
1084template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
1085 typename PivotType, typename SolutionType>
1087substitute_inplace(const MatrixType& A, const PivotType& piv, SolutionType& x) {
1088 using ctx =
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);
1099 } else {
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));
1107 } else {
1108 ctx::solver::template substitute_inplace<xt::is_internal, pa::is_internal,
1109 mt::is_internal>(
1110 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), xt::pointer(x),
1111 xt::stride(x));
1112 }
1113 }
1114}
1115
1123template<int pass_width = 0, typename MatrixType, typename PivotType, typename SolutionType>
1125substitute_inplace(const MatrixType& A, const PivotType& piv, SolutionType& x) {
1126 substitute_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, x);
1127}
1128
1150template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
1151 typename PivotType, typename SolutionType>
1153substitute_canonical(const MatrixType& A, const PivotType& piv, const int col, SolutionType& x) {
1154 using ctx =
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),
1169 xt::col_stride(x));
1170 } else {
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));
1177 }
1178 } else {
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));
1186 } else {
1187 ctx::solver::template substitute_canonical<xt::is_internal, pa::is_internal,
1188 mt::is_internal>(
1189 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), col,
1190 xt::pointer(x), xt::stride(x));
1191 }
1192 }
1193}
1194
1204template<int pass_width = 0, typename MatrixType, typename PivotType, typename SolutionType>
1206substitute_canonical(const MatrixType& A, const PivotType& piv, const int col, SolutionType& x) {
1207 substitute_canonical<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv,
1208 col, x);
1209}
1210
1211
1212
1213} // namespace tdls
1214
1215
1216
1217#endif // TDLS_TFEL_ADAPTORS_HPP
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
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
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