TDLS 0.4.0
Tiny Device-callable Linear Solvers
Loading...
Searching...
No Matches
solver_dynamic.hpp
Go to the documentation of this file.
1#ifndef TDLS_SOLVERS_ARROWHEAD_LUPP_SOLVER_DYNAMIC_HPP
2#define TDLS_SOLVERS_ARROWHEAD_LUPP_SOLVER_DYNAMIC_HPP
3
4
5
55
56
57
58#include <limits>
59
60#include <tdls/core/macros.hpp>
61#include <tdls/core/math.hpp>
65
66
67
68namespace tdls {
69
70
71
72/* Addressing macros. They expand inside member functions where the
73 parameter names (A, A_stride, y, rhs_stride, work, work_stride, piv,
74 piv_stride) are in scope. #undef'd at the end of this header. */
75
78#define TDLS_ARROW_LUPP_DYN_A(idx) A[TDLS_STRIDED(idx, A_stride)]
81#define TDLS_ARROW_LUPP_DYN_Y(i) y[TDLS_STRIDED(i, rhs_stride)]
87#define TDLS_ARROW_LUPP_DYN_YW(m, i) \
88 y[TDLS_STRIDED(i, rhs_stride) + \
89 (ycols == 1 ? std::size_t(0) : TDLS_STRIDED_WIDE(m, xcol_stride))]
92#define TDLS_ARROW_LUPP_DYN_W(e) work[TDLS_STRIDED(e, work_stride)]
96#define TDLS_ARROW_LUPP_DYN_F(r, c) TDLS_ARROW_LUPP_DYN_W((r) * front_cols + (c))
100#define TDLS_ARROW_LUPP_DYN_LU(i, j) TDLS_ARROW_LUPP_DYN_W(o.Dl + (i) * S + (j))
104#define TDLS_ARROW_LUPP_DYN_RC(c, i) TDLS_ARROW_LUPP_DYN_W(o.W + (c) * S + (i))
105
106
107
111template<typename T, ArrowheadLUppConfig<T> Config>
114 using storage_map = ArrowheadStorageMap<Config.storage, Config.coupling, Config.layout>;
115
116 private:
117 static constexpr bool col = Config.layout == MatrixLayout::ColMajor;
118 static constexpr bool diag = Config.coupling == ArrowheadCoupling::Diagonal;
122 static constexpr bool multiplier_update =
123 !col && !diag && Config.condensation == ArrowheadCondensation::MinFlops;
124 static constexpr bool in_place = Config.schur_in_place;
125
126 static_assert(Config.condensation_pivot_threshold.is_finite() &&
127 T(Config.condensation_pivot_threshold) >= T(0) &&
128 T(Config.condensation_pivot_threshold) <= T(1),
129 "ArrowheadLUppSolverDynamic: condensation_pivot_threshold must be in [0, 1]");
130
134 static constexpr T condensation_pivot_threshold = Config.condensation_pivot_threshold;
135 static_assert(Config.rhs_pass_width >= -1,
136 "ArrowheadLUppSolverDynamic: rhs_pass_width must be >= -1");
137
138 using BlockSolver = TiledLUppSolverDynamic<T, detail::arrowhead_block_config(Config)>;
139 using BorderSolver = TiledLUppSolverDynamic<T, detail::arrowhead_border_config(Config)>;
142 using PhysicalBlockSolver =
143 TiledLUppSolverDynamic<T, detail::arrowhead_physical_block_config(Config)>;
144
147 struct Sizes {
148 int B, S, NB, R, R0, front_rows, front_cols, extra_cols, pass_width;
151 bool reg;
152 };
155 template<typename I>
156 struct OffsetsOf {
157 I extra, Dl, U, W, X, Mb, m, a, acoef, Sacc, racc, Sbb, rb, w, s, refused, block_row, total;
158 };
160 using Offsets = OffsetsOf<int>;
163 template<typename I>
164 struct StreamedOffsetsOf {
165 I Abk, Dsave, Wsave, total;
166 };
168 using StreamedOffsets = StreamedOffsetsOf<int>;
169
176 [[nodiscard]] TDLS_HOST_DEVICE static constexpr Sizes
177 sizes(const ArrowheadShape& shape, const int nc, const int ny) noexcept {
178 Sizes z{};
179 z.B = shape.border;
180 z.S = shape.block;
181 z.NB = shape.blocks;
182 z.R = diag ? z.S : z.B;
183 z.R0 = diag ? shape.coupling_row : 0;
184 z.front_rows = z.S + z.B;
185 z.front_cols = z.S + z.B + z.R + nc;
186 z.extra_cols = z.R > z.S ? z.R - z.S : 0;
187 constexpr RowInterchange interchange = Config.block_row_interchange;
188 constexpr int width = Config.rhs_pass_width;
189 z.reg = detail::arrowhead_register_blocks_here<T>(interchange, z.S);
190 const int pass = detail::arrowhead_rhs_pass_width_here(width, z.reg);
191 z.pass_width = pass == 0 || pass > z.B + ny ? z.B + ny : pass;
192 return z;
193 }
194
203 template<typename I = int>
204 [[nodiscard]] TDLS_HOST_DEVICE static constexpr OffsetsOf<I>
205 offsets(const ArrowheadShape& shape, const int nc) noexcept {
206 const Sizes z = sizes(shape, nc, nc);
207 OffsetsOf<I> o{};
208 I e = I(z.front_rows) * z.front_cols;
209 o.extra = e;
210 e += I(z.NB) * z.S * z.extra_cols;
211 o.Dl = e;
212 e += I(z.S) * z.S;
213 o.U = e;
214 e += I(z.S) * z.S;
215 o.W = e;
216 e += (I(z.B) + nc) * z.S;
217 o.X = e;
218 e += (I(z.B) + nc) * z.S;
219 o.Mb = e;
220 e += I(z.R) * z.S;
221 o.m = e;
222 e += z.S;
223 o.a = e;
224 e += z.S;
225 o.acoef = e;
226 e += z.S;
227 o.Sacc = e;
228 e += I(z.R) * z.B;
229 o.racc = e;
230 e += I(z.R) * nc;
231 o.Sbb = e;
232 e += I(z.B) * z.B;
233 o.rb = e;
234 e += I(z.B) * nc;
235 o.w = e;
236 e += I(z.R) * nc;
237 o.s = e;
238 e += I(z.S) * nc;
239 o.refused = e;
240 e += z.NB;
241 o.block_row = e;
242 e += z.front_rows;
243 o.total = e;
244 return o;
245 }
246
255 template<typename I = int>
256 [[nodiscard]] TDLS_HOST_DEVICE static constexpr StreamedOffsetsOf<I>
257 streamed_offsets(const ArrowheadShape& shape, const int nc) noexcept {
258 const Sizes z = sizes(shape, nc, nc);
259 StreamedOffsetsOf<I> so{};
260 I e = offsets<I>(shape, nc).total;
261 so.Abk = e;
262 e += diag ? I(0) : I(z.B) * z.S;
263 so.Dsave = e;
264 e += I(z.S) * z.S;
265 so.Wsave = e;
266 e += (I(z.B) + 1) * z.S;
267 so.total = e;
268 return so;
269 }
270
271 public:
281 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool valid(const ArrowheadShape& shape,
282 const int nrhs = 1) noexcept {
283 return storage_map::valid(shape) && nrhs >= 1 && nrhs <= (1 << 30) &&
284 streamed_offsets<long long>(shape, nrhs).total < (1LL << 31);
285 }
286
291 [[nodiscard]] TDLS_HOST_DEVICE static constexpr int
292 workspace_size(const ArrowheadShape& shape) noexcept {
293 return valid(shape) ? offsets(shape, 1).total : 0;
294 }
295
303 [[nodiscard]] TDLS_HOST_DEVICE static constexpr int workspace_size(const ArrowheadShape& shape,
304 const int nrhs) noexcept {
305 return valid(shape, nrhs) ? offsets(shape, nrhs).total : 0;
306 }
307
314 [[nodiscard]] TDLS_HOST_DEVICE static constexpr int
316 return valid(shape) ? streamed_offsets(shape, 1).total : 0;
317 }
318
326 [[nodiscard]] TDLS_HOST_DEVICE static constexpr int
327 streamed_workspace_size(const ArrowheadShape& shape, const int nrhs) noexcept {
328 return valid(shape, nrhs) ? streamed_offsets(shape, nrhs).total : 0;
329 }
330
335 [[nodiscard]] TDLS_HOST_DEVICE static constexpr int
336 pivot_size(const ArrowheadShape& shape) noexcept {
337 if (!valid(shape)) return 0;
338 return shape.block > shape.border ? shape.block : shape.border;
339 }
340
357 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
358 solve_inplace(const ArrowheadShape& shape, T* TDLS_RESTRICT A, const int A_stride,
359 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT work,
360 const int work_stride, T* TDLS_RESTRICT y, const int rhs_stride,
361 int& refused_count) noexcept {
362 return solve_core<1, false>(shape, 1, A, A_stride, piv, piv_stride, work, work_stride, y,
363 rhs_stride, 0, 0, refused_count);
364 }
365
378 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
379 solve_inplace(const ArrowheadShape& shape, T* TDLS_RESTRICT A, const int A_stride,
380 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT work,
381 const int work_stride, T* TDLS_RESTRICT y, const int rhs_stride) noexcept {
382 int refused_count = 0;
383 return solve_inplace(shape, A, A_stride, piv, piv_stride, work, work_stride, y, rhs_stride,
384 refused_count);
385 }
386
405 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
406 solve(const ArrowheadShape& shape, T* TDLS_RESTRICT A, const int A_stride,
407 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT work,
408 const int work_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
409 int& refused_count) noexcept {
410 if (!valid(shape)) {
411 refused_count = 0;
412 return false;
413 }
414 const int n = shape.n();
415 for (int i = 0; i < n; ++i)
416 x[TDLS_STRIDED(i, rhs_stride)] = b[TDLS_STRIDED(i, rhs_stride)];
417 return solve_inplace(shape, A, A_stride, piv, piv_stride, work, work_stride, x, rhs_stride,
418 refused_count);
419 }
420
434 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
435 solve(const ArrowheadShape& shape, T* TDLS_RESTRICT A, const int A_stride,
436 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT work,
437 const int work_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
438 const int rhs_stride) noexcept {
439 int refused_count = 0;
440 return solve(shape, A, A_stride, piv, piv_stride, work, work_stride, b, x, rhs_stride,
441 refused_count);
442 }
443
469 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
470 solve_inplace_multirhs(const ArrowheadShape& shape, const int nrhs, T* TDLS_RESTRICT A,
471 const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
472 T* TDLS_RESTRICT work, const int work_stride, T* TDLS_RESTRICT y,
473 const int rhs_stride, const std::size_t xcol_stride,
474 int& refused_count) noexcept {
475 return solve_core<0, false>(shape, nrhs, A, A_stride, piv, piv_stride, work, work_stride, y,
476 rhs_stride, xcol_stride, 0, refused_count);
477 }
478
495 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
496 solve_inplace_multirhs(const ArrowheadShape& shape, const int nrhs, T* TDLS_RESTRICT A,
497 const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
498 T* TDLS_RESTRICT work, const int work_stride, T* TDLS_RESTRICT y,
499 const int rhs_stride, const std::size_t xcol_stride) noexcept {
500 int refused_count = 0;
501 return solve_inplace_multirhs(shape, nrhs, A, A_stride, piv, piv_stride, work, work_stride,
502 y, rhs_stride, xcol_stride, refused_count);
503 }
504
526 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
527 solve_multirhs(const ArrowheadShape& shape, const int nrhs, T* TDLS_RESTRICT A,
528 const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
529 T* TDLS_RESTRICT work, const int work_stride, const T* TDLS_RESTRICT b,
530 T* TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride,
531 int& refused_count) noexcept {
532 if (!valid(shape, nrhs)) {
533 refused_count = 0;
534 return false;
535 }
536 const int n = shape.n();
537 for (int m = 0; m < nrhs; ++m)
538 for (int i = 0; i < n; ++i)
539 x[TDLS_STRIDED(i, rhs_stride) + TDLS_STRIDED_WIDE(m, xcol_stride)] =
540 b[TDLS_STRIDED(i, rhs_stride) + TDLS_STRIDED_WIDE(m, xcol_stride)];
541 return solve_inplace_multirhs(shape, nrhs, A, A_stride, piv, piv_stride, work, work_stride,
542 x, rhs_stride, xcol_stride, refused_count);
543 }
544
561 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
562 solve_multirhs(const ArrowheadShape& shape, const int nrhs, T* TDLS_RESTRICT A,
563 const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
564 T* TDLS_RESTRICT work, const int work_stride, const T* TDLS_RESTRICT b,
565 T* TDLS_RESTRICT x, const int rhs_stride,
566 const std::size_t xcol_stride) noexcept {
567 int refused_count = 0;
568 return solve_multirhs(shape, nrhs, A, A_stride, piv, piv_stride, work, work_stride, b, x,
569 rhs_stride, xcol_stride, refused_count);
570 }
571
594 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool solve_border_canonical_multirhs(
595 const ArrowheadShape& shape, const int nrhs, T* TDLS_RESTRICT A, const int A_stride,
596 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT work, const int work_stride,
597 const int col0, T* TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride,
598 int& refused_count) noexcept {
599 return solve_core<0, true>(shape, nrhs, A, A_stride, piv, piv_stride, work, work_stride, x,
600 rhs_stride, xcol_stride, col0, refused_count);
601 }
602
621 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
623 const int A_stride, int* TDLS_RESTRICT piv,
624 const int piv_stride, T* TDLS_RESTRICT work,
625 const int work_stride, const int col0, T* TDLS_RESTRICT x,
626 const int rhs_stride, const std::size_t xcol_stride) noexcept {
627 int refused_count = 0;
628 return solve_border_canonical_multirhs(shape, nrhs, A, A_stride, piv, piv_stride, work,
629 work_stride, col0, x, rhs_stride, xcol_stride,
630 refused_count);
631 }
632
643 T* Dl;
644 T* W;
645 T* Abk;
646 int S;
647 int B;
648 int stride;
652 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL constexpr T& d(const int i,
653 const int j) const noexcept {
654 return Dl[TDLS_STRIDED(col ? j * S + i : i * S + j, stride)];
655 }
660 const int c) const noexcept {
661 return W[TDLS_STRIDED(c * S + i, stride)];
662 }
666 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL constexpr T& y(const int i) const noexcept {
667 return W[TDLS_STRIDED(B * S + i, stride)];
668 }
671 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL constexpr T& coupling(const int i) const noexcept {
672 return Abk[TDLS_STRIDED(i, stride)];
673 }
678 const int i) const noexcept {
679 return Abk[TDLS_STRIDED(r * S + i, stride)];
680 }
681 };
682
694 bool recovery = false;
695 };
696
698 enum class StreamedOutcome {
699 solved,
700 tangent,
702 singular
703 };
704
731 template<typename Provider>
732 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
733 solve_inplace_streamed(const ArrowheadShape& shape, Provider&& provider, T* TDLS_RESTRICT A,
734 const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
735 T* TDLS_RESTRICT work, const int work_stride, T* TDLS_RESTRICT y,
736 const int rhs_stride, int& refused_count) noexcept {
737 return solve_core<1, false>(shape, 1, A, A_stride, piv, piv_stride, work, work_stride, y,
738 rhs_stride, 0, 0, refused_count, &provider);
739 }
740
765 template<typename Provider>
767 const ArrowheadShape& shape, const int nrhs, Provider&& provider, T* TDLS_RESTRICT A,
768 const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT work,
769 const int work_stride, const int col0, T* TDLS_RESTRICT x, const int rhs_stride,
770 const std::size_t xcol_stride, int& refused_count) noexcept {
771 return solve_core<0, true>(shape, nrhs, A, A_stride, piv, piv_stride, work, work_stride, x,
772 rhs_stride, xcol_stride, col0, refused_count, &provider);
773 }
774
818 template<typename Provider, typename Converged>
819 [[nodiscard]] TDLS_HOST_DEVICE static constexpr StreamedOutcome
821 Provider&& provider, Converged&& converged,
822 T* TDLS_RESTRICT A, const int A_stride,
823 int* TDLS_RESTRICT piv, const int piv_stride,
824 T* TDLS_RESTRICT work, const int work_stride,
825 T* TDLS_RESTRICT y, const int rhs_stride, const int col0,
826 T* TDLS_RESTRICT x, const int x_stride,
827 const std::size_t xcol_stride, int& refused_count) noexcept {
828 if (!valid(shape, ntan)) {
829 refused_count = 0;
831 }
832 // The pass carries one right-hand side through the blocks (y, with
833 // one column); the workspace holds both phases, ntan columns in the
834 // front and in the border right-hand sides.
835 [[maybe_unused]] constexpr int ycols = 1;
836 constexpr int y_xcol_stride = 0;
837 const Sizes z = sizes(shape, 1, 1);
838 const Offsets o = offsets(shape, ntan);
839 const StreamedOffsets so = streamed_offsets(shape, ntan);
840 const int B = z.B, R = z.R, R0 = z.R0, NB = z.NB;
841 if constexpr (!in_place) {
842 for (int r = 0; r < R; ++r) {
843 for (int c = 0; c < B; ++c)
844 TDLS_ARROW_LUPP_DYN_W(o.Sacc + r * B + c) =
845 TDLS_ARROW_LUPP_DYN_A(storage_map::abb(shape, R0 + r, c));
846 TDLS_ARROW_LUPP_DYN_W(o.racc + r) = TDLS_ARROW_LUPP_DYN_Y(R0 + r);
847 }
848 }
849 refused_count = 0;
850 for (int k = 0; k < NB; ++k) {
851 StreamedBlock blk = streamed_block(z, o, so, work, work_stride);
852 provider(k, blk);
853 const bool ok = condense_or_store_streamed<1, false>(shape, z, o, so, A, A_stride, piv,
854 piv_stride, work, work_stride, y,
855 rhs_stride, y_xcol_stride, k, 1);
856 TDLS_ARROW_LUPP_DYN_W(o.refused + k) = ok ? T(0) : T(1);
857 if (!ok) ++refused_count;
858 }
859 // The rows as stored and the accumulated rows in separate loops, as
860 // in solve_core.
861 const int acc0 = in_place ? B : R0;
862 const int acc1 = in_place ? B : R0 + R;
863 for (int r = 0; r < acc0; ++r)
864 for (int c = 0; c < B; ++c)
865 TDLS_ARROW_LUPP_DYN_W(o.Sbb + r * B + c) =
867 for (int r = acc0; r < acc1; ++r)
868 for (int c = 0; c < B; ++c)
869 TDLS_ARROW_LUPP_DYN_W(o.Sbb + r * B + c) =
870 TDLS_ARROW_LUPP_DYN_W(o.Sacc + (r - R0) * B + c);
871 for (int r = acc1; r < B; ++r)
872 for (int c = 0; c < B; ++c)
873 TDLS_ARROW_LUPP_DYN_W(o.Sbb + r * B + c) =
875 T* const Sbb = work + TDLS_STRIDED(o.Sbb, work_stride);
876 T* const rb = work + TDLS_STRIDED(o.rb, work_stride);
877 if (converged()) {
878 const Sizes zt = sizes(shape, ntan, 0);
879 for (int m = 0; m < ntan; ++m)
880 for (int r = 0; r < B; ++r)
881 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r) = r == col0 + m ? T(1) : T(0);
882 // The refused blocks, stored as given, by the second pass of the
883 // canonical solve, which writes neither x nor y.
884 if (refused_count > 0 &&
885 !second_pass<0, true>(shape, zt, o, A, A_stride, work, work_stride, y, rhs_stride,
886 y_xcol_stride, ntan))
888 const bool ok = ntan == 1 ? BorderSolver::solve_inplace(B, Sbb, work_stride, piv,
889 piv_stride, rb, work_stride)
891 B, ntan, Sbb, work_stride, piv, piv_stride, rb,
892 work_stride, TDLS_STRIDED(B, work_stride));
893 if (!ok) return StreamedOutcome::singular;
894 for (int m = 0; m < ntan; ++m)
895 for (int r = 0; r < B; ++r)
896 x[TDLS_STRIDED(r, x_stride) + TDLS_STRIDED_WIDE(m, xcol_stride)] =
897 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r);
899 }
900 for (int r = 0; r < acc0; ++r)
902 for (int r = acc0; r < acc1; ++r)
903 TDLS_ARROW_LUPP_DYN_W(o.rb + r) = TDLS_ARROW_LUPP_DYN_W(o.racc + r - R0);
904 for (int r = acc1; r < B; ++r)
906 if (refused_count > 0 && !second_pass<1, false>(shape, z, o, A, A_stride, work, work_stride,
907 y, rhs_stride, y_xcol_stride, 1))
909 if (!BorderSolver::solve_inplace(B, Sbb, work_stride, piv, piv_stride, rb, work_stride))
911 for (int r = 0; r < B; ++r)
913 recover<1>(shape, z, o, A, A_stride, work, work_stride, y, rhs_stride, y_xcol_stride, 1);
915 }
916
919 enum class RecomputedOutcome {
920 solved,
921 tangent,
923 singular,
925 deferred
928 };
929
939 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
940 valid_recomputed(const ArrowheadShape& shape, const int capacity, const int nrhs = 1) noexcept {
941 constexpr int digits =
942 std::numeric_limits<T>::digits < 62 ? std::numeric_limits<T>::digits : 62;
943 return valid(shape, nrhs) && capacity >= 0 && capacity <= shape.blocks &&
944 static_cast<long long>(shape.blocks) <= (1LL << digits) &&
945 recomputed_offsets<long long>(shape, capacity, nrhs).total < (1LL << 31);
946 }
947
960 [[nodiscard]] TDLS_HOST_DEVICE static constexpr int
961 recomputed_workspace_size(const ArrowheadShape& shape, const int capacity,
962 const int nrhs = 1) noexcept {
963 return valid_recomputed(shape, capacity, nrhs)
964 ? recomputed_offsets(shape, capacity, nrhs).total
965 : 0;
966 }
967
1002 template<typename Provider>
1003 [[nodiscard]] TDLS_HOST_DEVICE static constexpr RecomputedOutcome
1004 solve_inplace_recomputed(const ArrowheadShape& shape, const int capacity, Provider&& provider,
1005 const T* TDLS_RESTRICT Abb, const int Abb_stride,
1006 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT work,
1007 const int work_stride, T* TDLS_RESTRICT y, const int rhs_stride,
1008 int& refused_count) noexcept {
1009 return recomputed_core(shape, capacity, 0, provider, Abb, Abb_stride, piv, piv_stride, work,
1010 work_stride, y, rhs_stride, 0, nullptr, 1, 0, refused_count);
1011 }
1012
1054 template<typename Provider, typename Converged>
1055 [[nodiscard]] TDLS_HOST_DEVICE static constexpr RecomputedOutcome
1057 const ArrowheadShape& shape, const int ntan, const int capacity, Provider&& provider,
1058 Converged&& converged, const T* TDLS_RESTRICT Abb, const int Abb_stride,
1059 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT work, const int work_stride,
1060 T* TDLS_RESTRICT y, const int rhs_stride, const int col0, T* TDLS_RESTRICT x,
1061 const int x_stride, const std::size_t xcol_stride, int& refused_count) noexcept {
1062 if (ntan < 1) {
1063 refused_count = 0;
1065 }
1066 return recomputed_core(shape, capacity, ntan, provider, Abb, Abb_stride, piv, piv_stride,
1067 work, work_stride, y, rhs_stride, col0, x, x_stride, xcol_stride,
1068 refused_count, &converged);
1069 }
1070
1071 private:
1076 template<typename I>
1077 struct RecomputedOffsetsOf {
1078 I store, ystore, slots, total;
1079 };
1081 using RecomputedOffsets = RecomputedOffsetsOf<int>;
1082
1087 [[nodiscard]] TDLS_HOST_DEVICE static constexpr ArrowheadShape
1088 store_shape(const ArrowheadShape& shape, const int capacity) noexcept {
1089 return {shape.border, shape.block, capacity, shape.coupling_row};
1090 }
1091
1100 template<typename I = int>
1101 [[nodiscard]] TDLS_HOST_DEVICE static constexpr RecomputedOffsetsOf<I>
1102 recomputed_offsets(const ArrowheadShape& shape, const int capacity, const int nc) noexcept {
1103 const ArrowheadShape st = store_shape(shape, capacity);
1104 RecomputedOffsetsOf<I> ro{};
1105 I e = streamed_offsets<I>(st, nc).total;
1106 ro.store = e;
1107 e += capacity > 0 ? I(storage_map::size(st)) : I(0);
1108 ro.ystore = e;
1109 e += I(shape.border) + I(capacity) * shape.block;
1110 ro.slots = e;
1111 e += capacity;
1112 ro.total = e;
1113 return ro;
1114 }
1115
1124 [[nodiscard]] TDLS_HOST_DEVICE static constexpr T border_entry(const T* TDLS_RESTRICT Abb,
1125 const int Abb_stride,
1126 const int B, const int r,
1127 const int c) noexcept {
1128 return Abb[TDLS_STRIDED(col ? c * B + r : r * B + c, Abb_stride)];
1129 }
1130
1159 template<typename Provider, typename... Converged>
1160 TDLS_HOST_DEVICE static constexpr RecomputedOutcome recomputed_core(
1161 const ArrowheadShape& shape, const int capacity, const int ntan, Provider& provider,
1162 const T* TDLS_RESTRICT Abb, const int Abb_stride, int* TDLS_RESTRICT piv,
1163 const int piv_stride, T* TDLS_RESTRICT work, const int work_stride, T* TDLS_RESTRICT y,
1164 const int rhs_stride, [[maybe_unused]] const int col0, [[maybe_unused]] T* TDLS_RESTRICT x,
1165 [[maybe_unused]] const int x_stride, [[maybe_unused]] const std::size_t xcol_stride,
1166 int& refused_count, Converged*... converged) noexcept {
1167 static_assert(!in_place, "ArrowheadLUppSolverDynamic: the recomputed solve accumulates "
1168 "the Schur complement apart (schur_in_place off)");
1169 // Columns of the front: the tangent columns, one at least.
1170 const int fnc = ntan > 1 ? ntan : 1;
1171 refused_count = 0;
1172 if (!valid_recomputed(shape, capacity, fnc) || ntan > shape.border)
1174 [[maybe_unused]] constexpr int ycols = 1;
1175 const ArrowheadShape st = store_shape(shape, capacity);
1176 const Sizes z = sizes(shape, 1, 1);
1177 const Offsets o = offsets(st, fnc);
1178 const StreamedOffsets so = streamed_offsets(st, fnc);
1179 const RecomputedOffsets ro = recomputed_offsets(shape, capacity, fnc);
1180 const int B = z.B, S = z.S, R = z.R, R0 = z.R0, NB = z.NB;
1181 for (int r = 0; r < R; ++r) {
1182 for (int c = 0; c < B; ++c)
1183 TDLS_ARROW_LUPP_DYN_W(o.Sacc + r * B + c) =
1184 border_entry(Abb, Abb_stride, B, R0 + r, c);
1185 TDLS_ARROW_LUPP_DYN_W(o.racc + r) = TDLS_ARROW_LUPP_DYN_Y(R0 + r);
1186 }
1187 int count = 0;
1188 for (int k = 0; k < NB; ++k) {
1189 RecomputedBlock blk{streamed_block(z, o, so, work, work_stride), false};
1190 provider(k, blk);
1191 const int verdict = condense_or_keep_recomputed(
1192 shape, z, o, so, ro, capacity, piv, piv_stride, work, work_stride, k, count);
1193 refused_count = count;
1194 if (verdict == 2) return RecomputedOutcome::deferred;
1195 }
1196 // The flags of the places of the store, for its second pass.
1197 for (int j = 0; j < capacity; ++j)
1198 TDLS_ARROW_LUPP_DYN_W(o.refused + j) = j < count ? T(1) : T(0);
1199 T* const store = work + TDLS_STRIDED(ro.store, work_stride);
1200 T* const ystore = work + TDLS_STRIDED(ro.ystore, work_stride);
1201 for (int r = 0; r < B; ++r) {
1202 const bool acc = r >= R0 && r < R0 + R;
1203 for (int c = 0; c < B; ++c)
1204 TDLS_ARROW_LUPP_DYN_W(o.Sbb + r * B + c) =
1205 acc ? TDLS_ARROW_LUPP_DYN_W(o.Sacc + (r - R0) * B + c)
1206 : border_entry(Abb, Abb_stride, B, r, c);
1207 }
1208 T* const Sbb = work + TDLS_STRIDED(o.Sbb, work_stride);
1209 T* const rb = work + TDLS_STRIDED(o.rb, work_stride);
1210 if constexpr (sizeof...(Converged) > 0) {
1211 if ((*streamed_arg(converged...))()) {
1212 // The tangent: unit columns, the refused blocks of the store
1213 // by the second pass of the canonical solve.
1214 const Sizes zt = sizes(st, fnc, 0);
1215 for (int m = 0; m < fnc; ++m)
1216 for (int r = 0; r < B; ++r)
1217 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r) = r == col0 + m ? T(1) : T(0);
1218 if (count > 0 && !second_pass<0, true>(st, zt, o, store, work_stride, work,
1219 work_stride, ystore, work_stride, 0, fnc))
1221 const bool ok = fnc == 1 ? BorderSolver::solve_inplace(B, Sbb, work_stride, piv,
1222 piv_stride, rb, work_stride)
1223 : BorderSolver::solve_inplace_multirhs(
1224 B, fnc, Sbb, work_stride, piv, piv_stride, rb,
1225 work_stride, TDLS_STRIDED(B, work_stride));
1226 if (!ok) return RecomputedOutcome::singular;
1227 for (int m = 0; m < fnc; ++m)
1228 for (int r = 0; r < B; ++r)
1229 x[TDLS_STRIDED(r, x_stride) + TDLS_STRIDED_WIDE(m, xcol_stride)] =
1230 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r);
1232 }
1233 }
1234 for (int r = 0; r < B; ++r) {
1235 const bool acc = r >= R0 && r < R0 + R;
1236 TDLS_ARROW_LUPP_DYN_W(o.rb + r) =
1237 acc ? TDLS_ARROW_LUPP_DYN_W(o.racc + r - R0) : TDLS_ARROW_LUPP_DYN_Y(r);
1238 }
1239 if (count > 0 && !second_pass<1, false>(st, sizes(st, 1, 1), o, store, work_stride, work,
1240 work_stride, ystore, work_stride, 0, 1))
1242 if (!BorderSolver::solve_inplace(B, Sbb, work_stride, piv, piv_stride, rb, work_stride))
1244 // The accepted blocks, recomputed, the last one first: the places
1245 // of the store hold the refused ones in increasing order.
1246 int j = count - 1;
1247 for (int k = NB - 1; k >= 0; --k) {
1248 if (j >= 0 && TDLS_ARROW_LUPP_DYN_W(ro.slots + j) == T(k)) {
1249 --j;
1250 continue;
1251 }
1252 RecomputedBlock blk{streamed_block(z, o, so, work, work_stride), true};
1253 provider(k, blk);
1254 // y_k - A_kb x_b, then the column solved in place of A_kb(:, 0)
1255 for (int i = 0; i < S; ++i) {
1256 T v = TDLS_ARROW_LUPP_DYN_W(o.W + B * S + i);
1257 for (int c = 0; c < B; ++c)
1258 v -= TDLS_ARROW_LUPP_DYN_W(o.W + c * S + i) * TDLS_ARROW_LUPP_DYN_W(o.rb + c);
1259 TDLS_ARROW_LUPP_DYN_W(o.s + i) = v;
1260 }
1261 for (int i = 0; i < S; ++i)
1262 TDLS_ARROW_LUPP_DYN_W(o.W + i) = TDLS_ARROW_LUPP_DYN_W(o.s + i);
1263 if (z.reg) {
1264 (void)physical_factor(z, o, work, work_stride, piv, piv_stride);
1265 physical_solve(z, o, work, work_stride, piv, piv_stride, 1, true);
1266 } else {
1267 T* const Dl = work + TDLS_STRIDED(o.Dl, work_stride);
1268 (void)BlockSolver::factorize(S, Dl, work_stride, piv, piv_stride);
1269 BlockSolver::substitute_inplace(S, Dl, work_stride, piv, piv_stride,
1270 work + TDLS_STRIDED(o.W, work_stride), work_stride);
1271 }
1272 for (int i = 0; i < S; ++i)
1273 TDLS_ARROW_LUPP_DYN_Y(B + k * S + i) = TDLS_ARROW_LUPP_DYN_W(o.W + i);
1274 }
1275 if (count > 0)
1276 recover_refused(shape, st, z, o, ro, work, work_stride, y, rhs_stride, count);
1277 for (int r = 0; r < B; ++r)
1280 }
1281
1299 TDLS_HOST_DEVICE static constexpr int
1300 condense_or_keep_recomputed(const ArrowheadShape& shape, const Sizes& z, const Offsets& o,
1301 const StreamedOffsets& so, const RecomputedOffsets& ro,
1302 const int capacity, int* TDLS_RESTRICT piv, const int piv_stride,
1303 T* TDLS_RESTRICT work, const int work_stride, const int k,
1304 int& count) noexcept {
1305 const int B = z.B, S = z.S;
1306 const bool forward_check = !z.reg && col && z.pass_width == B + 1;
1307 for (int e = 0; e < S * S; ++e)
1308 TDLS_ARROW_LUPP_DYN_W(so.Dsave + e) = TDLS_ARROW_LUPP_DYN_W(o.Dl + e);
1309 if (forward_check)
1310 for (int e = 0; e < (B + 1) * S; ++e)
1311 TDLS_ARROW_LUPP_DYN_W(so.Wsave + e) = TDLS_ARROW_LUPP_DYN_W(o.W + e);
1312 if (condense<1, false, true>(shape, z, o, nullptr, 1, piv, piv_stride, work, work_stride,
1313 nullptr, 1, 0, k, 1, &so))
1314 return 0;
1315 ++count;
1316 if (count > capacity) return 2;
1317 // The block as given, in place j of the store (W is intact unless the
1318 // forward stage ran).
1319 const int j = count - 1;
1320 const ArrowheadShape st = store_shape(shape, capacity);
1321 const int ow = forward_check ? so.Wsave : o.W;
1322 for (int i = 0; i < S; ++i) {
1323 for (int jj = 0; jj < S; ++jj)
1324 TDLS_ARROW_LUPP_DYN_W(ro.store + storage_map::d(st, j, i, jj)) =
1325 TDLS_ARROW_LUPP_DYN_W(so.Dsave + (col ? jj * S + i : i * S + jj));
1326 for (int c = 0; c < B; ++c)
1327 TDLS_ARROW_LUPP_DYN_W(ro.store + storage_map::akb(st, j, i, c)) =
1328 TDLS_ARROW_LUPP_DYN_W(ow + c * S + i);
1329 TDLS_ARROW_LUPP_DYN_W(ro.ystore + B + j * S + i) =
1330 TDLS_ARROW_LUPP_DYN_W(ow + B * S + i);
1331 }
1332 if constexpr (diag) {
1333 for (int i = 0; i < S; ++i)
1334 TDLS_ARROW_LUPP_DYN_W(ro.store + storage_map::abk_diag(st, j, i)) =
1335 TDLS_ARROW_LUPP_DYN_W(o.acoef + i);
1336 } else {
1337 for (int r = 0; r < B; ++r)
1338 for (int i = 0; i < S; ++i)
1339 TDLS_ARROW_LUPP_DYN_W(ro.store + storage_map::abk(st, j, r, i)) =
1340 TDLS_ARROW_LUPP_DYN_W(so.Abk + r * S + i);
1341 }
1342 TDLS_ARROW_LUPP_DYN_W(ro.slots + j) = T(k);
1343 return 1;
1344 }
1345
1361 TDLS_HOST_DEVICE static constexpr void
1362 recover_refused(const ArrowheadShape& shape, const ArrowheadShape& st, const Sizes& z,
1363 const Offsets& o, const RecomputedOffsets& ro, T* TDLS_RESTRICT work,
1364 const int work_stride, T* TDLS_RESTRICT y, const int rhs_stride,
1365 const int count) noexcept {
1366 (void)shape;
1367 const int B = z.B, S = z.S, R = z.R;
1368 for (int a = 0; a < R; ++a)
1369 TDLS_ARROW_LUPP_DYN_W(o.w + a) = T(0);
1370 for (int j = count - 1; j >= 0; --j) {
1371 const int k = static_cast<int>(TDLS_ARROW_LUPP_DYN_W(ro.slots + j));
1372 for (int i = 0; i < S; ++i) {
1373 T v = TDLS_ARROW_LUPP_DYN_W(ro.ystore + B + j * S + i);
1374 for (int c = 0; c < B; ++c)
1375 v -= TDLS_ARROW_LUPP_DYN_W(ro.store + storage_map::akb(st, j, i, c)) *
1376 TDLS_ARROW_LUPP_DYN_W(o.rb + c);
1377 for (int a = 0; a < R; ++a) {
1378 const T h =
1379 a < S ? TDLS_ARROW_LUPP_DYN_W(ro.store + storage_map::d(st, j, i, a))
1380 : TDLS_ARROW_LUPP_DYN_W(o.extra + (j * S + i) * z.extra_cols + a - S);
1381 v -= h * TDLS_ARROW_LUPP_DYN_W(o.w + a);
1382 }
1383 TDLS_ARROW_LUPP_DYN_W(o.s + i) = v;
1384 }
1385 if constexpr (diag) {
1386 for (int i = 0; i < S; ++i)
1387 TDLS_ARROW_LUPP_DYN_W(o.w + i) +=
1388 TDLS_ARROW_LUPP_DYN_W(ro.store + storage_map::abk_diag(st, j, i)) *
1389 TDLS_ARROW_LUPP_DYN_W(o.s + i);
1390 } else {
1391 for (int r = 0; r < B; ++r)
1392 for (int i = 0; i < S; ++i)
1393 TDLS_ARROW_LUPP_DYN_W(o.w + r) +=
1394 TDLS_ARROW_LUPP_DYN_W(ro.store + storage_map::abk(st, j, r, i)) *
1395 TDLS_ARROW_LUPP_DYN_W(o.s + i);
1396 }
1397 for (int i = 0; i < S; ++i)
1398 TDLS_ARROW_LUPP_DYN_Y(B + k * S + i) = TDLS_ARROW_LUPP_DYN_W(o.s + i);
1399 }
1400 }
1401
1402 private:
1410 template<typename P>
1411 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL static constexpr P*
1412 streamed_arg(P* p) noexcept {
1413 return p;
1414 }
1415
1424 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL static constexpr StreamedBlock
1425 streamed_block(const Sizes& z, const Offsets& o, const StreamedOffsets& so,
1426 T* TDLS_RESTRICT work, const int work_stride) noexcept {
1427 StreamedBlock blk{};
1428 blk.Dl = work + TDLS_STRIDED(o.Dl, work_stride);
1429 blk.W = work + TDLS_STRIDED(o.W, work_stride);
1430 blk.Abk = work + TDLS_STRIDED(diag ? o.acoef : so.Abk, work_stride);
1431 blk.S = z.S;
1432 blk.B = z.B;
1433 blk.stride = work_stride;
1434 return blk;
1435 }
1436
1462 template<int NC, bool canonical, typename... Provider>
1463 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL static constexpr bool
1464 solve_core(const ArrowheadShape& shape, const int nrhs, T* TDLS_RESTRICT A, const int A_stride,
1465 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT work,
1466 const int work_stride, T* TDLS_RESTRICT y, const int rhs_stride,
1467 [[maybe_unused]] const std::size_t xcol_stride, [[maybe_unused]] const int col0,
1468 int& refused_count, Provider*... provider) noexcept {
1469 const int nc = NC > 0 ? NC : nrhs;
1470 if (!valid(shape, nc)) {
1471 refused_count = 0;
1472 return false;
1473 }
1474 const int ny = canonical ? 0 : nc;
1475 // Columns of y when fixed, for TDLS_ARROW_LUPP_DYN_YW.
1476 [[maybe_unused]] constexpr int ycols = NC;
1477 const Sizes z = sizes(shape, nc, ny);
1478 const Offsets o = offsets(shape, nc);
1479 const int B = z.B, R = z.R, R0 = z.R0, NB = z.NB;
1480 if constexpr (!in_place) {
1481 for (int r = 0; r < R; ++r) {
1482 for (int c = 0; c < B; ++c)
1483 TDLS_ARROW_LUPP_DYN_W(o.Sacc + r * B + c) =
1484 TDLS_ARROW_LUPP_DYN_A(storage_map::abb(shape, R0 + r, c));
1485 for (int m = 0; m < ny; ++m)
1486 TDLS_ARROW_LUPP_DYN_W(o.racc + m * R + r) = TDLS_ARROW_LUPP_DYN_YW(m, R0 + r);
1487 }
1488 }
1489 refused_count = 0;
1490 if constexpr (sizeof...(Provider) > 0) {
1491 // Streamed: the provider writes block k in the workspace.
1492 const StreamedOffsets so = streamed_offsets(shape, nc);
1493 for (int k = 0; k < NB; ++k) {
1494 StreamedBlock blk = streamed_block(z, o, so, work, work_stride);
1495 (*streamed_arg(provider...))(k, blk);
1496 const bool ok = condense_or_store_streamed<NC, canonical>(
1497 shape, z, o, so, A, A_stride, piv, piv_stride, work, work_stride, y, rhs_stride,
1498 xcol_stride, k, nc);
1499 TDLS_ARROW_LUPP_DYN_W(o.refused + k) = ok ? T(0) : T(1);
1500 if (!ok) ++refused_count;
1501 }
1502 } else {
1503 for (int k = 0; k < NB; ++k) {
1504 const bool ok =
1505 condense<NC, canonical>(shape, z, o, A, A_stride, piv, piv_stride, work,
1506 work_stride, y, rhs_stride, xcol_stride, k, nc);
1507 TDLS_ARROW_LUPP_DYN_W(o.refused + k) = ok ? T(0) : T(1);
1508 if (!ok) ++refused_count;
1509 }
1510 }
1511 // The border system: the accumulated rows, the others as stored;
1512 // column m of its right-hand sides at m * B.
1513 // The rows as stored and the accumulated rows in separate loops:
1514 // GCC 11 and 12 vectorize a choice of source per entry wrongly (-O3
1515 // with AVX2 or AVX-512).
1516 const int acc0 = in_place ? B : R0;
1517 const int acc1 = in_place ? B : R0 + R;
1518 for (int r = 0; r < acc0; ++r) {
1519 for (int c = 0; c < B; ++c)
1520 TDLS_ARROW_LUPP_DYN_W(o.Sbb + r * B + c) =
1522 for (int m = 0; m < nc; ++m) {
1523 if constexpr (canonical)
1524 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r) = r == col0 + m ? T(1) : T(0);
1525 else
1526 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r) = TDLS_ARROW_LUPP_DYN_YW(m, r);
1527 }
1528 }
1529 for (int r = acc0; r < acc1; ++r) {
1530 for (int c = 0; c < B; ++c)
1531 TDLS_ARROW_LUPP_DYN_W(o.Sbb + r * B + c) =
1532 TDLS_ARROW_LUPP_DYN_W(o.Sacc + (r - R0) * B + c);
1533 for (int m = 0; m < nc; ++m) {
1534 if constexpr (canonical)
1535 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r) = r == col0 + m ? T(1) : T(0);
1536 else
1537 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r) =
1538 TDLS_ARROW_LUPP_DYN_W(o.racc + m * R + r - R0);
1539 }
1540 }
1541 for (int r = acc1; r < B; ++r) {
1542 for (int c = 0; c < B; ++c)
1543 TDLS_ARROW_LUPP_DYN_W(o.Sbb + r * B + c) =
1545 for (int m = 0; m < nc; ++m) {
1546 if constexpr (canonical)
1547 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r) = r == col0 + m ? T(1) : T(0);
1548 else
1549 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r) = TDLS_ARROW_LUPP_DYN_YW(m, r);
1550 }
1551 }
1552 if (refused_count > 0 &&
1553 !second_pass<NC, canonical>(shape, z, o, A, A_stride, work, work_stride, y, rhs_stride,
1554 xcol_stride, nc))
1555 return false;
1556 if (nc == 1) {
1557 if (!BorderSolver::solve_inplace(B, work + TDLS_STRIDED(o.Sbb, work_stride),
1558 work_stride, piv, piv_stride,
1559 work + TDLS_STRIDED(o.rb, work_stride), work_stride))
1560 return false;
1561 } else {
1563 B, nc, work + TDLS_STRIDED(o.Sbb, work_stride), work_stride, piv, piv_stride,
1564 work + TDLS_STRIDED(o.rb, work_stride), work_stride,
1565 TDLS_STRIDED(B, work_stride)))
1566 return false;
1567 }
1568 for (int m = 0; m < nc; ++m)
1569 for (int r = 0; r < B; ++r) {
1570 if constexpr (canonical)
1571 y[TDLS_STRIDED(r, rhs_stride) + TDLS_STRIDED_WIDE(m, xcol_stride)] =
1572 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r);
1573 else
1574 TDLS_ARROW_LUPP_DYN_YW(m, r) = TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r);
1575 }
1576 if constexpr (!canonical)
1577 recover<NC>(shape, z, o, A, A_stride, work, work_stride, y, rhs_stride, xcol_stride,
1578 nc);
1579 return true;
1580 }
1581
1610 template<int NC, bool canonical, bool recompute = false, typename... Streamed>
1611 TDLS_HOST_DEVICE static constexpr bool
1612 condense(const ArrowheadShape& shape, const Sizes& z, const Offsets& o, T* TDLS_RESTRICT A,
1613 const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
1614 T* TDLS_RESTRICT work, const int work_stride, T* TDLS_RESTRICT y, const int rhs_stride,
1615 [[maybe_unused]] const std::size_t xcol_stride, const int k,
1616 [[maybe_unused]] const int nrhs, const Streamed*... so) noexcept {
1617 constexpr bool streamed = sizeof...(Streamed) > 0;
1618 const int B = z.B, S = z.S;
1619 // Right-hand sides carried through the blocks.
1620 const int ny = canonical ? 0 : NC > 0 ? NC : nrhs;
1621 [[maybe_unused]] constexpr int ycols = NC;
1622 const T u = condensation_pivot_threshold;
1623 T* const Dl = work + TDLS_STRIDED(o.Dl, work_stride);
1624 T* const W = work + TDLS_STRIDED(o.W, work_stride);
1625 // The block, in the order of its storage (a streamed block is there
1626 // already).
1627 if constexpr (!streamed) {
1628 for (int i = 0; i < S; ++i)
1629 for (int j = 0; j < S; ++j)
1630 TDLS_ARROW_LUPP_DYN_W(o.Dl + (col ? j * S + i : i * S + j)) =
1631 TDLS_ARROW_LUPP_DYN_A(storage_map::d(shape, k, i, j));
1632 }
1633 // Factored by TiledLUpp, or by physical interchanges
1634 // (blocks in registers): Dl then holds U in logical row order and the
1635 // check of L is made.
1636 bool ok = true;
1637 if (z.reg) {
1638 const int verdict = physical_factor(z, o, work, work_stride, piv, piv_stride);
1639 if (verdict == 2) return false;
1640 ok = verdict == 0;
1641 } else {
1642 if (!BlockSolver::factorize(S, Dl, work_stride, piv, piv_stride)) return false;
1643 }
1644 if constexpr (diag && !streamed) {
1645 for (int i = 0; i < S; ++i)
1646 TDLS_ARROW_LUPP_DYN_W(o.acoef + i) =
1648 }
1649 const bool forward_check = !z.reg && col && z.pass_width == B + ny;
1650 // The multipliers L, and U in logical row order unless the forward
1651 // stage of the solve gives the border multipliers.
1652 if (z.reg) {
1653 for (int l = 0; l < S; ++l)
1654 for (int j = l; j < S; ++j)
1655 TDLS_ARROW_LUPP_DYN_W(o.U + l * S + j) =
1656 TDLS_ARROW_LUPP_DYN_W(o.Dl + l * S + j);
1657 } else {
1658 for (int l = 0; l < S; ++l) {
1659 const int p = piv[TDLS_STRIDED(l, piv_stride)];
1660 for (int j = 0; j < S; ++j) {
1661 if (j < l) {
1662 ok = ok && u * detail::abs(TDLS_ARROW_LUPP_DYN_W(o.Dl + p * S + j)) <= T(1);
1663 } else if (!forward_check) {
1664 TDLS_ARROW_LUPP_DYN_W(o.U + l * S + j) =
1665 TDLS_ARROW_LUPP_DYN_W(o.Dl + p * S + j);
1666 }
1667 }
1668 }
1669 }
1670 if (forward_check) {
1671 if constexpr (!streamed) {
1672 for (int i = 0; i < S; ++i) {
1673 for (int c = 0; c < B; ++c)
1674 TDLS_ARROW_LUPP_DYN_W(o.W + c * S + i) =
1675 TDLS_ARROW_LUPP_DYN_A(storage_map::akb(shape, k, i, c));
1676 for (int m = 0; m < ny; ++m)
1677 TDLS_ARROW_LUPP_DYN_W(o.W + (B + m) * S + i) =
1678 TDLS_ARROW_LUPP_DYN_YW(m, B + k * S + i);
1679 }
1680 }
1681 BlockSolver::forward_inplace_multirhs(S, B + ny, Dl, work_stride, piv, piv_stride, W,
1682 work_stride, TDLS_STRIDED(S, work_stride));
1683 // Column c of U^-T A_kb is row c of the border multipliers.
1684 for (int e = 0; e < B * S; ++e)
1685 ok = ok && u * detail::abs(TDLS_ARROW_LUPP_DYN_W(o.W + e)) <= T(1);
1686 if (!ok) return false;
1687 BlockSolver::backward_inplace_multirhs(S, B + ny, Dl, work_stride, piv, piv_stride, W,
1688 work_stride, TDLS_STRIDED(S, work_stride));
1689 condense_pass<NC, canonical, recompute>(shape, z, o, A, A_stride, piv, piv_stride, work,
1690 work_stride, y, rhs_stride, xcol_stride, k, 0,
1691 B + ny, true, so...);
1692 } else {
1693 ok = border_multipliers(shape, z, o, A, A_stride, work, work_stride, k, so...) && ok;
1694 if (!ok) return false;
1695 // U, recovery data: not for the canonical solve.
1696 if constexpr (multiplier_update && !canonical && !recompute) {
1697 for (int l = 0; l < S; ++l)
1698 for (int j = l; j < S; ++j)
1699 TDLS_ARROW_LUPP_DYN_A(storage_map::d(shape, k, l, j)) =
1700 TDLS_ARROW_LUPP_DYN_W(o.U + l * S + j);
1701 }
1702 for (int c0 = 0; c0 < B + ny; c0 += z.pass_width) {
1703 const int w = B + ny - c0 < z.pass_width ? B + ny - c0 : z.pass_width;
1704 condense_pass<NC, canonical, recompute>(shape, z, o, A, A_stride, piv, piv_stride,
1705 work, work_stride, y, rhs_stride,
1706 xcol_stride, k, c0, w, false, so...);
1707 }
1708 }
1709 return true;
1710 }
1711
1739 template<int NC, bool canonical>
1740 TDLS_HOST_DEVICE static constexpr bool condense_or_store_streamed(
1741 const ArrowheadShape& shape, const Sizes& z, const Offsets& o, const StreamedOffsets& so,
1742 T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
1743 T* TDLS_RESTRICT work, const int work_stride, T* TDLS_RESTRICT y, const int rhs_stride,
1744 [[maybe_unused]] const std::size_t xcol_stride, const int k, const int nrhs) noexcept {
1745 const int B = z.B, S = z.S;
1746 const int ny = canonical ? 0 : NC > 0 ? NC : nrhs;
1747 [[maybe_unused]] constexpr int ycols = NC;
1748 const bool forward_check = !z.reg && col && z.pass_width == B + ny;
1749 for (int e = 0; e < S * S; ++e)
1750 TDLS_ARROW_LUPP_DYN_W(so.Dsave + e) = TDLS_ARROW_LUPP_DYN_W(o.Dl + e);
1751 if (forward_check)
1752 for (int e = 0; e < (B + ny) * S; ++e)
1753 TDLS_ARROW_LUPP_DYN_W(so.Wsave + e) = TDLS_ARROW_LUPP_DYN_W(o.W + e);
1754 if (condense<NC, canonical>(shape, z, o, A, A_stride, piv, piv_stride, work, work_stride, y,
1755 rhs_stride, xcol_stride, k, nrhs, &so))
1756 return true;
1757 // Refused: the block as given, where the stored path leaves a
1758 // refused block (W is intact unless the forward stage ran).
1759 const int ow = forward_check ? so.Wsave : o.W;
1760 for (int i = 0; i < S; ++i) {
1761 for (int j = 0; j < S; ++j)
1762 TDLS_ARROW_LUPP_DYN_A(storage_map::d(shape, k, i, j)) =
1763 TDLS_ARROW_LUPP_DYN_W(so.Dsave + (col ? j * S + i : i * S + j));
1764 for (int c = 0; c < B; ++c)
1765 TDLS_ARROW_LUPP_DYN_A(storage_map::akb(shape, k, i, c)) =
1766 TDLS_ARROW_LUPP_DYN_W(ow + c * S + i);
1767 for (int m = 0; m < ny; ++m)
1768 TDLS_ARROW_LUPP_DYN_YW(m, B + k * S + i) = TDLS_ARROW_LUPP_DYN_W(ow + B * S + i);
1769 }
1770 if constexpr (diag) {
1771 for (int i = 0; i < S; ++i)
1773 TDLS_ARROW_LUPP_DYN_W(o.acoef + i);
1774 } else {
1775 for (int r = 0; r < B; ++r)
1776 for (int i = 0; i < S; ++i)
1777 TDLS_ARROW_LUPP_DYN_A(storage_map::abk(shape, k, r, i)) =
1778 TDLS_ARROW_LUPP_DYN_W(so.Abk + r * S + i);
1779 }
1780 return false;
1781 }
1782
1799 template<typename... Streamed>
1800 TDLS_HOST_DEVICE static constexpr bool
1801 border_multipliers(const ArrowheadShape& shape, const Sizes& z, const Offsets& o,
1802 const T* TDLS_RESTRICT A, const int A_stride, T* TDLS_RESTRICT work,
1803 const int work_stride, const int k, const Streamed*... so) noexcept {
1804 constexpr bool streamed = sizeof...(Streamed) > 0;
1805 const int B = z.B, S = z.S;
1806 bool ok = true;
1807 if constexpr (col) {
1808 for (int c = 0; c < B; ++c) {
1809 for (int i = 0; i < S; ++i) {
1810 if constexpr (streamed)
1811 TDLS_ARROW_LUPP_DYN_W(o.m + i) = TDLS_ARROW_LUPP_DYN_W(o.W + c * S + i);
1812 else
1813 TDLS_ARROW_LUPP_DYN_W(o.m + i) =
1814 TDLS_ARROW_LUPP_DYN_A(storage_map::akb(shape, k, i, c));
1815 }
1816 ok = multipliers_pass(z, o, work, work_stride, o.m, 0) && ok;
1817 }
1818 } else if constexpr (diag) {
1819 // Border row coupling_row + i holds one entry, in column i.
1820 for (int i = 0; i < S; ++i) {
1821 for (int j = 0; j < S; ++j)
1822 TDLS_ARROW_LUPP_DYN_W(o.m + j) =
1823 j == i ? TDLS_ARROW_LUPP_DYN_W(o.acoef + i) : T(0);
1824 ok = multipliers_pass(z, o, work, work_stride, o.m, i) && ok;
1825 }
1826 } else {
1827 for (int r = 0; r < B; ++r) {
1828 for (int j = 0; j < S; ++j) {
1829 if constexpr (streamed)
1830 TDLS_ARROW_LUPP_DYN_W(o.Mb + r * S + j) =
1831 TDLS_ARROW_LUPP_DYN_W(streamed_arg(so...)->Abk + r * S + j);
1832 else
1833 TDLS_ARROW_LUPP_DYN_W(o.Mb + r * S + j) =
1834 TDLS_ARROW_LUPP_DYN_A(storage_map::abk(shape, k, r, j));
1835 }
1836 ok = multipliers_pass(z, o, work, work_stride, o.Mb + r * S, 0) && ok;
1837 }
1838 }
1839 return ok;
1840 }
1841
1860 TDLS_HOST_DEVICE static constexpr int
1861 physical_factor(const Sizes& z, const Offsets& o, T* TDLS_RESTRICT work, const int work_stride,
1862 int* TDLS_RESTRICT piv, const int piv_stride) noexcept {
1864 z.S, work + TDLS_STRIDED(o.Dl, work_stride), work_stride, piv, piv_stride,
1865 condensation_pivot_threshold);
1866 }
1867
1885 TDLS_HOST_DEVICE static constexpr void
1886 physical_solve(const Sizes& z, const Offsets& o, T* TDLS_RESTRICT work, const int work_stride,
1887 const int* TDLS_RESTRICT piv, const int piv_stride, const int w,
1888 const bool backward = !multiplier_update) noexcept {
1889 const int S = z.S;
1890 const T* TDLS_RESTRICT F = work + TDLS_STRIDED(o.Dl, work_stride);
1891 T* TDLS_RESTRICT W = work + TDLS_STRIDED(o.W, work_stride);
1892 const std::size_t cs = TDLS_STRIDED(S, work_stride);
1893 PhysicalBlockSolver::physical_forward(S, w, F, work_stride, piv, piv_stride, W, work_stride,
1894 cs);
1895 if (col || backward)
1896 PhysicalBlockSolver::physical_backward(S, w, F, work_stride, piv, piv_stride, W,
1897 work_stride, cs);
1898 }
1899
1909 TDLS_HOST_DEVICE static constexpr bool multipliers_pass(const Sizes& z, const Offsets& o,
1910 T* TDLS_RESTRICT work,
1911 const int work_stride, const int m,
1912 const int first) noexcept {
1913 const int S = z.S;
1914 const T u = condensation_pivot_threshold;
1915 bool ok = true;
1916 for (int l = first; l < S; ++l) {
1917 TDLS_ARROW_LUPP_DYN_W(m + l) *= TDLS_ARROW_LUPP_DYN_W(o.U + l * S + l);
1918 ok = ok && u * detail::abs(TDLS_ARROW_LUPP_DYN_W(m + l)) <= T(1);
1919 for (int j = l + 1; j < S; ++j)
1920 TDLS_ARROW_LUPP_DYN_W(m + j) -=
1921 TDLS_ARROW_LUPP_DYN_W(m + l) * TDLS_ARROW_LUPP_DYN_W(o.U + l * S + j);
1922 }
1923 return ok;
1924 }
1925
1941 template<int NC>
1942 TDLS_HOST_DEVICE static constexpr void
1943 schur_sub([[maybe_unused]] const ArrowheadShape& shape, const Sizes& z,
1944 [[maybe_unused]] const Offsets& o, [[maybe_unused]] T* TDLS_RESTRICT A,
1945 [[maybe_unused]] const int A_stride, [[maybe_unused]] T* TDLS_RESTRICT work,
1946 [[maybe_unused]] const int work_stride, [[maybe_unused]] T* TDLS_RESTRICT y,
1947 [[maybe_unused]] const int rhs_stride, [[maybe_unused]] const std::size_t xcol_stride,
1948 const int r, const int c, const T v) noexcept {
1949 [[maybe_unused]] constexpr int ycols = NC;
1950 if constexpr (in_place) {
1951 if (c < z.B)
1952 TDLS_ARROW_LUPP_DYN_A(storage_map::abb(shape, z.R0 + r, c)) -= v;
1953 else
1954 TDLS_ARROW_LUPP_DYN_YW(c - z.B, z.R0 + r) -= v;
1955 } else {
1956 if (c < z.B)
1957 TDLS_ARROW_LUPP_DYN_W(o.Sacc + r * z.B + c) -= v;
1958 else
1959 TDLS_ARROW_LUPP_DYN_W(o.racc + (NC == 1 ? 0 : (c - z.B) * z.R) + r) -= v;
1960 }
1961 }
1962
1991 template<int NC, bool canonical, bool recompute = false, typename... Streamed>
1992 TDLS_HOST_DEVICE static constexpr void
1993 condense_pass(const ArrowheadShape& shape, const Sizes& z, const Offsets& o, T* TDLS_RESTRICT A,
1994 const int A_stride, const int* TDLS_RESTRICT piv, const int piv_stride,
1995 T* TDLS_RESTRICT work, const int work_stride, T* TDLS_RESTRICT y,
1996 const int rhs_stride, [[maybe_unused]] const std::size_t xcol_stride, const int k,
1997 const int c0, const int w, const bool solved, const Streamed*... so) noexcept {
1998 constexpr bool streamed = sizeof...(Streamed) > 0;
1999 [[maybe_unused]] constexpr int ycols = NC;
2000 const int B = z.B, S = z.S, R = z.R;
2001 const T* const Dl = work + TDLS_STRIDED(o.Dl, work_stride);
2002 T* const W = work + TDLS_STRIDED(o.W, work_stride);
2003 if (!solved) {
2004 if constexpr (streamed) {
2005 // Columns c0 .. c0 + w of the streamed block move to the
2006 // front of W: the columns before c0 are solved and stored
2007 // already, and c0 >= w when c0 > 0, so the two ranges do
2008 // not overlap.
2009 if (c0 > 0)
2010 for (int i = 0; i < S; ++i)
2011 for (int j = 0; j < w; ++j)
2012 TDLS_ARROW_LUPP_DYN_W(o.W + j * S + i) =
2013 TDLS_ARROW_LUPP_DYN_W(o.W + (c0 + j) * S + i);
2014 } else {
2015 for (int i = 0; i < S; ++i)
2016 for (int j = 0; j < w; ++j)
2017 TDLS_ARROW_LUPP_DYN_W(o.W + j * S + i) =
2018 c0 + j < B
2019 ? TDLS_ARROW_LUPP_DYN_A(storage_map::akb(shape, k, i, c0 + j))
2020 : TDLS_ARROW_LUPP_DYN_YW(c0 + j - B, B + k * S + i);
2021 }
2022 if (z.reg) {
2023 physical_solve(z, o, work, work_stride, piv, piv_stride, w);
2024 } else if constexpr (multiplier_update) {
2025 BlockSolver::forward_inplace_multirhs(S, w, Dl, work_stride, piv, piv_stride, W,
2026 work_stride, TDLS_STRIDED(S, work_stride));
2027 } else if constexpr (Config.gather_rhs) {
2028 T* const X = work + TDLS_STRIDED(o.X, work_stride);
2029 BlockSolver::substitute_multirhs(S, w, Dl, work_stride, piv, piv_stride, W, X,
2030 work_stride, TDLS_STRIDED(S, work_stride));
2031 for (int e = 0; e < w * S; ++e)
2032 TDLS_ARROW_LUPP_DYN_W(o.W + e) = TDLS_ARROW_LUPP_DYN_W(o.X + e);
2033 } else {
2034 BlockSolver::substitute_inplace_multirhs(S, w, Dl, work_stride, piv, piv_stride, W,
2035 work_stride, TDLS_STRIDED(S, work_stride));
2036 }
2037 }
2038 if constexpr (!canonical && !recompute) {
2039 for (int i = 0; i < S; ++i)
2040 for (int j = 0; j < w; ++j) {
2041 if (c0 + j < B)
2042 TDLS_ARROW_LUPP_DYN_A(storage_map::akb(shape, k, i, c0 + j)) =
2043 TDLS_ARROW_LUPP_DYN_W(o.W + j * S + i);
2044 else
2045 TDLS_ARROW_LUPP_DYN_YW(c0 + j - B, B + k * S + i) =
2046 TDLS_ARROW_LUPP_DYN_W(o.W + j * S + i);
2047 }
2048 }
2049 if constexpr (multiplier_update) {
2050 for (int r = 0; r < R; ++r)
2051 for (int j = 0; j < w; ++j) {
2052 T sum = T(0);
2053 for (int l = 0; l < S; ++l)
2054 sum += TDLS_ARROW_LUPP_DYN_W(o.Mb + r * S + l) *
2055 TDLS_ARROW_LUPP_DYN_W(o.W + j * S + l);
2056 schur_sub<NC>(shape, z, o, A, A_stride, work, work_stride, y, rhs_stride,
2057 xcol_stride, r, c0 + j, sum);
2058 }
2059 } else if constexpr (diag) {
2060 for (int i = 0; i < S; ++i)
2061 for (int j = 0; j < w; ++j)
2062 schur_sub<NC>(shape, z, o, A, A_stride, work, work_stride, y, rhs_stride,
2063 xcol_stride, i, c0 + j,
2064 TDLS_ARROW_LUPP_DYN_W(o.acoef + i) *
2065 TDLS_ARROW_LUPP_DYN_W(o.W + j * S + i));
2066 } else {
2067 for (int r = 0; r < B; ++r) {
2068 for (int i = 0; i < S; ++i) {
2069 if constexpr (streamed)
2070 TDLS_ARROW_LUPP_DYN_W(o.a + i) =
2071 TDLS_ARROW_LUPP_DYN_W(streamed_arg(so...)->Abk + r * S + i);
2072 else
2073 TDLS_ARROW_LUPP_DYN_W(o.a + i) =
2074 TDLS_ARROW_LUPP_DYN_A(storage_map::abk(shape, k, r, i));
2075 }
2076 for (int j = 0; j < w; ++j) {
2077 T sum = T(0);
2078 for (int i = 0; i < S; ++i)
2079 sum +=
2080 TDLS_ARROW_LUPP_DYN_W(o.a + i) * TDLS_ARROW_LUPP_DYN_W(o.W + j * S + i);
2081 schur_sub<NC>(shape, z, o, A, A_stride, work, work_stride, y, rhs_stride,
2082 xcol_stride, r, c0 + j, sum);
2083 }
2084 }
2085 }
2086 }
2087
2103 template<int NC>
2104 TDLS_HOST_DEVICE static constexpr void
2105 recover(const ArrowheadShape& shape, const Sizes& z, const Offsets& o, const T* TDLS_RESTRICT A,
2106 const int A_stride, T* TDLS_RESTRICT work, const int work_stride, T* TDLS_RESTRICT y,
2107 const int rhs_stride, [[maybe_unused]] const std::size_t xcol_stride,
2108 [[maybe_unused]] const int nrhs) noexcept {
2109 const int B = z.B, S = z.S, NB = z.NB, R = z.R;
2110 const int nc = NC > 0 ? NC : nrhs;
2111 [[maybe_unused]] constexpr int ycols = NC;
2112 // w and s: column m at m * R and m * S.
2113 for (int e = 0; e < R * nc; ++e)
2114 TDLS_ARROW_LUPP_DYN_W(o.w + e) = T(0);
2115 for (int k = NB - 1; k >= 0; --k) {
2116 for (int i = 0; i < S; ++i) {
2117 for (int m = 0; m < nc; ++m)
2118 TDLS_ARROW_LUPP_DYN_W(o.s + m * S + i) =
2119 TDLS_ARROW_LUPP_DYN_YW(m, B + k * S + i);
2120 for (int c = 0; c < B; ++c) {
2121 const T a = TDLS_ARROW_LUPP_DYN_A(storage_map::akb(shape, k, i, c));
2122 for (int m = 0; m < nc; ++m)
2123 TDLS_ARROW_LUPP_DYN_W(o.s + m * S + i) -=
2124 a * TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + c);
2125 }
2126 }
2127 if (TDLS_ARROW_LUPP_DYN_W(o.refused + k) != T(0)) {
2128 // x_k = c_k - G_k xb - H_k w, then w += A_bk x_k.
2129 for (int i = 0; i < S; ++i)
2130 for (int a = 0; a < R; ++a) {
2131 const T h = a < S ? TDLS_ARROW_LUPP_DYN_A(storage_map::d(shape, k, i, a))
2133 o.extra + (k * S + i) * z.extra_cols + a - S);
2134 for (int m = 0; m < nc; ++m)
2135 TDLS_ARROW_LUPP_DYN_W(o.s + m * S + i) -=
2136 h * TDLS_ARROW_LUPP_DYN_W(o.w + m * R + a);
2137 }
2138 if constexpr (diag) {
2139 for (int i = 0; i < S; ++i) {
2140 const T coef = TDLS_ARROW_LUPP_DYN_A(storage_map::abk_diag(shape, k, i));
2141 for (int m = 0; m < nc; ++m)
2142 TDLS_ARROW_LUPP_DYN_W(o.w + m * R + i) +=
2143 coef * TDLS_ARROW_LUPP_DYN_W(o.s + m * S + i);
2144 }
2145 } else {
2146 for (int r = 0; r < B; ++r)
2147 for (int i = 0; i < S; ++i) {
2148 const T coef = TDLS_ARROW_LUPP_DYN_A(storage_map::abk(shape, k, r, i));
2149 for (int m = 0; m < nc; ++m)
2150 TDLS_ARROW_LUPP_DYN_W(o.w + m * R + r) +=
2151 coef * TDLS_ARROW_LUPP_DYN_W(o.s + m * S + i);
2152 }
2153 }
2154 } else if constexpr (multiplier_update) {
2155 // x_k = U^-1 s, U in the upper triangle of D_k.
2156 for (int l = S - 1; l >= 0; --l) {
2157 for (int j = l + 1; j < S; ++j) {
2158 const T ulj = TDLS_ARROW_LUPP_DYN_A(storage_map::d(shape, k, l, j));
2159 for (int m = 0; m < nc; ++m)
2160 TDLS_ARROW_LUPP_DYN_W(o.s + m * S + l) -=
2161 ulj * TDLS_ARROW_LUPP_DYN_W(o.s + m * S + j);
2162 }
2163 const T ull = TDLS_ARROW_LUPP_DYN_A(storage_map::d(shape, k, l, l));
2164 for (int m = 0; m < nc; ++m)
2165 TDLS_ARROW_LUPP_DYN_W(o.s + m * S + l) *= ull;
2166 }
2167 }
2168 for (int i = 0; i < S; ++i)
2169 for (int m = 0; m < nc; ++m)
2170 TDLS_ARROW_LUPP_DYN_YW(m, B + k * S + i) =
2171 TDLS_ARROW_LUPP_DYN_W(o.s + m * S + i);
2172 }
2173 }
2174
2195 template<int NC, bool canonical>
2196 TDLS_HOST_DEVICE TDLS_NO_INLINE static constexpr bool
2197 second_pass(const ArrowheadShape& shape, const Sizes& z, const Offsets& o, T* TDLS_RESTRICT A,
2198 const int A_stride, T* TDLS_RESTRICT work, const int work_stride,
2199 T* TDLS_RESTRICT y, const int rhs_stride,
2200 [[maybe_unused]] const std::size_t xcol_stride,
2201 [[maybe_unused]] const int nrhs) noexcept {
2202 const int B = z.B, S = z.S, NB = z.NB, R = z.R, R0 = z.R0;
2203 const int front_rows = z.front_rows, front_cols = z.front_cols;
2204 const int nc = NC > 0 ? NC : nrhs;
2205 const int rhs0 = S + B + R;
2206 [[maybe_unused]] constexpr int ycols = NC;
2207 const T u = condensation_pivot_threshold;
2208 // The singularity floor of the block solver: a smaller pivot, zero
2209 // or subnormal, is singular there, and its reciprocal could overflow.
2210 constexpr T singular_floor = BlockSolver::singular_floor;
2211 for (int r = 0; r < B; ++r) {
2212 for (int c = 0; c < B; ++c)
2213 TDLS_ARROW_LUPP_DYN_F(S + r, S + c) = TDLS_ARROW_LUPP_DYN_W(o.Sbb + r * B + c);
2214 for (int a = 0; a < R; ++a)
2215 TDLS_ARROW_LUPP_DYN_F(S + r, S + B + a) = r == R0 + a ? T(1) : T(0);
2216 for (int m = 0; m < nc; ++m)
2217 TDLS_ARROW_LUPP_DYN_F(S + r, rhs0 + m) = TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r);
2218 }
2219 for (int k = 0; k < NB; ++k) {
2220 if (TDLS_ARROW_LUPP_DYN_W(o.refused + k) == T(0)) continue;
2221 for (int i = 0; i < S; ++i) {
2222 for (int j = 0; j < S; ++j)
2223 TDLS_ARROW_LUPP_DYN_F(i, j) =
2224 TDLS_ARROW_LUPP_DYN_A(storage_map::d(shape, k, i, j));
2225 for (int c = 0; c < B; ++c)
2226 TDLS_ARROW_LUPP_DYN_F(i, S + c) =
2227 TDLS_ARROW_LUPP_DYN_A(storage_map::akb(shape, k, i, c));
2228 for (int a = 0; a < R; ++a)
2229 TDLS_ARROW_LUPP_DYN_F(i, S + B + a) = T(0);
2230 for (int m = 0; m < nc; ++m) {
2231 if constexpr (canonical)
2232 TDLS_ARROW_LUPP_DYN_F(i, rhs0 + m) = T(0);
2233 else
2234 TDLS_ARROW_LUPP_DYN_F(i, rhs0 + m) =
2235 TDLS_ARROW_LUPP_DYN_YW(m, B + k * S + i);
2236 }
2237 }
2238 for (int r = 0; r < B; ++r) {
2239 for (int j = 0; j < S; ++j) {
2240 T v = T(0);
2241 if constexpr (diag) {
2242 v = TDLS_ARROW_LUPP_DYN_F(S + r, S + B + j) *
2244 } else {
2245 for (int a = 0; a < B; ++a)
2246 v += TDLS_ARROW_LUPP_DYN_F(S + r, S + B + a) *
2247 TDLS_ARROW_LUPP_DYN_A(storage_map::abk(shape, k, a, j));
2248 }
2249 TDLS_ARROW_LUPP_DYN_F(S + r, j) = v;
2250 }
2251 }
2252 for (int e = 0; e < front_rows; ++e)
2253 TDLS_ARROW_LUPP_DYN_W(o.block_row + e) = e < S ? T(1) : T(0);
2254 for (int j = 0; j < S; ++j) {
2255 int pmax = j, pblk = -1;
2256 T vmax = T(0), vblk = T(0);
2257 for (int e = j; e < front_rows; ++e) {
2258 const T v = detail::abs(TDLS_ARROW_LUPP_DYN_F(e, j));
2259 if (v > vmax) {
2260 vmax = v;
2261 pmax = e;
2262 }
2263 if (TDLS_ARROW_LUPP_DYN_W(o.block_row + e) != T(0) && v > vblk) {
2264 vblk = v;
2265 pblk = e;
2266 }
2267 }
2268 if (!(vmax >= singular_floor)) return false;
2269 const int p = pblk >= 0 && vblk >= u * vmax && vblk >= singular_floor ? pblk : pmax;
2270 if (p != j) {
2271 for (int c = j; c < front_cols; ++c) {
2272 const T t = TDLS_ARROW_LUPP_DYN_F(p, c);
2274 TDLS_ARROW_LUPP_DYN_F(j, c) = t;
2275 }
2276 const T t = TDLS_ARROW_LUPP_DYN_W(o.block_row + p);
2277 TDLS_ARROW_LUPP_DYN_W(o.block_row + p) = TDLS_ARROW_LUPP_DYN_W(o.block_row + j);
2278 TDLS_ARROW_LUPP_DYN_W(o.block_row + j) = t;
2279 }
2280 const T inv = T(1) / TDLS_ARROW_LUPP_DYN_F(j, j);
2281 TDLS_ARROW_LUPP_DYN_F(j, j) = inv;
2282 for (int e = j + 1; e < front_rows; ++e) {
2283 const T l = TDLS_ARROW_LUPP_DYN_F(e, j) * inv;
2284 if (l == T(0)) continue;
2285 for (int c = j + 1; c < front_cols; ++c)
2287 }
2288 }
2289 if constexpr (!canonical) {
2290 for (int c = S; c < front_cols; ++c) {
2291 for (int i = S - 1; i >= 0; --i) {
2292 T t = TDLS_ARROW_LUPP_DYN_F(i, c);
2293 for (int m = i + 1; m < S; ++m)
2296 }
2297 }
2298 for (int i = 0; i < S; ++i) {
2299 for (int c = 0; c < B; ++c)
2300 TDLS_ARROW_LUPP_DYN_A(storage_map::akb(shape, k, i, c)) =
2301 TDLS_ARROW_LUPP_DYN_F(i, S + c);
2302 for (int a = 0; a < R; ++a) {
2303 if (a < S)
2304 TDLS_ARROW_LUPP_DYN_A(storage_map::d(shape, k, i, a)) =
2305 TDLS_ARROW_LUPP_DYN_F(i, S + B + a);
2306 else
2307 TDLS_ARROW_LUPP_DYN_W(o.extra + (k * S + i) * z.extra_cols + a - S) =
2308 TDLS_ARROW_LUPP_DYN_F(i, S + B + a);
2309 }
2310 for (int m = 0; m < nc; ++m)
2311 TDLS_ARROW_LUPP_DYN_YW(m, B + k * S + i) =
2312 TDLS_ARROW_LUPP_DYN_F(i, rhs0 + m);
2313 }
2314 }
2315 }
2316 for (int r = 0; r < B; ++r) {
2317 for (int c = 0; c < B; ++c)
2318 TDLS_ARROW_LUPP_DYN_W(o.Sbb + r * B + c) = TDLS_ARROW_LUPP_DYN_F(S + r, S + c);
2319 for (int m = 0; m < nc; ++m)
2320 TDLS_ARROW_LUPP_DYN_W(o.rb + m * B + r) = TDLS_ARROW_LUPP_DYN_F(S + r, rhs0 + m);
2321 }
2322 return true;
2323 }
2324};
2325
2326
2327
2328#undef TDLS_ARROW_LUPP_DYN_A
2329#undef TDLS_ARROW_LUPP_DYN_Y
2330#undef TDLS_ARROW_LUPP_DYN_YW
2331#undef TDLS_ARROW_LUPP_DYN_W
2332#undef TDLS_ARROW_LUPP_DYN_F
2333#undef TDLS_ARROW_LUPP_DYN_LU
2334#undef TDLS_ARROW_LUPP_DYN_RC
2335
2336
2337
2338} // namespace tdls
2339
2340
2341
2342#endif // TDLS_SOLVERS_ARROWHEAD_LUPP_SOLVER_DYNAMIC_HPP
Shape and compile-time configuration of the ArrowheadLUpp solver family.
#define TDLS_ARROW_LUPP_DYN_A(idx)
Strided stored entry at index idx of the matrix.
Definition solver_dynamic.hpp:78
#define TDLS_ARROW_LUPP_DYN_Y(i)
Strided entry i of the right-hand side and solution.
Definition solver_dynamic.hpp:81
#define TDLS_ARROW_LUPP_DYN_F(r, c)
Entry (r, c) of the front of the second pass, row-major at the start of the workspace.
Definition solver_dynamic.hpp:96
#define TDLS_ARROW_LUPP_DYN_W(e)
Strided entry e of the workspace.
Definition solver_dynamic.hpp:92
#define TDLS_ARROW_LUPP_DYN_YW(m, i)
Strided entry i of column m of the right-hand sides and solutions. The function declares ycols,...
Definition solver_dynamic.hpp:87
Toolchain detection and portability macros.
#define TDLS_FORCE_INLINE_ALL
Inlining forced on every backend, CPU included: TDLS_FORCE_INLINE_GPU on GPU, __forceinline under MSV...
Definition macros.hpp:101
#define TDLS_RESTRICT
Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC; nothing in the code that nvc++ gene...
Definition macros.hpp:174
#define TDLS_STRIDED_WIDE(i, s)
Offset of entry i of a strided object, i times stride s, for a 64-bit stride s (std::size_t); i is a ...
Definition macros.hpp:227
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:50
#define TDLS_STRIDED(i, s)
Offset of entry i of a strided object, i times stride s, in 64-bit arithmetic (std::size_t); i and s ...
Definition macros.hpp:206
#define TDLS_NO_INLINE
Keeps a function out of line on every backend.
Definition macros.hpp:146
Scalar math helpers usable in constant expressions.
RowInterchange
Row interchanges of a factorization with partial pivoting.
Definition options.hpp:59
Position of every entry of an arrowhead matrix in its storage.
Block k of a recomputed solve (solve_inplace_recomputed, solve_inplace_recomputed_or_tangent): a Stre...
Definition solver_dynamic.hpp:690
bool recovery
Definition solver_dynamic.hpp:694
Block k of a streamed solve: a view of the areas of the workspace where the solver factors and solves...
Definition solver_dynamic.hpp:642
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL constexpr T & y(const int i) const noexcept
Definition solver_dynamic.hpp:666
int stride
Definition solver_dynamic.hpp:648
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL constexpr T & coupling(const int i) const noexcept
Definition solver_dynamic.hpp:671
T * Dl
D_k in the layout of the factorization.
Definition solver_dynamic.hpp:643
int S
block size
Definition solver_dynamic.hpp:646
T * Abk
A_bk: coupled rows (diagonal) or B x S (dense)
Definition solver_dynamic.hpp:645
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL constexpr T & akb(const int i, const int c) const noexcept
Definition solver_dynamic.hpp:659
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL constexpr T & abk(const int r, const int i) const noexcept
Definition solver_dynamic.hpp:677
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL constexpr T & d(const int i, const int j) const noexcept
Definition solver_dynamic.hpp:652
T * W
[A_kb | y_k], column c at c * S
Definition solver_dynamic.hpp:644
int B
border size
Definition solver_dynamic.hpp:647
Arrowhead solver with a runtime shape.
Definition solver_dynamic.hpp:112
static TDLS_HOST_DEVICE constexpr int pivot_size(const ArrowheadShape &shape) noexcept
Entries of the pivot array: the pivots of a block, then of the border system, reuse it.
Definition solver_dynamic.hpp:336
StreamedOutcome
Outcome of solve_inplace_streamed_or_tangent.
Definition solver_dynamic.hpp:698
@ singular
the matrix is singular, or the shape rejected
@ solved
y holds the solution (converged() returned false)
static TDLS_HOST_DEVICE constexpr bool solve_inplace(const ArrowheadShape &shape, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, T *TDLS_RESTRICT y, const int rhs_stride, int &refused_count) noexcept
Solves A x = y in place, counting the refused blocks.
Definition solver_dynamic.hpp:358
static TDLS_HOST_DEVICE constexpr bool valid_recomputed(const ArrowheadShape &shape, const int capacity, const int nrhs=1) noexcept
Whether the recomputed entry points accept a shape, a capacity and nrhs tangent columns: valid(shape,...
Definition solver_dynamic.hpp:940
static TDLS_HOST_DEVICE constexpr int workspace_size(const ArrowheadShape &shape) noexcept
Entries of the workspace of solve_inplace and solve.
Definition solver_dynamic.hpp:292
static TDLS_HOST_DEVICE constexpr RecomputedOutcome solve_inplace_recomputed_or_tangent(const ArrowheadShape &shape, const int ntan, const int capacity, Provider &&provider, Converged &&converged, const T *TDLS_RESTRICT Abb, const int Abb_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, T *TDLS_RESTRICT y, const int rhs_stride, const int col0, T *TDLS_RESTRICT x, const int x_stride, const std::size_t xcol_stride, int &refused_count) noexcept
One recomputed pass of a Newton iteration: the runtime counterpart of ArrowheadLUppSolverStatic::solv...
Definition solver_dynamic.hpp:1056
static TDLS_HOST_DEVICE constexpr bool solve_border_canonical_multirhs(const ArrowheadShape &shape, const int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, const int col0, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, int &refused_count) noexcept
Columns col0 .. col0 + nrhs - 1 of the border block of A^-1 (col0 + nrhs <= border),...
Definition solver_dynamic.hpp:594
static TDLS_HOST_DEVICE constexpr int workspace_size(const ArrowheadShape &shape, const int nrhs) noexcept
Entries of the workspace of the entry points with nrhs right-hand sides: the temporaries of the right...
Definition solver_dynamic.hpp:303
static TDLS_HOST_DEVICE constexpr bool solve_multirhs(const ArrowheadShape &shape, const int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, int &refused_count) noexcept
Solves A X = B for nrhs right-hand sides, B left untouched, counting the refused blocks: B is copied ...
Definition solver_dynamic.hpp:527
static TDLS_HOST_DEVICE constexpr bool solve_inplace(const ArrowheadShape &shape, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, T *TDLS_RESTRICT y, const int rhs_stride) noexcept
Solves A x = y in place.
Definition solver_dynamic.hpp:379
static TDLS_HOST_DEVICE constexpr bool valid(const ArrowheadShape &shape, const int nrhs=1) noexcept
Whether the solver accepts a shape, with nrhs right-hand sides (unit columns, tangent columns): stora...
Definition solver_dynamic.hpp:281
static TDLS_HOST_DEVICE constexpr bool solve_inplace_multirhs(const ArrowheadShape &shape, const int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, T *TDLS_RESTRICT y, const int rhs_stride, const std::size_t xcol_stride, int &refused_count) noexcept
Solves A X = Y in place for nrhs right-hand sides, counting the refused blocks. The right-hand sides ...
Definition solver_dynamic.hpp:470
static TDLS_HOST_DEVICE constexpr bool solve_inplace_streamed(const ArrowheadShape &shape, Provider &&provider, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, T *TDLS_RESTRICT y, const int rhs_stride, int &refused_count) noexcept
Solves A x = y in place, the blocks streamed: provider(k, blk) fills block k (StreamedBlock) when the...
Definition solver_dynamic.hpp:733
static TDLS_HOST_DEVICE constexpr int streamed_workspace_size(const ArrowheadShape &shape) noexcept
Entries of the workspace of solve_inplace_streamed: those of solve_inplace, then A_bk of a block unde...
Definition solver_dynamic.hpp:315
static TDLS_HOST_DEVICE constexpr int recomputed_workspace_size(const ArrowheadShape &shape, const int capacity, const int nrhs=1) noexcept
Entries of the workspace of solve_inplace_recomputed with a store of capacity refused blocks,...
Definition solver_dynamic.hpp:961
static TDLS_HOST_DEVICE constexpr int streamed_workspace_size(const ArrowheadShape &shape, const int nrhs) noexcept
Entries of the workspace of solve_border_canonical_multirhs_streamed with nrhs unit columns,...
Definition solver_dynamic.hpp:327
static TDLS_HOST_DEVICE constexpr bool solve_inplace_multirhs(const ArrowheadShape &shape, const int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, T *TDLS_RESTRICT y, const int rhs_stride, const std::size_t xcol_stride) noexcept
Solves A X = Y in place for nrhs right-hand sides.
Definition solver_dynamic.hpp:496
static TDLS_HOST_DEVICE constexpr bool solve_border_canonical_multirhs(const ArrowheadShape &shape, const int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, const int col0, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Columns col0 .. col0 + nrhs - 1 of the border block of A^-1 (col0 + nrhs <= border).
Definition solver_dynamic.hpp:622
static TDLS_HOST_DEVICE constexpr bool solve_multirhs(const ArrowheadShape &shape, const int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Solves A X = B for nrhs right-hand sides, B left untouched.
Definition solver_dynamic.hpp:562
static TDLS_HOST_DEVICE constexpr bool solve(const ArrowheadShape &shape, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Solves A x = b, b left untouched.
Definition solver_dynamic.hpp:435
RecomputedOutcome
Outcome of solve_inplace_recomputed and solve_inplace_recomputed_or_tangent.
Definition solver_dynamic.hpp:919
static TDLS_HOST_DEVICE constexpr bool solve(const ArrowheadShape &shape, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, int &refused_count) noexcept
Solves A x = b, b left untouched, counting the refused blocks: b is copied into x,...
Definition solver_dynamic.hpp:406
static TDLS_HOST_DEVICE constexpr bool solve_border_canonical_multirhs_streamed(const ArrowheadShape &shape, const int nrhs, Provider &&provider, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, const int col0, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, int &refused_count) noexcept
solve_border_canonical_multirhs with the blocks streamed (see solve_inplace_streamed; blk....
Definition solver_dynamic.hpp:766
static TDLS_HOST_DEVICE constexpr StreamedOutcome solve_inplace_streamed_or_tangent(const ArrowheadShape &shape, const int ntan, Provider &&provider, Converged &&converged, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, T *TDLS_RESTRICT y, const int rhs_stride, const int col0, T *TDLS_RESTRICT x, const int x_stride, const std::size_t xcol_stride, int &refused_count) noexcept
One streamed pass of a Newton iteration: the blocks condensed as solve_inplace_streamed does,...
Definition solver_dynamic.hpp:820
static TDLS_HOST_DEVICE constexpr RecomputedOutcome solve_inplace_recomputed(const ArrowheadShape &shape, const int capacity, Provider &&provider, const T *TDLS_RESTRICT Abb, const int Abb_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT work, const int work_stride, T *TDLS_RESTRICT y, const int rhs_stride, int &refused_count) noexcept
Solves A x = y in place, the blocks recomputed: the runtime counterpart of ArrowheadLUppSolverStatic:...
Definition solver_dynamic.hpp:1004
Shape of an arrowhead matrix: a template argument of the compile-time solver, a function argument of ...
Definition config.hpp:111
int blocks
Number of diagonal blocks.
Definition config.hpp:117
Index of every entry of an arrowhead matrix in its storage.
Definition storage_map.hpp:56
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int abb(const ArrowheadShape &s, const int r, const int c) noexcept
Border-border entry (r, c), r and c < border.
Definition storage_map.hpp:121
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int size(const ArrowheadShape &s) noexcept
Number of stored entries.
Definition storage_map.hpp:111
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int abk(const ArrowheadShape &s, const int k, const int r, const int i) noexcept
Coupling entry of border row r and unknown i of block k (dense coupling).
Definition storage_map.hpp:145
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int akb(const ArrowheadShape &s, const int k, const int i, const int c) noexcept
Entry of row i of block k in border column c.
Definition storage_map.hpp:189
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int d(const ArrowheadShape &s, const int k, const int i, const int j) noexcept
Entry (i, j) of the diagonal block k.
Definition storage_map.hpp:213
static TDLS_HOST_DEVICE constexpr bool valid(const ArrowheadShape &s) noexcept
Whether a shape suits this storage and coupling: positive sizes, coupled rows inside the border (diag...
Definition storage_map.hpp:69
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int abk_diag(const ArrowheadShape &s, const int k, const int i) noexcept
Coupling entry of border row coupling_row + i and unknown i of block k (diagonal coupling).
Definition storage_map.hpp:167
Runtime-size tiled dense LU factorization with logical partial pivoting and out-of-tile pivot recover...
Definition solver_dynamic.hpp:176
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool factorize(int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, int &oot_count, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
Factor A := P*L*U in place.
Definition solver_dynamic.hpp:1329
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_multirhs(int n, int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Solve nrhs right-hand-side columns from a prior factorize: X := U^-1 L^-1 P B (column-major layout: X...
Definition solver_dynamic.hpp:2114
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int physical_factorize_checked(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T u) noexcept
Physical factorization with its multipliers bounded, the check fused into the elimination (the conden...
Definition solver_dynamic.hpp:1167
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void forward_inplace_multirhs(int n, int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
First of the two stages of substitute_inplace_multirhs, in passes of pass_width columns: x := L^-1 P ...
Definition solver_dynamic.hpp:2347
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_forward(const int n, const int w, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
First stage of the physical substitution of w columns: x := L^-1 P x under the row-major layout,...
Definition solver_dynamic.hpp:1212
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace(int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride, int &oot_count) noexcept
Factorization with the forward substitution folded in.
Definition solver_dynamic.hpp:2558
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace_multirhs(int n, int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride, const std::size_t xcol_stride, int &oot_count) noexcept
factorize + substitute_inplace_multirhs in one call: y holds the nrhs unpermuted right-hand-side colu...
Definition solver_dynamic.hpp:2632
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_backward(const int n, const int w, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Second stage of the physical substitution of w columns: x := U^-1 x under the row-major layout,...
Definition solver_dynamic.hpp:1249
static constexpr T singular_floor
Singularity floor of the out-of-tile recovery, read once from the configuration (see TiledLUppConfig:...
Definition solver_dynamic.hpp:190
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_inplace_multirhs(int n, int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
nrhs columns solved in place: x holds the nrhs unpermuted right-hand-side columns on entry and the nr...
Definition solver_dynamic.hpp:2301
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void backward_inplace_multirhs(int n, int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Second of the two stages of substitute_inplace_multirhs, in passes of pass_width columns: x := U^-1 x...
Definition solver_dynamic.hpp:2387
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_inplace(int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Solve in place: x already holds the unpermuted RHS on entry and the solution on exit.
Definition solver_dynamic.hpp:2159
Runtime-size variant of the TiledLUpp solver.