1#ifndef TDLS_SOLVERS_TILED_LUPP_SOLVER_DYNAMIC_HPP
2#define TDLS_SOLVERS_TILED_LUPP_SOLVER_DYNAMIC_HPP
93#if defined(TDLS_DETAIL_COMPILER_GCC)
94#pragma GCC diagnostic push
95#pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
96#pragma GCC diagnostic ignored "-Warray-bounds"
111#define TDLS_LUPP_DYN_A(r, c) A[TDLS_STRIDED((r) * n + (c), A_stride)]
114#define TDLS_LUPP_DYN_PIV(i) piv[TDLS_STRIDED(i, piv_stride)]
117#define TDLS_LUPP_DYN_X(i) x[TDLS_STRIDED(i, rhs_stride)]
120#define TDLS_LUPP_DYN_B(i) b[TDLS_STRIDED(i, rhs_stride)]
123#define TDLS_LUPP_DYN_XW(w, i) x[TDLS_STRIDED(i, rhs_stride) + TDLS_STRIDED_WIDE(w, xcol_stride)]
127#define TDLS_LUPP_DYN_BW(w, i) b[TDLS_STRIDED(i, rhs_stride) + TDLS_STRIDED_WIDE(w, xcol_stride)]
130#define TDLS_LUPP_DYN_Y(i) y[TDLS_STRIDED(i, rhs_stride)]
175template<
typename T, TiledLUppConfig<T> Config = TiledLUppConfig<T>{}>
194 static constexpr bool transposed = Config.layout == MatrixLayout::ColMajor;
197 static constexpr bool physical = Config.row_interchange == RowInterchange::Physical;
199 static_assert(Config.row_interchange != RowInterchange::Auto,
200 "TiledLUppSolverDynamic: row_interchange must be Logical or Physical (Auto is "
201 "the default of the blocks of ArrowheadLUpp only)");
203 Config.oot_pivot_threshold.is_finite() && Config.singular_floor.is_finite(),
204 "TiledLUppSolverDynamic: oot_pivot_threshold and singular_floor must be finite (and fit a "
206 static_assert(Config.relative_pivot_threshold.is_finite(),
207 "TiledLUppSolverDynamic: relative_pivot_threshold must be finite (and fit a "
210 "TiledLUppSolverDynamic: relative_pivot_threshold must be in (0, 1]");
211 static_assert(
singular_floor > T(0),
"TiledLUppSolverDynamic: singular_floor must be positive");
213 "TiledLUppSolverDynamic: singular_floor must not exceed oot_pivot_threshold (the "
214 "floor applies to the out-of-tile recovery path)");
215 static_assert(
tile_size >= 1,
"TiledLUppSolverDynamic: tile size must be >= 1");
230 const int n)
noexcept {
247 for (
int j = 0; j < ke; ++j) {
265 const T inv_pivot = T(1) / t[k *
tile_size + k];
267 for (
int i = k + 1; i < re; ++i) {
269 for (
int j = k + 1; j < ce; ++j)
282 const int ce)
noexcept {
283 for (
int k = 0; k < kd; ++k) {
284 for (
int i = k + 1; i < kd; ++i) {
286 for (
int j = 0; j < ce; ++j)
303 const int re)
noexcept {
304 for (
int k = 0; k < kd; ++k) {
305 const T U_kk_inv = lu[k *
tile_size + k];
306 for (
int i = 0; i < re; ++i)
308 for (
int j = k + 1; j < kd; ++j) {
310 for (
int i = 0; i < re; ++i)
326 const int re,
const int ce,
const int kd)
noexcept {
327 for (
int i = 0; i < re; ++i) {
328 for (
int j = 0; j < ce; ++j) {
330 for (
int k = 0; k < kd; ++k)
353 const int ce)
noexcept {
354 for (
int i = 0; i < re; ++i) {
355 for (
int j = 0; j < ce; ++j)
371 const int col0,
const T*
TDLS_RESTRICT t,
const int re,
const int ce)
noexcept {
372 for (
int i = 0; i < re; ++i) {
373 for (
int j = 0; j < ce; ++j)
392 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int row0,
393 const int col0, T*
TDLS_RESTRICT t,
const int re,
const int ce)
noexcept {
394 for (
int i = 0; i < re; ++i) {
396 for (
int j = 0; j < ce; ++j)
415 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int row0,
416 const int col0,
const T*
TDLS_RESTRICT t,
const int re,
const int ce)
noexcept {
417 for (
int i = 0; i < re; ++i) {
419 for (
int j = 0; j < ce; ++j)
438 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int row0,
440 for (
int i = 1; i < re; ++i) {
442 for (
int j = 0; j < i; ++j)
461 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int row0,
463 for (
int i = 0; i < re; ++i) {
465 for (
int j = i; j < re; ++j)
495 template<
bool oot_diagnostics,
bool fuse_rhs>
498 const int piv_stride,
const int k0, T*
TDLS_RESTRICT tile,
int& oot_count,
500 const int rhs_stride)
noexcept {
502 const int gc = k0 + c;
506 T best = detail::abs(tile[c *
tile_size + c]);
507 for (
int r = c + 1; r < ke; ++r) {
508 const T v = detail::abs(tile[r *
tile_size + c]);
521 const T in_place = detail::abs(tile[c *
tile_size + c]);
529 piv_row = k0 + best_r;
535 if constexpr (oot_diagnostics) ++oot_count;
536 piv_row = k0 + best_r;
542 if constexpr (oot_diagnostics) ++oot_count;
544 int gbest_row = k0 + best_r;
546 for (
int row = k0 + ke; row < n; ++row) {
550 if constexpr (
schedule == Schedule::LeftLooking) {
551 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
559 for (
int t = 0; t < c; ++t) {
561 if constexpr (
schedule == Schedule::LeftLooking) {
562 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
567 for (
int p = 0; p < t; ++p)
568 a_t -= L_row[p] * tile[p *
tile_size + t];
569 L_row[t] = a_t * tile[t *
tile_size + t];
570 corrected -= L_row[t] * tile[t *
tile_size + c];
574 const T v = detail::abs(corrected);
583 if constexpr (Config.oot_first_acceptable)
607 if (piv_row < k0 + ke) {
614 for (
int j = 0; j < ke; ++j) {
616 if constexpr (
schedule == Schedule::LeftLooking) {
617 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
627 for (
int t = 0; t < c; ++t) {
629 for (
int p = 0; p < t; ++p)
630 a_t -= L_row[p] * tile[p *
tile_size + t];
631 L_row[t] = a_t * tile[t *
tile_size + t];
634 for (
int j = c; j < ke; ++j) {
635 for (
int t = 0; t < c; ++t)
665 template<
bool oot_diagnostics,
bool fuse_rhs>
668 const int piv_stride,
const int k0, T*
TDLS_RESTRICT tile,
int& oot_count,
670 const int rhs_stride = 1) noexcept {
671 for (
int c = 0; c < ke; ++c) {
673 n, A, A_stride, piv, piv_stride, k0, tile, oot_count, c, ke, y, rhs_stride))
695 const int ke,
const int je)
noexcept {
697 load_tile(n, A, A_stride, pk, j0, Akj, ke, je);
699 store_tile(n, A, A_stride, pk, j0, Akj, ke, je);
720 const int ke,
const int je,
const int k0, T*
TDLS_RESTRICT y,
721 const int rhs_stride)
noexcept {
723 load_tile(n, A, A_stride, pk, j0, Akj, ke, je);
725 store_tile(n, A, A_stride, pk, j0, Akj, ke, je);
744 const int ke,
const int ie,
const int je)
noexcept {
746 load_tile(n, A, A_stride, pi, j0, Aij, ie, je);
748 for (
int p = 0; p < ke; ++p) {
750 for (
int j = 0; j < je; ++j)
752 for (
int i = 0; i < ie; ++i) {
754 for (
int j = 0; j < je; ++j)
755 Aij[i *
tile_size + j] -= L_ip * Akj_row[j];
759 store_tile(n, A, A_stride, pi, j0, Aij, ie, je);
778 template<
bool fuse_rhs>
783 const int i0,
const int ke,
const int ie, T*
TDLS_RESTRICT y =
nullptr,
784 const int rhs_stride = 1) noexcept {
789 for (
int i = 0; i < ie; ++i)
794 load_tile(n, A, A_stride, pi, k0, Aik, ie, ke);
796 store_tile(n, A, A_stride, pi, k0, Aik, ie, ke);
802 for (
int r = 0; r < ie; ++r) {
804 for (
int j = 0; j < ke; ++j)
811 for (
int tj = k + 1; tj < nt; ++tj)
830 template<
bool oot_diagnostics,
bool fuse_rhs>
833 const int piv_stride,
const int k,
int& oot_count, T*
TDLS_RESTRICT y =
nullptr,
834 const int rhs_stride = 1) noexcept {
840 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
843 oot_count, ke, y, rhs_stride))
848 for (
int i = 0; i < ke; ++i)
851 store_tile(n, A, A_stride, pk, k0, tile, ke, ke);
858 }
else if constexpr (fuse_rhs) {
859 for (
int kk = 0; kk < ke; ++kk) {
860 for (
int i = kk + 1; i < ke; ++i)
867 for (
int tj = k + 1; tj < nt; ++tj) {
877 for (
int ti = k + 1; ti < nt; ++ti)
904 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int row0,
905 const int col0,
const int k0, T*
TDLS_RESTRICT t,
const int re,
906 const int ce)
noexcept {
907 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size) {
931 template<
bool fuse_rhs>
934 const int piv_stride,
const T*
TDLS_RESTRICT tile,
const int k0,
const int i0,
const int ke,
935 const int ie, T*
TDLS_RESTRICT y =
nullptr,
const int rhs_stride = 1) noexcept {
937 load_tile_piv(n, A, A_stride, piv, piv_stride, i0, k0, B, ie, ke);
938 ll_correct_tile(n, A, A_stride, piv, piv_stride, i0, k0, k0, B, ie, ke);
940 store_tile_piv(n, A, A_stride, piv, piv_stride, i0, k0, B, ie, ke);
946 for (
int r = 0; r < ie; ++r) {
948 for (
int j = 0; j < ke; ++j)
969 const T*
TDLS_RESTRICT tile,
const int k0,
const int j0,
const int ke,
970 const int je)
noexcept {
972 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
973 ll_correct_tile(n, A, A_stride, piv, piv_stride, k0, j0, k0, B, ke, je);
975 store_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
997 const T*
TDLS_RESTRICT tile,
const int k0,
const int j0,
const int ke,
998 const int je, T*
TDLS_RESTRICT y,
const int rhs_stride)
noexcept {
1000 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
1001 ll_correct_tile(n, A, A_stride, piv, piv_stride, k0, j0, k0, B, ke, je);
1003 store_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
1021 template<
bool oot_diagnostics,
bool fuse_rhs>
1024 const int piv_stride,
const int k,
int& oot_count, T*
TDLS_RESTRICT y =
nullptr,
1025 const int rhs_stride = 1) noexcept {
1031 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
1032 ll_correct_tile(n, A, A_stride, piv, piv_stride, k0, k0, k0, tile, ke, ke);
1035 oot_count, ke, y, rhs_stride))
1038 store_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
1045 }
else if constexpr (fuse_rhs) {
1046 for (
int kk = 0; kk < ke; ++kk) {
1047 for (
int i = kk + 1; i < ke; ++i)
1053 for (
int ti = k + 1; ti < nt; ++ti)
1059 for (
int tj = k + 1; tj < nt; ++tj) {
1093 template<
bool check>
1097 [[maybe_unused]]
const T u,
int& weak)
noexcept {
1099 bool singular =
false;
1101 for (
int j = 0; j < n; ++j) {
1104 for (
int i = j + 1; i < n; ++i) {
1117 singular = singular || !pivot_ok;
1121 for (
int c = 0; c < n; ++c) {
1129 for (
int i = j + 1; i < n; ++i) {
1132 if constexpr (check) ok = ok && u * detail::abs(l) <= T(1);
1133 for (
int c = j + 1; c < n; ++c)
1137 return singular ? 2 : ok ? 0 : 1;
1151 const int piv_stride,
int& weak)
noexcept {
1168 int*
TDLS_RESTRICT piv,
const int piv_stride,
const T u)
noexcept {
1185 T*
TDLS_RESTRICT x,
const int rhs_stride,
const std::size_t xcol_stride,
1186 const bool reverse)
noexcept {
1187 for (
int t = 0; t < n; ++t) {
1188 const int j = reverse ? n - 1 - t : t;
1191 for (
int c = 0; c < w; ++c) {
1214 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
1216 physical_permute(n, w, piv, piv_stride, x, rhs_stride, xcol_stride,
false);
1217 for (
int j = 0; j < n; ++j)
1218 for (
int i = j + 1; i < n; ++i) {
1220 for (
int c = 0; c < w; ++c)
1224 for (
int j = 0; j < n; ++j) {
1225 for (
int c = 0; c < w; ++c)
1227 for (
int q = j + 1; q < n; ++q) {
1229 for (
int c = 0; c < w; ++c)
1251 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
1253 for (
int j = n - 1; j >= 0; --j)
1254 for (
int c = 0; c < w; ++c) {
1256 for (
int q = j + 1; q < n; ++q)
1261 for (
int j = n - 1; j >= 0; --j)
1262 for (
int q = j + 1; q < n; ++q) {
1264 for (
int c = 0; c < w; ++c)
1284 template<
int pass_w
idth,
int stages>
1288 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
1289 const int cpp = (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
1290 for (
int c0 = 0; c0 < nrhs; c0 += cpp) {
1291 const int w = nrhs - c0 < cpp ? nrhs - c0 : cpp;
1293 if constexpr ((stages & 1) != 0)
1294 physical_forward(n, w, A, A_stride, piv, piv_stride, xp, rhs_stride, xcol_stride);
1295 if constexpr ((stages & 2) != 0)
1296 physical_backward(n, w, A, A_stride, piv, piv_stride, xp, rhs_stride, xcol_stride);
1327 template<
bool oot_diagnostics = true,
bool fuse_rhs = false>
1330 const int piv_stride,
int& oot_count, T*
TDLS_RESTRICT y =
nullptr,
1331 const int rhs_stride = 1) noexcept {
1332 n = detail::runtime_dimension(n);
1336 if constexpr (oot_diagnostics) oot_count = weak;
1337 if (!ok)
return false;
1338 if constexpr (fuse_rhs)
1343 if constexpr (oot_diagnostics) oot_count = 0;
1345 for (
int i = 0; i < n; ++i)
1349 for (
int k = 0; k < nt; ++k) {
1350 if constexpr (
schedule == Schedule::RightLooking) {
1352 oot_count, y, rhs_stride))
1356 oot_count, y, rhs_stride))
1374 const int piv_stride)
noexcept {
1375 n = detail::runtime_dimension(n);
1402 const int rhs_stride,
const std::size_t xcol_stride,
const int k0,
const int m0,
1403 const int ke,
const int me)
noexcept {
1405 load_tile_piv(n, A, A_stride, piv, piv_stride, m0, k0, Lmk, me, ke);
1406 for (
int i = 0; i < me; ++i) {
1407 for (
int w = 0; w < nrhs; ++w) {
1409 for (
int j = 0; j < ke; ++j)
1431 const int rhs_stride,
const std::size_t xcol_stride,
const int k)
noexcept {
1440 for (
int kk = 0; kk < ke; ++kk) {
1441 for (
int i = kk + 1; i < ke; ++i) {
1442 for (
int w = 0; w < nrhs; ++w)
1449 for (
int m = k + 1; m < nt; ++m)
1450 fwd_push_one(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k0,
1472 const int rhs_stride,
const std::size_t xcol_stride,
const int k0,
const int m0,
1473 const int ke,
const int me)
noexcept {
1475 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, m0, Ukm, ke, me);
1476 for (
int i = 0; i < ke; ++i) {
1477 for (
int w = 0; w < nrhs; ++w) {
1479 for (
int j = 0; j < me; ++j)
1501 const int rhs_stride,
const std::size_t xcol_stride,
const int k)
noexcept {
1507 for (
int m = k + 1; m < nt; ++m)
1508 bwd_pull_one(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k0,
1515 for (
int kk = ke - 1; kk >= 0; --kk) {
1516 for (
int w = 0; w < nrhs; ++w)
1518 for (
int i = 0; i < kk; ++i) {
1519 for (
int w = 0; w < nrhs; ++w)
1540 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
1541 for (
int k =
num_tiles(n) - 1; k >= 0; --k)
1542 bwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1561 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
1563 for (
int k = 0; k < nt; ++k)
1564 fwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1566 bwd_only(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1590 unsigned long long visited = 0ull;
1591 for (
int s = 0; s < n; ++s) {
1592 if ((visited >> s) & 1ull)
continue;
1598 visited |= 1ull << cur;
1603 visited |= 1ull << cur;
1609 for (
int s = 0; s < n; ++s) {
1613 if (probe != s)
continue;
1646 const std::size_t xcol_stride,
const int nrhs,
const int k0,
1647 const int ke)
noexcept {
1648 for (
int kk = 0; kk < ke; ++kk) {
1649 for (
int w = 0; w < nrhs; ++w)
1651 for (
int i = kk + 1; i < ke; ++i) {
1652 for (
int w = 0; w < nrhs; ++w)
1671 const std::size_t xcol_stride,
const int nrhs,
const int k0,
const int m0,
1672 const int ke,
const int me)
noexcept {
1673 for (
int i = 0; i < me; ++i) {
1674 for (
int w = 0; w < nrhs; ++w) {
1676 for (
int j = 0; j < ke; ++j)
1695 const std::size_t xcol_stride,
const int nrhs,
const int k0,
const int m0,
1696 const int ke,
const int me)
noexcept {
1697 for (
int i = 0; i < ke; ++i) {
1698 for (
int w = 0; w < nrhs; ++w) {
1700 for (
int j = 0; j < me; ++j)
1717 const int rhs_stride,
const std::size_t xcol_stride,
1718 const int nrhs,
const int k0,
const int ke)
noexcept {
1719 for (
int kk = ke - 1; kk > 0; --kk) {
1720 for (
int i = 0; i < kk; ++i) {
1721 for (
int w = 0; w < nrhs; ++w)
1743 const int rhs_stride,
const std::size_t xcol_stride,
const int k)
noexcept {
1750 for (
int m = k + 1; m < nt; ++m) {
1773 const int rhs_stride,
const std::size_t xcol_stride,
const int k)
noexcept {
1777 for (
int m = k + 1; m < nt; ++m) {
1802 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
1803 for (
int k =
num_tiles(n) - 1; k >= 0; --k)
1823 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
1825 for (
int k = 0; k < nt; ++k)
1844 unsigned long long visited = 0ull;
1845 for (
int s = 0; s < n; ++s) {
1846 if ((visited >> s) & 1ull)
continue;
1853 visited |= 1ull << dst;
1857 visited |= 1ull << s;
1860 for (
int s = 0; s < n; ++s) {
1864 if (probe != s)
continue;
1897 const int rhs_stride)
noexcept {
1898 n = detail::runtime_dimension(n);
1900 for (
int i = 0; i < n; ++i)
1907 for (
int i = 0; i < n; ++i)
1912 for (
int i = 0; i < n; ++i)
1914 fwd_bwd(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
1934 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int col,
1936 n = detail::runtime_dimension(n);
1938 for (
int i = 0; i < n; ++i)
1965 const int n,
const int nrhs,
const T*
TDLS_RESTRICT A,
const int A_stride,
1967 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
1970 for (
int i = 0; i < n; ++i)
1971 for (
int w = 0; w < nrhs; ++w)
1973 transposed_fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1974 for (
int w = 0; w < nrhs; ++w)
1978 for (
int i = 0; i < n; ++i) {
1980 for (
int w = 0; w < nrhs; ++w)
1983 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2011 template<
int pass_w
idth = 0>
2016 const std::size_t xcol_stride)
noexcept {
2017 n = detail::runtime_dimension(n);
2018 nrhs = detail::runtime_dimension(nrhs);
2020 for (
int c = 0; c < nrhs; ++c)
2021 for (
int i = 0; i < n; ++i)
2027 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
2028 if constexpr (pass_width == 0) {
2030 rhs_stride, xcol_stride);
2034 for (
int c0 = 0; c0 < nrhs; c0 += pass_width) {
2035 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
2038 x + off, rhs_stride, xcol_stride);
2062 const int A_stride,
const int*
TDLS_RESTRICT piv,
const int piv_stride,
2064 const std::size_t xcol_stride)
noexcept {
2067 for (
int i = 0; i < n; ++i)
2068 for (
int w = 0; w < nrhs; ++w)
2070 transposed_fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2071 for (
int w = 0; w < nrhs; ++w)
2075 for (
int i = 0; i < n; ++i) {
2077 for (
int w = 0; w < nrhs; ++w)
2080 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2112 template<
int pass_w
idth = 0>
2117 const std::size_t xcol_stride)
noexcept {
2118 n = detail::runtime_dimension(n);
2119 nrhs = detail::runtime_dimension(nrhs);
2121 for (
int c = 0; c < nrhs; ++c)
2122 for (
int i = 0; i < n; ++i)
2128 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
2129 if constexpr (pass_width == 0) {
2135 for (
int c0 = 0; c0 < nrhs; c0 += pass_width) {
2136 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
2139 rhs_stride, xcol_stride);
2161 const int rhs_stride)
noexcept {
2162 n = detail::runtime_dimension(n);
2174 fwd_bwd(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2198 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
2199 const std::size_t xcol_stride)
noexcept {
2202 transposed_fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2203 for (
int w = 0; w < nrhs; ++w)
2207 for (
int w = 0; w < nrhs; ++w)
2210 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2230 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
2231 const std::size_t xcol_stride)
noexcept {
2234 for (
int k = 0; k < nt; ++k)
2238 for (
int w = 0; w < nrhs; ++w)
2241 for (
int k = 0; k < nt; ++k)
2242 fwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
2262 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
2263 const std::size_t xcol_stride)
noexcept {
2266 for (
int w = 0; w < nrhs; ++w)
2270 bwd_only(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2299 template<
int pass_w
idth = 0>
2304 const std::size_t xcol_stride)
noexcept {
2305 n = detail::runtime_dimension(n);
2306 nrhs = detail::runtime_dimension(nrhs);
2312 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
2313 if constexpr (pass_width == 0) {
2319 for (
int c0 = 0; c0 < nrhs; c0 += pass_width) {
2320 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
2323 rhs_stride, xcol_stride);
2345 template<
int pass_w
idth = 0>
2349 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
2350 n = detail::runtime_dimension(n);
2351 nrhs = detail::runtime_dimension(nrhs);
2357 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
2358 if constexpr (pass_width == 0) {
2362 for (
int c0 = 0; c0 < nrhs; c0 += pass_width) {
2363 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
2366 rhs_stride, xcol_stride);
2385 template<
int pass_w
idth = 0>
2390 const std::size_t xcol_stride)
noexcept {
2391 n = detail::runtime_dimension(n);
2392 nrhs = detail::runtime_dimension(nrhs);
2398 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
2399 if constexpr (pass_width == 0) {
2403 for (
int c0 = 0; c0 < nrhs; c0 += pass_width) {
2404 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
2407 rhs_stride, xcol_stride);
2431 template<
bool oot_diagnostics = true>
2435 int& oot_count)
noexcept {
2436 n = detail::runtime_dimension(n);
2438 substitute(n, A, A_stride, piv, piv_stride, b, x, rhs_stride);
2457 const int rhs_stride)
noexcept {
2458 n = detail::runtime_dimension(n);
2460 return solve<false>(n, A, A_stride, piv, piv_stride, b, x, rhs_stride, unused);
2485 template<
int pass_w
idth = 0,
bool oot_diagnostics = true>
2489 const int rhs_stride,
const std::size_t xcol_stride,
int& oot_count)
noexcept {
2490 n = detail::runtime_dimension(n);
2491 nrhs = detail::runtime_dimension(nrhs);
2517 template<
int pass_w
idth = 0>
2521 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
2522 n = detail::runtime_dimension(n);
2523 nrhs = detail::runtime_dimension(nrhs);
2526 rhs_stride, xcol_stride, unused);
2556 template<
bool oot_diagnostics = true>
2559 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride,
2560 int& oot_count)
noexcept {
2561 n = detail::runtime_dimension(n);
2565 if constexpr (oot_diagnostics) oot_count = weak;
2566 if (!ok)
return false;
2580 bwd_only(n, 1, A, A_stride, piv, piv_stride, y, rhs_stride, 0);
2598 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride)
noexcept {
2599 n = detail::runtime_dimension(n);
2630 template<
int pass_w
idth = 0,
bool oot_diagnostics = true>
2634 const int rhs_stride,
const std::size_t xcol_stride,
2635 int& oot_count)
noexcept {
2636 n = detail::runtime_dimension(n);
2637 nrhs = detail::runtime_dimension(nrhs);
2640 rhs_stride, xcol_stride);
2662 template<
int pass_w
idth = 0>
2666 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
2667 n = detail::runtime_dimension(n);
2668 nrhs = detail::runtime_dimension(nrhs);
2671 rhs_stride, xcol_stride, unused);
2677#undef TDLS_LUPP_DYN_A
2678#undef TDLS_LUPP_DYN_PIV
2679#undef TDLS_LUPP_DYN_X
2680#undef TDLS_LUPP_DYN_B
2681#undef TDLS_LUPP_DYN_XW
2682#undef TDLS_LUPP_DYN_BW
2683#undef TDLS_LUPP_DYN_Y
2689#if defined(TDLS_DETAIL_COMPILER_GCC)
2690#pragma GCC diagnostic pop
#define TDLS_FORCE_INLINE_GPU
Inlining forced on GPU only: __forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS...
Definition macros.hpp:77
#define TDLS_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
Scalar math helpers usable in constant expressions.
Schedule
Elimination schedule of a factorization.
Definition options.hpp:51
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 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:832
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_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 std::size_t xcol_stride) noexcept
Physical substitution of nrhs columns in passes of pass_width columns: the first stage,...
Definition solver_dynamic.hpp:1286
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 (the backward one under the...
Definition solver_dynamic.hpp:437
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 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:1585
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t 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:1470
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t 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:1429
static constexpr T relative_pivot_threshold
Relative threshold of the in-tile pivot choice, read once from the configuration (see TiledLUppConfig...
Definition solver_dynamic.hpp:187
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:903
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ll_update_right_push_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, T *TDLS_RESTRICT y, const int rhs_stride) noexcept
LL, column-major layout, solve_inplace: ll_update_right_one, then the fused U^T forward push....
Definition solver_dynamic.hpp:995
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:281
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 bool solve_multirhs(int n, 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 std::size_t xcol_stride, int &oot_count) noexcept
factorize + substitute_multirhs in one call.
Definition solver_dynamic.hpp:2487
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 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) noexcept
Diagnostics-free solve_inplace overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:2597
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_push_one(const T *TDLS_RESTRICT u, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int nrhs, const int k0, const int m0, const int ke, const int me) noexcept
Forward push with U^T: x_m segment -= U(k,m)^T * x_k.
Definition solver_dynamic.hpp:1670
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t 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:2196
static constexpr Schedule schedule
elimination schedule (RightLooking or LeftLooking)
Definition solver_dynamic.hpp:179
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_pull_one(const T *TDLS_RESTRICT l, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int nrhs, const int k0, const int m0, const int ke, const int me) noexcept
Backward pull with L^T: x_k segment -= L(m,k)^T * x_m.
Definition solver_dynamic.hpp:1694
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 (the forward one under the...
Definition solver_dynamic.hpp:460
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t xcol_stride, const int k) noexcept
Backward step: pull the trailing contributions, then upper-solve the diagonal tile's segment.
Definition solver_dynamic.hpp:1499
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_multirhs(int n, 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 std::size_t xcol_stride) noexcept
Diagnostics-free solve_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:2519
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_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 std::size_t xcol_stride, const int k) noexcept
Transposed forward step: U^T-solve the diagonal tile's segment, then push it into the following segme...
Definition solver_dynamic.hpp:1741
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) noexcept
Diagnostics-free solve_inplace_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:2664
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void unpermute_inplace(const int n, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
x := P^T x in place for one strided column (x[piv[i]] receives x[i]): the inverse rotation of permute...
Definition solver_dynamic.hpp:1841
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int num_tiles(const int n) noexcept
Number of tiles per dimension (last one possibly partial).
Definition solver_dynamic.hpp:220
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 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:667
static constexpr bool transposed
Whether the solve uses the transpose of the factored matrix: S = A^T under the column-major layout (s...
Definition solver_dynamic.hpp:194
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) noexcept
Diagnostics-free factorize overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:1373
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void rl_trsm_right_push_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, const int k0, T *TDLS_RESTRICT y, const int rhs_stride) noexcept
RL, column-major layout, solve_inplace: rl_trsm_right_one, then the fused U^T forward push....
Definition solver_dynamic.hpp:718
static constexpr bool physical
Whether the rows are interchanged physically, the whole matrix a single tile (see TiledLUppConfig::ro...
Definition solver_dynamic.hpp:197
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:693
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 void transposed_solve_tile_lower_unit(const T *TDLS_RESTRICT l, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int nrhs, const int k0, const int ke) noexcept
In-tile backward solve with L^T (unit diagonal).
Definition solver_dynamic.hpp:1716
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_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 std::size_t xcol_stride) noexcept
Transposed triangular sweeps on unpermuted column(s) x: U^T z = x, then L^T w = z,...
Definition solver_dynamic.hpp:1821
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:967
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void forward_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 std::size_t xcol_stride) noexcept
Forward stage of one pass of substitute_inplace_multirhs: the permutation and the sweep with L under ...
Definition solver_dynamic.hpp:2228
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
static constexpr T oot_pivot_threshold
Acceptable-pivot threshold of the out-of-tile search, read once from the configuration (see TiledLUpp...
Definition solver_dynamic.hpp:184
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t 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:1538
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:264
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:780
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:391
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 substitute_canonical_multirhs(int n, 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 std::size_t 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:2013
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:351
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:742
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t 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:1400
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_permute(const int n, const int w, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const bool reverse) noexcept
Applies the interchanges to w columns of x, in their order (P x) or in reverse order (P^T x).
Definition solver_dynamic.hpp:1184
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:414
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:302
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
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool physical_factorize(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, int &weak) noexcept
Physical factorization (see physical_factorize_impl).
Definition solver_dynamic.hpp:1150
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void backward_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 std::size_t xcol_stride) noexcept
Backward stage of one pass of substitute_inplace_multirhs: the sweep with U under the row-major layou...
Definition solver_dynamic.hpp:2260
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_canonical(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:1933
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:325
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_solve_tile_upper(const T *TDLS_RESTRICT t, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int nrhs, const int k0, const int ke) noexcept
In-tile forward solve with U^T (the diagonal of the tile holds the pivot reciprocals).
Definition solver_dynamic.hpp:1645
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_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 std::size_t xcol_stride) noexcept
Transposed backward pass alone, used by transposed_fwd_bwd and by solve_inplace (whose forward pass h...
Definition solver_dynamic.hpp:1800
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 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 std::size_t xcol_stride) noexcept
Triangular solves on already-permuted column(s) x. nrhs columns are processed per tile visit,...
Definition solver_dynamic.hpp:1559
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute(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 (column-major layout: x := P^T L^-T U^-T b)....
Definition solver_dynamic.hpp:1895
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t xcol_stride) noexcept
One pass of the canonical multi right-hand-side substitution: nrhs canonical columns e_col0 ....
Definition solver_dynamic.hpp:1964
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:497
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t 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:2061
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:246
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:370
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve(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:2455
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 transposed_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 std::size_t xcol_stride, const int k) noexcept
Transposed backward step: pull the trailing contributions with L^T, then L^T-solve the diagonal tile'...
Definition solver_dynamic.hpp:1771
static constexpr int tile_size
tile size (int)
Definition solver_dynamic.hpp:178
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve(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:2433
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:932
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:1023
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
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int physical_factorize_impl(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T u, int &weak) noexcept
Physical factorization: P S = L U in place, L below the diagonal and U above it in the pivoted row or...
Definition solver_dynamic.hpp:1095
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:229
Detection of the compiler, of the language and of the side of the compilation pass: the one place whe...
Compile-time configuration of the TiledLUpp solver family.
#define TDLS_LUPP_DYN_X(i)
Strided entry i of the solution vector.
Definition solver_dynamic.hpp:117
#define TDLS_LUPP_DYN_PIV(i)
Strided pivot entry i.
Definition solver_dynamic.hpp:114
#define TDLS_LUPP_DYN_Y(i)
Strided entry i of the fused right-hand side of solve_inplace.
Definition solver_dynamic.hpp:130
#define TDLS_LUPP_DYN_A(r, c)
Strided element (r, c) of S, the stored matrix read row-major (A or A^T, see the file documentation).
Definition solver_dynamic.hpp:111
#define TDLS_LUPP_DYN_XW(w, i)
Strided entry i of column w of a multi right-hand-side block.
Definition solver_dynamic.hpp:123
#define TDLS_LUPP_DYN_B(i)
Strided entry i of the right-hand side.
Definition solver_dynamic.hpp:120
#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:127