1#ifndef TDLS_SOLVERS_TILED_LUPP_SOLVER_DYNAMIC_HPP
2#define TDLS_SOLVERS_TILED_LUPP_SOLVER_DYNAMIC_HPP
79#if defined(__GNUC__) && !defined(__clang__)
80#pragma GCC diagnostic push
81#pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
82#pragma GCC diagnostic ignored "-Warray-bounds"
97#define TDLS_LUPP_DYN_A(r, c) A[TDLS_LAYOUT_INDEX(r, c, n) * unsigned(A_stride)]
100#define TDLS_LUPP_DYN_PIV(i) piv[unsigned(i) * unsigned(piv_stride)]
103#define TDLS_LUPP_DYN_X(i) x[unsigned(i) * unsigned(rhs_stride)]
106#define TDLS_LUPP_DYN_B(i) b[unsigned(i) * unsigned(rhs_stride)]
109#define TDLS_LUPP_DYN_XW(w, i) \
110 x[unsigned(i) * unsigned(rhs_stride) + unsigned(w) * unsigned(xcol_stride)]
114#define TDLS_LUPP_DYN_BW(w, i) \
115 b[unsigned(i) * unsigned(rhs_stride) + unsigned(w) * unsigned(xcol_stride)]
118#define TDLS_LUPP_DYN_Y(i) y[unsigned(i) * unsigned(rhs_stride)]
158template<
typename T, TiledLUppConfig<T> Config = TiledLUppConfig<T>{}>
173 Config.oot_threshold.is_finite() && Config.singular_floor.is_finite(),
174 "TiledLUppSolverDynamic: oot_threshold and singular_floor must be finite (and fit a "
176 static_assert(
singular_floor > T(0),
"TiledLUppSolverDynamic: singular_floor must be positive");
178 "TiledLUppSolverDynamic: singular_floor must not exceed oot_threshold (the "
179 "floor applies to the out-of-tile recovery path)");
180 static_assert(
tile_size >= 1,
"TiledLUppSolverDynamic: tile size must be >= 1");
193 const int n)
noexcept {
210 for (
int j = 0; j < ke; ++j) {
228 const T inv_pivot = T(1) / t[k *
tile_size + k];
230 for (
int i = k + 1; i < re; ++i) {
232 for (
int j = k + 1; j < ce; ++j)
245 const int ce)
noexcept {
246 for (
int k = 0; k < kd; ++k) {
247 for (
int i = k + 1; i < kd; ++i) {
249 for (
int j = 0; j < ce; ++j)
266 const int re)
noexcept {
267 for (
int k = 0; k < kd; ++k) {
268 const T U_kk_inv = lu[k *
tile_size + k];
269 for (
int i = 0; i < re; ++i)
271 for (
int j = k + 1; j < kd; ++j) {
273 for (
int i = 0; i < re; ++i)
289 const int re,
const int ce,
const int kd)
noexcept {
290 for (
int i = 0; i < re; ++i) {
291 for (
int j = 0; j < ce; ++j) {
293 for (
int k = 0; k < kd; ++k)
316 const int ce)
noexcept {
317 for (
int i = 0; i < re; ++i) {
318 for (
int j = 0; j < ce; ++j)
334 const int col0,
const T*
TDLS_RESTRICT t,
const int re,
const int ce)
noexcept {
335 for (
int i = 0; i < re; ++i) {
336 for (
int j = 0; j < ce; ++j)
355 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int row0,
356 const int col0, T*
TDLS_RESTRICT t,
const int re,
const int ce)
noexcept {
357 for (
int i = 0; i < re; ++i) {
359 for (
int j = 0; j < ce; ++j)
378 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int row0,
379 const int col0,
const T*
TDLS_RESTRICT t,
const int re,
const int ce)
noexcept {
380 for (
int i = 0; i < re; ++i) {
382 for (
int j = 0; j < ce; ++j)
400 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int row0,
402 for (
int i = 1; i < re; ++i) {
404 for (
int j = 0; j < i; ++j)
423 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int row0,
425 for (
int i = 0; i < re; ++i) {
427 for (
int j = i; j < re; ++j)
456 template<
bool oot_diagnostics,
bool fuse_rhs>
459 const int piv_stride,
const int k0, T*
TDLS_RESTRICT tile,
int& oot_count,
461 const int rhs_stride)
noexcept {
463 const int gc = k0 + c;
467 T best = detail::abs(tile[c *
tile_size + c]);
468 for (
int r = c + 1; r < ke; ++r) {
469 const T v = detail::abs(tile[r *
tile_size + c]);
479 piv_row = k0 + best_r;
485 if constexpr (oot_diagnostics) ++oot_count;
486 piv_row = k0 + best_r;
492 if constexpr (oot_diagnostics) ++oot_count;
494 int gbest_row = k0 + best_r;
496 for (
int row = k0 + ke; row < n; ++row) {
500 if constexpr (
schedule == Schedule::LeftLooking) {
501 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
509 for (
int t = 0; t < c; ++t) {
511 if constexpr (
schedule == Schedule::LeftLooking) {
512 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
517 for (
int p = 0; p < t; ++p)
518 a_t -= L_row[p] * tile[p *
tile_size + t];
519 L_row[t] = a_t * tile[t *
tile_size + t];
520 corrected -= L_row[t] * tile[t *
tile_size + c];
524 const T v = detail::abs(corrected);
533 if constexpr (Config.oot_first_acceptable)
548 if constexpr (fuse_rhs) {
555 if (piv_row < k0 + ke) {
562 for (
int j = 0; j < ke; ++j) {
564 if constexpr (
schedule == Schedule::LeftLooking) {
565 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
575 for (
int t = 0; t < c; ++t) {
577 for (
int p = 0; p < t; ++p)
578 a_t -= L_row[p] * tile[p *
tile_size + t];
579 L_row[t] = a_t * tile[t *
tile_size + t];
582 for (
int j = c; j < ke; ++j) {
583 for (
int t = 0; t < c; ++t)
612 template<
bool oot_diagnostics,
bool fuse_rhs>
615 const int piv_stride,
const int k0, T*
TDLS_RESTRICT tile,
int& oot_count,
617 const int rhs_stride = 1) noexcept {
618 for (
int c = 0; c < ke; ++c) {
620 n, A, A_stride, piv, piv_stride, k0, tile, oot_count, c, ke, y, rhs_stride))
642 const int ke,
const int je)
noexcept {
644 load_tile(n, A, A_stride, pk, j0, Akj, ke, je);
646 store_tile(n, A, A_stride, pk, j0, Akj, ke, je);
664 const int ke,
const int ie,
const int je)
noexcept {
666 load_tile(n, A, A_stride, pi, j0, Aij, ie, je);
668 for (
int p = 0; p < ke; ++p) {
670 for (
int j = 0; j < je; ++j)
672 for (
int i = 0; i < ie; ++i) {
674 for (
int j = 0; j < je; ++j)
675 Aij[i *
tile_size + j] -= L_ip * Akj_row[j];
679 store_tile(n, A, A_stride, pi, j0, Aij, ie, je);
698 template<
bool fuse_rhs>
703 const int i0,
const int ke,
const int ie, T*
TDLS_RESTRICT y =
nullptr,
704 const int rhs_stride = 1) noexcept {
709 for (
int i = 0; i < ie; ++i)
714 load_tile(n, A, A_stride, pi, k0, Aik, ie, ke);
716 store_tile(n, A, A_stride, pi, k0, Aik, ie, ke);
720 if constexpr (fuse_rhs) {
721 for (
int r = 0; r < ie; ++r) {
723 for (
int j = 0; j < ke; ++j)
730 for (
int tj = k + 1; tj < nt; ++tj)
749 template<
bool oot_diagnostics,
bool fuse_rhs>
752 const int piv_stride,
const int k,
int& oot_count, T*
TDLS_RESTRICT y =
nullptr,
753 const int rhs_stride = 1) noexcept {
759 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
762 oot_count, ke, y, rhs_stride))
767 for (
int i = 0; i < ke; ++i)
770 store_tile(n, A, A_stride, pk, k0, tile, ke, ke);
774 if constexpr (fuse_rhs) {
775 for (
int kk = 0; kk < ke; ++kk) {
776 for (
int i = kk + 1; i < ke; ++i)
782 for (
int tj = k + 1; tj < nt; ++tj)
787 for (
int ti = k + 1; ti < nt; ++ti)
814 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int row0,
815 const int col0,
const int k0, T*
TDLS_RESTRICT t,
const int re,
816 const int ce)
noexcept {
817 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size) {
841 template<
bool fuse_rhs>
844 const int piv_stride,
const T*
TDLS_RESTRICT tile,
const int k0,
const int i0,
const int ke,
845 const int ie, T*
TDLS_RESTRICT y =
nullptr,
const int rhs_stride = 1) noexcept {
847 load_tile_piv(n, A, A_stride, piv, piv_stride, i0, k0, B, ie, ke);
848 ll_correct_tile(n, A, A_stride, piv, piv_stride, i0, k0, k0, B, ie, ke);
850 store_tile_piv(n, A, A_stride, piv, piv_stride, i0, k0, B, ie, ke);
854 if constexpr (fuse_rhs) {
855 for (
int r = 0; r < ie; ++r) {
857 for (
int j = 0; j < ke; ++j)
878 const T*
TDLS_RESTRICT tile,
const int k0,
const int j0,
const int ke,
879 const int je)
noexcept {
881 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
882 ll_correct_tile(n, A, A_stride, piv, piv_stride, k0, j0, k0, B, ke, je);
884 store_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
901 template<
bool oot_diagnostics,
bool fuse_rhs>
904 const int piv_stride,
const int k,
int& oot_count, T*
TDLS_RESTRICT y =
nullptr,
905 const int rhs_stride = 1) noexcept {
911 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
912 ll_correct_tile(n, A, A_stride, piv, piv_stride, k0, k0, k0, tile, ke, ke);
915 oot_count, ke, y, rhs_stride))
918 store_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
922 if constexpr (fuse_rhs) {
923 for (
int kk = 0; kk < ke; ++kk) {
924 for (
int i = kk + 1; i < ke; ++i)
930 for (
int ti = k + 1; ti < nt; ++ti)
935 for (
int tj = k + 1; tj < nt; ++tj)
965 template<
bool oot_diagnostics = true,
bool fuse_rhs = false>
968 const int piv_stride,
int& oot_count, T*
TDLS_RESTRICT y =
nullptr,
969 const int rhs_stride = 1) noexcept {
971 if constexpr (oot_diagnostics) oot_count = 0;
973 for (
int i = 0; i < n; ++i)
977 for (
int k = 0; k < nt; ++k) {
978 if constexpr (
schedule == Schedule::RightLooking) {
980 oot_count, y, rhs_stride))
984 oot_count, y, rhs_stride))
1002 const int piv_stride)
noexcept {
1029 const int rhs_stride,
const int xcol_stride,
const int k0,
const int m0,
1030 const int ke,
const int me)
noexcept {
1032 load_tile_piv(n, A, A_stride, piv, piv_stride, m0, k0, Lmk, me, ke);
1033 for (
int i = 0; i < me; ++i) {
1034 for (
int w = 0; w < nrhs; ++w) {
1036 for (
int j = 0; j < ke; ++j)
1058 const int rhs_stride,
const int xcol_stride,
const int k)
noexcept {
1067 for (
int kk = 0; kk < ke; ++kk) {
1068 for (
int i = kk + 1; i < ke; ++i) {
1069 for (
int w = 0; w < nrhs; ++w)
1076 for (
int m = k + 1; m < nt; ++m)
1077 fwd_push_one(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k0,
1099 const int rhs_stride,
const int xcol_stride,
const int k0,
const int m0,
1100 const int ke,
const int me)
noexcept {
1102 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, m0, Ukm, ke, me);
1103 for (
int i = 0; i < ke; ++i) {
1104 for (
int w = 0; w < nrhs; ++w) {
1106 for (
int j = 0; j < me; ++j)
1128 const int rhs_stride,
const int xcol_stride,
const int k)
noexcept {
1134 for (
int m = k + 1; m < nt; ++m)
1135 bwd_pull_one(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k0,
1142 for (
int kk = ke - 1; kk >= 0; --kk) {
1143 for (
int w = 0; w < nrhs; ++w)
1145 for (
int i = 0; i < kk; ++i) {
1146 for (
int w = 0; w < nrhs; ++w)
1167 const int rhs_stride,
const int xcol_stride)
noexcept {
1168 for (
int k =
num_tiles(n) - 1; k >= 0; --k)
1169 bwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1188 const int rhs_stride,
const int xcol_stride)
noexcept {
1190 for (
int k = 0; k < nt; ++k)
1191 fwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1193 bwd_only(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1217 unsigned long long visited = 0ull;
1218 for (
int s = 0; s < n; ++s) {
1219 if ((visited >> s) & 1ull)
continue;
1225 visited |= 1ull << cur;
1230 visited |= 1ull << cur;
1236 for (
int s = 0; s < n; ++s) {
1240 if (probe != s)
continue;
1274 for (
int i = 0; i < n; ++i)
1276 fwd_bwd(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
1295 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int col,
1320 const int rhs_stride,
const int xcol_stride)
noexcept {
1321 for (
int i = 0; i < n; ++i) {
1323 for (
int w = 0; w < nrhs; ++w)
1326 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1352 template<
int pass_w
idth = 0>
1357 const int rhs_stride,
const int xcol_stride)
noexcept {
1358 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
1359 if constexpr (pass_width == 0) {
1361 rhs_stride, xcol_stride);
1365 for (
int c0 = 0; c0 < nrhs; c0 += pass_width) {
1366 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
1367 const unsigned off = unsigned(c0) * unsigned(xcol_stride);
1369 x + off, rhs_stride, xcol_stride);
1391 const int A_stride,
const int*
TDLS_RESTRICT piv,
const int piv_stride,
1393 const int xcol_stride)
noexcept {
1394 for (
int i = 0; i < n; ++i) {
1396 for (
int w = 0; w < nrhs; ++w)
1399 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1428 template<
int pass_w
idth = 0>
1433 const int xcol_stride)
noexcept {
1434 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
1435 if constexpr (pass_width == 0) {
1441 for (
int c0 = 0; c0 < nrhs; c0 += pass_width) {
1442 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
1443 const unsigned off = unsigned(c0) * unsigned(xcol_stride);
1445 rhs_stride, xcol_stride);
1465 const int rhs_stride)
noexcept {
1467 fwd_bwd(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
1488 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1489 const int xcol_stride)
noexcept {
1490 for (
int w = 0; w < nrhs; ++w)
1491 permute_inplace(n, piv, piv_stride, x +
unsigned(w) *
unsigned(xcol_stride),
1493 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1519 template<
int pass_w
idth = 0>
1523 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1524 const int xcol_stride)
noexcept {
1525 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
1526 if constexpr (pass_width == 0) {
1532 for (
int c0 = 0; c0 < nrhs; c0 += pass_width) {
1533 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
1534 const unsigned off = unsigned(c0) * unsigned(xcol_stride);
1536 rhs_stride, xcol_stride);
1560 template<
bool oot_diagnostics = true>
1564 int& oot_count)
noexcept {
1566 substitute(n, A, A_stride, piv, piv_stride, b, x, rhs_stride);
1585 const int rhs_stride)
noexcept {
1587 return solve<false>(n, A, A_stride, piv, piv_stride, b, x, rhs_stride, unused);
1612 template<
int pass_w
idth = 0,
bool oot_diagnostics = true>
1616 T*
TDLS_RESTRICT x,
const int rhs_stride,
const int xcol_stride,
1617 int& oot_count)
noexcept {
1643 template<
int pass_w
idth = 0>
1647 T*
TDLS_RESTRICT x,
const int rhs_stride,
const int xcol_stride)
noexcept {
1650 rhs_stride, xcol_stride, unused);
1674 template<
bool oot_diagnostics = true>
1677 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride,
1678 int& oot_count)
noexcept {
1684 bwd_only(n, 1, A, A_stride, piv, piv_stride, y, rhs_stride, 0);
1701 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride)
noexcept {
1732 template<
int pass_w
idth = 0,
bool oot_diagnostics = true>
1736 const int rhs_stride,
const int xcol_stride,
int& oot_count)
noexcept {
1739 rhs_stride, xcol_stride);
1761 template<
int pass_w
idth = 0>
1765 const int rhs_stride,
const int xcol_stride)
noexcept {
1768 rhs_stride, xcol_stride, unused);
1774#undef TDLS_LUPP_DYN_A
1775#undef TDLS_LUPP_DYN_PIV
1776#undef TDLS_LUPP_DYN_X
1777#undef TDLS_LUPP_DYN_B
1778#undef TDLS_LUPP_DYN_XW
1779#undef TDLS_LUPP_DYN_BW
1780#undef TDLS_LUPP_DYN_Y
1786#if defined(__GNUC__) && !defined(__clang__)
1787#pragma GCC diagnostic pop
Compile-time configuration of the TiledLUpp solver family.
#define TDLS_RESTRICT
Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC).
Definition macros.hpp:80
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:45
#define TDLS_FORCEINLINE
__forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS_HOST_DEVICE for why),...
Definition macros.hpp:67
Scalar math helpers usable in constant expressions.
Schedule
Elimination schedule of a factorization.
Definition options.hpp:37
#define TDLS_LUPP_DYN_X(i)
Strided entry i of the solution vector.
Definition solver_dynamic.hpp:103
#define TDLS_LUPP_DYN_PIV(i)
Strided pivot entry i.
Definition solver_dynamic.hpp:100
#define TDLS_LUPP_DYN_Y(i)
Strided entry i of the fused right-hand side of solve_inplace.
Definition solver_dynamic.hpp:118
#define TDLS_LUPP_DYN_A(r, c)
Strided element (r, c) of the factor matrix, flat index remapped by Config.layout.
Definition solver_dynamic.hpp:97
#define TDLS_LUPP_DYN_XW(w, i)
Strided entry i of column w of a multi right-hand-side block.
Definition solver_dynamic.hpp:109
#define TDLS_LUPP_DYN_B(i)
Strided entry i of the right-hand side.
Definition solver_dynamic.hpp:106
#define TDLS_LUPP_DYN_BW(w, i)
Strided entry i of column w of a multi right-hand-side block of b, addressed exactly as TDLS_LUPP_DYN...
Definition solver_dynamic.hpp:114
Runtime-size tiled dense LU factorization with logical partial pivoting and out-of-tile pivot recover...
Definition solver_dynamic.hpp:159
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void bwd_only(const int n, const 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 int xcol_stride) noexcept
Backward pass alone, used by fwd_bwd and by solve_inplace (whose forward pass happens inside the fact...
Definition solver_dynamic.hpp:1165
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool ll_step(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k, int &oot_count, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
LL: one full factorization step (correct + factor the diagonal tile, then its L and U panels).
Definition solver_dynamic.hpp:903
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile_piv(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, T *TDLS_RESTRICT t, const int re, const int ce) noexcept
Load a re x ce tile, reading the permutation inline (one read per row).
Definition solver_dynamic.hpp:354
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace(const 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) noexcept
Diagnostics-free solve_inplace overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:1700
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile_piv_lower(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, T *TDLS_RESTRICT t, const int re) noexcept
Triangular variant of the diagonal-tile load for the forward substitution: only the strict lower tria...
Definition solver_dynamic.hpp:399
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void rl_schur_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT pk, const int *TDLS_RESTRICT pi, const T *TDLS_RESTRICT Aik, const int j0, const int ke, const int ie, const int je) noexcept
RL: one Schur-complement tile update, Aij -= Aik * Akj, streaming the factored Akj row by row from re...
Definition solver_dynamic.hpp:662
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ll_update_right_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT tile, const int k0, const int j0, const int ke, const int je) noexcept
LL: correct + TRSM one U-panel tile right of the diagonal.
Definition solver_dynamic.hpp:876
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factor_diag_tile(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k0, T *TDLS_RESTRICT tile, int &oot_count, const int ke, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
Factor the ke x ke diagonal tile in registers, with out-of-tile pivot recovery (drives the per-column...
Definition solver_dynamic.hpp:614
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace_multirhs(const int n, const 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 int xcol_stride) noexcept
Diagnostics-free solve_inplace_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:1763
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factorize(const 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:967
static constexpr Schedule schedule
elimination schedule (RightLooking or LeftLooking)
Definition solver_dynamic.hpp:162
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ops_eliminate_column(T *TDLS_RESTRICT t, const int k, const int re, const int ce) noexcept
Gaussian elimination of column k inside the diagonal tile.
Definition solver_dynamic.hpp:227
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void fwd_bwd(const int n, const 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 int xcol_stride) noexcept
Triangular solves on already-permuted column(s) x. nrhs columns are processed per tile visit,...
Definition solver_dynamic.hpp:1186
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void bwd_pull_one(const int n, const 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 int xcol_stride, const int k0, const int m0, const int ke, const int me) noexcept
Backward pull: subtract U(k,m) * x_m from the x_k segment, for nrhs columns at once.
Definition solver_dynamic.hpp:1097
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT prow, const int col0, T *TDLS_RESTRICT t, const int re, const int ce) noexcept
Load a re x ce tile through a cached physical-row segment.
Definition solver_dynamic.hpp:314
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical_multirhs_pass(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int col0, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
One pass of the canonical multi right-hand-side substitution: nrhs canonical columns e_col0 ....
Definition solver_dynamic.hpp:1317
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool rl_step(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k, int &oot_count, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
RL: one full factorization step (diagonal tile + trailing updates).
Definition solver_dynamic.hpp:751
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int tile_size_at(const int t0, const int n) noexcept
Extent of the tile starting at row/column t0: tile_size, or less for the last tile.
Definition solver_dynamic.hpp:192
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace_multirhs_pass(const int n, const 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 int xcol_stride) noexcept
One pass of the in-place multi right-hand-side substitution: nrhs columns of x permuted by the cycle ...
Definition solver_dynamic.hpp:1486
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_multirhs(const int n, const int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
Diagnostics-free solve_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:1645
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace(const 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:1463
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void permute_inplace(const int n, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
In-place application of the permutation to one strided column x, by cycle decomposition: the permute ...
Definition solver_dynamic.hpp:1212
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void fwd_push_one(const int n, const 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 int xcol_stride, const int k0, const int m0, const int ke, const int me) noexcept
Forward push: subtract L(m,k) * x_k from the x_m segment, for nrhs columns at once.
Definition solver_dynamic.hpp:1027
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factorize(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride) noexcept
Diagnostics-free factorize overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:1001
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace_multirhs(const int n, const 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 int 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:1734
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ops_trsm_right(const T *TDLS_RESTRICT lu, T *TDLS_RESTRICT B, const int kd, const int re) noexcept
B := B U^-1, with U the upper part of the factored diagonal tile. U is kd x kd, B is re x kd.
Definition solver_dynamic.hpp:265
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void rl_update_row_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int *TDLS_RESTRICT pk, const T *TDLS_RESTRICT tile, const int k, const int i0, const int ke, const int ie, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
RL: TRSM down + Schur sweep of one row block below the diagonal, with the optional fused forward-subs...
Definition solver_dynamic.hpp:700
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ll_correct_tile(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, const int k0, T *TDLS_RESTRICT t, const int re, const int ce) noexcept
LL: t (re x ce, rows row0.., cols col0..) -= sum over prior tiles bj < k0/tile_size of L(row0....
Definition solver_dynamic.hpp:813
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int col, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Solve with b = e_col generated on the fly: the consistent-tangent-operator path.
Definition solver_dynamic.hpp:1294
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void bwd_step(const int n, const 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 int xcol_stride, const int k) noexcept
Backward step: pull the trailing contributions, then upper-solve the diagonal tile's segment.
Definition solver_dynamic.hpp:1126
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_multirhs(const int n, const int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride, int &oot_count) noexcept
factorize + substitute_multirhs in one call.
Definition solver_dynamic.hpp:1614
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int num_tiles(const int n) noexcept
Number of tiles per dimension (last one possibly partial).
Definition solver_dynamic.hpp:184
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace_multirhs(const int n, const 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 int 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:1521
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, int &oot_count) noexcept
factorize + substitute in one call.
Definition solver_dynamic.hpp:1562
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Diagnostics-free solve overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:1583
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ll_update_below_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT tile, const int k0, const int i0, const int ke, const int ie, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
LL: correct + TRSM one L-panel tile below the diagonal, with the optional fused forward-substitution ...
Definition solver_dynamic.hpp:842
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace(const 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:1676
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ops_trsm_left_unit(const T *TDLS_RESTRICT lu, T *TDLS_RESTRICT B, const int kd, const int ce) noexcept
B := L^-1 B, with L the unit lower part of the factored diagonal tile. L is kd x kd,...
Definition solver_dynamic.hpp:244
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void store_tile(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT prow, const int col0, const T *TDLS_RESTRICT t, const int re, const int ce) noexcept
Store a re x ce tile through a cached physical-row segment.
Definition solver_dynamic.hpp:333
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_multirhs(const int n, const 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 int xcol_stride) noexcept
Solve nrhs right-hand-side columns from a prior factorize: X := U^-1 L^-1 P B, column by column....
Definition solver_dynamic.hpp:1430
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void rl_trsm_right_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT pk, const T *TDLS_RESTRICT tile, const int j0, const int ke, const int je) noexcept
RL: one row-panel tile update, Akj := L^-1 Akj.
Definition solver_dynamic.hpp:640
static constexpr T singular_floor
Singularity floor of the out-of-tile recovery, read once from the configuration (see TiledLUppConfig:...
Definition solver_dynamic.hpp:170
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factor_diag_column(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k0, T *TDLS_RESTRICT tile, int &oot_count, const int c, const int ke, T *TDLS_RESTRICT y, const int rhs_stride) noexcept
One column step of the diagonal-tile factorization: pivot search (in-tile, then out-of-tile recovery)...
Definition solver_dynamic.hpp:458
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute(const int n, 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) noexcept
Solve x := U^-1 L^-1 P b from a prior factorize. b and x must not alias (use substitute_inplace for t...
Definition solver_dynamic.hpp:1271
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile_piv_upper(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, T *TDLS_RESTRICT t, const int re) noexcept
Triangular variant of the diagonal-tile load for the backward substitution: only the upper triangle i...
Definition solver_dynamic.hpp:422
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void store_tile_piv(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, const T *TDLS_RESTRICT t, const int re, const int ce) noexcept
Store a re x ce tile, reading the permutation inline (one read per row).
Definition solver_dynamic.hpp:377
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ops_swap_rows(T *TDLS_RESTRICT t, const int k, const int r, const int ke) noexcept
Row swap k <-> r inside the tile (direct, first ke columns).
Definition solver_dynamic.hpp:209
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical_multirhs(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int col0, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
nrhs canonical columns e_col0 .. e_{col0+nrhs-1} solved from a prior factorize: the columns of the in...
Definition solver_dynamic.hpp:1354
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ops_gemm_sub(T *TDLS_RESTRICT Ct, const T *TDLS_RESTRICT At, const T *TDLS_RESTRICT Bt, const int re, const int ce, const int kd) noexcept
Ct -= At*Bt with per-element dot-product accumulation. At is re x kd, Bt is kd x ce,...
Definition solver_dynamic.hpp:288
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void fwd_step(const int n, const 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 int xcol_stride, const int k) noexcept
Forward step: unit-lower solve the diagonal tile's segment, then push it into the tiles below.
Definition solver_dynamic.hpp:1056
static constexpr int tile_size
tile size (int)
Definition solver_dynamic.hpp:161
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_multirhs_pass(const int n, const 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 int xcol_stride) noexcept
One pass of the multi right-hand-side substitution: nrhs columns of b gathered in permuted order into...
Definition solver_dynamic.hpp:1390
static constexpr T oot_threshold
Acceptable-pivot threshold of the out-of-tile search, read once from the configuration (see TiledLUpp...
Definition solver_dynamic.hpp:167