TDLS 0.4.0
Tiny Device-callable Linear Solvers
Loading...
Searching...
No Matches
solver_static.hpp
Go to the documentation of this file.
1#ifndef TDLS_SOLVERS_TILED_LUPP_SOLVER_STATIC_HPP
2#define TDLS_SOLVERS_TILED_LUPP_SOLVER_STATIC_HPP
3
4
5
92
93
94
95#include <tdls/core/math.hpp>
96#include <tdls/core/target.hpp>
97#include <type_traits>
98
101
102
103
104// The flow analysis of older gcc releases (observed on the 11.4 of
105// Ubuntu 22.04) cannot prove that the replay row of the out-of-tile
106// search is written before being read: L_row[p] is only read for
107// p < t, and slot t is written at the end of iteration t, a
108// loop-carried property validated by the oracle suites. The resulting
109// -Wmaybe-uninitialized reports are spurious.
110//
111// Under UBSan instrumentation (-fsanitize=undefined) at -O2 and
112// above, gcc additionally loses the value ranges of the loop indices
113// and reports impossible subscripts through -Warray-bounds: the
114// unsigned cast of a provably non-negative index read as a wrapped
115// negative, or a loop counter assumed below its own loop guard.
116// Observed from gcc 12 to gcc 16; plain builds and clang are clean,
117// and the exercised paths are certified free of out-of-bounds
118// accesses by the constexpr suite.
119//
120// Both suppressions are scoped to this header and to those warnings
121// only, exactly as in solver_dynamic.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"
126#endif
127
128
129
130// clang reports a forced unrolling that the optimizer could not perform
131// through -Wpass-failed. The unrolling requested by TDLS_UNROLL_FORCE is
132// a performance hint: a failed hint does not affect correctness. The
133// suppression is scoped to this header and to that warning only.
134#if defined(TDLS_DETAIL_COMPILER_CLANG)
135#pragma clang diagnostic push
136#pragma clang diagnostic ignored "-Wpass-failed"
137#endif
138
139namespace tdls {
140
141
142
143/* Addressing macros. They expand inside member functions where the
144 parameter names (A, A_stride, piv, piv_stride, x, b, y, rhs_stride,
145 xcol_stride), the residency template booleans and the Config value
146 are in scope. #undef'd at the end of this header. */
147
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) { \
190 TDLS_UNROLL_FORCE \
191 for (int i = (first); i < (last); ++i) { \
192 __VA_ARGS__ \
193 } \
194 } else { \
195 for (int i = (first); i < (last); ++i) { \
196 __VA_ARGS__ \
197 } \
198 }
199
200
201
254template<typename T, int N, TiledLUppConfig<T> Config = TiledLUppConfig<T>{}>
256
257 static constexpr int tile_size = Config.tile_size;
258 static constexpr Schedule schedule =
259 Config.schedule;
260
263 static constexpr T oot_pivot_threshold = Config.oot_pivot_threshold;
266 static constexpr T relative_pivot_threshold = Config.relative_pivot_threshold;
269 static constexpr T singular_floor = Config.singular_floor;
276 static constexpr bool transposed = Config.layout == MatrixLayout::ColMajor;
279 static constexpr bool physical = Config.row_interchange == RowInterchange::Physical;
280
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)");
285 static_assert(
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 "
288 "63-bit mantissa)");
289 static_assert(Config.relative_pivot_threshold.is_finite(),
290 "TiledLUppSolverStatic: relative_pivot_threshold must be finite (and fit a "
291 "63-bit mantissa)");
292 static_assert(relative_pivot_threshold > T(0) && relative_pivot_threshold <= T(1),
293 "TiledLUppSolverStatic: relative_pivot_threshold must be in (0, 1]");
294 static_assert(singular_floor > T(0), "TiledLUppSolverStatic: singular_floor must be positive");
295 static_assert(singular_floor <= oot_pivot_threshold,
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");
299
300 static constexpr int full_tiles = N / tile_size;
302 static constexpr int last_tile_tail = N - full_tiles * tile_size;
304 static constexpr int num_tiles = full_tiles + (last_tile_tail > 0 ? 1 : 0);
305
307 using Operations = TiledLUppTileOperations<T, tile_size, Config.unroll_loops>;
308
309 /* =====================================================================
310 Remote <-> register tile movement.
311 `load/store_tile` go through a cached physical-row segment (the RL
312 schedule keeps one per tile); `load/store_tile_piv` read the
313 permutation inline, once per row (LL schedule and the substitution
314 do this; it adds no integer registers).
315 ===================================================================== */
316
326 template<int row_extent, int col_extent, bool internal_matrix>
327 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
328 load_tile(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT prow,
329 const int col0, T* TDLS_RESTRICT t) noexcept {
330 if constexpr (Config.unroll_loops) {
332 for (int i = 0; i < row_extent; ++i) {
334 for (int j = 0; j < col_extent; ++j)
335 t[i * tile_size + j] = TDLS_LUPP_A(prow[i], col0 + j);
336 }
337 } else {
338 for (int i = 0; i < row_extent; ++i) {
339 for (int j = 0; j < col_extent; ++j)
340 t[i * tile_size + j] = TDLS_LUPP_A(prow[i], col0 + j);
341 }
342 }
343 }
344
354 template<int row_extent, int col_extent, bool internal_matrix>
355 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
356 store_tile(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT prow,
357 const int col0, const T* TDLS_RESTRICT t) noexcept {
358 if constexpr (Config.unroll_loops) {
360 for (int i = 0; i < row_extent; ++i) {
362 for (int j = 0; j < col_extent; ++j)
363 TDLS_LUPP_A(prow[i], col0 + j) = t[i * tile_size + j];
364 }
365 } else {
366 for (int i = 0; i < row_extent; ++i) {
367 for (int j = 0; j < col_extent; ++j)
368 TDLS_LUPP_A(prow[i], col0 + j) = t[i * tile_size + j];
369 }
370 }
371 }
372
386 template<int row_extent, int col_extent, bool internal_piv, bool internal_matrix>
387 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
388 load_tile_piv(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
389 const int piv_stride, const int row0, const int col0,
390 T* TDLS_RESTRICT t) noexcept {
391 if constexpr (Config.unroll_loops) {
393 for (int i = 0; i < row_extent; ++i) {
394 const int phys = TDLS_LUPP_PIV(row0 + i);
396 for (int j = 0; j < col_extent; ++j)
397 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
398 }
399 } else {
400 for (int i = 0; i < row_extent; ++i) {
401 const int phys = TDLS_LUPP_PIV(row0 + i);
402 for (int j = 0; j < col_extent; ++j)
403 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
404 }
405 }
406 }
407
421 template<int row_extent, int col_extent, bool internal_piv, bool internal_matrix>
422 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
423 store_tile_piv(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
424 const int piv_stride, const int row0, const int col0,
425 const T* TDLS_RESTRICT t) noexcept {
426 if constexpr (Config.unroll_loops) {
428 for (int i = 0; i < row_extent; ++i) {
429 const int phys = TDLS_LUPP_PIV(row0 + i);
431 for (int j = 0; j < col_extent; ++j)
432 TDLS_LUPP_A(phys, col0 + j) = t[i * tile_size + j];
433 }
434 } else {
435 for (int i = 0; i < row_extent; ++i) {
436 const int phys = TDLS_LUPP_PIV(row0 + i);
437 for (int j = 0; j < col_extent; ++j)
438 TDLS_LUPP_A(phys, col0 + j) = t[i * tile_size + j];
439 }
440 }
441 }
442
459 template<int row_extent, bool internal_piv, bool internal_matrix>
460 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
461 load_tile_piv_lower(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
462 const int piv_stride, const int row0, const int col0,
463 T* TDLS_RESTRICT t) noexcept {
464 if constexpr (Config.unroll_loops) {
466 for (int i = 1; i < row_extent; ++i) {
467 const int phys = TDLS_LUPP_PIV(row0 + i);
469 for (int j = 0; j < i; ++j)
470 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
471 }
472 } else {
473 for (int i = 1; i < row_extent; ++i) {
474 const int phys = TDLS_LUPP_PIV(row0 + i);
475 for (int j = 0; j < i; ++j)
476 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
477 }
478 }
479 }
480
494 template<int row_extent, bool internal_piv, bool internal_matrix>
495 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
496 load_tile_piv_upper(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
497 const int piv_stride, const int row0, const int col0,
498 T* TDLS_RESTRICT t) noexcept {
499 if constexpr (Config.unroll_loops) {
501 for (int i = 0; i < row_extent; ++i) {
502 const int phys = TDLS_LUPP_PIV(row0 + i);
504 for (int j = i; j < row_extent; ++j)
505 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
506 }
507 } else {
508 for (int i = 0; i < row_extent; ++i) {
509 const int phys = TDLS_LUPP_PIV(row0 + i);
510 for (int j = i; j < row_extent; ++j)
511 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
512 }
513 }
514 }
515
516 /* =====================================================================
517 Diagonal-tile factorization with out-of-tile pivoting.
518 Shared by both schedules; the only schedule-dependent part is the
519 replay a candidate row needs before it can be compared:
520 RL: the trailing matrix is already Schur-updated, so a candidate
521 only misses the current tile's columns t < c.
522 LL: rows below additionally miss every prior tile's L*U
523 contribution (columns [0, k0)), replayed term by term.
524 ===================================================================== */
525
550 template<int k_extent, bool internal_piv, bool internal_matrix, bool oot_diagnostics,
551 bool fuse_rhs, bool internal_rhs>
552 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
553 factor_diag_column(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
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 {
556
557 const int gc = k0 + c; // global column
558
559 // In-tile pivot search (rows c..k_extent of the register tile)
560 int best_r = 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]);
566 if (v > best) {
567 best = v;
568 best_r = r;
569 }
570 }
571 } else {
572 for (int r = c + 1; r < k_extent; ++r) {
573 const T v = detail::abs(tile[r * tile_size + c]);
574 if (v > best) {
575 best = v;
576 best_r = r;
577 }
578 }
579 }
580
581 // Threshold pivoting inside the tile: the row in place keeps the
582 // pivot when it reaches both relative_pivot_threshold times the
583 // best in-tile candidate and oot_pivot_threshold. The out-of-tile
584 // recovery below keeps its own absolute trigger. A threshold of 1
585 // compiles the test out: the best in-tile candidate wins.
586 if constexpr (relative_pivot_threshold < T(1)) {
587 const T in_place = detail::abs(tile[c * tile_size + c]);
588 if (in_place >= relative_pivot_threshold * best && in_place >= oot_pivot_threshold)
589 best_r = c;
590 }
591
592 int piv_row; // winning global (logical) row
593
594 if (best >= oot_pivot_threshold) {
595 piv_row = k0 + best_r;
596 } else if constexpr (k_extent < tile_size) {
597 // Trailing tile: no rows below to recover from. Diagnostic
598 // order: singularity verdict first, then count the weak pivot
599 // (full tiles count before the verdict).
600 if (best < singular_floor) return false;
601 if constexpr (oot_diagnostics) ++oot_count;
602 piv_row = k0 + best_r;
603 } else {
604 // Out-of-tile recovery: scan the rows below the tile and
605 // evaluate each candidate as if it had received the
606 // eliminations it is missing, keeping the best (or, with
607 // Config.oot_first_acceptable, the first to reach the threshold).
608 // Cold path, never unroll-annotated.
609 if constexpr (oot_diagnostics) ++oot_count;
610 T gbest = best;
611 int gbest_row = k0 + best_r;
612
613 for (int row = k0 + k_extent; row < N; ++row) {
614 const int phys = TDLS_LUPP_PIV(row);
615
616 T corrected = TDLS_LUPP_A(phys, gc);
617 if constexpr (schedule == Schedule::LeftLooking) {
618 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
619 for (int p = 0; p < tile_size; ++p)
620 corrected -= TDLS_LUPP_A(phys, bj0 + p) *
621 TDLS_LUPP_A(TDLS_LUPP_PIV(bj0 + p), gc);
622 }
623
624 if (c > 0) {
625 T L_row[tile_size];
626 for (int t = 0; t < c; ++t) {
627 T a_t = TDLS_LUPP_A(phys, k0 + t);
628 if constexpr (schedule == Schedule::LeftLooking) {
629 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
630 for (int p = 0; p < tile_size; ++p)
631 a_t -= TDLS_LUPP_A(phys, bj0 + p) *
632 TDLS_LUPP_A(TDLS_LUPP_PIV(bj0 + p), k0 + t);
633 }
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]; // diag holds 1/pivot
637 corrected -= L_row[t] * tile[t * tile_size + c];
638 }
639 }
640
641 const T v = detail::abs(corrected);
642 if (v > gbest) {
643 gbest = v;
644 gbest_row = row;
645 }
646
647 // First-acceptable out-of-tile pivot: a candidate that
648 // reaches the threshold already beats the sub-threshold
649 // in-tile pivot, so stop scanning. The running max above is
650 // kept as the fallback when no candidate is acceptable.
651 if constexpr (Config.oot_first_acceptable)
652 if (v >= oot_pivot_threshold) break;
653 }
654
655 if (gbest < singular_floor) return false;
656 piv_row = gbest_row;
657 }
658
659 {
660 // Swap the permutation entries unconditionally: when
661 // piv_row == gc the exchanges write back the same values
662 // (bitwise no-ops), and removing the comparison removes a
663 // divergent branch from the hot path (the lanes of a warp can
664 // pick different pivots).
665 const int tmp = TDLS_LUPP_PIV(gc);
666 TDLS_LUPP_PIV(gc) = TDLS_LUPP_PIV(piv_row);
667 TDLS_LUPP_PIV(piv_row) = tmp;
668
669 if constexpr (fuse_rhs && !transposed) {
670 // The fused RHS follows the rows through pivoting. For an
671 // internal y the swap is predicated over the unrolled index
672 // range: y must never be dynamically indexed or it is
673 // demoted to local memory. Under the column-major layout y
674 // is indexed by the columns of S, which pivoting never moves.
675 if constexpr (internal_rhs) {
676 if constexpr (Config.unroll_loops) {
678 for (int r = 0; r < N; ++r) {
679 if (r == piv_row) {
680 const T ty = TDLS_LUPP_Y(gc);
681 TDLS_LUPP_Y(gc) = TDLS_LUPP_Y(r);
682 TDLS_LUPP_Y(r) = ty;
683 }
684 }
685 } else {
686 for (int r = 0; r < N; ++r) {
687 if (r == piv_row) {
688 const T ty = TDLS_LUPP_Y(gc);
689 TDLS_LUPP_Y(gc) = TDLS_LUPP_Y(r);
690 TDLS_LUPP_Y(r) = ty;
691 }
692 }
693 }
694 } else {
695 const T ty = TDLS_LUPP_Y(gc);
696 TDLS_LUPP_Y(gc) = TDLS_LUPP_Y(piv_row);
697 TDLS_LUPP_Y(piv_row) = ty;
698 }
699 }
700
701 if (piv_row < k0 + k_extent) {
702 Operations::template swap_rows<k_extent>(tile, c, piv_row - k0);
703 } else {
704 // Cross-tile swap: pull the new row into the tile and
705 // replay everything it missed: prior tiles (LL), then
706 // the current tile's factored columns, on the FULL row.
707 const int phys = TDLS_LUPP_PIV(gc);
708 if constexpr (Config.unroll_loops) {
710 for (int j = 0; j < k_extent; ++j) {
711 tile[c * tile_size + j] = TDLS_LUPP_A(phys, k0 + j);
712 if constexpr (schedule == Schedule::LeftLooking) {
713 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
714 for (int p = 0; p < tile_size; ++p)
715 tile[c * tile_size + j] -=
716 TDLS_LUPP_A(phys, bj0 + p) *
717 TDLS_LUPP_A(TDLS_LUPP_PIV(bj0 + p), k0 + j);
718 }
719 }
720 } else {
721 for (int j = 0; j < k_extent; ++j) {
722 tile[c * tile_size + j] = TDLS_LUPP_A(phys, k0 + j);
723 if constexpr (schedule == Schedule::LeftLooking) {
724 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
725 for (int p = 0; p < tile_size; ++p)
726 tile[c * tile_size + j] -=
727 TDLS_LUPP_A(phys, bj0 + p) *
728 TDLS_LUPP_A(TDLS_LUPP_PIV(bj0 + p), k0 + j);
729 }
730 }
731 }
732
733 if (c > 0) {
734 T L_row[tile_size];
735 for (int t = 0; t < c; ++t) {
736 T a_t = tile[c * tile_size + 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]; // diag holds 1/pivot
740 tile[c * tile_size + t] = L_row[t];
741 }
742 for (int j = c; j < k_extent; ++j) {
743 for (int t = 0; t < c; ++t)
744 tile[c * tile_size + j] -= L_row[t] * tile[t * tile_size + j];
745 }
746 }
747 }
748 }
749
750 Operations::template eliminate_column<k_extent, k_extent>(tile, c);
751 return true;
752 }
753
775 template<int k_extent, bool internal_piv, bool internal_matrix, bool oot_diagnostics,
776 bool fuse_rhs = false, bool internal_rhs = true>
777 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
778 factor_diag_tile(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
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) {
784 if (!factor_diag_column<k_extent, internal_piv, internal_matrix, oot_diagnostics,
785 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k0,
786 tile, oot_count, c, y, rhs_stride))
787 return false;
788 }
789 } else {
790 for (int c = 0; c < k_extent; ++c) {
791 if (!factor_diag_column<k_extent, internal_piv, internal_matrix, oot_diagnostics,
792 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k0,
793 tile, oot_count, c, y, rhs_stride))
794 return false;
795 }
796 }
797 return true;
798 }
799
800 /* =====================================================================
801 RIGHT-LOOKING schedule. Step k: factor the diagonal tile, then push
802 its factors into the trailing matrix (TRSM right over the row panel,
803 TRSM down + Schur complement over the rows below). The physical-row
804 segments pk/pi are cached per tile. Their entries are known at run
805 time only, an internal permutation included: a caller-local matrix
806 is indexed through them, so on GPU it lives in local memory.
807 ===================================================================== */
808
818 template<int k_extent, int j_extent, bool internal_matrix>
819 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
820 rl_trsm_right_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT pk,
821 const T* TDLS_RESTRICT tile, const int j0) noexcept {
822 T Akj[tile_size * tile_size];
823 load_tile<k_extent, j_extent, internal_matrix>(A, A_stride, pk, j0, Akj);
824 Operations::template trsm_left_unit<k_extent, j_extent>(tile, Akj);
825 store_tile<k_extent, j_extent, internal_matrix>(A, A_stride, pk, j0, Akj);
826 }
827
844 template<int k_extent, int j_extent, bool internal_matrix, bool internal_rhs>
845 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
846 rl_trsm_right_push_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT pk,
847 const T* TDLS_RESTRICT tile, const int j0, const int k0,
848 T* TDLS_RESTRICT y, const int rhs_stride) noexcept {
849 T Akj[tile_size * tile_size];
850 load_tile<k_extent, j_extent, internal_matrix>(A, A_stride, pk, j0, Akj);
851 Operations::template trsm_left_unit<k_extent, j_extent>(tile, Akj);
852 store_tile<k_extent, j_extent, internal_matrix>(A, A_stride, pk, j0, Akj);
854 }
855
869 template<int k_extent, int i_extent, int j_extent, bool internal_matrix>
870 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
871 rl_schur_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT pk,
872 const int* TDLS_RESTRICT pi, const T* TDLS_RESTRICT Aik, const int j0) noexcept {
873 T Aij[tile_size * tile_size];
874 load_tile<i_extent, j_extent, internal_matrix>(A, A_stride, pi, j0, Aij);
875
876 if constexpr (Config.unroll_loops) {
878 for (int p = 0; p < k_extent; ++p) {
879 T Akj_row[tile_size];
881 for (int j = 0; j < j_extent; ++j)
882 Akj_row[j] = TDLS_LUPP_A(pk[p], j0 + j);
884 for (int i = 0; i < i_extent; ++i) {
885 const T L_ip = Aik[i * tile_size + p];
887 for (int j = 0; j < j_extent; ++j)
888 Aij[i * tile_size + j] -= L_ip * Akj_row[j];
889 }
890 }
891 } else {
892 for (int p = 0; p < k_extent; ++p) {
893 T Akj_row[tile_size];
894 for (int j = 0; j < j_extent; ++j)
895 Akj_row[j] = TDLS_LUPP_A(pk[p], j0 + j);
896 for (int i = 0; i < i_extent; ++i) {
897 const T L_ip = Aik[i * tile_size + p];
898 for (int j = 0; j < j_extent; ++j)
899 Aij[i * tile_size + j] -= L_ip * Akj_row[j];
900 }
901 }
902 }
903
904 store_tile<i_extent, j_extent, internal_matrix>(A, A_stride, pi, j0, Aij);
905 }
906
925 template<int k_extent, int i_extent, bool internal_piv, bool internal_matrix,
926 bool fuse_rhs = false, bool internal_rhs = true>
927 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
928 rl_update_row_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
929 const int piv_stride, const int* TDLS_RESTRICT pk,
930 const T* TDLS_RESTRICT tile, const int k, const int i0,
931 T* TDLS_RESTRICT y = nullptr, const int rhs_stride = 1) noexcept {
932 const int k0 = k * tile_size;
933
934 int pi[tile_size];
935 if constexpr (Config.unroll_loops) {
937 for (int i = 0; i < i_extent; ++i)
938 pi[i] = TDLS_LUPP_PIV(i0 + i);
939 } else {
940 for (int i = 0; i < i_extent; ++i)
941 pi[i] = TDLS_LUPP_PIV(i0 + i);
942 }
943
944 // TRSM down: Aik := Aik * U^-1
945 T Aik[tile_size * tile_size];
946 load_tile<i_extent, k_extent, internal_matrix>(A, A_stride, pi, k0, Aik);
947 Operations::template trsm_right<k_extent, i_extent>(tile, Aik);
948 store_tile<i_extent, k_extent, internal_matrix>(A, A_stride, pi, k0, Aik);
949
950 // Fused forward substitution: push the solved y_k segment into this
951 // row block while its L panel sits in registers (this is the whole
952 // point of solve_inplace: the separate forward pass reloads vanish).
953 // Under the column-major layout the forward pass runs on U^T and is
954 // pushed from the U row panel instead (rl_trsm_right_push_one).
955 if constexpr (fuse_rhs && !transposed) {
956 if constexpr (Config.unroll_loops) {
958 for (int r = 0; r < i_extent; ++r) {
959 T sum = T(0);
961 for (int j = 0; j < k_extent; ++j)
962 sum += Aik[r * tile_size + j] * TDLS_LUPP_Y(k0 + j);
963 TDLS_LUPP_Y(i0 + r) -= sum;
964 }
965 } else {
966 for (int r = 0; r < i_extent; ++r) {
967 T sum = T(0);
968 for (int j = 0; j < k_extent; ++j)
969 sum += Aik[r * tile_size + j] * TDLS_LUPP_Y(k0 + j);
970 TDLS_LUPP_Y(i0 + r) -= sum;
971 }
972 }
973 }
974
975 // Schur sweep over the trailing columns
976 for (int j = k + 1; j < full_tiles; ++j)
978 j * tile_size);
979 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
981 A, A_stride, pk, pi, Aik, full_tiles * tile_size);
982 }
983
1001 template<int k_extent, bool internal_piv, bool internal_matrix, bool oot_diagnostics,
1002 bool fuse_rhs = false, bool internal_rhs = true>
1003 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1004 rl_step(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
1005 const int k, int& oot_count, T* TDLS_RESTRICT y = nullptr,
1006 const int rhs_stride = 1) noexcept {
1007 const int k0 = k * tile_size;
1008
1009 T tile[tile_size * tile_size];
1011 piv_stride, k0, k0, tile);
1012
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,
1015 rhs_stride))
1016 return false;
1017
1018 // Physical rows of the tile after the swaps of this step
1019 int pk[tile_size];
1020 if constexpr (Config.unroll_loops) {
1022 for (int i = 0; i < k_extent; ++i)
1023 pk[i] = TDLS_LUPP_PIV(k0 + i);
1024 } else {
1025 for (int i = 0; i < k_extent; ++i)
1026 pk[i] = TDLS_LUPP_PIV(k0 + i);
1027 }
1028
1029 store_tile<k_extent, k_extent, internal_matrix>(A, A_stride, pk, k0, tile);
1030
1031 // Fused forward substitution: this tile's y segment is final from
1032 // here on: solve it while the tile is in registers, unit-lower
1033 // (row-major layout) or with U^T (column-major layout).
1034 if constexpr (fuse_rhs) {
1035 if constexpr (transposed) {
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)
1042 TDLS_LUPP_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_Y(k0 + kk);
1043 }
1044 } else {
1045 for (int kk = 0; kk < k_extent; ++kk) {
1046 for (int i = kk + 1; i < k_extent; ++i)
1047 TDLS_LUPP_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_Y(k0 + kk);
1048 }
1049 }
1050 }
1051
1052 // TRSM right over the row panel; under the column-major layout,
1053 // solve_inplace pushes the solved z_k segment along
1054 for (int j = k + 1; j < full_tiles; ++j) {
1055 if constexpr (fuse_rhs && transposed)
1057 A, A_stride, pk, tile, j * tile_size, k0, y, rhs_stride);
1058 else
1060 j * tile_size);
1061 }
1062 if constexpr (last_tile_tail > 0 && k_extent == tile_size) {
1063 if constexpr (fuse_rhs && transposed)
1065 A, A_stride, pk, tile, full_tiles * tile_size, k0, y, rhs_stride);
1066 else
1068 A, A_stride, pk, tile, full_tiles * tile_size);
1069 }
1070
1071 // TRSM down + Schur over the rows below
1072 for (int i = k + 1; i < full_tiles; ++i)
1073 rl_update_row_one<k_extent, tile_size, internal_piv, internal_matrix, fuse_rhs,
1074 internal_rhs>(A, A_stride, piv, piv_stride, pk, tile, k,
1075 i * tile_size, y, rhs_stride);
1076 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
1077 rl_update_row_one<k_extent, last_tile_tail, internal_piv, internal_matrix, fuse_rhs,
1078 internal_rhs>(A, A_stride, piv, piv_stride, pk, tile, k,
1079 full_tiles * tile_size, y, rhs_stride);
1080 return true;
1081 }
1082
1083 /* =====================================================================
1084 LEFT-LOOKING schedule. Step k: pull every prior tile's L*U
1085 contribution into the current block column/row, then factor the
1086 diagonal tile and triangular-solve its panels. Each tile of the
1087 trailing matrix is written once per factorization (vs once per step
1088 for RL) at the cost of replaying prior tiles on each visit. The
1089 permutation is read inline (one read per row), never cached in
1090 segments.
1091 ===================================================================== */
1092
1107 template<int row_extent, int col_extent, bool internal_piv, bool internal_matrix>
1108 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1109 ll_correct_tile(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1110 const int piv_stride, const int row0, const int col0, const int k0,
1111 T* TDLS_RESTRICT t) noexcept {
1112 for (int bj0 = 0; bj0 < k0; bj0 += tile_size) {
1113 T Lt[tile_size * tile_size];
1115 A, A_stride, piv, piv_stride, row0, bj0, Lt);
1116 T Ut[tile_size * tile_size];
1118 A, A_stride, piv, piv_stride, bj0, col0, Ut);
1119 Operations::template gemm_sub<row_extent, col_extent, tile_size>(t, Lt, Ut);
1120 }
1121 }
1122
1140 template<int k_extent, int i_extent, bool internal_piv, bool internal_matrix,
1141 bool fuse_rhs = false, bool internal_rhs = true>
1142 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1143 ll_update_below_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1144 const int piv_stride, const T* TDLS_RESTRICT tile, const int k0,
1145 const int i0, T* TDLS_RESTRICT y = nullptr,
1146 const int rhs_stride = 1) noexcept {
1147 T B[tile_size * tile_size];
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);
1155
1156 // Fused forward substitution: B is the final L(i,k) panel, so push the
1157 // solved y_k segment into this row block while it is in registers.
1158 // Under the column-major layout the forward pass runs on U^T and is
1159 // pushed from the U row panel instead (ll_update_right_push_one).
1160 if constexpr (fuse_rhs && !transposed) {
1161 if constexpr (Config.unroll_loops) {
1163 for (int r = 0; r < i_extent; ++r) {
1164 T sum = T(0);
1166 for (int j = 0; j < k_extent; ++j)
1167 sum += B[r * tile_size + j] * TDLS_LUPP_Y(k0 + j);
1168 TDLS_LUPP_Y(i0 + r) -= sum;
1169 }
1170 } else {
1171 for (int r = 0; r < i_extent; ++r) {
1172 T sum = T(0);
1173 for (int j = 0; j < k_extent; ++j)
1174 sum += B[r * tile_size + j] * TDLS_LUPP_Y(k0 + j);
1175 TDLS_LUPP_Y(i0 + r) -= sum;
1176 }
1177 }
1178 }
1179 }
1180
1193 template<int k_extent, int j_extent, bool internal_piv, bool internal_matrix>
1194 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1195 ll_update_right_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1196 const int piv_stride, const T* TDLS_RESTRICT tile, const int k0,
1197 const int j0) noexcept {
1198 T B[tile_size * tile_size];
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);
1206 }
1207
1226 template<int k_extent, int j_extent, bool internal_piv, bool internal_matrix, bool internal_rhs>
1227 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1228 ll_update_right_push_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
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 {
1231 T B[tile_size * tile_size];
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);
1240 }
1241
1259 template<int k_extent, bool internal_piv, bool internal_matrix, bool oot_diagnostics,
1260 bool fuse_rhs = false, bool internal_rhs = true>
1261 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1262 ll_step(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
1263 const int k, int& oot_count, T* TDLS_RESTRICT y = nullptr,
1264 const int rhs_stride = 1) noexcept {
1265 const int k0 = k * tile_size;
1266
1267 T tile[tile_size * tile_size];
1269 piv_stride, k0, k0, tile);
1271 A, A_stride, piv, piv_stride, k0, k0, k0, tile);
1272
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,
1275 rhs_stride))
1276 return false;
1277
1279 piv_stride, k0, k0, tile);
1280
1281 // Fused forward substitution: this tile's y segment is final from
1282 // here on: solve it while the tile is in registers, unit-lower
1283 // (row-major layout) or with U^T (column-major layout).
1284 if constexpr (fuse_rhs) {
1285 if constexpr (transposed) {
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)
1292 TDLS_LUPP_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_Y(k0 + kk);
1293 }
1294 } else {
1295 for (int kk = 0; kk < k_extent; ++kk) {
1296 for (int i = kk + 1; i < k_extent; ++i)
1297 TDLS_LUPP_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_Y(k0 + kk);
1298 }
1299 }
1300 }
1301
1302 // L panel below the diagonal
1303 for (int i = k + 1; i < full_tiles; ++i)
1304 ll_update_below_one<k_extent, tile_size, internal_piv, internal_matrix, fuse_rhs,
1305 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0, i * tile_size,
1306 y, rhs_stride);
1307 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
1308 ll_update_below_one<k_extent, last_tile_tail, internal_piv, internal_matrix, fuse_rhs,
1309 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0,
1310 full_tiles * tile_size, y, rhs_stride);
1311
1312 // U row panel right of the diagonal; under the column-major layout,
1313 // solve_inplace pushes the solved z_k segment along
1314 for (int j = k + 1; j < full_tiles; ++j) {
1315 if constexpr (fuse_rhs && transposed)
1316 ll_update_right_push_one<k_extent, tile_size, internal_piv, internal_matrix,
1317 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0,
1318 j * tile_size, y, rhs_stride);
1319 else
1321 A, A_stride, piv, piv_stride, tile, k0, j * tile_size);
1322 }
1323 if constexpr (last_tile_tail > 0 && k_extent == tile_size) {
1324 if constexpr (fuse_rhs && transposed)
1325 ll_update_right_push_one<k_extent, last_tile_tail, internal_piv, internal_matrix,
1326 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0,
1327 full_tiles * tile_size, y, rhs_stride);
1328 else
1330 A, A_stride, piv, piv_stride, tile, k0, full_tiles * tile_size);
1331 }
1332 return true;
1333 }
1334
1335 /* =====================================================================
1336 Physical row interchanges (Config.row_interchange == Physical): the
1337 whole matrix is one tile, the rows exchanged as in LAPACK. Every
1338 loop runs over a full range with the triangle selected inside: nvcc
1339 does not unroll a loop whose trip count depends on an outer index,
1340 and a run-time index would send a caller-local array to the stack.
1341 ===================================================================== */
1342
1352 template<int w, bool internal_rhs>
1353 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1354 physical_interchange_rhs(T* TDLS_RESTRICT x, const int rhs_stride,
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) {
1359 TDLS_LUPP_FOR(c, 0, w, const T t = TDLS_LUPP_XW(c, j);
1360 TDLS_LUPP_XW(c, j) = TDLS_LUPP_XW(c, i); TDLS_LUPP_XW(c, i) = t;)
1361 })
1362 }
1363
1377 template<int n>
1378 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1379 physical_interchange_local(T* TDLS_RESTRICT X, const int j, const int r, const int cs,
1380 const int rs) noexcept {
1382 i, 0, N, if (i > j && i == r) {
1383 TDLS_LUPP_FOR(c, 0, n, const T t = X[j * rs + c * cs];
1384 X[j * rs + c * cs] = X[i * rs + c * cs]; X[i * rs + c * cs] = t;)
1385 })
1386 }
1387
1399 template<int w, bool internal_rhs, bool internal_piv, bool reverse>
1400 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1401 physical_permute(const int* TDLS_RESTRICT piv, [[maybe_unused]] const int piv_stride,
1402 T* TDLS_RESTRICT x, const int rhs_stride,
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)
1408 else
1409 physical_interchange_rhs<w, internal_rhs>(x, rhs_stride, xcol_stride, j,
1410 TDLS_LUPP_PIV(j));
1411 })
1412 }
1413
1435 template<bool internal_piv, bool internal_matrix, bool check>
1436 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1437 physical_step(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1438 const int piv_stride, const int j, [[maybe_unused]] const T u, bool& singular,
1439 int& weak) noexcept {
1440 T best = detail::abs(TDLS_LUPP_A(j, j));
1441 int r = j;
1443 i, 0, N, const T a = detail::abs(TDLS_LUPP_A(i, j)); if (i > j && a > best) {
1444 best = a;
1445 r = i;
1446 })
1447 if constexpr (relative_pivot_threshold < T(1)) {
1448 const T in_place = detail::abs(TDLS_LUPP_A(j, j));
1449 if (in_place >= relative_pivot_threshold * best && in_place >= oot_pivot_threshold)
1450 r = j;
1451 }
1452 const bool pivot_ok = best >= singular_floor;
1453 singular = singular || !pivot_ok;
1454 weak += best < oot_pivot_threshold ? 1 : 0;
1455 TDLS_LUPP_PIV(j) = r;
1456 // Rarely taken on matrices that keep their rows: uniform across a warp.
1457 if (r != j) {
1459 i, 0, N, if (i > j && i == r) {
1460 TDLS_LUPP_FOR(c, 0, N, const T t = TDLS_LUPP_A(j, c);
1461 TDLS_LUPP_A(j, c) = TDLS_LUPP_A(i, c); TDLS_LUPP_A(i, c) = t;)
1462 })
1463 }
1464 const T inv = pivot_ok ? T(1) / TDLS_LUPP_A(j, j) : T(0);
1465 TDLS_LUPP_A(j, j) = inv;
1466 bool ok = true;
1468 i, 0, N, if (i > j) {
1469 const T l = TDLS_LUPP_A(i, j) * inv;
1470 TDLS_LUPP_A(i, j) = l;
1471 if constexpr (check) ok = ok && u * detail::abs(l) <= T(1);
1472 TDLS_LUPP_FOR(c, 0, N, if (c > j) TDLS_LUPP_A(i, c) -= l * TDLS_LUPP_A(j, c);)
1473 })
1474 return ok;
1475 }
1476
1489 template<bool internal_piv, bool internal_matrix>
1490 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1491 physical_factorize(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1492 const int piv_stride, int& weak) noexcept {
1493 bool singular = false;
1494 weak = 0;
1495 TDLS_LUPP_FOR(j, 0, N,
1497 A, A_stride, piv, piv_stride, j, T(0), singular, weak);)
1498 return !singular;
1499 }
1500
1514 template<bool internal_piv, bool internal_matrix>
1515 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
1516 physical_factorize_checked(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1517 const int piv_stride, const T u) noexcept {
1518 bool ok = true, singular = false;
1519 int weak = 0;
1520 TDLS_LUPP_FOR(j, 0, N,
1522 A, A_stride, piv, piv_stride, j, u, singular, weak) &&
1523 ok;)
1524 return singular ? 2 : ok ? 0 : 1;
1525 }
1526
1541 template<int w, bool internal_rhs, bool internal_piv, bool internal_matrix>
1542 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1543 physical_forward(const T* TDLS_RESTRICT A, const int A_stride,
1544 [[maybe_unused]] const int* TDLS_RESTRICT piv,
1545 [[maybe_unused]] const int piv_stride, T* TDLS_RESTRICT x,
1546 const int rhs_stride, const std::size_t xcol_stride) noexcept {
1547 if constexpr (!transposed) {
1548 physical_permute<w, internal_rhs, internal_piv, false>(piv, piv_stride, x, rhs_stride,
1549 xcol_stride);
1550 TDLS_LUPP_FOR(j, 0, N,
1552 i, 0, N, if (i > j) {
1553 const T l = TDLS_LUPP_A(i, j);
1554 TDLS_LUPP_FOR(c, 0, w,
1555 TDLS_LUPP_XW(c, i) -= l * TDLS_LUPP_XW(c, j);)
1556 }))
1557 } else {
1559 j, 0, N,
1560 TDLS_LUPP_FOR(c, 0, w, TDLS_LUPP_XW(c, j) *= TDLS_LUPP_A(j, j);) TDLS_LUPP_FOR(
1561 q, 0, N, if (q > j) {
1562 const T uq = TDLS_LUPP_A(j, q);
1563 TDLS_LUPP_FOR(c, 0, w, TDLS_LUPP_XW(c, q) -= uq * TDLS_LUPP_XW(c, j);)
1564 }))
1565 }
1566 }
1567
1583 template<int w, bool internal_rhs, bool internal_piv, bool internal_matrix>
1584 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1585 physical_backward(const T* TDLS_RESTRICT A, const int A_stride,
1586 [[maybe_unused]] const int* TDLS_RESTRICT piv,
1587 [[maybe_unused]] const int piv_stride, T* TDLS_RESTRICT x,
1588 const int rhs_stride, const std::size_t xcol_stride) noexcept {
1589 if constexpr (!transposed) {
1591 t, 0, N, const int j = N - 1 - t; TDLS_LUPP_FOR(
1592 c, 0, w, T v = TDLS_LUPP_XW(c, j);
1593 TDLS_LUPP_FOR(q, 0, N, if (q > j) v -= TDLS_LUPP_A(j, q) * TDLS_LUPP_XW(c, q);)
1594 TDLS_LUPP_XW(c, j) = v * TDLS_LUPP_A(j, j);))
1595 } else {
1596 TDLS_LUPP_FOR(t, 0, N, const int j = N - 1 - t; TDLS_LUPP_FOR(
1597 q, 0, N, if (q > j) {
1598 const T l = TDLS_LUPP_A(q, j);
1599 TDLS_LUPP_FOR(c, 0, w,
1600 TDLS_LUPP_XW(c, j) -= l * TDLS_LUPP_XW(c, q);)
1601 }))
1602 physical_permute<w, internal_rhs, internal_piv, true>(piv, piv_stride, x, rhs_stride,
1603 xcol_stride);
1604 }
1605 }
1606
1622 template<int nrhs, int pass_width, int stages, bool internal_rhs, bool internal_piv,
1623 bool internal_matrix>
1624 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1625 physical_multirhs(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
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 =
1634 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
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);
1641 }
1642 if constexpr (tail > 0) {
1643 constexpr int c0 = nrhs - tail;
1644 const std::size_t off =
1645 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
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);
1652 }
1653 }
1654
1655 /* =====================================================================
1656 FACTORIZE - public entry point.
1657 ===================================================================== */
1658
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,
1697 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
1698 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1699 factorize(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_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),
1706 oot_count, y,
1707 detail::stride_or<rhs_static_stride>(rhs_stride));
1708 } else {
1709 if constexpr (physical) {
1710 int weak = 0;
1711 const bool ok = physical_factorize<internal_piv, internal_matrix>(A, A_stride, piv,
1712 piv_stride, weak);
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);
1718 return true;
1719 }
1720
1721 if constexpr (oot_diagnostics) oot_count = 0;
1722
1723 if constexpr (Config.unroll_loops) {
1725 for (int i = 0; i < N; ++i)
1726 TDLS_LUPP_PIV(i) = i;
1727 } else {
1728 for (int i = 0; i < N; ++i)
1729 TDLS_LUPP_PIV(i) = i;
1730 }
1731
1732 if constexpr (schedule == Schedule::RightLooking) {
1733 for (int k = 0; k < full_tiles; ++k)
1734 if (!rl_step<tile_size, internal_piv, internal_matrix, oot_diagnostics,
1735 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k, oot_count,
1736 y, rhs_stride))
1737 return false;
1738 if constexpr (last_tile_tail > 0)
1739 if (!rl_step<last_tile_tail, internal_piv, internal_matrix, oot_diagnostics,
1740 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, full_tiles,
1741 oot_count, y, rhs_stride))
1742 return false;
1743 } else {
1744 for (int k = 0; k < full_tiles; ++k)
1745 if (!ll_step<tile_size, internal_piv, internal_matrix, oot_diagnostics,
1746 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k, oot_count,
1747 y, rhs_stride))
1748 return false;
1749 if constexpr (last_tile_tail > 0)
1750 if (!ll_step<last_tile_tail, internal_piv, internal_matrix, oot_diagnostics,
1751 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, full_tiles,
1752 oot_count, y, rhs_stride))
1753 return false;
1754 }
1755
1756 return true;
1757 }
1758 }
1759
1773 template<bool internal_piv, bool internal_matrix, int A_static_stride = dynamic_stride,
1774 int piv_static_stride = dynamic_stride>
1775 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1776 factorize(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1777 const int piv_stride) noexcept {
1778 int unused = 0;
1779 return factorize<internal_piv, internal_matrix, false, false, true, A_static_stride,
1780 piv_static_stride>(A, A_stride, piv, piv_stride, unused);
1781 }
1782
1783 /* =====================================================================
1784 SUBSTITUTION - schedule-independent.
1785 Forward: unit-lower solve tile by tile, each solved segment pushed
1786 into the tiles below it. Backward: each segment first pulls the
1787 trailing contributions, then upper-solves its diagonal tile.
1788 ===================================================================== */
1789
1807 template<int k_extent, int m_extent, int pass_width, bool internal_rhs, bool internal_piv,
1808 bool internal_matrix>
1809 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1810 fwd_push_one(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
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 {
1813 T Lmk[tile_size * tile_size];
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) {
1821 T sum = T(0);
1823 for (int j = 0; j < k_extent; ++j)
1824 sum += Lmk[i * tile_size + j] * TDLS_LUPP_XW(w, k0 + j);
1825 TDLS_LUPP_XW(w, m0 + i) -= sum;
1826 }
1827 }
1828 } else {
1829 for (int i = 0; i < m_extent; ++i) {
1830 for (int w = 0; w < pass_width; ++w) {
1831 T sum = T(0);
1832 for (int j = 0; j < k_extent; ++j)
1833 sum += Lmk[i * tile_size + j] * TDLS_LUPP_XW(w, k0 + j);
1834 TDLS_LUPP_XW(w, m0 + i) -= sum;
1835 }
1836 }
1837 }
1838 }
1839
1855 template<int k_extent, int pass_width, bool internal_rhs, bool internal_piv,
1856 bool internal_matrix>
1857 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1858 fwd_step(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1859 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
1860 const std::size_t xcol_stride, const int k) noexcept {
1861 const int k0 = k * tile_size;
1862
1863 T Lkk[tile_size * tile_size];
1865 k0, k0, Lkk);
1866
1867 // In-tile unit-lower solve
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)
1875 TDLS_LUPP_XW(w, k0 + i) -=
1876 Lkk[i * tile_size + kk] * TDLS_LUPP_XW(w, k0 + kk);
1877 }
1878 }
1879 } else {
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)
1883 TDLS_LUPP_XW(w, k0 + i) -=
1884 Lkk[i * tile_size + kk] * TDLS_LUPP_XW(w, k0 + kk);
1885 }
1886 }
1887 }
1888
1889 // Push into the tiles below
1890 for (int m = k + 1; m < full_tiles; ++m)
1891 fwd_push_one<k_extent, tile_size, pass_width, internal_rhs, internal_piv,
1892 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1893 k0, m * tile_size);
1894 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
1895 fwd_push_one<k_extent, last_tile_tail, pass_width, internal_rhs, internal_piv,
1896 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1897 k0, full_tiles * tile_size);
1898 }
1899
1917 template<int k_extent, int m_extent, int pass_width, bool internal_rhs, bool internal_piv,
1918 bool internal_matrix>
1919 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1920 bwd_pull_one(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
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 {
1923 T Ukm[tile_size * tile_size];
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) {
1931 T sum = T(0);
1933 for (int j = 0; j < m_extent; ++j)
1934 sum += Ukm[i * tile_size + j] * TDLS_LUPP_XW(w, m0 + j);
1935 TDLS_LUPP_XW(w, k0 + i) -= sum;
1936 }
1937 }
1938 } else {
1939 for (int i = 0; i < k_extent; ++i) {
1940 for (int w = 0; w < pass_width; ++w) {
1941 T sum = T(0);
1942 for (int j = 0; j < m_extent; ++j)
1943 sum += Ukm[i * tile_size + j] * TDLS_LUPP_XW(w, m0 + j);
1944 TDLS_LUPP_XW(w, k0 + i) -= sum;
1945 }
1946 }
1947 }
1948 }
1949
1965 template<int k_extent, int pass_width, bool internal_rhs, bool internal_piv,
1966 bool internal_matrix>
1967 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1968 bwd_step(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1969 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
1970 const std::size_t xcol_stride, const int k) noexcept {
1971 const int k0 = k * tile_size;
1972
1973 // Pull the trailing contributions
1974 for (int m = k + 1; m < full_tiles; ++m)
1975 bwd_pull_one<k_extent, tile_size, pass_width, internal_rhs, internal_piv,
1976 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1977 k0, m * tile_size);
1978 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
1979 bwd_pull_one<k_extent, last_tile_tail, pass_width, internal_rhs, internal_piv,
1980 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1981 k0, full_tiles * tile_size);
1982
1983 // In-tile upper solve
1984 T Ukk[tile_size * tile_size];
1986 k0, k0, Ukk);
1987
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)
1993 TDLS_LUPP_XW(w, k0 + kk) *= Ukk[kk * tile_size + kk]; // diag holds 1/pivot
1995 for (int i = 0; i < kk; ++i) {
1997 for (int w = 0; w < pass_width; ++w)
1998 TDLS_LUPP_XW(w, k0 + i) -=
1999 Ukk[i * tile_size + kk] * TDLS_LUPP_XW(w, k0 + kk);
2000 }
2001 }
2002 } else {
2003 for (int kk = k_extent - 1; kk >= 0; --kk) {
2004 for (int w = 0; w < pass_width; ++w)
2005 TDLS_LUPP_XW(w, k0 + kk) *= Ukk[kk * tile_size + kk]; // diag holds 1/pivot
2006 for (int i = 0; i < kk; ++i) {
2007 for (int w = 0; w < pass_width; ++w)
2008 TDLS_LUPP_XW(w, k0 + i) -=
2009 Ukk[i * tile_size + kk] * TDLS_LUPP_XW(w, k0 + kk);
2010 }
2011 }
2012 }
2013 }
2014
2028 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
2029 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2030 bwd_only(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
2031 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
2032 const std::size_t xcol_stride) noexcept {
2033 if constexpr (last_tile_tail > 0)
2035 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, full_tiles);
2036 for (int k = full_tiles - 1; k >= 0; --k)
2038 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
2039 }
2040
2056 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
2057 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2058 fwd_bwd(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
2059 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
2060 const std::size_t xcol_stride) noexcept {
2061 for (int k = 0; k < full_tiles; ++k)
2063 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
2064 if constexpr (last_tile_tail > 0)
2066 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, full_tiles);
2067
2069 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2070 }
2071
2072 /* =====================================================================
2073 TRANSPOSED SUBSTITUTION - column-major layout. With P S = L U and
2074 S = A^T, A = U^T L^T P, so A x = b is U^T z = b (forward, the U
2075 tiles read transposed, the diagonal holding 1/pivot), then
2076 L^T w = z (backward, the L tiles read transposed, unit diagonal),
2077 then x = P^T w. The right-hand side is indexed by the columns of S
2078 and is never permuted before the sweeps. The tile kernels below are
2079 shared by the fused forward pass of solve_inplace and by the
2080 separate passes, which therefore agree bitwise when the compiler
2081 contracts multiply-adds alike in both (-ffp-contract=off, or
2082 --fmad=false for nvcc).
2083 ===================================================================== */
2084
2095 template<int k_extent, int pass_width, bool internal_rhs>
2096 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2097 transposed_solve_tile_upper(const T* TDLS_RESTRICT t, T* TDLS_RESTRICT x, const int rhs_stride,
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)
2104 TDLS_LUPP_XW(w, k0 + kk) *= t[kk * tile_size + kk]; // diag holds 1/pivot
2106 for (int i = kk + 1; i < k_extent; ++i) {
2108 for (int w = 0; w < pass_width; ++w)
2109 TDLS_LUPP_XW(w, k0 + i) -= t[kk * tile_size + i] * TDLS_LUPP_XW(w, k0 + kk);
2110 }
2111 }
2112 } else {
2113 for (int kk = 0; kk < k_extent; ++kk) {
2114 for (int w = 0; w < pass_width; ++w)
2115 TDLS_LUPP_XW(w, k0 + kk) *= t[kk * tile_size + kk]; // diag holds 1/pivot
2116 for (int i = kk + 1; i < k_extent; ++i) {
2117 for (int w = 0; w < pass_width; ++w)
2118 TDLS_LUPP_XW(w, k0 + i) -= t[kk * tile_size + i] * TDLS_LUPP_XW(w, k0 + kk);
2119 }
2120 }
2121 }
2122 }
2123
2135 template<int k_extent, int m_extent, int pass_width, bool internal_rhs>
2136 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2137 transposed_push_one(const T* TDLS_RESTRICT u, T* TDLS_RESTRICT x, const int rhs_stride,
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) {
2144 T sum = T(0);
2146 for (int j = 0; j < k_extent; ++j)
2147 sum += u[j * tile_size + i] * TDLS_LUPP_XW(w, k0 + j);
2148 TDLS_LUPP_XW(w, m0 + i) -= sum;
2149 }
2150 }
2151 } else {
2152 for (int i = 0; i < m_extent; ++i) {
2153 for (int w = 0; w < pass_width; ++w) {
2154 T sum = T(0);
2155 for (int j = 0; j < k_extent; ++j)
2156 sum += u[j * tile_size + i] * TDLS_LUPP_XW(w, k0 + j);
2157 TDLS_LUPP_XW(w, m0 + i) -= sum;
2158 }
2159 }
2160 }
2161 }
2162
2174 template<int k_extent, int m_extent, int pass_width, bool internal_rhs>
2175 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2176 transposed_pull_one(const T* TDLS_RESTRICT l, T* TDLS_RESTRICT x, const int rhs_stride,
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) {
2183 T sum = T(0);
2185 for (int j = 0; j < m_extent; ++j)
2186 sum += l[j * tile_size + i] * TDLS_LUPP_XW(w, m0 + j);
2187 TDLS_LUPP_XW(w, k0 + i) -= sum;
2188 }
2189 }
2190 } else {
2191 for (int i = 0; i < k_extent; ++i) {
2192 for (int w = 0; w < pass_width; ++w) {
2193 T sum = T(0);
2194 for (int j = 0; j < m_extent; ++j)
2195 sum += l[j * tile_size + i] * TDLS_LUPP_XW(w, m0 + j);
2196 TDLS_LUPP_XW(w, k0 + i) -= sum;
2197 }
2198 }
2199 }
2200 }
2201
2212 template<int k_extent, int pass_width, bool internal_rhs>
2213 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
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)
2224 TDLS_LUPP_XW(w, k0 + i) -= l[kk * tile_size + i] * TDLS_LUPP_XW(w, k0 + kk);
2225 }
2226 }
2227 } else {
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)
2231 TDLS_LUPP_XW(w, k0 + i) -= l[kk * tile_size + i] * TDLS_LUPP_XW(w, k0 + kk);
2232 }
2233 }
2234 }
2235 }
2236
2252 template<int k_extent, int pass_width, bool internal_rhs, bool internal_piv,
2253 bool internal_matrix>
2254 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2255 transposed_fwd_step(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
2256 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
2257 const std::size_t xcol_stride, const int k) noexcept {
2258 const int k0 = k * tile_size;
2259 T Ukk[tile_size * tile_size];
2261 k0, k0, Ukk);
2263 xcol_stride, k0);
2264 for (int m = k + 1; m < full_tiles; ++m) {
2265 T Ukm[tile_size * tile_size];
2267 A, A_stride, piv, piv_stride, k0, m * tile_size, Ukm);
2269 Ukm, x, rhs_stride, xcol_stride, k0, m * tile_size);
2270 }
2271 if constexpr (last_tile_tail > 0 && k_extent == tile_size) {
2272 T Ukm[tile_size * tile_size];
2274 A, A_stride, piv, piv_stride, k0, full_tiles * tile_size, Ukm);
2276 Ukm, x, rhs_stride, xcol_stride, k0, full_tiles * tile_size);
2277 }
2278 }
2279
2295 template<int k_extent, int pass_width, bool internal_rhs, bool internal_piv,
2296 bool internal_matrix>
2297 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2298 transposed_bwd_step(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
2299 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
2300 const std::size_t xcol_stride, const int k) noexcept {
2301 const int k0 = k * tile_size;
2302 for (int m = k + 1; m < full_tiles; ++m) {
2303 T Lmk[tile_size * tile_size];
2305 A, A_stride, piv, piv_stride, m * tile_size, k0, Lmk);
2307 Lmk, x, rhs_stride, xcol_stride, k0, m * tile_size);
2308 }
2309 if constexpr (last_tile_tail > 0 && k_extent == tile_size) {
2310 T Lmk[tile_size * tile_size];
2312 A, A_stride, piv, piv_stride, full_tiles * tile_size, k0, Lmk);
2314 Lmk, x, rhs_stride, xcol_stride, k0, full_tiles * tile_size);
2315 }
2316 T Lkk[tile_size * tile_size];
2318 k0, k0, Lkk);
2320 xcol_stride, k0);
2321 }
2322
2336 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
2337 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2338 transposed_bwd_only(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
2339 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
2340 const std::size_t xcol_stride) noexcept {
2341 if constexpr (last_tile_tail > 0)
2342 transposed_bwd_step<last_tile_tail, pass_width, internal_rhs, internal_piv,
2343 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride,
2344 xcol_stride, full_tiles);
2345 for (int k = full_tiles - 1; k >= 0; --k)
2347 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
2348 }
2349
2364 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
2365 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2366 transposed_fwd_bwd(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
2367 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
2368 const std::size_t xcol_stride) noexcept {
2369 for (int k = 0; k < full_tiles; ++k)
2371 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
2372 if constexpr (last_tile_tail > 0)
2373 transposed_fwd_step<last_tile_tail, pass_width, internal_rhs, internal_piv,
2374 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride,
2375 xcol_stride, full_tiles);
2376
2378 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2379 }
2380
2394 template<int pass_width, bool internal_rhs, bool internal_piv>
2395 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2396 unpermute_cycles(const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
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)
2408 carry[w] = TDLS_LUPP_XW(w, s);
2409 int dst = TDLS_LUPP_PIV(s);
2410 while (dst != s) {
2412 for (int w = 0; w < pass_width; ++w) {
2413 const T displaced = TDLS_LUPP_XW(w, dst);
2414 TDLS_LUPP_XW(w, dst) = carry[w];
2415 carry[w] = displaced;
2416 }
2417 visited |= mask_t(1) << dst;
2418 dst = TDLS_LUPP_PIV(dst);
2419 }
2421 for (int w = 0; w < pass_width; ++w)
2422 TDLS_LUPP_XW(w, s) = carry[w];
2423 visited |= mask_t(1) << s;
2424 }
2425 } else {
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)
2430 carry[w] = TDLS_LUPP_XW(w, s);
2431 int dst = TDLS_LUPP_PIV(s);
2432 while (dst != s) {
2433 for (int w = 0; w < pass_width; ++w) {
2434 const T displaced = TDLS_LUPP_XW(w, dst);
2435 TDLS_LUPP_XW(w, dst) = carry[w];
2436 carry[w] = displaced;
2437 }
2438 visited |= mask_t(1) << dst;
2439 dst = TDLS_LUPP_PIV(dst);
2440 }
2441 for (int w = 0; w < pass_width; ++w)
2442 TDLS_LUPP_XW(w, s) = carry[w];
2443 visited |= mask_t(1) << s;
2444 }
2445 }
2446 } else {
2447 for (int s = 0; s < N; ++s) {
2448 int probe = TDLS_LUPP_PIV(s);
2449 while (probe > s)
2450 probe = TDLS_LUPP_PIV(probe);
2451 if (probe != s) continue;
2452
2453 T carry[pass_width];
2454 for (int w = 0; w < pass_width; ++w)
2455 carry[w] = TDLS_LUPP_XW(w, s);
2456 int dst = TDLS_LUPP_PIV(s);
2457 while (dst != s) {
2458 for (int w = 0; w < pass_width; ++w) {
2459 const T displaced = TDLS_LUPP_XW(w, dst);
2460 TDLS_LUPP_XW(w, dst) = carry[w];
2461 carry[w] = displaced;
2462 }
2463 dst = TDLS_LUPP_PIV(dst);
2464 }
2465 for (int w = 0; w < pass_width; ++w)
2466 TDLS_LUPP_XW(w, s) = carry[w];
2467 }
2468 }
2469 }
2470
2482 template<int pass_width, bool internal_piv>
2483 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2484 unpermute_predicated(const int* TDLS_RESTRICT piv, const int piv_stride,
2485 T* TDLS_RESTRICT x) noexcept {
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];
2490 int p[N];
2491 if constexpr (Config.unroll_loops) {
2493 for (int i = 0; i < pass_width * N; ++i)
2494 w_copy[i] = x[i];
2496 for (int i = 0; i < N; ++i)
2497 p[i] = TDLS_LUPP_PIV(i);
2499 for (int j = 0; j < N; ++j) {
2501 for (int w = 0; w < pass_width; ++w) {
2502 T v = T(0);
2504 for (int i = 0; i < N; ++i)
2505 if (p[i] == j) v = w_copy[w * N + i];
2506 TDLS_LUPP_XW(w, j) = v;
2507 }
2508 }
2509 } else {
2510 for (int i = 0; i < pass_width * N; ++i)
2511 w_copy[i] = x[i];
2512 for (int i = 0; i < N; ++i)
2513 p[i] = TDLS_LUPP_PIV(i);
2514 for (int j = 0; j < N; ++j) {
2515 for (int w = 0; w < pass_width; ++w) {
2516 T v = T(0);
2517 for (int i = 0; i < N; ++i)
2518 if (p[i] == j) v = w_copy[w * N + i];
2519 TDLS_LUPP_XW(w, j) = v;
2520 }
2521 }
2522 }
2523 }
2524
2535 template<int pass_width, bool internal_rhs, bool internal_piv>
2536 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2537 unpermute(const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
2538 const int rhs_stride, const std::size_t xcol_stride) noexcept {
2539 if constexpr (internal_rhs)
2541 else
2542 unpermute_cycles<pass_width, internal_rhs, internal_piv>(piv, piv_stride, x, rhs_stride,
2543 xcol_stride);
2544 }
2545
2546 /* =====================================================================
2547 SUBSTITUTION - public entry points.
2548 ===================================================================== */
2549
2571 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
2572 int A_static_stride = dynamic_stride, int piv_static_stride = dynamic_stride,
2573 int rhs_static_stride = dynamic_stride>
2574 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2575 substitute(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
2576 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
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));
2584 } else {
2585 if constexpr (physical) {
2586 TDLS_LUPP_FOR(i, 0, N, TDLS_LUPP_X(i) = TDLS_LUPP_B(i);)
2588 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2589 return;
2590 }
2591 if constexpr (transposed) {
2592 // Column-major layout: b is consumed unpermuted, the solution is
2593 // permuted at the end (x = P^T w).
2594 if constexpr (Config.unroll_loops) {
2596 for (int i = 0; i < N; ++i)
2597 TDLS_LUPP_X(i) = TDLS_LUPP_B(i);
2598 } else {
2599 for (int i = 0; i < N; ++i)
2600 TDLS_LUPP_X(i) = TDLS_LUPP_B(i);
2601 }
2603 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2604 unpermute<1, internal_rhs, internal_piv>(piv, piv_stride, x, rhs_stride, 0);
2605 } else {
2606 if constexpr (internal_rhs) {
2607 // Predicated gather: b[piv[i]] would index the internal b with
2608 // a runtime value and demote it to local memory. The equality
2609 // sweep keeps every index compile-time (measured: local memory
2610 // eliminated, identical values).
2611 if constexpr (Config.unroll_loops) {
2613 for (int i = 0; i < N; ++i) {
2614 const int p = TDLS_LUPP_PIV(i);
2615 T v = T(0);
2617 for (int j = 0; j < N; ++j)
2618 if (p == j) v = b[j];
2619 TDLS_LUPP_X(i) = v;
2620 }
2621 } else {
2622 for (int i = 0; i < N; ++i) {
2623 const int p = TDLS_LUPP_PIV(i);
2624 T v = T(0);
2625 for (int j = 0; j < N; ++j)
2626 if (p == j) v = b[j];
2627 TDLS_LUPP_X(i) = v;
2628 }
2629 }
2630 } else {
2631 if constexpr (Config.unroll_loops) {
2633 for (int i = 0; i < N; ++i)
2635 } else {
2636 for (int i = 0; i < N; ++i)
2638 }
2639 }
2641 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2642 }
2643 }
2644 }
2645
2668 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
2669 int A_static_stride = dynamic_stride, int piv_static_stride = dynamic_stride,
2670 int rhs_static_stride = dynamic_stride>
2671 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2672 substitute_canonical(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
2673 const int piv_stride, const int col, T* TDLS_RESTRICT x,
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));
2681 } else {
2682 if constexpr (physical) {
2683 TDLS_LUPP_FOR(i, 0, N, TDLS_LUPP_X(i) = i == col ? T(1) : T(0);)
2685 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2686 return;
2687 }
2689 A, A_stride, piv, piv_stride, col, x, rhs_stride, 0);
2690 }
2691 }
2692
2713 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
2714 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2716 const int* TDLS_RESTRICT piv, const int piv_stride,
2717 const int col0, T* TDLS_RESTRICT x, const int rhs_stride,
2718 const std::size_t xcol_stride) noexcept {
2719 if constexpr (transposed) {
2720 // Column-major layout: e_col unpermuted, the solution permuted at
2721 // the end (x = P^T w).
2722 if constexpr (Config.unroll_loops) {
2724 for (int i = 0; i < N; ++i) {
2726 for (int w = 0; w < pass_width; ++w)
2727 TDLS_LUPP_XW(w, i) = (i == col0 + w) ? T(1) : T(0);
2728 }
2729 } else {
2730 for (int i = 0; i < N; ++i) {
2731 for (int w = 0; w < pass_width; ++w)
2732 TDLS_LUPP_XW(w, i) = (i == col0 + w) ? T(1) : T(0);
2733 }
2734 }
2736 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2737 unpermute<pass_width, internal_rhs, internal_piv>(piv, piv_stride, x, rhs_stride,
2738 xcol_stride);
2739 } else {
2740 if constexpr (Config.unroll_loops) {
2742 for (int i = 0; i < N; ++i) {
2743 const int p = TDLS_LUPP_PIV(i);
2745 for (int w = 0; w < pass_width; ++w)
2746 TDLS_LUPP_XW(w, i) = (p == col0 + w) ? T(1) : T(0);
2747 }
2748 } else {
2749 for (int i = 0; i < N; ++i) {
2750 const int p = TDLS_LUPP_PIV(i);
2751 for (int w = 0; w < pass_width; ++w)
2752 TDLS_LUPP_XW(w, i) = (p == col0 + w) ? T(1) : T(0);
2753 }
2754 }
2756 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2757 }
2758 }
2759
2794 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
2795 int pass_width = 0, int A_static_stride = dynamic_stride,
2796 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
2797 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2798 substitute_canonical_multirhs(const T* TDLS_RESTRICT A, const int A_stride,
2799 const int* TDLS_RESTRICT piv, const int piv_stride,
2800 const int col0, T* TDLS_RESTRICT x, const int rhs_stride,
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>) {
2804 return substitute_canonical_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix,
2805 pass_width>(
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);
2809 } else {
2810 if constexpr (physical) {
2812 c, 0, nrhs,
2813 TDLS_LUPP_FOR(i, 0, N, TDLS_LUPP_XW(c, i) = i == col0 + c ? T(1) : T(0);))
2815 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2816 return;
2817 }
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 =
2824 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
2825 substitute_canonical_multirhs_pass<columns_per_pass, internal_rhs, internal_piv,
2826 internal_matrix>(
2827 A, A_stride, piv, piv_stride, col0 + c0, x + off, rhs_stride, xcol_stride);
2828 }
2829 if constexpr (nrhs % columns_per_pass > 0) {
2830 constexpr int c0 = nrhs - nrhs % columns_per_pass;
2831 const std::size_t off =
2832 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
2833 substitute_canonical_multirhs_pass<nrhs % columns_per_pass, internal_rhs,
2834 internal_piv, internal_matrix>(
2835 A, A_stride, piv, piv_stride, col0 + c0, x + off, rhs_stride, xcol_stride);
2836 }
2837 }
2838 }
2839
2860 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
2861 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2862 substitute_multirhs_pass(const T* TDLS_RESTRICT A, const int A_stride,
2863 const int* TDLS_RESTRICT piv, const int piv_stride,
2864 const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
2865 const std::size_t xcol_stride) noexcept {
2866 if constexpr (transposed) {
2867 // Column-major layout, as in substitute.
2868 if constexpr (Config.unroll_loops) {
2870 for (int i = 0; i < N; ++i) {
2872 for (int w = 0; w < pass_width; ++w)
2873 TDLS_LUPP_XW(w, i) = TDLS_LUPP_BW(w, i);
2874 }
2875 } else {
2876 for (int i = 0; i < N; ++i) {
2877 for (int w = 0; w < pass_width; ++w)
2878 TDLS_LUPP_XW(w, i) = TDLS_LUPP_BW(w, i);
2879 }
2880 }
2882 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2883 unpermute<pass_width, internal_rhs, internal_piv>(piv, piv_stride, x, rhs_stride,
2884 xcol_stride);
2885 } else {
2886 if constexpr (internal_rhs) {
2887 // Predicated gather, exactly as in substitute: b[piv[i]] would
2888 // index the internal block with a runtime value and demote it
2889 // to local memory.
2890 if constexpr (Config.unroll_loops) {
2892 for (int i = 0; i < N; ++i) {
2893 const int p = TDLS_LUPP_PIV(i);
2895 for (int w = 0; w < pass_width; ++w) {
2896 T v = T(0);
2898 for (int j = 0; j < N; ++j)
2899 if (p == j) v = TDLS_LUPP_BW(w, j);
2900 TDLS_LUPP_XW(w, i) = v;
2901 }
2902 }
2903 } else {
2904 for (int i = 0; i < N; ++i) {
2905 const int p = TDLS_LUPP_PIV(i);
2906 for (int w = 0; w < pass_width; ++w) {
2907 T v = T(0);
2908 for (int j = 0; j < N; ++j)
2909 if (p == j) v = TDLS_LUPP_BW(w, j);
2910 TDLS_LUPP_XW(w, i) = v;
2911 }
2912 }
2913 }
2914 } else {
2915 if constexpr (Config.unroll_loops) {
2917 for (int i = 0; i < N; ++i) {
2918 const int p = TDLS_LUPP_PIV(i);
2920 for (int w = 0; w < pass_width; ++w)
2921 TDLS_LUPP_XW(w, i) = TDLS_LUPP_BW(w, p);
2922 }
2923 } else {
2924 for (int i = 0; i < N; ++i) {
2925 const int p = TDLS_LUPP_PIV(i);
2926 for (int w = 0; w < pass_width; ++w)
2927 TDLS_LUPP_XW(w, i) = TDLS_LUPP_BW(w, p);
2928 }
2929 }
2930 }
2932 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2933 }
2934 }
2935
2974 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
2975 int pass_width = 0, int A_static_stride = dynamic_stride,
2976 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
2977 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2978 substitute_multirhs(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
2979 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
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>) {
2983 return substitute_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix,
2984 pass_width>(
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);
2988 } else {
2989 if constexpr (physical) {
2990 TDLS_LUPP_FOR(c, 0, nrhs,
2991 TDLS_LUPP_FOR(i, 0, N, TDLS_LUPP_XW(c, i) = TDLS_LUPP_BW(c, i);))
2993 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2994 return;
2995 }
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 =
3002 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
3003 substitute_multirhs_pass<columns_per_pass, internal_rhs, internal_piv,
3004 internal_matrix>(A, A_stride, piv, piv_stride, b + off,
3005 x + off, rhs_stride, xcol_stride);
3006 }
3007 if constexpr (nrhs % columns_per_pass > 0) {
3008 constexpr int c0 = nrhs - nrhs % columns_per_pass;
3009 const std::size_t off =
3010 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
3011 substitute_multirhs_pass<nrhs % columns_per_pass, internal_rhs, internal_piv,
3012 internal_matrix>(A, A_stride, piv, piv_stride, b + off,
3013 x + off, rhs_stride, xcol_stride);
3014 }
3015 }
3016 }
3017
3043 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
3044 int A_static_stride = dynamic_stride, int piv_static_stride = dynamic_stride,
3045 int rhs_static_stride = dynamic_stride>
3046 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
3047 substitute_inplace(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
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));
3055 } else {
3056 if constexpr (physical) {
3058 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
3059 return;
3060 }
3061 if constexpr (transposed) {
3062 // Column-major layout: the sweeps on the unpermuted x, then the
3063 // inverse cycle rotation (x = P^T w).
3065 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
3066 unpermute_cycles<1, internal_rhs, internal_piv>(piv, piv_stride, x, rhs_stride, 0);
3067 } else {
3068 if constexpr (N <= 64) {
3069 // Bitmask cycle decomposition: one visited bit per row, the
3070 // whole state in a single 32- or 64-bit register.
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;
3077 const T tmp = TDLS_LUPP_X(s);
3078 int cur = s;
3079 int nxt = TDLS_LUPP_PIV(cur);
3080 while (nxt != s) {
3081 TDLS_LUPP_X(cur) = TDLS_LUPP_X(nxt);
3082 visited |= mask_t(1) << cur;
3083 cur = nxt;
3084 nxt = TDLS_LUPP_PIV(cur);
3085 }
3086 TDLS_LUPP_X(cur) = tmp;
3087 visited |= mask_t(1) << cur;
3088 }
3089 } else {
3090 for (int s = 0; s < N; ++s) {
3091 if ((visited >> s) & mask_t(1)) continue;
3092 const T tmp = TDLS_LUPP_X(s);
3093 int cur = s;
3094 int nxt = TDLS_LUPP_PIV(cur);
3095 while (nxt != s) {
3096 TDLS_LUPP_X(cur) = TDLS_LUPP_X(nxt);
3097 visited |= mask_t(1) << cur;
3098 cur = nxt;
3099 nxt = TDLS_LUPP_PIV(cur);
3100 }
3101 TDLS_LUPP_X(cur) = tmp;
3102 visited |= mask_t(1) << cur;
3103 }
3104 }
3105 } else {
3106 // Cycle-leader scan (no visited storage): a cycle is rotated
3107 // only when reached from its smallest index, detected by
3108 // walking the orbit. Values and order match the mask path.
3109 for (int s = 0; s < N; ++s) {
3110 int probe = TDLS_LUPP_PIV(s);
3111 while (probe > s)
3112 probe = TDLS_LUPP_PIV(probe);
3113 if (probe != s) continue;
3114
3115 const T tmp = TDLS_LUPP_X(s);
3116 int cur = s;
3117 int nxt = TDLS_LUPP_PIV(cur);
3118 while (nxt != s) {
3119 TDLS_LUPP_X(cur) = TDLS_LUPP_X(nxt);
3120 cur = nxt;
3121 nxt = TDLS_LUPP_PIV(cur);
3122 }
3123 TDLS_LUPP_X(cur) = tmp;
3124 }
3125 }
3127 A, A_stride, piv, piv_stride, x, rhs_stride, 0);
3128 }
3129 }
3130 }
3131
3148 template<int pass_width, bool internal_rhs, bool internal_piv>
3149 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL static constexpr void
3150 permute_cycles(const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
3151 const int rhs_stride, const std::size_t xcol_stride) noexcept {
3152 if constexpr (N <= 64) {
3153 // Bitmask cycle decomposition, as in substitute_inplace; tmp
3154 // widens to one entry per column.
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;
3161 T tmp[pass_width];
3163 for (int w = 0; w < pass_width; ++w)
3164 tmp[w] = TDLS_LUPP_XW(w, s);
3165 int cur = s;
3166 int nxt = TDLS_LUPP_PIV(cur);
3167 while (nxt != s) {
3169 for (int w = 0; w < pass_width; ++w)
3170 TDLS_LUPP_XW(w, cur) = TDLS_LUPP_XW(w, nxt);
3171 visited |= mask_t(1) << cur;
3172 cur = nxt;
3173 nxt = TDLS_LUPP_PIV(cur);
3174 }
3176 for (int w = 0; w < pass_width; ++w)
3177 TDLS_LUPP_XW(w, cur) = tmp[w];
3178 visited |= mask_t(1) << cur;
3179 }
3180 } else {
3181 for (int s = 0; s < N; ++s) {
3182 if ((visited >> s) & mask_t(1)) continue;
3183 T tmp[pass_width];
3184 for (int w = 0; w < pass_width; ++w)
3185 tmp[w] = TDLS_LUPP_XW(w, s);
3186 int cur = s;
3187 int nxt = TDLS_LUPP_PIV(cur);
3188 while (nxt != s) {
3189 for (int w = 0; w < pass_width; ++w)
3190 TDLS_LUPP_XW(w, cur) = TDLS_LUPP_XW(w, nxt);
3191 visited |= mask_t(1) << cur;
3192 cur = nxt;
3193 nxt = TDLS_LUPP_PIV(cur);
3194 }
3195 for (int w = 0; w < pass_width; ++w)
3196 TDLS_LUPP_XW(w, cur) = tmp[w];
3197 visited |= mask_t(1) << cur;
3198 }
3199 }
3200 } else {
3201 // Cycle-leader scan, as in substitute_inplace.
3202 for (int s = 0; s < N; ++s) {
3203 int probe = TDLS_LUPP_PIV(s);
3204 while (probe > s)
3205 probe = TDLS_LUPP_PIV(probe);
3206 if (probe != s) continue;
3207
3208 T tmp[pass_width];
3209 for (int w = 0; w < pass_width; ++w)
3210 tmp[w] = TDLS_LUPP_XW(w, s);
3211 int cur = s;
3212 int nxt = TDLS_LUPP_PIV(cur);
3213 while (nxt != s) {
3214 for (int w = 0; w < pass_width; ++w)
3215 TDLS_LUPP_XW(w, cur) = TDLS_LUPP_XW(w, nxt);
3216 cur = nxt;
3217 nxt = TDLS_LUPP_PIV(cur);
3218 }
3219 for (int w = 0; w < pass_width; ++w)
3220 TDLS_LUPP_XW(w, cur) = tmp[w];
3221 }
3222 }
3223 }
3224
3245 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
3246 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
3247 substitute_inplace_multirhs_pass(const T* TDLS_RESTRICT A, const int A_stride,
3248 const int* TDLS_RESTRICT piv, const int piv_stride,
3249 T* TDLS_RESTRICT x, const int rhs_stride,
3250 const std::size_t xcol_stride) noexcept {
3251 if constexpr (transposed) {
3252 // Column-major layout, as in substitute_inplace.
3254 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3255 unpermute_cycles<pass_width, internal_rhs, internal_piv>(piv, piv_stride, x, rhs_stride,
3256 xcol_stride);
3257 } else {
3258 permute_cycles<pass_width, internal_rhs, internal_piv>(piv, piv_stride, x, rhs_stride,
3259 xcol_stride);
3261 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3262 }
3263 }
3264
3281 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
3282 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
3283 forward_inplace_multirhs_pass(const T* TDLS_RESTRICT A, const int A_stride,
3284 const int* TDLS_RESTRICT piv, const int piv_stride,
3285 T* TDLS_RESTRICT x, const int rhs_stride,
3286 const std::size_t xcol_stride) noexcept {
3287 if constexpr (transposed) {
3288 for (int k = 0; k < full_tiles; ++k)
3289 transposed_fwd_step<tile_size, pass_width, internal_rhs, internal_piv,
3290 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride,
3291 xcol_stride, k);
3292 if constexpr (last_tile_tail > 0)
3293 transposed_fwd_step<last_tile_tail, pass_width, internal_rhs, internal_piv,
3294 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride,
3295 xcol_stride, full_tiles);
3296 } else {
3297 permute_cycles<pass_width, internal_rhs, internal_piv>(piv, piv_stride, x, rhs_stride,
3298 xcol_stride);
3299 for (int k = 0; k < full_tiles; ++k)
3301 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
3302 if constexpr (last_tile_tail > 0)
3304 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, full_tiles);
3305 }
3306 }
3307
3324 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
3325 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
3326 backward_inplace_multirhs_pass(const T* TDLS_RESTRICT A, const int A_stride,
3327 const int* TDLS_RESTRICT piv, const int piv_stride,
3328 T* TDLS_RESTRICT x, const int rhs_stride,
3329 const std::size_t xcol_stride) noexcept {
3330 if constexpr (transposed) {
3332 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3333 unpermute_cycles<pass_width, internal_rhs, internal_piv>(piv, piv_stride, x, rhs_stride,
3334 xcol_stride);
3335 } else {
3337 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3338 }
3339 }
3340
3374 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
3375 int pass_width = 0, int A_static_stride = dynamic_stride,
3376 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
3377 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
3378 substitute_inplace_multirhs(const T* TDLS_RESTRICT A, const int A_stride,
3379 const int* TDLS_RESTRICT piv, const int piv_stride,
3380 T* TDLS_RESTRICT x, const int rhs_stride,
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>) {
3384 return substitute_inplace_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix,
3385 pass_width>(
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);
3389 } else {
3390 if constexpr (physical) {
3392 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3393 return;
3394 }
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 =
3401 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
3402 substitute_inplace_multirhs_pass<columns_per_pass, internal_rhs, internal_piv,
3403 internal_matrix>(A, A_stride, piv, piv_stride,
3404 x + off, rhs_stride, xcol_stride);
3405 }
3406 if constexpr (nrhs % columns_per_pass > 0) {
3407 constexpr int c0 = nrhs - nrhs % columns_per_pass;
3408 const std::size_t off =
3409 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
3410 substitute_inplace_multirhs_pass<nrhs % columns_per_pass, internal_rhs,
3411 internal_piv, internal_matrix>(
3412 A, A_stride, piv, piv_stride, x + off, rhs_stride, xcol_stride);
3413 }
3414 }
3415 }
3416
3443 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
3444 int pass_width = 0, int A_static_stride = dynamic_stride,
3445 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
3446 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
3447 forward_inplace_multirhs(const T* TDLS_RESTRICT A, const int A_stride,
3448 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
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>) {
3452 return forward_inplace_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix,
3453 pass_width>(
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);
3457 } else {
3458 if constexpr (physical) {
3460 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3461 return;
3462 }
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 =
3469 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
3470 forward_inplace_multirhs_pass<columns_per_pass, internal_rhs, internal_piv,
3471 internal_matrix>(A, A_stride, piv, piv_stride,
3472 x + off, rhs_stride, xcol_stride);
3473 }
3474 if constexpr (nrhs % columns_per_pass > 0) {
3475 constexpr int c0 = nrhs - nrhs % columns_per_pass;
3476 const std::size_t off =
3477 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
3478 forward_inplace_multirhs_pass<nrhs % columns_per_pass, internal_rhs, internal_piv,
3479 internal_matrix>(A, A_stride, piv, piv_stride,
3480 x + off, rhs_stride, xcol_stride);
3481 }
3482 }
3483 }
3484
3508 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
3509 int pass_width = 0, int A_static_stride = dynamic_stride,
3510 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
3511 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
3512 backward_inplace_multirhs(const T* TDLS_RESTRICT A, const int A_stride,
3513 const int* TDLS_RESTRICT piv, const int piv_stride,
3514 T* TDLS_RESTRICT x, const int rhs_stride,
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>) {
3518 return backward_inplace_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix,
3519 pass_width>(
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);
3523 } else {
3524 if constexpr (physical) {
3526 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
3527 return;
3528 }
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 =
3535 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
3536 backward_inplace_multirhs_pass<columns_per_pass, internal_rhs, internal_piv,
3537 internal_matrix>(A, A_stride, piv, piv_stride,
3538 x + off, rhs_stride, xcol_stride);
3539 }
3540 if constexpr (nrhs % columns_per_pass > 0) {
3541 constexpr int c0 = nrhs - nrhs % columns_per_pass;
3542 const std::size_t off =
3543 internal_rhs ? unsigned(c0) * unsigned(N) : TDLS_STRIDED_WIDE(c0, xcol_stride);
3544 backward_inplace_multirhs_pass<nrhs % columns_per_pass, internal_rhs, internal_piv,
3545 internal_matrix>(A, A_stride, piv, piv_stride,
3546 x + off, rhs_stride, xcol_stride);
3547 }
3548 }
3549 }
3550
3551 /* =====================================================================
3552 SOLVE - convenience wrappers.
3553 ===================================================================== */
3554
3578 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
3579 bool oot_diagnostics = true, int A_static_stride = dynamic_stride,
3580 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
3581 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
3582 solve(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
3583 const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_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);
3591 } else {
3593 piv_stride, oot_count))
3594 return false;
3595 substitute<internal_rhs, internal_piv, internal_matrix>(A, A_stride, piv, piv_stride, b,
3596 x, rhs_stride);
3597 return true;
3598 }
3599 }
3600
3621 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
3622 int A_static_stride = dynamic_stride, int piv_static_stride = dynamic_stride,
3623 int rhs_static_stride = dynamic_stride>
3624 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
3625 solve(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
3626 const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride) noexcept {
3627 int unused = 0;
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);
3631 }
3632
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,
3666 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
3667 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
3668 solve_multirhs(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
3669 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
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,
3674 oot_diagnostics>(
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);
3678 } else {
3680 piv_stride, oot_count))
3681 return false;
3683 A, A_stride, piv, piv_stride, b, x, rhs_stride, xcol_stride);
3684 return true;
3685 }
3686 }
3687
3716 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
3717 int pass_width = 0, int A_static_stride = dynamic_stride,
3718 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
3719 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
3720 solve_multirhs(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
3721 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
3722 const int rhs_stride, const std::size_t xcol_stride) noexcept {
3723 int unused = 0;
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);
3727 }
3728
3765 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
3766 bool oot_diagnostics = true, int A_static_stride = dynamic_stride,
3767 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
3768 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
3769 solve_inplace(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
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);
3778 } else {
3779 if constexpr (physical) {
3780 int weak = 0;
3781 const bool ok = physical_factorize<internal_piv, internal_matrix>(A, A_stride, piv,
3782 piv_stride, weak);
3783 if constexpr (oot_diagnostics) oot_count = weak;
3784 if (!ok) return false;
3786 A, A_stride, piv, piv_stride, y, rhs_stride, 0);
3787 return true;
3788 }
3790 A, A_stride, piv, piv_stride, oot_count, y, rhs_stride))
3791 return false;
3792
3793 // Backward pass only: the forward one happened inside factorize.
3794 if constexpr (transposed) {
3795 // Column-major layout: L^T w = z, then x = P^T w.
3797 A, A_stride, piv, piv_stride, y, rhs_stride, 0);
3798 unpermute<1, internal_rhs, internal_piv>(piv, piv_stride, y, rhs_stride, 0);
3799 } else {
3801 A, A_stride, piv, piv_stride, y, rhs_stride, 0);
3802 }
3803 return true;
3804 }
3805 }
3806
3826 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
3827 int A_static_stride = dynamic_stride, int piv_static_stride = dynamic_stride,
3828 int rhs_static_stride = dynamic_stride>
3829 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
3830 solve_inplace(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
3831 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride) noexcept {
3832 int unused = 0;
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);
3836 }
3837
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,
3875 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
3876 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
3877 solve_inplace_multirhs(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
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>) {
3882 return solve_inplace_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix,
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);
3887 } else {
3889 piv_stride, oot_count))
3890 return false;
3891 substitute_inplace_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix,
3892 pass_width>(A, A_stride, piv, piv_stride, y, rhs_stride,
3893 xcol_stride);
3894 return true;
3895 }
3896 }
3897
3925 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
3926 int pass_width = 0, int A_static_stride = dynamic_stride,
3927 int piv_static_stride = dynamic_stride, int rhs_static_stride = dynamic_stride>
3928 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
3929 solve_inplace_multirhs(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
3930 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride,
3931 const std::size_t xcol_stride) noexcept {
3932 int unused = 0;
3933 return solve_inplace_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix, pass_width,
3934 false, A_static_stride, piv_static_stride, rhs_static_stride>(
3935 A, A_stride, piv, piv_stride, y, rhs_stride, xcol_stride, unused);
3936 }
3937
3938 /* =====================================================================
3939 Note on pivot ownership: the pivot storage is ALWAYS caller-provided
3940 (internal array or remote scratch), for every entry point: one
3941 uniform calling convention. A solver-internal pivot would only be
3942 expressible for the single combination {solve/solve_inplace x internal
3943 pivot}; externalizing it there costs nothing (the caller's int
3944 piv[N] inlines to the exact same codegen), so no special case is
3945 kept. For one-shot solves the contents are simply treated as
3946 scratch; for factorize + substitute* (consistent tangent), the same
3947 storage carries the permutation across the calls.
3948 ===================================================================== */
3949};
3950
3951
3952
3953#undef TDLS_LUPP_A
3954#undef TDLS_LUPP_PIV
3955#undef TDLS_LUPP_X
3956#undef TDLS_LUPP_B
3957#undef TDLS_LUPP_XW
3958#undef TDLS_LUPP_BW
3959#undef TDLS_LUPP_Y
3960#undef TDLS_LUPP_FOR
3961
3962
3963
3964} // namespace tdls
3965
3966#if defined(TDLS_DETAIL_COMPILER_GCC)
3967#pragma GCC diagnostic pop
3968#endif
3969
3970#if defined(TDLS_DETAIL_COMPILER_CLANG)
3971#pragma clang diagnostic pop
3972#endif
3973
3974
3975
3976#endif // TDLS_SOLVERS_TILED_LUPP_SOLVER_STATIC_HPP
#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.