1#ifndef TDLS_SOLVERS_TILED_LUPP_SOLVER_STATIC_HPP
2#define TDLS_SOLVERS_TILED_LUPP_SOLVER_STATIC_HPP
91#if defined(__GNUC__) && !defined(__clang__)
92#pragma GCC diagnostic push
93#pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
94#pragma GCC diagnostic ignored "-Warray-bounds"
103#if defined(__clang__)
104#pragma clang diagnostic push
105#pragma clang diagnostic ignored "-Wpass-failed"
120#define TDLS_LUPP_A(r, c) \
121 A[internal_matrix ? TDLS_LAYOUT_INDEX(r, c, N) \
122 : TDLS_LAYOUT_INDEX(r, c, N) * unsigned(A_stride)]
126#define TDLS_LUPP_PIV(i) piv[internal_piv ? unsigned(i) : unsigned(i) * unsigned(piv_stride)]
130#define TDLS_LUPP_X(i) x[internal_rhs ? unsigned(i) : unsigned(i) * unsigned(rhs_stride)]
134#define TDLS_LUPP_B(i) b[internal_rhs ? unsigned(i) : unsigned(i) * unsigned(rhs_stride)]
140#define TDLS_LUPP_XW(w, i) \
141 x[internal_rhs ? unsigned((w) * N + (i)) \
142 : unsigned(i) * unsigned(rhs_stride) + unsigned(w) * unsigned(xcol_stride)]
146#define TDLS_LUPP_BW(w, i) \
147 b[internal_rhs ? unsigned((w) * N + (i)) \
148 : unsigned(i) * unsigned(rhs_stride) + unsigned(w) * unsigned(xcol_stride)]
152#define TDLS_LUPP_Y(i) y[internal_rhs ? unsigned(i) : unsigned(i) * unsigned(rhs_stride)]
199template<
typename T,
int N, TiledLUppConfig<T> Config = TiledLUppConfig<T>{}>
213 static_assert(N >= 1,
"TiledLUppSolverStatic: N must be >= 1");
215 Config.oot_threshold.is_finite() && Config.singular_floor.is_finite(),
216 "TiledLUppSolverStatic: oot_threshold and singular_floor must be finite (and fit a "
218 static_assert(
singular_floor > T(0),
"TiledLUppSolverStatic: singular_floor must be positive");
220 "TiledLUppSolverStatic: singular_floor must not exceed oot_threshold (the "
221 "floor applies to the out-of-tile recovery path)");
222 static_assert(
tile_size >= 1,
"TiledLUppSolverStatic: tile size must be >= 1");
250 template<
int row_extent,
int col_extent,
bool internal_matrix>
254 if constexpr (Config.unroll_inner) {
256 for (
int i = 0; i < row_extent; ++i) {
258 for (
int j = 0; j < col_extent; ++j)
262 for (
int i = 0; i < row_extent; ++i) {
263 for (
int j = 0; j < col_extent; ++j)
278 template<
int row_extent,
int col_extent,
bool internal_matrix>
282 if constexpr (Config.unroll_inner) {
284 for (
int i = 0; i < row_extent; ++i) {
286 for (
int j = 0; j < col_extent; ++j)
290 for (
int i = 0; i < row_extent; ++i) {
291 for (
int j = 0; j < col_extent; ++j)
310 template<
int row_extent,
int col_extent,
bool internal_piv,
bool internal_matrix>
313 const int piv_stride,
const int row0,
const int col0,
315 if constexpr (Config.unroll_inner) {
317 for (
int i = 0; i < row_extent; ++i) {
320 for (
int j = 0; j < col_extent; ++j)
324 for (
int i = 0; i < row_extent; ++i) {
326 for (
int j = 0; j < col_extent; ++j)
345 template<
int row_extent,
int col_extent,
bool internal_piv,
bool internal_matrix>
348 const int piv_stride,
const int row0,
const int col0,
350 if constexpr (Config.unroll_inner) {
352 for (
int i = 0; i < row_extent; ++i) {
355 for (
int j = 0; j < col_extent; ++j)
359 for (
int i = 0; i < row_extent; ++i) {
361 for (
int j = 0; j < col_extent; ++j)
382 template<
int row_extent,
bool internal_piv,
bool internal_matrix>
385 const int piv_stride,
const int row0,
const int col0,
387 if constexpr (Config.unroll_inner) {
389 for (
int i = 1; i < row_extent; ++i) {
392 for (
int j = 0; j < i; ++j)
396 for (
int i = 1; i < row_extent; ++i) {
398 for (
int j = 0; j < i; ++j)
416 template<
int row_extent,
bool internal_piv,
bool internal_matrix>
419 const int piv_stride,
const int row0,
const int col0,
421 if constexpr (Config.unroll_inner) {
423 for (
int i = 0; i < row_extent; ++i) {
426 for (
int j = i; j < row_extent; ++j)
430 for (
int i = 0; i < row_extent; ++i) {
432 for (
int j = i; j < row_extent; ++j)
471 template<
int k_extent,
bool internal_piv,
bool internal_matrix,
bool oot_diagnostics,
472 bool fuse_rhs,
bool internal_rhs>
475 const int piv_stride,
const int k0, T*
TDLS_RESTRICT tile,
int& oot_count,
476 const int c, T*
TDLS_RESTRICT y,
const int rhs_stride)
noexcept {
478 const int gc = k0 + c;
482 T best = detail::abs(tile[c *
tile_size + c]);
483 if constexpr (Config.unroll_inner) {
485 for (
int r = c + 1; r < k_extent; ++r) {
486 const T v = detail::abs(tile[r *
tile_size + c]);
493 for (
int r = c + 1; r < k_extent; ++r) {
494 const T v = detail::abs(tile[r *
tile_size + c]);
505 piv_row = k0 + best_r;
506 }
else if constexpr (k_extent <
tile_size) {
511 if constexpr (oot_diagnostics) ++oot_count;
512 piv_row = k0 + best_r;
519 if constexpr (oot_diagnostics) ++oot_count;
521 int gbest_row = k0 + best_r;
523 for (
int row = k0 + k_extent; row < N; ++row) {
527 if constexpr (
schedule == Schedule::LeftLooking) {
528 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
536 for (
int t = 0; t < c; ++t) {
538 if constexpr (
schedule == Schedule::LeftLooking) {
539 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
544 for (
int p = 0; p < t; ++p)
545 a_t -= L_row[p] * tile[p *
tile_size + t];
546 L_row[t] = a_t * tile[t *
tile_size + t];
547 corrected -= L_row[t] * tile[t *
tile_size + c];
551 const T v = detail::abs(corrected);
561 if constexpr (Config.oot_first_acceptable)
579 if constexpr (fuse_rhs) {
584 if constexpr (internal_rhs) {
585 if constexpr (Config.unroll_inner) {
587 for (
int r = 0; r < N; ++r) {
595 for (
int r = 0; r < N; ++r) {
610 if (piv_row < k0 + k_extent) {
611 Operations::template swap_rows<k_extent>(tile, c, piv_row - k0);
617 if constexpr (Config.unroll_inner) {
619 for (
int j = 0; j < k_extent; ++j) {
621 if constexpr (
schedule == Schedule::LeftLooking) {
622 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
630 for (
int j = 0; j < k_extent; ++j) {
632 if constexpr (
schedule == Schedule::LeftLooking) {
633 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
644 for (
int t = 0; t < c; ++t) {
646 for (
int p = 0; p < t; ++p)
647 a_t -= L_row[p] * tile[p *
tile_size + t];
648 L_row[t] = a_t * tile[t *
tile_size + t];
651 for (
int j = c; j < k_extent; ++j) {
652 for (
int t = 0; t < c; ++t)
659 Operations::template eliminate_column<k_extent, k_extent>(tile, c);
683 template<
int k_extent,
bool internal_piv,
bool internal_matrix,
bool oot_diagnostics,
684 bool fuse_rhs =
false,
bool internal_rhs =
true>
687 const int piv_stride,
const int k0, T*
TDLS_RESTRICT tile,
int& oot_count,
688 T*
TDLS_RESTRICT y =
nullptr,
const int rhs_stride = 1) noexcept {
689 if constexpr (Config.unroll_inner) {
691 for (
int c = 0; c < k_extent; ++c) {
693 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k0,
694 tile, oot_count, c, y, rhs_stride))
698 for (
int c = 0; c < k_extent; ++c) {
700 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k0,
701 tile, oot_count, c, y, rhs_stride))
725 template<
int k_extent,
int j_extent,
bool internal_matrix>
731 Operations::template trsm_left_unit<k_extent, j_extent>(tile, Akj);
748 template<
int k_extent,
int i_extent,
int j_extent,
bool internal_matrix>
755 if constexpr (Config.unroll_inner) {
757 for (
int p = 0; p < k_extent; ++p) {
760 for (
int j = 0; j < j_extent; ++j)
763 for (
int i = 0; i < i_extent; ++i) {
766 for (
int j = 0; j < j_extent; ++j)
767 Aij[i *
tile_size + j] -= L_ip * Akj_row[j];
771 for (
int p = 0; p < k_extent; ++p) {
773 for (
int j = 0; j < j_extent; ++j)
775 for (
int i = 0; i < i_extent; ++i) {
777 for (
int j = 0; j < j_extent; ++j)
778 Aij[i *
tile_size + j] -= L_ip * Akj_row[j];
804 template<
int k_extent,
int i_extent,
bool internal_piv,
bool internal_matrix,
805 bool fuse_rhs =
false,
bool internal_rhs =
true>
810 T*
TDLS_RESTRICT y =
nullptr,
const int rhs_stride = 1) noexcept {
814 if constexpr (Config.unroll_inner) {
816 for (
int i = 0; i < i_extent; ++i)
819 for (
int i = 0; i < i_extent; ++i)
826 Operations::template trsm_right<k_extent, i_extent>(tile, Aik);
832 if constexpr (fuse_rhs) {
833 if constexpr (Config.unroll_inner) {
835 for (
int r = 0; r < i_extent; ++r) {
838 for (
int j = 0; j < k_extent; ++j)
843 for (
int r = 0; r < i_extent; ++r) {
845 for (
int j = 0; j < k_extent; ++j)
878 template<
int k_extent,
bool internal_piv,
bool internal_matrix,
bool oot_diagnostics,
879 bool fuse_rhs =
false,
bool internal_rhs =
true>
883 const int rhs_stride = 1) noexcept {
888 piv_stride, k0, k0, tile);
890 if (!
factor_diag_tile<k_extent, internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
891 internal_rhs>(A, A_stride, piv, piv_stride, k0, tile, oot_count, y,
897 if constexpr (Config.unroll_inner) {
899 for (
int i = 0; i < k_extent; ++i)
902 for (
int i = 0; i < k_extent; ++i)
910 if constexpr (fuse_rhs) {
911 if constexpr (Config.unroll_inner) {
913 for (
int kk = 0; kk < k_extent; ++kk) {
915 for (
int i = kk + 1; i < k_extent; ++i)
919 for (
int kk = 0; kk < k_extent; ++kk) {
920 for (
int i = kk + 1; i < k_extent; ++i)
937 internal_rhs>(A, A_stride, piv, piv_stride, pk, tile, k,
941 internal_rhs>(A, A_stride, piv, piv_stride, pk, tile, k,
970 template<
int row_extent,
int col_extent,
bool internal_piv,
bool internal_matrix>
973 const int piv_stride,
const int row0,
const int col0,
const int k0,
975 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size) {
978 A, A_stride, piv, piv_stride, row0, bj0, Lt);
981 A, A_stride, piv, piv_stride, bj0, col0, Ut);
982 Operations::template gemm_sub<row_extent, col_extent, tile_size>(t, Lt, Ut);
1003 template<
int k_extent,
int i_extent,
bool internal_piv,
bool internal_matrix,
1004 bool fuse_rhs =
false,
bool internal_rhs =
true>
1007 const int piv_stride,
const T*
TDLS_RESTRICT tile,
const int k0,
1009 const int rhs_stride = 1) noexcept {
1012 piv_stride, i0, k0, B);
1014 A, A_stride, piv, piv_stride, i0, k0, k0, B);
1015 Operations::template trsm_right<k_extent, i_extent>(tile, B);
1017 piv_stride, i0, k0, B);
1021 if constexpr (fuse_rhs) {
1022 if constexpr (Config.unroll_inner) {
1024 for (
int r = 0; r < i_extent; ++r) {
1027 for (
int j = 0; j < k_extent; ++j)
1032 for (
int r = 0; r < i_extent; ++r) {
1034 for (
int j = 0; j < k_extent; ++j)
1054 template<
int k_extent,
int j_extent,
bool internal_piv,
bool internal_matrix>
1057 const int piv_stride,
const T*
TDLS_RESTRICT tile,
const int k0,
1058 const int j0)
noexcept {
1061 piv_stride, k0, j0, B);
1063 A, A_stride, piv, piv_stride, k0, j0, k0, B);
1064 Operations::template trsm_left_unit<k_extent, j_extent>(tile, B);
1066 piv_stride, k0, j0, B);
1086 template<
int k_extent,
bool internal_piv,
bool internal_matrix,
bool oot_diagnostics,
1087 bool fuse_rhs =
false,
bool internal_rhs =
true>
1091 const int rhs_stride = 1) noexcept {
1096 piv_stride, k0, k0, tile);
1098 A, A_stride, piv, piv_stride, k0, k0, k0, tile);
1100 if (!
factor_diag_tile<k_extent, internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
1101 internal_rhs>(A, A_stride, piv, piv_stride, k0, tile, oot_count, y,
1106 piv_stride, k0, k0, tile);
1110 if constexpr (fuse_rhs) {
1111 if constexpr (Config.unroll_inner) {
1113 for (
int kk = 0; kk < k_extent; ++kk) {
1115 for (
int i = kk + 1; i < k_extent; ++i)
1119 for (
int kk = 0; kk < k_extent; ++kk) {
1120 for (
int i = kk + 1; i < k_extent; ++i)
1129 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0, i *
tile_size,
1133 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0,
1139 A, A_stride, piv, piv_stride, tile, k0, j *
tile_size);
1177 template<
bool internal_piv,
bool internal_matrix,
bool oot_diagnostics =
true,
1178 bool fuse_rhs =
false,
bool internal_rhs =
true>
1181 int& oot_count, T*
TDLS_RESTRICT y =
nullptr,
const int rhs_stride = 1) noexcept {
1183 if constexpr (oot_diagnostics) oot_count = 0;
1185 if constexpr (Config.unroll_inner) {
1187 for (
int i = 0; i < N; ++i)
1190 for (
int i = 0; i < N; ++i)
1194 if constexpr (
schedule == Schedule::RightLooking) {
1196 if (!
rl_step<
tile_size, internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
1197 internal_rhs>(A, A_stride, piv, piv_stride, k, oot_count, y,
1202 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride,
full_tiles,
1203 oot_count, y, rhs_stride))
1207 if (!
ll_step<
tile_size, internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
1208 internal_rhs>(A, A_stride, piv, piv_stride, k, oot_count, y,
1213 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride,
full_tiles,
1214 oot_count, y, rhs_stride))
1230 template<
bool internal_piv,
bool internal_matrix>
1233 const int piv_stride)
noexcept {
1263 template<
int k_extent,
int m_extent,
int pass_width,
bool internal_rhs,
bool internal_piv,
1264 bool internal_matrix>
1267 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1268 const int xcol_stride,
const int k0,
const int m0)
noexcept {
1271 piv_stride, m0, k0, Lmk);
1272 if constexpr (Config.unroll_inner) {
1274 for (
int i = 0; i < m_extent; ++i) {
1276 for (
int w = 0; w < pass_width; ++w) {
1279 for (
int j = 0; j < k_extent; ++j)
1285 for (
int i = 0; i < m_extent; ++i) {
1286 for (
int w = 0; w < pass_width; ++w) {
1288 for (
int j = 0; j < k_extent; ++j)
1311 template<
int k_extent,
int pass_width,
bool internal_rhs,
bool internal_piv,
1312 bool internal_matrix>
1315 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
const int xcol_stride,
1316 const int k)
noexcept {
1324 if constexpr (Config.unroll_inner) {
1326 for (
int kk = 0; kk < k_extent; ++kk) {
1328 for (
int i = kk + 1; i < k_extent; ++i) {
1330 for (
int w = 0; w < pass_width; ++w)
1336 for (
int kk = 0; kk < k_extent; ++kk) {
1337 for (
int i = kk + 1; i < k_extent; ++i) {
1338 for (
int w = 0; w < pass_width; ++w)
1348 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1352 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1373 template<
int k_extent,
int m_extent,
int pass_width,
bool internal_rhs,
bool internal_piv,
1374 bool internal_matrix>
1377 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1378 const int xcol_stride,
const int k0,
const int m0)
noexcept {
1381 piv_stride, k0, m0, Ukm);
1382 if constexpr (Config.unroll_inner) {
1384 for (
int i = 0; i < k_extent; ++i) {
1386 for (
int w = 0; w < pass_width; ++w) {
1389 for (
int j = 0; j < m_extent; ++j)
1395 for (
int i = 0; i < k_extent; ++i) {
1396 for (
int w = 0; w < pass_width; ++w) {
1398 for (
int j = 0; j < m_extent; ++j)
1421 template<
int k_extent,
int pass_width,
bool internal_rhs,
bool internal_piv,
1422 bool internal_matrix>
1425 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
const int xcol_stride,
1426 const int k)
noexcept {
1432 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1436 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1444 if constexpr (Config.unroll_inner) {
1446 for (
int kk = k_extent - 1; kk >= 0; --kk) {
1448 for (
int w = 0; w < pass_width; ++w)
1451 for (
int i = 0; i < kk; ++i) {
1453 for (
int w = 0; w < pass_width; ++w)
1459 for (
int kk = k_extent - 1; kk >= 0; --kk) {
1460 for (
int w = 0; w < pass_width; ++w)
1462 for (
int i = 0; i < kk; ++i) {
1463 for (
int w = 0; w < pass_width; ++w)
1484 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
1487 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1488 const int xcol_stride)
noexcept {
1491 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
full_tiles);
1494 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1512 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
1515 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1516 const int xcol_stride)
noexcept {
1519 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1522 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
full_tiles);
1525 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1547 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
1551 const int rhs_stride)
noexcept {
1552 if constexpr (internal_rhs) {
1557 if constexpr (Config.unroll_inner) {
1559 for (
int i = 0; i < N; ++i) {
1563 for (
int j = 0; j < N; ++j)
1564 if (p == j) v = b[j];
1568 for (
int i = 0; i < N; ++i) {
1571 for (
int j = 0; j < N; ++j)
1572 if (p == j) v = b[j];
1577 if constexpr (Config.unroll_inner) {
1579 for (
int i = 0; i < N; ++i)
1582 for (
int i = 0; i < N; ++i)
1606 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
1610 const int rhs_stride)
noexcept {
1612 A, A_stride, piv, piv_stride, col, x, rhs_stride, 0);
1633 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
1638 const int xcol_stride)
noexcept {
1639 if constexpr (Config.unroll_inner) {
1641 for (
int i = 0; i < N; ++i) {
1644 for (
int w = 0; w < pass_width; ++w)
1648 for (
int i = 0; i < N; ++i) {
1650 for (
int w = 0; w < pass_width; ++w)
1655 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1685 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
1691 const int xcol_stride)
noexcept {
1692 static_assert(nrhs >= 1,
"tdls: nrhs must be at least 1");
1693 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
1694 constexpr int columns_per_pass =
1695 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
1696 for (
int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
1697 const unsigned off =
1698 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
1701 A, A_stride, piv, piv_stride, col0 + c0, x + off, rhs_stride, xcol_stride);
1703 if constexpr (nrhs % columns_per_pass > 0) {
1704 constexpr int c0 = nrhs - nrhs % columns_per_pass;
1705 const unsigned off =
1706 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
1709 A, A_stride, piv, piv_stride, col0 + c0, x + off, rhs_stride, xcol_stride);
1731 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
1736 const int xcol_stride)
noexcept {
1737 if constexpr (internal_rhs) {
1741 if constexpr (Config.unroll_inner) {
1743 for (
int i = 0; i < N; ++i) {
1746 for (
int w = 0; w < pass_width; ++w) {
1749 for (
int j = 0; j < N; ++j)
1755 for (
int i = 0; i < N; ++i) {
1757 for (
int w = 0; w < pass_width; ++w) {
1759 for (
int j = 0; j < N; ++j)
1766 if constexpr (Config.unroll_inner) {
1768 for (
int i = 0; i < N; ++i) {
1771 for (
int w = 0; w < pass_width; ++w)
1775 for (
int i = 0; i < N; ++i) {
1777 for (
int w = 0; w < pass_width; ++w)
1783 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1816 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
1821 const int rhs_stride,
const int xcol_stride)
noexcept {
1822 static_assert(nrhs >= 1,
"tdls: nrhs must be at least 1");
1823 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
1824 constexpr int columns_per_pass =
1825 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
1826 for (
int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
1827 const unsigned off =
1828 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
1830 A, A_stride, piv, piv_stride, b + off, x + off, rhs_stride, xcol_stride);
1832 if constexpr (nrhs % columns_per_pass > 0) {
1833 constexpr int c0 = nrhs - nrhs % columns_per_pass;
1834 const unsigned off =
1835 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
1837 internal_matrix>(A, A_stride, piv, piv_stride, b + off,
1838 x + off, rhs_stride, xcol_stride);
1860 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
1863 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride)
noexcept {
1864 if constexpr (N <= 64) {
1867 using mask_t = std::conditional_t<(N <= 32), unsigned, unsigned long long>;
1868 mask_t visited = mask_t(0);
1869 if constexpr (Config.unroll_inner) {
1871 for (
int s = 0; s < N; ++s) {
1872 if ((visited >> s) & mask_t(1))
continue;
1878 visited |= mask_t(1) << cur;
1883 visited |= mask_t(1) << cur;
1886 for (
int s = 0; s < N; ++s) {
1887 if ((visited >> s) & mask_t(1))
continue;
1893 visited |= mask_t(1) << cur;
1898 visited |= mask_t(1) << cur;
1905 for (
int s = 0; s < N; ++s) {
1909 if (probe != s)
continue;
1944 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
1949 const int xcol_stride)
noexcept {
1950 if constexpr (N <= 64) {
1953 using mask_t = std::conditional_t<(N <= 32), unsigned, unsigned long long>;
1954 mask_t visited = mask_t(0);
1955 if constexpr (Config.unroll_inner) {
1957 for (
int s = 0; s < N; ++s) {
1958 if ((visited >> s) & mask_t(1))
continue;
1961 for (
int w = 0; w < pass_width; ++w)
1967 for (
int w = 0; w < pass_width; ++w)
1969 visited |= mask_t(1) << cur;
1974 for (
int w = 0; w < pass_width; ++w)
1976 visited |= mask_t(1) << cur;
1979 for (
int s = 0; s < N; ++s) {
1980 if ((visited >> s) & mask_t(1))
continue;
1982 for (
int w = 0; w < pass_width; ++w)
1987 for (
int w = 0; w < pass_width; ++w)
1989 visited |= mask_t(1) << cur;
1993 for (
int w = 0; w < pass_width; ++w)
1995 visited |= mask_t(1) << cur;
2000 for (
int s = 0; s < N; ++s) {
2004 if (probe != s)
continue;
2007 for (
int w = 0; w < pass_width; ++w)
2012 for (
int w = 0; w < pass_width; ++w)
2017 for (
int w = 0; w < pass_width; ++w)
2022 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2050 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2056 const int xcol_stride)
noexcept {
2057 static_assert(nrhs >= 1,
"tdls: nrhs must be at least 1");
2058 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
2059 constexpr int columns_per_pass =
2060 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
2061 for (
int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
2062 const unsigned off =
2063 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
2065 internal_matrix>(A, A_stride, piv, piv_stride, x + off,
2066 rhs_stride, xcol_stride);
2068 if constexpr (nrhs % columns_per_pass > 0) {
2069 constexpr int c0 = nrhs - nrhs % columns_per_pass;
2070 const unsigned off =
2071 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
2073 internal_matrix>(A, A_stride, piv, piv_stride, x + off,
2074 rhs_stride, xcol_stride);
2099 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2100 bool oot_diagnostics =
true>
2104 int& oot_count)
noexcept {
2127 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
2133 A, A_stride, piv, piv_stride, b, x, rhs_stride, unused);
2161 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2162 int pass_width = 0,
bool oot_diagnostics =
true>
2166 const int rhs_stride,
const int xcol_stride,
int& oot_count)
noexcept {
2171 A, A_stride, piv, piv_stride, b, x, rhs_stride, xcol_stride);
2197 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2202 const int rhs_stride,
const int xcol_stride)
noexcept {
2205 A, A_stride, piv, piv_stride, b, x, rhs_stride, xcol_stride, unused);
2232 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2233 bool oot_diagnostics =
true>
2236 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride,
2237 int& oot_count)
noexcept {
2239 A, A_stride, piv, piv_stride, oot_count, y, rhs_stride))
2261 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
2264 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride)
noexcept {
2267 A, A_stride, piv, piv_stride, y, rhs_stride, unused);
2299 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2300 int pass_width = 0,
bool oot_diagnostics =
true>
2303 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride,
2304 const int xcol_stride,
int& oot_count)
noexcept {
2309 A, A_stride, piv, piv_stride, y, rhs_stride, xcol_stride);
2334 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2338 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride,
2339 const int xcol_stride)
noexcept {
2342 false>(A, A_stride, piv, piv_stride, y, rhs_stride,
2343 xcol_stride, unused);
2373#if defined(__GNUC__) && !defined(__clang__)
2374#pragma GCC diagnostic pop
2377#if defined(__clang__)
2378#pragma clang 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_UNROLL_FORCE
Full-unroll pragma in the local compiler dialect.
Definition macros.hpp:103
#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_XW(w, i)
Entry i of column w of a multi right-hand-side block: pass_width contiguous columns under internal re...
Definition solver_static.hpp:140
#define TDLS_LUPP_Y(i)
Entry i of the fused right-hand side of solve_inplace (y follows the matrix rows through pivoting).
Definition solver_static.hpp:152
#define TDLS_LUPP_PIV(i)
Pivot entry i: contiguous under internal residency, strided otherwise.
Definition solver_static.hpp:126
#define TDLS_LUPP_A(r, c)
Element (r, c) of the factor matrix: contiguous under internal residency, strided otherwise,...
Definition solver_static.hpp:120
#define TDLS_LUPP_BW(w, i)
Entry i of column w of a multi right-hand-side block of b, addressed exactly as TDLS_LUPP_XW (b and x...
Definition solver_static.hpp:146
#define TDLS_LUPP_X(i)
Entry i of the solution vector: contiguous under internal residency, strided otherwise.
Definition solver_static.hpp:130
#define TDLS_LUPP_B(i)
Entry i of the right-hand side: contiguous under internal residency, strided otherwise.
Definition solver_static.hpp:134
Tiled dense LU factorization with logical partial pivoting and out-of-tile pivot recovery,...
Definition solver_static.hpp:200
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical_multirhs(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_static.hpp:1688
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factor_diag_tile(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, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
Factor the KExKE diagonal tile in registers, with out-of-tile pivot recovery (drives the per-column l...
Definition solver_static.hpp:686
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ll_update_below_one(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, 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_static.hpp:1006
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void rl_trsm_right_one(T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT pk, const T *TDLS_RESTRICT tile, const int j0) noexcept
RL: one row-panel tile update, Akj := L^-1 Akj.
Definition solver_static.hpp:727
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace_multirhs(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_static.hpp:2302
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve(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_static.hpp:2102
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace(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_static.hpp:1862
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ll_correct_tile(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) noexcept
LL: t (RExCE, rows row0.., cols col0..) -= sum over prior tiles bj < k0/tile_size of L(row0....
Definition solver_static.hpp:972
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical(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_static.hpp:1608
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void store_tile(T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT prow, const int col0, const T *TDLS_RESTRICT t) noexcept
Store an RxC tile through a cached physical-row segment.
Definition solver_static.hpp:280
static constexpr int num_tiles
Tiles per dimension, the partial one included.
Definition solver_static.hpp:228
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factorize(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_static.hpp:1180
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_multirhs(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_static.hpp:2164
static constexpr T oot_threshold
Acceptable-pivot threshold of the out-of-tile search, read once from the configuration (see TiledLUpp...
Definition solver_static.hpp:208
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical_multirhs_pass(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: pass_width canonical columns e_col0 ....
Definition solver_static.hpp:1635
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve(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_static.hpp:2129
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile_piv_upper(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) noexcept
Triangular variant of the diagonal-tile load for the backward substitution: only the upper triangle i...
Definition solver_static.hpp:418
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void fwd_push_one(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) noexcept
Forward push: subtract L(m,k) * x_k from the x_m segment, for pass_width columns at once.
Definition solver_static.hpp:1266
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile_piv(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) noexcept
Load an RxC tile, reading the permutation inline (one read per row).
Definition solver_static.hpp:312
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool rl_step(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_static.hpp:881
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factorize(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_static.hpp:1232
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void bwd_only(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_static.hpp:1486
static constexpr int tile_size
tile size (int)
Definition solver_static.hpp:202
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_multirhs_pass(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: pass_width columns of b gathered in permuted orde...
Definition solver_static.hpp:1733
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void fwd_bwd(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. pass_width columns are processed per tile visit,...
Definition solver_static.hpp:1514
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT prow, const int col0, T *TDLS_RESTRICT t) noexcept
Load an RxC tile through a cached physical-row segment.
Definition solver_static.hpp:252
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace_multirhs_pass(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: pass_width columns of x permuted by the ...
Definition solver_static.hpp:1946
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute(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_static.hpp:1549
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void fwd_step(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_static.hpp:1314
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ll_update_right_one(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) noexcept
LL: correct + TRSM one U-panel tile right of the diagonal.
Definition solver_static.hpp:1056
static constexpr int full_tiles
Definition solver_static.hpp:224
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace_multirhs(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_static.hpp:2053
static constexpr Schedule schedule
elimination schedule (RightLooking or LeftLooking)
Definition solver_static.hpp:203
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void bwd_pull_one(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) noexcept
Backward pull: subtract U(k,m) * x_m from the x_k segment, for pass_width columns at once.
Definition solver_static.hpp:1376
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void rl_schur_one(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) noexcept
RL: one Schur-complement tile update, Aij -= Aik * Akj, streaming the factored Akj row by row from re...
Definition solver_static.hpp:750
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile_piv_lower(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) noexcept
Triangular variant of the diagonal-tile load for the forward substitution: only the strict lower tria...
Definition solver_static.hpp:384
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace(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_static.hpp:2235
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void rl_update_row_one(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, 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_static.hpp:807
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void bwd_step(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_static.hpp:1424
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factor_diag_column(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, 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_static.hpp:474
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_multirhs(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_static.hpp:1819
static constexpr T singular_floor
Singularity floor of the out-of-tile recovery, read once from the configuration (see TiledLUppConfig:...
Definition solver_static.hpp:211
static constexpr int last_tile_tail
Extent of the last, partial tile: 0 when N is a multiple of tile_size.
Definition solver_static.hpp:226
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void store_tile_piv(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) noexcept
Store an RxC tile, reading the permutation inline (one read per row).
Definition solver_static.hpp:347
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace_multirhs(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_static.hpp:2337
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace(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_static.hpp:2263
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_multirhs(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_static.hpp:2200
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool ll_step(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_static.hpp:1089
tile_size x tile_size register-tile micro-kernels of the TiledLUpp solvers: the shared kernels plus t...
Definition tile_operations.hpp:46
LU register-tile micro-kernel of the TiledLUpp solver family.