TDLS 0.4.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_FORCE_INLINE_GPU 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_FORCE_INLINE_GPU 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_FORCE_INLINE_GPU 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_FORCE_INLINE_GPU 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_FORCE_INLINE_GPU 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_FORCE_INLINE_GPU 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_FORCE_INLINE_GPU static constexpr int
324 stride(const DenseType& o) noexcept {
325 return policy_stride * runtime_view_stride(o);
326 }
327
328 // The three accessors below describe a matrix-like block as the
329 // _multirhs entry points of the raw API see it: a column count, a
330 // total distance between row-consecutive elements and a total
331 // distance between column-consecutive elements. Each reads the
332 // policy type when the extents are compile-time and the policy
333 // instance of the object otherwise, so a fixed-size block needs no
334 // getIndexingPolicy() member, as the structural contract states.
335
338 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
339 columns(const DenseType& o) noexcept {
340 if constexpr (has_runtime_extents) {
341 return runtime_extent1(o);
342 } else {
343 return extent1;
344 }
345 }
346
351 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
352 row_stride(const DenseType& o) noexcept {
353 if constexpr (has_runtime_extents) {
354 return runtime_row_stride(o) * runtime_view_stride(o);
355 } else {
356 return extent1 * stride(o);
357 }
358 }
359
364 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr std::size_t
365 col_stride(const DenseType& o) noexcept {
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));
369 } else {
370 return static_cast<std::size_t>(stride(o));
371 }
372 }
373};
374
375
376
377namespace detail {
378
381template<typename T, typename = void>
382struct has_storage_traits : std::false_type {};
383template<typename T>
384struct has_storage_traits<T, std::void_t<typename storage_traits<T>::value_type>> : std::true_type {
385};
386
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)");
398};
399
402template<typename MatrixType>
404
409template<typename ConfigType>
410struct is_solver_config : std::false_type {};
413template<typename Scalar>
414struct is_solver_config<TiledLUppConfig<Scalar>> : std::true_type {};
415
424template<typename ConfigType>
425concept solver_config = is_solver_config<std::remove_cvref_t<ConfigType>>::value;
426
432template<typename MatrixType, auto UserConfig>
433consteval auto checked_config() {
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");
437 return UserConfig;
438}
439
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");
455 static constexpr bool runtime_sized = mtraits::has_runtime_extents;
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;
469};
470
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");
487}
488
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");
507}
508
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)");
540}
541
548template<typename PivotType, int N>
551 static constexpr bool is_raw =
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");
558
561 static constexpr bool has_matching_shape = [] {
562 if constexpr (is_raw) {
563 return true;
564 } else {
566 return pt::arity == 1 && (N == 0 || pt::has_runtime_extents || pt::extent0 == N);
567 }
568 }();
569 static_assert(has_matching_shape, "tdls adaptors: the pivot must be vector-like, of the "
570 "system dimension when both are fixed-size");
571
574 static constexpr bool is_mutable = [] {
575 if constexpr (is_raw) {
576 return !std::is_const_v<std::remove_pointer_t<std::decay_t<PivotType>>>;
577 } else {
578 return !std::is_const_v<PivotType> &&
580 }
581 }();
582
584 static constexpr bool is_internal = [] {
585 if constexpr (is_raw) {
586 return true;
587 } else {
589 }
590 }();
591
594 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int*
595 pointer(const PivotType& p) noexcept {
596 if constexpr (is_raw) {
597 return const_cast<int*>(static_cast<const int*>(p));
598 } else {
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);
603 }
604 }
605
608 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
609 stride(const PivotType& p) noexcept {
610 if constexpr (is_raw) {
611 return 1;
612 } else {
614 }
615 }
616};
617
634template<typename Ctx, int pass_width, bool inplace, typename MatrixType, typename PivotType,
635 typename RhsType, typename SolutionType>
637substitute_multirhs_dispatch(const MatrixType& A, const PivotType& piv, const RhsType& b,
638 SolutionType& x) {
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);
653 } else {
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);
657 }
658 } else {
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),
667 rs, xcol);
668 } else {
669 Ctx::solver::template substitute_multirhs<M, false, pa::is_internal, mt::is_internal,
670 pass_width>(
671 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), bt::pointer(b),
672 xt::pointer(x), rs, xcol);
673 }
674 }
675}
676
686template<auto UserConfig, bool oot_diagnostics, typename MatrixType, typename PivotType>
687[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
688factorize_dispatch(MatrixType& A, PivotType& piv, int& oot_count) {
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);
701 } else {
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);
704 }
705}
706
718template<auto UserConfig, int pass_width, bool oot_diagnostics, typename MatrixType,
719 typename PivotType, typename RhsType, typename SolutionType>
720[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
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);
737 return true;
738 } else {
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);
748 } else {
749 return ctx::solver::template solve<xt::is_internal, pa::is_internal, mt::is_internal,
750 oot_diagnostics>(
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);
753 }
754 }
755}
756
768template<auto UserConfig, int pass_width, bool oot_diagnostics, typename MatrixType,
769 typename PivotType, typename VectorType>
770[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
771solve_inplace_dispatch(MatrixType& A, PivotType& piv, VectorType& y, int& oot_count) {
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);
786 return true;
787 } else {
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);
795 } else {
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);
800 }
801 }
802}
803
804} // namespace detail
805
806
807
814template<detail::solver_config auto UserConfig, typename MatrixType, typename PivotType>
815[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool factorize(MatrixType& A,
816 PivotType& piv) {
817 int unused = 0;
818 return detail::factorize_dispatch<UserConfig, false>(A, piv, unused);
819}
820
825template<typename MatrixType, typename PivotType>
826[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool factorize(MatrixType& A,
827 PivotType& piv) {
828 return factorize<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}>(A, piv);
829}
830
840template<detail::solver_config auto UserConfig, typename MatrixType, typename PivotType>
841[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
842factorize(MatrixType& A, PivotType& piv, int& oot_count) {
843 return detail::factorize_dispatch<UserConfig, true>(A, piv, oot_count);
844}
845
852template<typename MatrixType, typename PivotType>
853[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
854factorize(MatrixType& A, PivotType& piv, int& oot_count) {
855 return factorize<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}>(A, piv, oot_count);
856}
857
875template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
876 typename PivotType, typename RhsType, typename SolutionType>
877[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
878solve(MatrixType& A, PivotType& piv, const RhsType& b, SolutionType& x) {
879 int unused = 0;
880 return detail::solve_dispatch<UserConfig, pass_width, false>(A, piv, b, x, unused);
881}
882
892template<int pass_width = 0, typename MatrixType, typename PivotType, typename RhsType,
893 typename SolutionType>
894[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
895solve(MatrixType& A, PivotType& piv, const RhsType& b, SolutionType& x) {
896 return solve<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, b, x);
897}
898
912template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
913 typename PivotType, typename RhsType, typename SolutionType>
914[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
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);
917}
918
929template<int pass_width = 0, typename MatrixType, typename PivotType, typename RhsType,
930 typename SolutionType>
931[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
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,
934 oot_count);
935}
936
956template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
957 typename PivotType, typename VectorType>
958[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
959solve_inplace(MatrixType& A, PivotType& piv, VectorType& y) {
960 int unused = 0;
961 return detail::solve_inplace_dispatch<UserConfig, pass_width, false>(A, piv, y, unused);
962}
963
974template<int pass_width = 0, typename MatrixType, typename PivotType, typename VectorType>
975[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
976solve_inplace(MatrixType& A, PivotType& piv, VectorType& y) {
977 return solve_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv,
978 y);
979}
980
994template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
995 typename PivotType, typename VectorType>
996[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
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);
999}
1000
1011template<int pass_width = 0, typename MatrixType, typename PivotType, typename VectorType>
1012[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr bool
1013solve_inplace(MatrixType& A, PivotType& piv, VectorType& y, int& oot_count) {
1014 return solve_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(
1015 A, piv, y, oot_count);
1016}
1017
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) {
1038 using ctx =
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);
1050 } else {
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));
1060 } else {
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));
1064 }
1065 }
1066}
1067
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);
1081}
1082
1097template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
1098 typename PivotType, typename SolutionType>
1100substitute_inplace(const MatrixType& A, const PivotType& piv, SolutionType& x) {
1101 using ctx =
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);
1112 } else {
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));
1120 } else {
1121 ctx::solver::template substitute_inplace<xt::is_internal, pa::is_internal,
1122 mt::is_internal>(
1123 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), xt::pointer(x),
1124 xt::stride(x));
1125 }
1126 }
1127}
1128
1136template<int pass_width = 0, typename MatrixType, typename PivotType, typename SolutionType>
1138substitute_inplace(const MatrixType& A, const PivotType& piv, SolutionType& x) {
1139 substitute_inplace<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv, x);
1140}
1141
1163template<detail::solver_config auto UserConfig, int pass_width = 0, typename MatrixType,
1164 typename PivotType, typename SolutionType>
1166substitute_canonical(const MatrixType& A, const PivotType& piv, const int col, SolutionType& x) {
1167 using ctx =
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),
1182 xt::col_stride(x));
1183 } else {
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));
1190 }
1191 } else {
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));
1199 } else {
1200 ctx::solver::template substitute_canonical<xt::is_internal, pa::is_internal,
1201 mt::is_internal>(
1202 mt::pointer(A), mt::stride(A), pa::pointer(piv), pa::stride(piv), col,
1203 xt::pointer(x), xt::stride(x));
1204 }
1205 }
1206}
1207
1217template<int pass_width = 0, typename MatrixType, typename PivotType, typename SolutionType>
1219substitute_canonical(const MatrixType& A, const PivotType& piv, const int col, SolutionType& x) {
1220 substitute_canonical<TiledLUppConfig<detail::matrix_scalar<MatrixType>>{}, pass_width>(A, piv,
1221 col, x);
1222}
1223
1224
1225
1226} // namespace tdls
1227
1228
1229
1230#endif // TDLS_TFEL_ADAPTORS_HPP
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
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
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.