1#ifndef TDLS_SOLVERS_TILED_LUPP_SOLVER_STATIC_HPP
2#define TDLS_SOLVERS_TILED_LUPP_SOLVER_STATIC_HPP
122#if defined(TDLS_DETAIL_COMPILER_GCC)
123#pragma GCC diagnostic push
124#pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
125#pragma GCC diagnostic ignored "-Warray-bounds"
134#if defined(TDLS_DETAIL_COMPILER_CLANG)
135#pragma clang diagnostic push
136#pragma clang diagnostic ignored "-Wpass-failed"
152#define TDLS_LUPP_A(r, c) \
153 A[internal_matrix ? unsigned((r) * N + (c)) : TDLS_STRIDED((r) * N + (c), A_stride)]
157#define TDLS_LUPP_PIV(i) piv[internal_piv ? unsigned(i) : TDLS_STRIDED(i, piv_stride)]
161#define TDLS_LUPP_X(i) x[internal_rhs ? unsigned(i) : TDLS_STRIDED(i, rhs_stride)]
165#define TDLS_LUPP_B(i) b[internal_rhs ? unsigned(i) : TDLS_STRIDED(i, rhs_stride)]
171#define TDLS_LUPP_XW(w, i) \
172 x[internal_rhs ? unsigned((w) * N + (i)) \
173 : TDLS_STRIDED(i, rhs_stride) + TDLS_STRIDED_WIDE(w, xcol_stride)]
177#define TDLS_LUPP_BW(w, i) \
178 b[internal_rhs ? unsigned((w) * N + (i)) \
179 : TDLS_STRIDED(i, rhs_stride) + TDLS_STRIDED_WIDE(w, xcol_stride)]
183#define TDLS_LUPP_Y(i) y[internal_rhs ? unsigned(i) : TDLS_STRIDED(i, rhs_stride)]
188#define TDLS_LUPP_FOR(i, first, last, ...) \
189 if constexpr (Config.unroll_loops) { \
191 for (int i = (first); i < (last); ++i) { \
195 for (int i = (first); i < (last); ++i) { \
254template<
typename T,
int N, TiledLUppConfig<T> Config = TiledLUppConfig<T>{}>
276 static constexpr bool transposed = Config.layout == MatrixLayout::ColMajor;
279 static constexpr bool physical = Config.row_interchange == RowInterchange::Physical;
281 static_assert(N >= 1,
"TiledLUppSolverStatic: N must be >= 1");
282 static_assert(Config.row_interchange != RowInterchange::Auto,
283 "TiledLUppSolverStatic: row_interchange must be Logical or Physical (Auto is "
284 "the default of the blocks of ArrowheadLUpp only)");
286 Config.oot_pivot_threshold.is_finite() && Config.singular_floor.is_finite(),
287 "TiledLUppSolverStatic: oot_pivot_threshold and singular_floor must be finite (and fit a "
289 static_assert(Config.relative_pivot_threshold.is_finite(),
290 "TiledLUppSolverStatic: relative_pivot_threshold must be finite (and fit a "
293 "TiledLUppSolverStatic: relative_pivot_threshold must be in (0, 1]");
294 static_assert(
singular_floor > T(0),
"TiledLUppSolverStatic: singular_floor must be positive");
296 "TiledLUppSolverStatic: singular_floor must not exceed oot_pivot_threshold (the "
297 "floor applies to the out-of-tile recovery path)");
298 static_assert(
tile_size >= 1,
"TiledLUppSolverStatic: tile size must be >= 1");
326 template<
int row_extent,
int col_extent,
bool internal_matrix>
330 if constexpr (Config.unroll_loops) {
332 for (
int i = 0; i < row_extent; ++i) {
334 for (
int j = 0; j < col_extent; ++j)
338 for (
int i = 0; i < row_extent; ++i) {
339 for (
int j = 0; j < col_extent; ++j)
354 template<
int row_extent,
int col_extent,
bool internal_matrix>
358 if constexpr (Config.unroll_loops) {
360 for (
int i = 0; i < row_extent; ++i) {
362 for (
int j = 0; j < col_extent; ++j)
366 for (
int i = 0; i < row_extent; ++i) {
367 for (
int j = 0; j < col_extent; ++j)
386 template<
int row_extent,
int col_extent,
bool internal_piv,
bool internal_matrix>
389 const int piv_stride,
const int row0,
const int col0,
391 if constexpr (Config.unroll_loops) {
393 for (
int i = 0; i < row_extent; ++i) {
396 for (
int j = 0; j < col_extent; ++j)
400 for (
int i = 0; i < row_extent; ++i) {
402 for (
int j = 0; j < col_extent; ++j)
421 template<
int row_extent,
int col_extent,
bool internal_piv,
bool internal_matrix>
424 const int piv_stride,
const int row0,
const int col0,
426 if constexpr (Config.unroll_loops) {
428 for (
int i = 0; i < row_extent; ++i) {
431 for (
int j = 0; j < col_extent; ++j)
435 for (
int i = 0; i < row_extent; ++i) {
437 for (
int j = 0; j < col_extent; ++j)
459 template<
int row_extent,
bool internal_piv,
bool internal_matrix>
462 const int piv_stride,
const int row0,
const int col0,
464 if constexpr (Config.unroll_loops) {
466 for (
int i = 1; i < row_extent; ++i) {
469 for (
int j = 0; j < i; ++j)
473 for (
int i = 1; i < row_extent; ++i) {
475 for (
int j = 0; j < i; ++j)
494 template<
int row_extent,
bool internal_piv,
bool internal_matrix>
497 const int piv_stride,
const int row0,
const int col0,
499 if constexpr (Config.unroll_loops) {
501 for (
int i = 0; i < row_extent; ++i) {
504 for (
int j = i; j < row_extent; ++j)
508 for (
int i = 0; i < row_extent; ++i) {
510 for (
int j = i; j < row_extent; ++j)
550 template<
int k_extent,
bool internal_piv,
bool internal_matrix,
bool oot_diagnostics,
551 bool fuse_rhs,
bool internal_rhs>
554 const int piv_stride,
const int k0, T*
TDLS_RESTRICT tile,
int& oot_count,
555 const int c, T*
TDLS_RESTRICT y,
const int rhs_stride)
noexcept {
557 const int gc = k0 + c;
561 T best = detail::abs(tile[c *
tile_size + c]);
562 if constexpr (Config.unroll_loops) {
564 for (
int r = c + 1; r < k_extent; ++r) {
565 const T v = detail::abs(tile[r *
tile_size + c]);
572 for (
int r = c + 1; r < k_extent; ++r) {
573 const T v = detail::abs(tile[r *
tile_size + c]);
587 const T in_place = detail::abs(tile[c *
tile_size + c]);
595 piv_row = k0 + best_r;
596 }
else if constexpr (k_extent <
tile_size) {
601 if constexpr (oot_diagnostics) ++oot_count;
602 piv_row = k0 + best_r;
609 if constexpr (oot_diagnostics) ++oot_count;
611 int gbest_row = k0 + best_r;
613 for (
int row = k0 + k_extent; row < N; ++row) {
617 if constexpr (
schedule == Schedule::LeftLooking) {
618 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
626 for (
int t = 0; t < c; ++t) {
628 if constexpr (
schedule == Schedule::LeftLooking) {
629 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
634 for (
int p = 0; p < t; ++p)
635 a_t -= L_row[p] * tile[p *
tile_size + t];
636 L_row[t] = a_t * tile[t *
tile_size + t];
637 corrected -= L_row[t] * tile[t *
tile_size + c];
641 const T v = detail::abs(corrected);
651 if constexpr (Config.oot_first_acceptable)
675 if constexpr (internal_rhs) {
676 if constexpr (Config.unroll_loops) {
678 for (
int r = 0; r < N; ++r) {
686 for (
int r = 0; r < N; ++r) {
701 if (piv_row < k0 + k_extent) {
702 Operations::template swap_rows<k_extent>(tile, c, piv_row - k0);
708 if constexpr (Config.unroll_loops) {
710 for (
int j = 0; j < k_extent; ++j) {
712 if constexpr (
schedule == Schedule::LeftLooking) {
713 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
721 for (
int j = 0; j < k_extent; ++j) {
723 if constexpr (
schedule == Schedule::LeftLooking) {
724 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size)
735 for (
int t = 0; t < c; ++t) {
737 for (
int p = 0; p < t; ++p)
738 a_t -= L_row[p] * tile[p *
tile_size + t];
739 L_row[t] = a_t * tile[t *
tile_size + t];
742 for (
int j = c; j < k_extent; ++j) {
743 for (
int t = 0; t < c; ++t)
750 Operations::template eliminate_column<k_extent, k_extent>(tile, c);
775 template<
int k_extent,
bool internal_piv,
bool internal_matrix,
bool oot_diagnostics,
776 bool fuse_rhs =
false,
bool internal_rhs =
true>
779 const int piv_stride,
const int k0, T*
TDLS_RESTRICT tile,
int& oot_count,
780 T*
TDLS_RESTRICT y =
nullptr,
const int rhs_stride = 1) noexcept {
781 if constexpr (Config.unroll_loops) {
783 for (
int c = 0; c < k_extent; ++c) {
785 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k0,
786 tile, oot_count, c, y, rhs_stride))
790 for (
int c = 0; c < k_extent; ++c) {
792 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k0,
793 tile, oot_count, c, y, rhs_stride))
818 template<
int k_extent,
int j_extent,
bool internal_matrix>
824 Operations::template trsm_left_unit<k_extent, j_extent>(tile, Akj);
844 template<
int k_extent,
int j_extent,
bool internal_matrix,
bool internal_rhs>
851 Operations::template trsm_left_unit<k_extent, j_extent>(tile, Akj);
869 template<
int k_extent,
int i_extent,
int j_extent,
bool internal_matrix>
876 if constexpr (Config.unroll_loops) {
878 for (
int p = 0; p < k_extent; ++p) {
881 for (
int j = 0; j < j_extent; ++j)
884 for (
int i = 0; i < i_extent; ++i) {
887 for (
int j = 0; j < j_extent; ++j)
888 Aij[i *
tile_size + j] -= L_ip * Akj_row[j];
892 for (
int p = 0; p < k_extent; ++p) {
894 for (
int j = 0; j < j_extent; ++j)
896 for (
int i = 0; i < i_extent; ++i) {
898 for (
int j = 0; j < j_extent; ++j)
899 Aij[i *
tile_size + j] -= L_ip * Akj_row[j];
925 template<
int k_extent,
int i_extent,
bool internal_piv,
bool internal_matrix,
926 bool fuse_rhs =
false,
bool internal_rhs =
true>
931 T*
TDLS_RESTRICT y =
nullptr,
const int rhs_stride = 1) noexcept {
935 if constexpr (Config.unroll_loops) {
937 for (
int i = 0; i < i_extent; ++i)
940 for (
int i = 0; i < i_extent; ++i)
947 Operations::template trsm_right<k_extent, i_extent>(tile, Aik);
956 if constexpr (Config.unroll_loops) {
958 for (
int r = 0; r < i_extent; ++r) {
961 for (
int j = 0; j < k_extent; ++j)
966 for (
int r = 0; r < i_extent; ++r) {
968 for (
int j = 0; j < k_extent; ++j)
1001 template<
int k_extent,
bool internal_piv,
bool internal_matrix,
bool oot_diagnostics,
1002 bool fuse_rhs =
false,
bool internal_rhs =
true>
1006 const int rhs_stride = 1) noexcept {
1011 piv_stride, k0, k0, tile);
1013 if (!
factor_diag_tile<k_extent, internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
1014 internal_rhs>(A, A_stride, piv, piv_stride, k0, tile, oot_count, y,
1020 if constexpr (Config.unroll_loops) {
1022 for (
int i = 0; i < k_extent; ++i)
1025 for (
int i = 0; i < k_extent; ++i)
1034 if constexpr (fuse_rhs) {
1037 }
else if constexpr (Config.unroll_loops) {
1039 for (
int kk = 0; kk < k_extent; ++kk) {
1041 for (
int i = kk + 1; i < k_extent; ++i)
1045 for (
int kk = 0; kk < k_extent; ++kk) {
1046 for (
int i = kk + 1; i < k_extent; ++i)
1057 A, A_stride, pk, tile, j *
tile_size, k0, y, rhs_stride);
1074 internal_rhs>(A, A_stride, piv, piv_stride, pk, tile, k,
1078 internal_rhs>(A, A_stride, piv, piv_stride, pk, tile, k,
1107 template<
int row_extent,
int col_extent,
bool internal_piv,
bool internal_matrix>
1110 const int piv_stride,
const int row0,
const int col0,
const int k0,
1112 for (
int bj0 = 0; bj0 < k0; bj0 +=
tile_size) {
1115 A, A_stride, piv, piv_stride, row0, bj0, Lt);
1118 A, A_stride, piv, piv_stride, bj0, col0, Ut);
1119 Operations::template gemm_sub<row_extent, col_extent, tile_size>(t, Lt, Ut);
1140 template<
int k_extent,
int i_extent,
bool internal_piv,
bool internal_matrix,
1141 bool fuse_rhs =
false,
bool internal_rhs =
true>
1144 const int piv_stride,
const T*
TDLS_RESTRICT tile,
const int k0,
1146 const int rhs_stride = 1) noexcept {
1149 piv_stride, i0, k0, B);
1151 A, A_stride, piv, piv_stride, i0, k0, k0, B);
1152 Operations::template trsm_right<k_extent, i_extent>(tile, B);
1154 piv_stride, i0, k0, B);
1161 if constexpr (Config.unroll_loops) {
1163 for (
int r = 0; r < i_extent; ++r) {
1166 for (
int j = 0; j < k_extent; ++j)
1171 for (
int r = 0; r < i_extent; ++r) {
1173 for (
int j = 0; j < k_extent; ++j)
1193 template<
int k_extent,
int j_extent,
bool internal_piv,
bool internal_matrix>
1196 const int piv_stride,
const T*
TDLS_RESTRICT tile,
const int k0,
1197 const int j0)
noexcept {
1200 piv_stride, k0, j0, B);
1202 A, A_stride, piv, piv_stride, k0, j0, k0, B);
1203 Operations::template trsm_left_unit<k_extent, j_extent>(tile, B);
1205 piv_stride, k0, j0, B);
1226 template<
int k_extent,
int j_extent,
bool internal_piv,
bool internal_matrix,
bool internal_rhs>
1229 const int piv_stride,
const T*
TDLS_RESTRICT tile,
const int k0,
1230 const int j0, T*
TDLS_RESTRICT y,
const int rhs_stride)
noexcept {
1233 piv_stride, k0, j0, B);
1235 A, A_stride, piv, piv_stride, k0, j0, k0, B);
1236 Operations::template trsm_left_unit<k_extent, j_extent>(tile, B);
1238 piv_stride, k0, j0, B);
1259 template<
int k_extent,
bool internal_piv,
bool internal_matrix,
bool oot_diagnostics,
1260 bool fuse_rhs =
false,
bool internal_rhs =
true>
1264 const int rhs_stride = 1) noexcept {
1269 piv_stride, k0, k0, tile);
1271 A, A_stride, piv, piv_stride, k0, k0, k0, tile);
1273 if (!
factor_diag_tile<k_extent, internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
1274 internal_rhs>(A, A_stride, piv, piv_stride, k0, tile, oot_count, y,
1279 piv_stride, k0, k0, tile);
1284 if constexpr (fuse_rhs) {
1287 }
else if constexpr (Config.unroll_loops) {
1289 for (
int kk = 0; kk < k_extent; ++kk) {
1291 for (
int i = kk + 1; i < k_extent; ++i)
1295 for (
int kk = 0; kk < k_extent; ++kk) {
1296 for (
int i = kk + 1; i < k_extent; ++i)
1305 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0, i *
tile_size,
1309 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0,
1317 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0,
1321 A, A_stride, piv, piv_stride, tile, k0, j *
tile_size);
1326 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0,
1352 template<
int w,
bool internal_rhs>
1355 [[maybe_unused]]
const std::size_t xcol_stride,
const int j,
1356 const int r)
noexcept {
1358 i, 0, N,
if (i > j && i == r) {
1380 const int rs)
noexcept {
1382 i, 0, N,
if (i > j && i == r) {
1384 X[j * rs + c * cs] = X[i * rs + c * cs]; X[i * rs + c * cs] = t;)
1399 template<
int w,
bool internal_rhs,
bool internal_piv,
bool reverse>
1403 const std::size_t xcol_stride)
noexcept {
1405 t, 0, N,
const int j = reverse ? N - 1 - t : t;
if (
TDLS_LUPP_PIV(j) != j) {
1406 if constexpr (internal_rhs)
1435 template<
bool internal_piv,
bool internal_matrix,
bool check>
1438 const int piv_stride,
const int j, [[maybe_unused]]
const T u,
bool& singular,
1439 int& weak)
noexcept {
1443 i, 0, N,
const T a = detail::abs(
TDLS_LUPP_A(i, j));
if (i > j && a > best) {
1448 const T in_place = detail::abs(
TDLS_LUPP_A(j, j));
1453 singular = singular || !pivot_ok;
1459 i, 0, N,
if (i > j && i == r) {
1464 const T inv = pivot_ok ? T(1) /
TDLS_LUPP_A(j, j) : T(0);
1468 i, 0, N,
if (i > j) {
1471 if constexpr (check) ok = ok && u * detail::abs(l) <= T(1);
1489 template<
bool internal_piv,
bool internal_matrix>
1492 const int piv_stride,
int& weak)
noexcept {
1493 bool singular =
false;
1497 A, A_stride, piv, piv_stride, j, T(0), singular, weak);)
1514 template<
bool internal_piv,
bool internal_matrix>
1517 const int piv_stride,
const T u)
noexcept {
1518 bool ok =
true, singular =
false;
1522 A, A_stride, piv, piv_stride, j, u, singular, weak) &&
1524 return singular ? 2 : ok ? 0 : 1;
1541 template<
int w,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
1546 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
1552 i, 0, N,
if (i > j) {
1561 q, 0, N,
if (q > j) {
1583 template<
int w,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
1588 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
1597 q, 0, N,
if (q > j) {
1622 template<
int nrhs,
int pass_width,
int stages,
bool internal_rhs,
bool internal_piv,
1623 bool internal_matrix>
1626 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1627 const std::size_t xcol_stride)
noexcept {
1628 constexpr int cpp = (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
1629 constexpr int passes = nrhs / cpp;
1630 constexpr int tail = nrhs % cpp;
1631 for (
int p = 0; p < passes; ++p) {
1632 const int c0 = p * cpp;
1633 const std::size_t off =
1635 if constexpr ((stages & 1) != 0)
1637 A, A_stride, piv, piv_stride, x + off, rhs_stride, xcol_stride);
1638 if constexpr ((stages & 2) != 0)
1640 A, A_stride, piv, piv_stride, x + off, rhs_stride, xcol_stride);
1642 if constexpr (tail > 0) {
1643 constexpr int c0 = nrhs - tail;
1644 const std::size_t off =
1646 if constexpr ((stages & 1) != 0)
1648 A, A_stride, piv, piv_stride, x + off, rhs_stride, xcol_stride);
1649 if constexpr ((stages & 2) != 0)
1651 A, A_stride, piv, piv_stride, x + off, rhs_stride, xcol_stride);
1695 template<
bool internal_piv,
bool internal_matrix,
bool oot_diagnostics =
true,
1696 bool fuse_rhs =
false,
bool internal_rhs =
true,
int A_static_stride =
dynamic_stride,
1700 int& oot_count, T*
TDLS_RESTRICT y =
nullptr,
const int rhs_stride = 1) noexcept {
1701 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
1702 rhs_static_stride>) {
1703 return factorize<internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
1704 internal_rhs>(A, detail::stride_or<A_static_stride>(A_stride), piv,
1705 detail::stride_or<piv_static_stride>(piv_stride),
1707 detail::stride_or<rhs_static_stride>(rhs_stride));
1713 if constexpr (oot_diagnostics) oot_count = weak;
1714 if (!ok)
return false;
1715 if constexpr (fuse_rhs)
1717 A, A_stride, piv, piv_stride, y, rhs_stride, 0);
1721 if constexpr (oot_diagnostics) oot_count = 0;
1723 if constexpr (Config.unroll_loops) {
1725 for (
int i = 0; i < N; ++i)
1728 for (
int i = 0; i < N; ++i)
1732 if constexpr (
schedule == Schedule::RightLooking) {
1735 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k, oot_count,
1740 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride,
full_tiles,
1741 oot_count, y, rhs_stride))
1746 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k, oot_count,
1751 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride,
full_tiles,
1752 oot_count, y, rhs_stride))
1773 template<
bool internal_piv,
bool internal_matrix,
int A_static_stride =
dynamic_stride,
1777 const int piv_stride)
noexcept {
1779 return factorize<internal_piv, internal_matrix,
false,
false,
true, A_static_stride,
1780 piv_static_stride>(A, A_stride, piv, piv_stride, unused);
1807 template<
int k_extent,
int m_extent,
int pass_width,
bool internal_rhs,
bool internal_piv,
1808 bool internal_matrix>
1811 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1812 const std::size_t xcol_stride,
const int k0,
const int m0)
noexcept {
1815 piv_stride, m0, k0, Lmk);
1816 if constexpr (Config.unroll_loops) {
1818 for (
int i = 0; i < m_extent; ++i) {
1820 for (
int w = 0; w < pass_width; ++w) {
1823 for (
int j = 0; j < k_extent; ++j)
1829 for (
int i = 0; i < m_extent; ++i) {
1830 for (
int w = 0; w < pass_width; ++w) {
1832 for (
int j = 0; j < k_extent; ++j)
1855 template<
int k_extent,
int pass_width,
bool internal_rhs,
bool internal_piv,
1856 bool internal_matrix>
1859 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1860 const std::size_t xcol_stride,
const int k)
noexcept {
1868 if constexpr (Config.unroll_loops) {
1870 for (
int kk = 0; kk < k_extent; ++kk) {
1872 for (
int i = kk + 1; i < k_extent; ++i) {
1874 for (
int w = 0; w < pass_width; ++w)
1880 for (
int kk = 0; kk < k_extent; ++kk) {
1881 for (
int i = kk + 1; i < k_extent; ++i) {
1882 for (
int w = 0; w < pass_width; ++w)
1892 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1896 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1917 template<
int k_extent,
int m_extent,
int pass_width,
bool internal_rhs,
bool internal_piv,
1918 bool internal_matrix>
1921 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1922 const std::size_t xcol_stride,
const int k0,
const int m0)
noexcept {
1925 piv_stride, k0, m0, Ukm);
1926 if constexpr (Config.unroll_loops) {
1928 for (
int i = 0; i < k_extent; ++i) {
1930 for (
int w = 0; w < pass_width; ++w) {
1933 for (
int j = 0; j < m_extent; ++j)
1939 for (
int i = 0; i < k_extent; ++i) {
1940 for (
int w = 0; w < pass_width; ++w) {
1942 for (
int j = 0; j < m_extent; ++j)
1965 template<
int k_extent,
int pass_width,
bool internal_rhs,
bool internal_piv,
1966 bool internal_matrix>
1969 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
1970 const std::size_t xcol_stride,
const int k)
noexcept {
1976 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1980 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1988 if constexpr (Config.unroll_loops) {
1990 for (
int kk = k_extent - 1; kk >= 0; --kk) {
1992 for (
int w = 0; w < pass_width; ++w)
1995 for (
int i = 0; i < kk; ++i) {
1997 for (
int w = 0; w < pass_width; ++w)
2003 for (
int kk = k_extent - 1; kk >= 0; --kk) {
2004 for (
int w = 0; w < pass_width; ++w)
2006 for (
int i = 0; i < kk; ++i) {
2007 for (
int w = 0; w < pass_width; ++w)
2028 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
2031 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
2032 const std::size_t xcol_stride)
noexcept {
2035 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
full_tiles);
2038 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
2056 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
2059 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
2060 const std::size_t xcol_stride)
noexcept {
2063 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
2066 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
full_tiles);
2069 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2095 template<
int k_extent,
int pass_w
idth,
bool internal_rhs>
2098 const std::size_t xcol_stride,
const int k0)
noexcept {
2099 if constexpr (Config.unroll_loops) {
2101 for (
int kk = 0; kk < k_extent; ++kk) {
2103 for (
int w = 0; w < pass_width; ++w)
2106 for (
int i = kk + 1; i < k_extent; ++i) {
2108 for (
int w = 0; w < pass_width; ++w)
2113 for (
int kk = 0; kk < k_extent; ++kk) {
2114 for (
int w = 0; w < pass_width; ++w)
2116 for (
int i = kk + 1; i < k_extent; ++i) {
2117 for (
int w = 0; w < pass_width; ++w)
2135 template<
int k_extent,
int m_extent,
int pass_w
idth,
bool internal_rhs>
2138 const std::size_t xcol_stride,
const int k0,
const int m0)
noexcept {
2139 if constexpr (Config.unroll_loops) {
2141 for (
int i = 0; i < m_extent; ++i) {
2143 for (
int w = 0; w < pass_width; ++w) {
2146 for (
int j = 0; j < k_extent; ++j)
2152 for (
int i = 0; i < m_extent; ++i) {
2153 for (
int w = 0; w < pass_width; ++w) {
2155 for (
int j = 0; j < k_extent; ++j)
2174 template<
int k_extent,
int m_extent,
int pass_w
idth,
bool internal_rhs>
2177 const std::size_t xcol_stride,
const int k0,
const int m0)
noexcept {
2178 if constexpr (Config.unroll_loops) {
2180 for (
int i = 0; i < k_extent; ++i) {
2182 for (
int w = 0; w < pass_width; ++w) {
2185 for (
int j = 0; j < m_extent; ++j)
2191 for (
int i = 0; i < k_extent; ++i) {
2192 for (
int w = 0; w < pass_width; ++w) {
2194 for (
int j = 0; j < m_extent; ++j)
2212 template<
int k_extent,
int pass_w
idth,
bool internal_rhs>
2215 const int rhs_stride,
const std::size_t xcol_stride,
2216 const int k0)
noexcept {
2217 if constexpr (Config.unroll_loops) {
2219 for (
int kk = k_extent - 1; kk > 0; --kk) {
2221 for (
int i = 0; i < kk; ++i) {
2223 for (
int w = 0; w < pass_width; ++w)
2228 for (
int kk = k_extent - 1; kk > 0; --kk) {
2229 for (
int i = 0; i < kk; ++i) {
2230 for (
int w = 0; w < pass_width; ++w)
2252 template<
int k_extent,
int pass_width,
bool internal_rhs,
bool internal_piv,
2253 bool internal_matrix>
2256 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
2257 const std::size_t xcol_stride,
const int k)
noexcept {
2267 A, A_stride, piv, piv_stride, k0, m *
tile_size, Ukm);
2269 Ukm, x, rhs_stride, xcol_stride, k0, m *
tile_size);
2295 template<
int k_extent,
int pass_width,
bool internal_rhs,
bool internal_piv,
2296 bool internal_matrix>
2299 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
2300 const std::size_t xcol_stride,
const int k)
noexcept {
2305 A, A_stride, piv, piv_stride, m *
tile_size, k0, Lmk);
2307 Lmk, x, rhs_stride, xcol_stride, k0, m *
tile_size);
2336 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
2339 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
2340 const std::size_t xcol_stride)
noexcept {
2343 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride,
2347 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
2364 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
2367 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride,
2368 const std::size_t xcol_stride)
noexcept {
2371 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
2374 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride,
2378 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2394 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv>
2397 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
2398 if constexpr (N <= 64) {
2399 using mask_t = std::conditional_t<(N <= 32), unsigned, unsigned long long>;
2400 mask_t visited = mask_t(0);
2401 if constexpr (Config.unroll_loops) {
2403 for (
int s = 0; s < N; ++s) {
2404 if ((visited >> s) & mask_t(1))
continue;
2405 T carry[pass_width];
2407 for (
int w = 0; w < pass_width; ++w)
2412 for (
int w = 0; w < pass_width; ++w) {
2415 carry[w] = displaced;
2417 visited |= mask_t(1) << dst;
2421 for (
int w = 0; w < pass_width; ++w)
2423 visited |= mask_t(1) << s;
2426 for (
int s = 0; s < N; ++s) {
2427 if ((visited >> s) & mask_t(1))
continue;
2428 T carry[pass_width];
2429 for (
int w = 0; w < pass_width; ++w)
2433 for (
int w = 0; w < pass_width; ++w) {
2436 carry[w] = displaced;
2438 visited |= mask_t(1) << dst;
2441 for (
int w = 0; w < pass_width; ++w)
2443 visited |= mask_t(1) << s;
2447 for (
int s = 0; s < N; ++s) {
2451 if (probe != s)
continue;
2453 T carry[pass_width];
2454 for (
int w = 0; w < pass_width; ++w)
2458 for (
int w = 0; w < pass_width; ++w) {
2461 carry[w] = displaced;
2465 for (
int w = 0; w < pass_width; ++w)
2482 template<
int pass_w
idth,
bool internal_piv>
2486 constexpr bool internal_rhs =
true;
2487 constexpr int rhs_stride = 1;
2488 constexpr std::size_t xcol_stride = N;
2489 T w_copy[pass_width * N];
2491 if constexpr (Config.unroll_loops) {
2493 for (
int i = 0; i < pass_width * N; ++i)
2496 for (
int i = 0; i < N; ++i)
2499 for (
int j = 0; j < N; ++j) {
2501 for (
int w = 0; w < pass_width; ++w) {
2504 for (
int i = 0; i < N; ++i)
2505 if (p[i] == j) v = w_copy[w * N + i];
2510 for (
int i = 0; i < pass_width * N; ++i)
2512 for (
int i = 0; i < N; ++i)
2514 for (
int j = 0; j < N; ++j) {
2515 for (
int w = 0; w < pass_width; ++w) {
2517 for (
int i = 0; i < N; ++i)
2518 if (p[i] == j) v = w_copy[w * N + i];
2535 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv>
2538 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
2539 if constexpr (internal_rhs)
2571 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2577 const int rhs_stride)
noexcept {
2578 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
2579 rhs_static_stride>) {
2581 A, detail::stride_or<A_static_stride>(A_stride), piv,
2582 detail::stride_or<piv_static_stride>(piv_stride), b, x,
2583 detail::stride_or<rhs_static_stride>(rhs_stride));
2588 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2594 if constexpr (Config.unroll_loops) {
2596 for (
int i = 0; i < N; ++i)
2599 for (
int i = 0; i < N; ++i)
2603 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2606 if constexpr (internal_rhs) {
2611 if constexpr (Config.unroll_loops) {
2613 for (
int i = 0; i < N; ++i) {
2617 for (
int j = 0; j < N; ++j)
2618 if (p == j) v = b[j];
2622 for (
int i = 0; i < N; ++i) {
2625 for (
int j = 0; j < N; ++j)
2626 if (p == j) v = b[j];
2631 if constexpr (Config.unroll_loops) {
2633 for (
int i = 0; i < N; ++i)
2636 for (
int i = 0; i < N; ++i)
2641 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2668 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2674 const int rhs_stride)
noexcept {
2675 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
2676 rhs_static_stride>) {
2678 A, detail::stride_or<A_static_stride>(A_stride), piv,
2679 detail::stride_or<piv_static_stride>(piv_stride), col, x,
2680 detail::stride_or<rhs_static_stride>(rhs_stride));
2685 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2689 A, A_stride, piv, piv_stride, col, x, rhs_stride, 0);
2713 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
2718 const std::size_t xcol_stride)
noexcept {
2722 if constexpr (Config.unroll_loops) {
2724 for (
int i = 0; i < N; ++i) {
2726 for (
int w = 0; w < pass_width; ++w)
2730 for (
int i = 0; i < N; ++i) {
2731 for (
int w = 0; w < pass_width; ++w)
2736 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2740 if constexpr (Config.unroll_loops) {
2742 for (
int i = 0; i < N; ++i) {
2745 for (
int w = 0; w < pass_width; ++w)
2749 for (
int i = 0; i < N; ++i) {
2751 for (
int w = 0; w < pass_width; ++w)
2756 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2794 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2801 const std::size_t xcol_stride)
noexcept {
2802 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
2803 rhs_static_stride>) {
2806 A, detail::stride_or<A_static_stride>(A_stride), piv,
2807 detail::stride_or<piv_static_stride>(piv_stride), col0, x,
2808 detail::stride_or<rhs_static_stride>(rhs_stride), xcol_stride);
2815 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2818 static_assert(nrhs >= 1,
"tdls: nrhs must be at least 1");
2819 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
2820 constexpr int columns_per_pass =
2821 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
2822 for (
int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
2823 const std::size_t off =
2827 A, A_stride, piv, piv_stride, col0 + c0, x + off, rhs_stride, xcol_stride);
2829 if constexpr (nrhs % columns_per_pass > 0) {
2830 constexpr int c0 = nrhs - nrhs % columns_per_pass;
2831 const std::size_t off =
2834 internal_piv, internal_matrix>(
2835 A, A_stride, piv, piv_stride, col0 + c0, x + off, rhs_stride, xcol_stride);
2860 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
2865 const std::size_t xcol_stride)
noexcept {
2868 if constexpr (Config.unroll_loops) {
2870 for (
int i = 0; i < N; ++i) {
2872 for (
int w = 0; w < pass_width; ++w)
2876 for (
int i = 0; i < N; ++i) {
2877 for (
int w = 0; w < pass_width; ++w)
2882 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2886 if constexpr (internal_rhs) {
2890 if constexpr (Config.unroll_loops) {
2892 for (
int i = 0; i < N; ++i) {
2895 for (
int w = 0; w < pass_width; ++w) {
2898 for (
int j = 0; j < N; ++j)
2904 for (
int i = 0; i < N; ++i) {
2906 for (
int w = 0; w < pass_width; ++w) {
2908 for (
int j = 0; j < N; ++j)
2915 if constexpr (Config.unroll_loops) {
2917 for (
int i = 0; i < N; ++i) {
2920 for (
int w = 0; w < pass_width; ++w)
2924 for (
int i = 0; i < N; ++i) {
2926 for (
int w = 0; w < pass_width; ++w)
2932 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2974 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
2980 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
2981 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
2982 rhs_static_stride>) {
2985 A, detail::stride_or<A_static_stride>(A_stride), piv,
2986 detail::stride_or<piv_static_stride>(piv_stride), b, x,
2987 detail::stride_or<rhs_static_stride>(rhs_stride), xcol_stride);
2993 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2996 static_assert(nrhs >= 1,
"tdls: nrhs must be at least 1");
2997 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
2998 constexpr int columns_per_pass =
2999 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
3000 for (
int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
3001 const std::size_t off =
3004 internal_matrix>(A, A_stride, piv, piv_stride, b + off,
3005 x + off, rhs_stride, xcol_stride);
3007 if constexpr (nrhs % columns_per_pass > 0) {
3008 constexpr int c0 = nrhs - nrhs % columns_per_pass;
3009 const std::size_t off =
3012 internal_matrix>(A, A_stride, piv, piv_stride, b + off,
3013 x + off, rhs_stride, xcol_stride);
3043 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3048 const int piv_stride, T*
TDLS_RESTRICT x,
const int rhs_stride)
noexcept {
3049 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
3050 rhs_static_stride>) {
3052 A, detail::stride_or<A_static_stride>(A_stride), piv,
3053 detail::stride_or<piv_static_stride>(piv_stride), x,
3054 detail::stride_or<rhs_static_stride>(rhs_stride));
3058 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
3065 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
3068 if constexpr (N <= 64) {
3071 using mask_t = std::conditional_t<(N <= 32), unsigned, unsigned long long>;
3072 mask_t visited = mask_t(0);
3073 if constexpr (Config.unroll_loops) {
3075 for (
int s = 0; s < N; ++s) {
3076 if ((visited >> s) & mask_t(1))
continue;
3082 visited |= mask_t(1) << cur;
3087 visited |= mask_t(1) << cur;
3090 for (
int s = 0; s < N; ++s) {
3091 if ((visited >> s) & mask_t(1))
continue;
3097 visited |= mask_t(1) << cur;
3102 visited |= mask_t(1) << cur;
3109 for (
int s = 0; s < N; ++s) {
3113 if (probe != s)
continue;
3127 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
3148 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv>
3151 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
3152 if constexpr (N <= 64) {
3155 using mask_t = std::conditional_t<(N <= 32), unsigned, unsigned long long>;
3156 mask_t visited = mask_t(0);
3157 if constexpr (Config.unroll_loops) {
3159 for (
int s = 0; s < N; ++s) {
3160 if ((visited >> s) & mask_t(1))
continue;
3163 for (
int w = 0; w < pass_width; ++w)
3169 for (
int w = 0; w < pass_width; ++w)
3171 visited |= mask_t(1) << cur;
3176 for (
int w = 0; w < pass_width; ++w)
3178 visited |= mask_t(1) << cur;
3181 for (
int s = 0; s < N; ++s) {
3182 if ((visited >> s) & mask_t(1))
continue;
3184 for (
int w = 0; w < pass_width; ++w)
3189 for (
int w = 0; w < pass_width; ++w)
3191 visited |= mask_t(1) << cur;
3195 for (
int w = 0; w < pass_width; ++w)
3197 visited |= mask_t(1) << cur;
3202 for (
int s = 0; s < N; ++s) {
3206 if (probe != s)
continue;
3209 for (
int w = 0; w < pass_width; ++w)
3214 for (
int w = 0; w < pass_width; ++w)
3219 for (
int w = 0; w < pass_width; ++w)
3245 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
3250 const std::size_t xcol_stride)
noexcept {
3254 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3261 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3281 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
3286 const std::size_t xcol_stride)
noexcept {
3290 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride,
3294 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride,
3301 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
3304 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
full_tiles);
3324 template<
int pass_w
idth,
bool internal_rhs,
bool internal_piv,
bool internal_matrix>
3329 const std::size_t xcol_stride)
noexcept {
3332 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3337 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3374 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3381 const std::size_t xcol_stride)
noexcept {
3382 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
3383 rhs_static_stride>) {
3386 A, detail::stride_or<A_static_stride>(A_stride), piv,
3387 detail::stride_or<piv_static_stride>(piv_stride), x,
3388 detail::stride_or<rhs_static_stride>(rhs_stride), xcol_stride);
3392 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3395 static_assert(nrhs >= 1,
"tdls: nrhs must be at least 1");
3396 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
3397 constexpr int columns_per_pass =
3398 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
3399 for (
int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
3400 const std::size_t off =
3403 internal_matrix>(A, A_stride, piv, piv_stride,
3404 x + off, rhs_stride, xcol_stride);
3406 if constexpr (nrhs % columns_per_pass > 0) {
3407 constexpr int c0 = nrhs - nrhs % columns_per_pass;
3408 const std::size_t off =
3411 internal_piv, internal_matrix>(
3412 A, A_stride, piv, piv_stride, x + off, rhs_stride, xcol_stride);
3443 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3449 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
3450 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
3451 rhs_static_stride>) {
3454 A, detail::stride_or<A_static_stride>(A_stride), piv,
3455 detail::stride_or<piv_static_stride>(piv_stride), x,
3456 detail::stride_or<rhs_static_stride>(rhs_stride), xcol_stride);
3460 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3463 static_assert(nrhs >= 1,
"tdls: nrhs must be at least 1");
3464 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
3465 constexpr int columns_per_pass =
3466 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
3467 for (
int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
3468 const std::size_t off =
3471 internal_matrix>(A, A_stride, piv, piv_stride,
3472 x + off, rhs_stride, xcol_stride);
3474 if constexpr (nrhs % columns_per_pass > 0) {
3475 constexpr int c0 = nrhs - nrhs % columns_per_pass;
3476 const std::size_t off =
3479 internal_matrix>(A, A_stride, piv, piv_stride,
3480 x + off, rhs_stride, xcol_stride);
3508 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3515 const std::size_t xcol_stride)
noexcept {
3516 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
3517 rhs_static_stride>) {
3520 A, detail::stride_or<A_static_stride>(A_stride), piv,
3521 detail::stride_or<piv_static_stride>(piv_stride), x,
3522 detail::stride_or<rhs_static_stride>(rhs_stride), xcol_stride);
3526 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3529 static_assert(nrhs >= 1,
"tdls: nrhs must be at least 1");
3530 static_assert(pass_width >= 0,
"tdls: pass_width must not be negative");
3531 constexpr int columns_per_pass =
3532 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
3533 for (
int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
3534 const std::size_t off =
3537 internal_matrix>(A, A_stride, piv, piv_stride,
3538 x + off, rhs_stride, xcol_stride);
3540 if constexpr (nrhs % columns_per_pass > 0) {
3541 constexpr int c0 = nrhs - nrhs % columns_per_pass;
3542 const std::size_t off =
3545 internal_matrix>(A, A_stride, piv, piv_stride,
3546 x + off, rhs_stride, xcol_stride);
3578 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3579 bool oot_diagnostics =
true,
int A_static_stride =
dynamic_stride,
3584 int& oot_count)
noexcept {
3585 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
3586 rhs_static_stride>) {
3588 A, detail::stride_or<A_static_stride>(A_stride), piv,
3589 detail::stride_or<piv_static_stride>(piv_stride), b, x,
3590 detail::stride_or<rhs_static_stride>(rhs_stride), oot_count);
3593 piv_stride, oot_count))
3621 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3628 return solve<internal_rhs, internal_piv, internal_matrix,
false, A_static_stride,
3629 piv_static_stride, rhs_static_stride>(A, A_stride, piv, piv_stride, b, x,
3630 rhs_stride, unused);
3664 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3665 int pass_width = 0,
bool oot_diagnostics =
true,
int A_static_stride =
dynamic_stride,
3670 const int rhs_stride,
const std::size_t xcol_stride,
int& oot_count)
noexcept {
3671 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
3672 rhs_static_stride>) {
3673 return solve_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix, pass_width,
3675 A, detail::stride_or<A_static_stride>(A_stride), piv,
3676 detail::stride_or<piv_static_stride>(piv_stride), b, x,
3677 detail::stride_or<rhs_static_stride>(rhs_stride), xcol_stride, oot_count);
3680 piv_stride, oot_count))
3683 A, A_stride, piv, piv_stride, b, x, rhs_stride, xcol_stride);
3716 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3722 const int rhs_stride,
const std::size_t xcol_stride)
noexcept {
3724 return solve_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix, pass_width,
false,
3725 A_static_stride, piv_static_stride, rhs_static_stride>(
3726 A, A_stride, piv, piv_stride, b, x, rhs_stride, xcol_stride, unused);
3765 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3766 bool oot_diagnostics =
true,
int A_static_stride =
dynamic_stride,
3770 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride,
3771 int& oot_count)
noexcept {
3772 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
3773 rhs_static_stride>) {
3775 A, detail::stride_or<A_static_stride>(A_stride), piv,
3776 detail::stride_or<piv_static_stride>(piv_stride), y,
3777 detail::stride_or<rhs_static_stride>(rhs_stride), oot_count);
3783 if constexpr (oot_diagnostics) oot_count = weak;
3784 if (!ok)
return false;
3786 A, A_stride, piv, piv_stride, y, rhs_stride, 0);
3790 A, A_stride, piv, piv_stride, oot_count, y, rhs_stride))
3797 A, A_stride, piv, piv_stride, y, rhs_stride, 0);
3801 A, A_stride, piv, piv_stride, y, rhs_stride, 0);
3826 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3831 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride)
noexcept {
3833 return solve_inplace<internal_rhs, internal_piv, internal_matrix,
false, A_static_stride,
3834 piv_static_stride, rhs_static_stride>(A, A_stride, piv, piv_stride, y,
3835 rhs_stride, unused);
3873 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3874 int pass_width = 0,
bool oot_diagnostics =
true,
int A_static_stride =
dynamic_stride,
3878 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride,
3879 const std::size_t xcol_stride,
int& oot_count)
noexcept {
3880 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
3881 rhs_static_stride>) {
3883 pass_width, oot_diagnostics>(
3884 A, detail::stride_or<A_static_stride>(A_stride), piv,
3885 detail::stride_or<piv_static_stride>(piv_stride), y,
3886 detail::stride_or<rhs_static_stride>(rhs_stride), xcol_stride, oot_count);
3889 piv_stride, oot_count))
3892 pass_width>(A, A_stride, piv, piv_stride, y, rhs_stride,
3925 template<
int nrhs,
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
3930 const int piv_stride, T*
TDLS_RESTRICT y,
const int rhs_stride,
3931 const std::size_t xcol_stride)
noexcept {
3934 false, A_static_stride, piv_static_stride, rhs_static_stride>(
3935 A, A_stride, piv, piv_stride, y, rhs_stride, xcol_stride, unused);
3966#if defined(TDLS_DETAIL_COMPILER_GCC)
3967#pragma GCC diagnostic pop
3970#if defined(TDLS_DETAIL_COMPILER_CLANG)
3971#pragma clang 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_FORCE_INLINE_ALL
Inlining forced on every backend, CPU included: TDLS_FORCE_INLINE_GPU on GPU, __forceinline under MSV...
Definition macros.hpp:101
#define TDLS_RESTRICT
Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC; nothing in the code that nvc++ gene...
Definition macros.hpp:174
#define TDLS_STRIDED_WIDE(i, s)
Offset of entry i of a strided object, i times stride s, for a 64-bit stride s (std::size_t); i is a ...
Definition macros.hpp:227
#define TDLS_UNROLL_FORCE
Full-unroll pragma in the local compiler dialect.
Definition macros.hpp:250
#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
constexpr int dynamic_stride
The compile-time stride of an operand whose stride is known at run time only: the default of the stri...
Definition options.hpp:181
Tiled dense LU factorization with logical partial pivoting and out-of-tile pivot recovery,...
Definition solver_static.hpp:255
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t 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:1920
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t 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:3247
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:3769
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 (the backward one under the...
Definition solver_static.hpp:461
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void unpermute(const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
x := P^T x, predicated for an internal x, by cycles for a remote one.
Definition solver_static.hpp:2537
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 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_static.hpp:1858
static constexpr T relative_pivot_threshold
Relative threshold of the in-tile pivot choice, read once from the configuration (see TiledLUppConfig...
Definition solver_static.hpp:266
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void rl_trsm_right_push_one(T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT pk, const T *TDLS_RESTRICT tile, const int j0, 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_static.hpp:846
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_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 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_static.hpp:2366
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t 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:1810
static constexpr bool physical
Whether the rows are interchanged physically, the whole matrix a single tile (see TiledLUppConfig::ro...
Definition solver_static.hpp:279
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool physical_factorize(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, 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_static.hpp:1491
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void unpermute_predicated(const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x) noexcept
x := P^T x for an internal x, predicated over the unrolled index range: x must never be dynamically i...
Definition solver_static.hpp:2484
static constexpr int num_tiles
Tiles per dimension, the partial one included.
Definition solver_static.hpp:304
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:820
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:1776
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 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_static.hpp:2798
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:388
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 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_static.hpp:3877
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:1699
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:1195
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t xcol_stride) noexcept
One pass of the canonical multi right-hand-side substitution: pass_width canonical columns e_col0 ....
Definition solver_static.hpp:2715
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int physical_factorize_checked(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_static.hpp:1516
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void backward_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 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_static.hpp:3326
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void backward_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 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_static.hpp:3512
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void forward_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 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_static.hpp:3283
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:553
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:3625
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_ALL constexpr void permute_cycles(const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
x := P x in place for pass_width columns (row i receives row piv[i]), by the cycle decomposition of s...
Definition solver_static.hpp:3150
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:928
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 (column-major layout: x := P^T L^-T U^-T b)....
Definition solver_static.hpp:2575
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 k0, const int m0) noexcept
Forward push with U^T: x_m segment -= U(k,m)^T * x_k.
Definition solver_static.hpp:2137
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 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_static.hpp:2030
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 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_static.hpp:1968
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:3047
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t xcol_stride, int &oot_count) noexcept
factorize + substitute_multirhs in one call.
Definition solver_static.hpp:3668
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:871
static constexpr int tile_size
tile size (int)
Definition solver_static.hpp:257
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void forward_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 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_static.hpp:3447
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:1004
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 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_static.hpp:2978
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t xcol_stride) noexcept
Triangular solves on already-permuted column(s) x. pass_width columns are processed per tile visit,...
Definition solver_static.hpp:2058
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 k0) noexcept
In-tile forward solve with U^T: x_k segment := U(k,k)^-T x_k (the diagonal of the tile holds the pivo...
Definition solver_static.hpp:2097
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_interchange_local(T *TDLS_RESTRICT X, const int j, const int r, const int cs, const int rs) noexcept
Interchanges rows j and r > j of the n columns of X, entry (i, c) at i * rs + c * cs,...
Definition solver_static.hpp:1379
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:1109
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t 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:2862
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void unpermute_cycles(const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
x := P^T x in place (x[piv[i]] receives x[i]), the last step of the transposed solve,...
Definition solver_static.hpp:2396
static constexpr T oot_pivot_threshold
Acceptable-pivot threshold of the out-of-tile search, read once from the configuration (see TiledLUpp...
Definition solver_static.hpp:263
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t xcol_stride) noexcept
Diagnostics-free solve_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_static.hpp:3720
static constexpr int full_tiles
Definition solver_static.hpp:300
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_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 std::size_t xcol_stride) noexcept
Physical substitution of nrhs columns in passes of pass_width columns: the first stage,...
Definition solver_static.hpp:1625
static constexpr Schedule schedule
elimination schedule (RightLooking or LeftLooking)
Definition solver_static.hpp:258
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:2672
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool physical_step(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int j, const T u, bool &singular, int &weak) noexcept
Step j of the physical factorization: the pivot of column j by the rule of a single tile (threshold p...
Definition solver_static.hpp:1437
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:3830
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:356
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_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 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_static.hpp:2298
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 k0) noexcept
In-tile backward solve with L^T (unit diagonal: no division and no multiplication by the diagonal).
Definition solver_static.hpp:2214
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 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_static.hpp:3378
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ll_update_right_push_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, 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_static.hpp:1228
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 k0, const int m0) noexcept
Backward pull with L^T: x_k segment -= L(m,k)^T * x_m.
Definition solver_static.hpp:2176
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_backward(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_static.hpp:1585
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:1143
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:778
static constexpr T singular_floor
Singularity floor of the out-of-tile recovery, read once from the configuration (see TiledLUppConfig:...
Definition solver_static.hpp:269
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_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 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_static.hpp:2338
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:302
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:328
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:1262
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 (the forward one under the...
Definition solver_static.hpp:496
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_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 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_static.hpp:2255
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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 std::size_t xcol_stride) noexcept
Diagnostics-free solve_inplace_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_static.hpp:3929
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_forward(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 (the ...
Definition solver_static.hpp:1543
static constexpr bool transposed
Whether the solve uses the transpose of the factored matrix. The solver factors S,...
Definition solver_static.hpp:276
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:423
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_interchange_rhs(T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int j, const int r) noexcept
Interchanges rows j and r > j of the w columns of x (physical scheme): a test of every unrolled row i...
Definition solver_static.hpp:1354
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:3582
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_permute(const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Applies the interchanges to w columns of x (physical scheme): in their order (P x),...
Definition solver_static.hpp:1401
tile_size x tile_size register-tile micro-kernels of the TiledLUpp solvers: the shared kernels plus t...
Definition tile_operations.hpp:47
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_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:171
#define TDLS_LUPP_Y(i)
Entry i of the fused right-hand side of solve_inplace (under the row-major layout y follows the matri...
Definition solver_static.hpp:183
#define TDLS_LUPP_PIV(i)
Pivot entry i: contiguous under internal residency, strided otherwise.
Definition solver_static.hpp:157
#define TDLS_LUPP_FOR(i, first, last,...)
Loop of the physical row interchanges, unrolled under Config.unroll_loops: the bounds are compile-tim...
Definition solver_static.hpp:188
#define TDLS_LUPP_A(r, c)
Element (r, c) of S, the stored matrix read row-major (A or A^T, see the file documentation): contigu...
Definition solver_static.hpp:152
#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:177
#define TDLS_LUPP_X(i)
Entry i of the solution vector: contiguous under internal residency, strided otherwise.
Definition solver_static.hpp:161
#define TDLS_LUPP_B(i)
Entry i of the right-hand side: contiguous under internal residency, strided otherwise.
Definition solver_static.hpp:165
LU register-tile micro-kernel of the TiledLUpp solver family.