TDLS 0.2.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
62
63
64
65#include <tdls/core/math.hpp>
66#include <type_traits>
67
70
71
72
73// The flow analysis of older gcc releases (observed on the 11.4 of
74// Ubuntu 22.04) cannot prove that the replay row of the out-of-tile
75// search is written before being read: L_row[p] is only read for
76// p < t, and slot t is written at the end of iteration t, a
77// loop-carried property validated by the oracle suites. The resulting
78// -Wmaybe-uninitialized reports are spurious.
79//
80// Under UBSan instrumentation (-fsanitize=undefined) at -O2 and
81// above, gcc additionally loses the value ranges of the loop indices
82// and reports impossible subscripts through -Warray-bounds: the
83// unsigned cast of a provably non-negative index read as a wrapped
84// negative, or a loop counter assumed below its own loop guard.
85// Observed from gcc 12 to gcc 16; plain builds and clang are clean,
86// and the exercised paths are certified free of out-of-bounds
87// accesses by the constexpr suite.
88//
89// Both suppressions are scoped to this header and to those warnings
90// only, exactly as in solver_dynamic.hpp.
91#if defined(__GNUC__) && !defined(__clang__)
92#pragma GCC diagnostic push
93#pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
94#pragma GCC diagnostic ignored "-Warray-bounds"
95#endif
96
97
98
99// clang reports a forced unrolling that the optimizer could not perform
100// through -Wpass-failed. The unrolling requested by TDLS_UNROLL_FORCE is
101// a performance hint: a failed hint does not affect correctness. The
102// suppression is scoped to this header and to that warning only.
103#if defined(__clang__)
104#pragma clang diagnostic push
105#pragma clang diagnostic ignored "-Wpass-failed"
106#endif
107
108namespace tdls {
109
110
111
112/* Addressing macros. They expand inside member functions where the
113 parameter names (A, A_stride, piv, piv_stride, x, b, y, rhs_stride,
114 xcol_stride), the residency template booleans and the Config value
115 are in scope. #undef'd at the end of this header. */
116
120#define TDLS_LUPP_A(r, c) \
121 A[internal_matrix ? TDLS_LAYOUT_INDEX(r, c, N) \
122 : TDLS_LAYOUT_INDEX(r, c, N) * unsigned(A_stride)]
126#define TDLS_LUPP_PIV(i) piv[internal_piv ? unsigned(i) : unsigned(i) * unsigned(piv_stride)]
130#define TDLS_LUPP_X(i) x[internal_rhs ? unsigned(i) : unsigned(i) * unsigned(rhs_stride)]
134#define TDLS_LUPP_B(i) b[internal_rhs ? unsigned(i) : unsigned(i) * unsigned(rhs_stride)]
140#define TDLS_LUPP_XW(w, i) \
141 x[internal_rhs ? unsigned((w) * N + (i)) \
142 : unsigned(i) * unsigned(rhs_stride) + unsigned(w) * unsigned(xcol_stride)]
146#define TDLS_LUPP_BW(w, i) \
147 b[internal_rhs ? unsigned((w) * N + (i)) \
148 : unsigned(i) * unsigned(rhs_stride) + unsigned(w) * unsigned(xcol_stride)]
152#define TDLS_LUPP_Y(i) y[internal_rhs ? unsigned(i) : unsigned(i) * unsigned(rhs_stride)]
153
154
155
199template<typename T, int N, TiledLUppConfig<T> Config = TiledLUppConfig<T>{}>
201
202 static constexpr int tile_size = Config.tile_size;
203 static constexpr Schedule schedule =
204 Config.schedule;
205
208 static constexpr T oot_threshold = Config.oot_threshold;
211 static constexpr T singular_floor = Config.singular_floor;
212
213 static_assert(N >= 1, "TiledLUppSolverStatic: N must be >= 1");
214 static_assert(
215 Config.oot_threshold.is_finite() && Config.singular_floor.is_finite(),
216 "TiledLUppSolverStatic: oot_threshold and singular_floor must be finite (and fit a "
217 "63-bit mantissa)");
218 static_assert(singular_floor > T(0), "TiledLUppSolverStatic: singular_floor must be positive");
219 static_assert(singular_floor <= oot_threshold,
220 "TiledLUppSolverStatic: singular_floor must not exceed oot_threshold (the "
221 "floor applies to the out-of-tile recovery path)");
222 static_assert(tile_size >= 1, "TiledLUppSolverStatic: tile size must be >= 1");
223
224 static constexpr int full_tiles = N / tile_size;
226 static constexpr int last_tile_tail = N - full_tiles * tile_size;
228 static constexpr int num_tiles = full_tiles + (last_tile_tail > 0 ? 1 : 0);
229
231 using Operations = TiledLUppTileOperations<T, tile_size, Config.unroll_inner>;
232
233 /* =====================================================================
234 Remote <-> register tile movement.
235 `load/store_tile` go through a cached physical-row segment (the RL
236 schedule keeps one per tile); `load/store_tile_piv` read the
237 permutation inline, once per row (LL schedule and the substitution
238 do this; it adds no integer registers).
239 ===================================================================== */
240
250 template<int row_extent, int col_extent, bool internal_matrix>
251 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
252 load_tile(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT prow,
253 const int col0, T* TDLS_RESTRICT t) noexcept {
254 if constexpr (Config.unroll_inner) {
256 for (int i = 0; i < row_extent; ++i) {
258 for (int j = 0; j < col_extent; ++j)
259 t[i * tile_size + j] = TDLS_LUPP_A(prow[i], col0 + j);
260 }
261 } else {
262 for (int i = 0; i < row_extent; ++i) {
263 for (int j = 0; j < col_extent; ++j)
264 t[i * tile_size + j] = TDLS_LUPP_A(prow[i], col0 + j);
265 }
266 }
267 }
268
278 template<int row_extent, int col_extent, bool internal_matrix>
279 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
280 store_tile(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT prow,
281 const int col0, const T* TDLS_RESTRICT t) noexcept {
282 if constexpr (Config.unroll_inner) {
284 for (int i = 0; i < row_extent; ++i) {
286 for (int j = 0; j < col_extent; ++j)
287 TDLS_LUPP_A(prow[i], col0 + j) = t[i * tile_size + j];
288 }
289 } else {
290 for (int i = 0; i < row_extent; ++i) {
291 for (int j = 0; j < col_extent; ++j)
292 TDLS_LUPP_A(prow[i], col0 + j) = t[i * tile_size + j];
293 }
294 }
295 }
296
310 template<int row_extent, int col_extent, bool internal_piv, bool internal_matrix>
311 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
312 load_tile_piv(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
313 const int piv_stride, const int row0, const int col0,
314 T* TDLS_RESTRICT t) noexcept {
315 if constexpr (Config.unroll_inner) {
317 for (int i = 0; i < row_extent; ++i) {
318 const int phys = TDLS_LUPP_PIV(row0 + i);
320 for (int j = 0; j < col_extent; ++j)
321 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
322 }
323 } else {
324 for (int i = 0; i < row_extent; ++i) {
325 const int phys = TDLS_LUPP_PIV(row0 + i);
326 for (int j = 0; j < col_extent; ++j)
327 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
328 }
329 }
330 }
331
345 template<int row_extent, int col_extent, bool internal_piv, bool internal_matrix>
346 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
347 store_tile_piv(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
348 const int piv_stride, const int row0, const int col0,
349 const T* TDLS_RESTRICT t) noexcept {
350 if constexpr (Config.unroll_inner) {
352 for (int i = 0; i < row_extent; ++i) {
353 const int phys = TDLS_LUPP_PIV(row0 + i);
355 for (int j = 0; j < col_extent; ++j)
356 TDLS_LUPP_A(phys, col0 + j) = t[i * tile_size + j];
357 }
358 } else {
359 for (int i = 0; i < row_extent; ++i) {
360 const int phys = TDLS_LUPP_PIV(row0 + i);
361 for (int j = 0; j < col_extent; ++j)
362 TDLS_LUPP_A(phys, col0 + j) = t[i * tile_size + j];
363 }
364 }
365 }
366
382 template<int row_extent, bool internal_piv, bool internal_matrix>
383 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
384 load_tile_piv_lower(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
385 const int piv_stride, const int row0, const int col0,
386 T* TDLS_RESTRICT t) noexcept {
387 if constexpr (Config.unroll_inner) {
389 for (int i = 1; i < row_extent; ++i) {
390 const int phys = TDLS_LUPP_PIV(row0 + i);
392 for (int j = 0; j < i; ++j)
393 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
394 }
395 } else {
396 for (int i = 1; i < row_extent; ++i) {
397 const int phys = TDLS_LUPP_PIV(row0 + i);
398 for (int j = 0; j < i; ++j)
399 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
400 }
401 }
402 }
403
416 template<int row_extent, bool internal_piv, bool internal_matrix>
417 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
418 load_tile_piv_upper(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
419 const int piv_stride, const int row0, const int col0,
420 T* TDLS_RESTRICT t) noexcept {
421 if constexpr (Config.unroll_inner) {
423 for (int i = 0; i < row_extent; ++i) {
424 const int phys = TDLS_LUPP_PIV(row0 + i);
426 for (int j = i; j < row_extent; ++j)
427 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
428 }
429 } else {
430 for (int i = 0; i < row_extent; ++i) {
431 const int phys = TDLS_LUPP_PIV(row0 + i);
432 for (int j = i; j < row_extent; ++j)
433 t[i * tile_size + j] = TDLS_LUPP_A(phys, col0 + j);
434 }
435 }
436 }
437
438 /* =====================================================================
439 Diagonal-tile factorization with out-of-tile pivoting.
440 Shared by both schedules; the only schedule-dependent part is the
441 replay a candidate row needs before it can be compared:
442 RL: the trailing matrix is already Schur-updated, so a candidate
443 only misses the current tile's columns t < c.
444 LL: rows below additionally miss every prior tile's L*U
445 contribution (columns [0, k0)), replayed term by term.
446 ===================================================================== */
447
471 template<int k_extent, bool internal_piv, bool internal_matrix, bool oot_diagnostics,
472 bool fuse_rhs, bool internal_rhs>
473 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
474 factor_diag_column(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
475 const int piv_stride, const int k0, T* TDLS_RESTRICT tile, int& oot_count,
476 const int c, T* TDLS_RESTRICT y, const int rhs_stride) noexcept {
477
478 const int gc = k0 + c; // global column
479
480 // In-tile pivot search (rows c..k_extent of the register tile)
481 int best_r = c;
482 T best = detail::abs(tile[c * tile_size + c]);
483 if constexpr (Config.unroll_inner) {
485 for (int r = c + 1; r < k_extent; ++r) {
486 const T v = detail::abs(tile[r * tile_size + c]);
487 if (v > best) {
488 best = v;
489 best_r = r;
490 }
491 }
492 } else {
493 for (int r = c + 1; r < k_extent; ++r) {
494 const T v = detail::abs(tile[r * tile_size + c]);
495 if (v > best) {
496 best = v;
497 best_r = r;
498 }
499 }
500 }
501
502 int piv_row; // winning global (logical) row
503
504 if (best >= oot_threshold) {
505 piv_row = k0 + best_r;
506 } else if constexpr (k_extent < tile_size) {
507 // Trailing tile: no rows below to recover from. Diagnostic
508 // order: singularity verdict first, then count the weak pivot
509 // (full tiles count before the verdict).
510 if (best < singular_floor) return false;
511 if constexpr (oot_diagnostics) ++oot_count;
512 piv_row = k0 + best_r;
513 } else {
514 // Out-of-tile recovery: scan the rows below the tile and
515 // evaluate each candidate as if it had received the
516 // eliminations it is missing, keeping the best (or, with
517 // Config.oot_first_acceptable, the first to reach the threshold).
518 // Cold path, never unroll-annotated.
519 if constexpr (oot_diagnostics) ++oot_count;
520 T gbest = best;
521 int gbest_row = k0 + best_r;
522
523 for (int row = k0 + k_extent; row < N; ++row) {
524 const int phys = TDLS_LUPP_PIV(row);
525
526 T corrected = TDLS_LUPP_A(phys, gc);
527 if constexpr (schedule == Schedule::LeftLooking) {
528 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
529 for (int p = 0; p < tile_size; ++p)
530 corrected -= TDLS_LUPP_A(phys, bj0 + p) *
531 TDLS_LUPP_A(TDLS_LUPP_PIV(bj0 + p), gc);
532 }
533
534 if (c > 0) {
535 T L_row[tile_size];
536 for (int t = 0; t < c; ++t) {
537 T a_t = TDLS_LUPP_A(phys, k0 + t);
538 if constexpr (schedule == Schedule::LeftLooking) {
539 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
540 for (int p = 0; p < tile_size; ++p)
541 a_t -= TDLS_LUPP_A(phys, bj0 + p) *
542 TDLS_LUPP_A(TDLS_LUPP_PIV(bj0 + p), k0 + t);
543 }
544 for (int p = 0; p < t; ++p)
545 a_t -= L_row[p] * tile[p * tile_size + t];
546 L_row[t] = a_t * tile[t * tile_size + t]; // diag holds 1/pivot
547 corrected -= L_row[t] * tile[t * tile_size + c];
548 }
549 }
550
551 const T v = detail::abs(corrected);
552 if (v > gbest) {
553 gbest = v;
554 gbest_row = row;
555 }
556
557 // First-acceptable out-of-tile pivot: a candidate that
558 // reaches the threshold already beats the sub-threshold
559 // in-tile pivot, so stop scanning. The running max above is
560 // kept as the fallback when no candidate is acceptable.
561 if constexpr (Config.oot_first_acceptable)
562 if (v >= oot_threshold) break;
563 }
564
565 if (gbest < singular_floor) return false;
566 piv_row = gbest_row;
567 }
568
569 {
570 // Swap the permutation entries unconditionally: when
571 // piv_row == gc the exchanges write back the same values
572 // (bitwise no-ops), and removing the comparison removes a
573 // divergent branch from the hot path (lanes of a warp pick
574 // different pivots almost every column).
575 const int tmp = TDLS_LUPP_PIV(gc);
576 TDLS_LUPP_PIV(gc) = TDLS_LUPP_PIV(piv_row);
577 TDLS_LUPP_PIV(piv_row) = tmp;
578
579 if constexpr (fuse_rhs) {
580 // The fused RHS follows the rows through pivoting. For an
581 // internal y the swap is predicated over the unrolled index
582 // range: y must never be dynamically indexed or it is
583 // demoted to local memory.
584 if constexpr (internal_rhs) {
585 if constexpr (Config.unroll_inner) {
587 for (int r = 0; r < N; ++r) {
588 if (r == piv_row) {
589 const T ty = TDLS_LUPP_Y(gc);
590 TDLS_LUPP_Y(gc) = TDLS_LUPP_Y(r);
591 TDLS_LUPP_Y(r) = ty;
592 }
593 }
594 } else {
595 for (int r = 0; r < N; ++r) {
596 if (r == piv_row) {
597 const T ty = TDLS_LUPP_Y(gc);
598 TDLS_LUPP_Y(gc) = TDLS_LUPP_Y(r);
599 TDLS_LUPP_Y(r) = ty;
600 }
601 }
602 }
603 } else {
604 const T ty = TDLS_LUPP_Y(gc);
605 TDLS_LUPP_Y(gc) = TDLS_LUPP_Y(piv_row);
606 TDLS_LUPP_Y(piv_row) = ty;
607 }
608 }
609
610 if (piv_row < k0 + k_extent) {
611 Operations::template swap_rows<k_extent>(tile, c, piv_row - k0);
612 } else {
613 // Cross-tile swap: pull the new row into the tile and
614 // replay everything it missed: prior tiles (LL), then
615 // the current tile's factored columns, on the FULL row.
616 const int phys = TDLS_LUPP_PIV(gc);
617 if constexpr (Config.unroll_inner) {
619 for (int j = 0; j < k_extent; ++j) {
620 tile[c * tile_size + j] = TDLS_LUPP_A(phys, k0 + j);
621 if constexpr (schedule == Schedule::LeftLooking) {
622 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
623 for (int p = 0; p < tile_size; ++p)
624 tile[c * tile_size + j] -=
625 TDLS_LUPP_A(phys, bj0 + p) *
626 TDLS_LUPP_A(TDLS_LUPP_PIV(bj0 + p), k0 + j);
627 }
628 }
629 } else {
630 for (int j = 0; j < k_extent; ++j) {
631 tile[c * tile_size + j] = TDLS_LUPP_A(phys, k0 + j);
632 if constexpr (schedule == Schedule::LeftLooking) {
633 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
634 for (int p = 0; p < tile_size; ++p)
635 tile[c * tile_size + j] -=
636 TDLS_LUPP_A(phys, bj0 + p) *
637 TDLS_LUPP_A(TDLS_LUPP_PIV(bj0 + p), k0 + j);
638 }
639 }
640 }
641
642 if (c > 0) {
643 T L_row[tile_size];
644 for (int t = 0; t < c; ++t) {
645 T a_t = tile[c * tile_size + t];
646 for (int p = 0; p < t; ++p)
647 a_t -= L_row[p] * tile[p * tile_size + t];
648 L_row[t] = a_t * tile[t * tile_size + t]; // diag holds 1/pivot
649 tile[c * tile_size + t] = L_row[t];
650 }
651 for (int j = c; j < k_extent; ++j) {
652 for (int t = 0; t < c; ++t)
653 tile[c * tile_size + j] -= L_row[t] * tile[t * tile_size + j];
654 }
655 }
656 }
657 }
658
659 Operations::template eliminate_column<k_extent, k_extent>(tile, c);
660 return true;
661 }
662
683 template<int k_extent, bool internal_piv, bool internal_matrix, bool oot_diagnostics,
684 bool fuse_rhs = false, bool internal_rhs = true>
685 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
686 factor_diag_tile(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
687 const int piv_stride, const int k0, T* TDLS_RESTRICT tile, int& oot_count,
688 T* TDLS_RESTRICT y = nullptr, const int rhs_stride = 1) noexcept {
689 if constexpr (Config.unroll_inner) {
691 for (int c = 0; c < k_extent; ++c) {
692 if (!factor_diag_column<k_extent, internal_piv, internal_matrix, oot_diagnostics,
693 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k0,
694 tile, oot_count, c, y, rhs_stride))
695 return false;
696 }
697 } else {
698 for (int c = 0; c < k_extent; ++c) {
699 if (!factor_diag_column<k_extent, internal_piv, internal_matrix, oot_diagnostics,
700 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, k0,
701 tile, oot_count, c, y, rhs_stride))
702 return false;
703 }
704 }
705 return true;
706 }
707
708 /* =====================================================================
709 RIGHT-LOOKING schedule. Step k: factor the diagonal tile, then push
710 its factors into the trailing matrix (TRSM right over the row panel,
711 TRSM down + Schur complement over the rows below). The physical-row
712 segments pk/pi are cached per tile (they fold to nothing when the
713 permutation itself is internal).
714 ===================================================================== */
715
725 template<int k_extent, int j_extent, bool internal_matrix>
726 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
727 rl_trsm_right_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT pk,
728 const T* TDLS_RESTRICT tile, const int j0) noexcept {
729 T Akj[tile_size * tile_size];
730 load_tile<k_extent, j_extent, internal_matrix>(A, A_stride, pk, j0, Akj);
731 Operations::template trsm_left_unit<k_extent, j_extent>(tile, Akj);
732 store_tile<k_extent, j_extent, internal_matrix>(A, A_stride, pk, j0, Akj);
733 }
734
748 template<int k_extent, int i_extent, int j_extent, bool internal_matrix>
749 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
750 rl_schur_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT pk,
751 const int* TDLS_RESTRICT pi, const T* TDLS_RESTRICT Aik, const int j0) noexcept {
752 T Aij[tile_size * tile_size];
753 load_tile<i_extent, j_extent, internal_matrix>(A, A_stride, pi, j0, Aij);
754
755 if constexpr (Config.unroll_inner) {
757 for (int p = 0; p < k_extent; ++p) {
758 T Akj_row[tile_size];
760 for (int j = 0; j < j_extent; ++j)
761 Akj_row[j] = TDLS_LUPP_A(pk[p], j0 + j);
763 for (int i = 0; i < i_extent; ++i) {
764 const T L_ip = Aik[i * tile_size + p];
766 for (int j = 0; j < j_extent; ++j)
767 Aij[i * tile_size + j] -= L_ip * Akj_row[j];
768 }
769 }
770 } else {
771 for (int p = 0; p < k_extent; ++p) {
772 T Akj_row[tile_size];
773 for (int j = 0; j < j_extent; ++j)
774 Akj_row[j] = TDLS_LUPP_A(pk[p], j0 + j);
775 for (int i = 0; i < i_extent; ++i) {
776 const T L_ip = Aik[i * tile_size + p];
777 for (int j = 0; j < j_extent; ++j)
778 Aij[i * tile_size + j] -= L_ip * Akj_row[j];
779 }
780 }
781 }
782
783 store_tile<i_extent, j_extent, internal_matrix>(A, A_stride, pi, j0, Aij);
784 }
785
804 template<int k_extent, int i_extent, bool internal_piv, bool internal_matrix,
805 bool fuse_rhs = false, bool internal_rhs = true>
806 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
807 rl_update_row_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
808 const int piv_stride, const int* TDLS_RESTRICT pk,
809 const T* TDLS_RESTRICT tile, const int k, const int i0,
810 T* TDLS_RESTRICT y = nullptr, const int rhs_stride = 1) noexcept {
811 const int k0 = k * tile_size;
812
813 int pi[tile_size];
814 if constexpr (Config.unroll_inner) {
816 for (int i = 0; i < i_extent; ++i)
817 pi[i] = TDLS_LUPP_PIV(i0 + i);
818 } else {
819 for (int i = 0; i < i_extent; ++i)
820 pi[i] = TDLS_LUPP_PIV(i0 + i);
821 }
822
823 // TRSM down: Aik := Aik * U^-1
824 T Aik[tile_size * tile_size];
825 load_tile<i_extent, k_extent, internal_matrix>(A, A_stride, pi, k0, Aik);
826 Operations::template trsm_right<k_extent, i_extent>(tile, Aik);
827 store_tile<i_extent, k_extent, internal_matrix>(A, A_stride, pi, k0, Aik);
828
829 // Fused forward substitution: push the solved y_k segment into this
830 // row block while its L panel sits in registers (this is the whole
831 // point of solve_inplace: the separate forward pass reloads vanish).
832 if constexpr (fuse_rhs) {
833 if constexpr (Config.unroll_inner) {
835 for (int r = 0; r < i_extent; ++r) {
836 T sum = T(0);
838 for (int j = 0; j < k_extent; ++j)
839 sum += Aik[r * tile_size + j] * TDLS_LUPP_Y(k0 + j);
840 TDLS_LUPP_Y(i0 + r) -= sum;
841 }
842 } else {
843 for (int r = 0; r < i_extent; ++r) {
844 T sum = T(0);
845 for (int j = 0; j < k_extent; ++j)
846 sum += Aik[r * tile_size + j] * TDLS_LUPP_Y(k0 + j);
847 TDLS_LUPP_Y(i0 + r) -= sum;
848 }
849 }
850 }
851
852 // Schur sweep over the trailing columns
853 for (int j = k + 1; j < full_tiles; ++j)
855 j * tile_size);
856 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
858 A, A_stride, pk, pi, Aik, full_tiles * tile_size);
859 }
860
878 template<int k_extent, bool internal_piv, bool internal_matrix, bool oot_diagnostics,
879 bool fuse_rhs = false, bool internal_rhs = true>
880 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
881 rl_step(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
882 const int k, int& oot_count, T* TDLS_RESTRICT y = nullptr,
883 const int rhs_stride = 1) noexcept {
884 const int k0 = k * tile_size;
885
886 T tile[tile_size * tile_size];
888 piv_stride, k0, k0, tile);
889
890 if (!factor_diag_tile<k_extent, internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
891 internal_rhs>(A, A_stride, piv, piv_stride, k0, tile, oot_count, y,
892 rhs_stride))
893 return false;
894
895 // Physical rows of the tile after the swaps of this step
896 int pk[tile_size];
897 if constexpr (Config.unroll_inner) {
899 for (int i = 0; i < k_extent; ++i)
900 pk[i] = TDLS_LUPP_PIV(k0 + i);
901 } else {
902 for (int i = 0; i < k_extent; ++i)
903 pk[i] = TDLS_LUPP_PIV(k0 + i);
904 }
905
906 store_tile<k_extent, k_extent, internal_matrix>(A, A_stride, pk, k0, tile);
907
908 // Fused forward substitution: this tile's y segment is final from
909 // here on: unit-lower-solve it while the tile is in registers.
910 if constexpr (fuse_rhs) {
911 if constexpr (Config.unroll_inner) {
913 for (int kk = 0; kk < k_extent; ++kk) {
915 for (int i = kk + 1; i < k_extent; ++i)
916 TDLS_LUPP_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_Y(k0 + kk);
917 }
918 } else {
919 for (int kk = 0; kk < k_extent; ++kk) {
920 for (int i = kk + 1; i < k_extent; ++i)
921 TDLS_LUPP_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_Y(k0 + kk);
922 }
923 }
924 }
925
926 // TRSM right over the row panel
927 for (int j = k + 1; j < full_tiles; ++j)
929 j * tile_size);
930 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
933
934 // TRSM down + Schur over the rows below
935 for (int i = k + 1; i < full_tiles; ++i)
936 rl_update_row_one<k_extent, tile_size, internal_piv, internal_matrix, fuse_rhs,
937 internal_rhs>(A, A_stride, piv, piv_stride, pk, tile, k,
938 i * tile_size, y, rhs_stride);
939 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
940 rl_update_row_one<k_extent, last_tile_tail, internal_piv, internal_matrix, fuse_rhs,
941 internal_rhs>(A, A_stride, piv, piv_stride, pk, tile, k,
942 full_tiles * tile_size, y, rhs_stride);
943 return true;
944 }
945
946 /* =====================================================================
947 LEFT-LOOKING schedule. Step k: pull every prior tile's L*U
948 contribution into the current block column/row, then factor the
949 diagonal tile and triangular-solve its panels. Each tile of the
950 trailing matrix is written once per factorization (vs once per step
951 for RL) at the cost of replaying prior tiles on each visit. The
952 permutation is read inline (one read per row), never cached in
953 segments.
954 ===================================================================== */
955
970 template<int row_extent, int col_extent, bool internal_piv, bool internal_matrix>
971 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
972 ll_correct_tile(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
973 const int piv_stride, const int row0, const int col0, const int k0,
974 T* TDLS_RESTRICT t) noexcept {
975 for (int bj0 = 0; bj0 < k0; bj0 += tile_size) {
976 T Lt[tile_size * tile_size];
978 A, A_stride, piv, piv_stride, row0, bj0, Lt);
979 T Ut[tile_size * tile_size];
981 A, A_stride, piv, piv_stride, bj0, col0, Ut);
982 Operations::template gemm_sub<row_extent, col_extent, tile_size>(t, Lt, Ut);
983 }
984 }
985
1003 template<int k_extent, int i_extent, bool internal_piv, bool internal_matrix,
1004 bool fuse_rhs = false, bool internal_rhs = true>
1005 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1006 ll_update_below_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1007 const int piv_stride, const T* TDLS_RESTRICT tile, const int k0,
1008 const int i0, T* TDLS_RESTRICT y = nullptr,
1009 const int rhs_stride = 1) noexcept {
1010 T B[tile_size * tile_size];
1012 piv_stride, i0, k0, B);
1014 A, A_stride, piv, piv_stride, i0, k0, k0, B);
1015 Operations::template trsm_right<k_extent, i_extent>(tile, B);
1017 piv_stride, i0, k0, B);
1018
1019 // Fused forward substitution: B is the final L(i,k) panel, so push the
1020 // solved y_k segment into this row block while it is in registers.
1021 if constexpr (fuse_rhs) {
1022 if constexpr (Config.unroll_inner) {
1024 for (int r = 0; r < i_extent; ++r) {
1025 T sum = T(0);
1027 for (int j = 0; j < k_extent; ++j)
1028 sum += B[r * tile_size + j] * TDLS_LUPP_Y(k0 + j);
1029 TDLS_LUPP_Y(i0 + r) -= sum;
1030 }
1031 } else {
1032 for (int r = 0; r < i_extent; ++r) {
1033 T sum = T(0);
1034 for (int j = 0; j < k_extent; ++j)
1035 sum += B[r * tile_size + j] * TDLS_LUPP_Y(k0 + j);
1036 TDLS_LUPP_Y(i0 + r) -= sum;
1037 }
1038 }
1039 }
1040 }
1041
1054 template<int k_extent, int j_extent, bool internal_piv, bool internal_matrix>
1055 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1056 ll_update_right_one(T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1057 const int piv_stride, const T* TDLS_RESTRICT tile, const int k0,
1058 const int j0) noexcept {
1059 T B[tile_size * tile_size];
1061 piv_stride, k0, j0, B);
1063 A, A_stride, piv, piv_stride, k0, j0, k0, B);
1064 Operations::template trsm_left_unit<k_extent, j_extent>(tile, B);
1066 piv_stride, k0, j0, B);
1067 }
1068
1086 template<int k_extent, bool internal_piv, bool internal_matrix, bool oot_diagnostics,
1087 bool fuse_rhs = false, bool internal_rhs = true>
1088 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1089 ll_step(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
1090 const int k, int& oot_count, T* TDLS_RESTRICT y = nullptr,
1091 const int rhs_stride = 1) noexcept {
1092 const int k0 = k * tile_size;
1093
1094 T tile[tile_size * tile_size];
1096 piv_stride, k0, k0, tile);
1098 A, A_stride, piv, piv_stride, k0, k0, k0, tile);
1099
1100 if (!factor_diag_tile<k_extent, internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
1101 internal_rhs>(A, A_stride, piv, piv_stride, k0, tile, oot_count, y,
1102 rhs_stride))
1103 return false;
1104
1106 piv_stride, k0, k0, tile);
1107
1108 // Fused forward substitution: this tile's y segment is final from
1109 // here on: unit-lower-solve it while the tile is in registers.
1110 if constexpr (fuse_rhs) {
1111 if constexpr (Config.unroll_inner) {
1113 for (int kk = 0; kk < k_extent; ++kk) {
1115 for (int i = kk + 1; i < k_extent; ++i)
1116 TDLS_LUPP_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_Y(k0 + kk);
1117 }
1118 } else {
1119 for (int kk = 0; kk < k_extent; ++kk) {
1120 for (int i = kk + 1; i < k_extent; ++i)
1121 TDLS_LUPP_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_Y(k0 + kk);
1122 }
1123 }
1124 }
1125
1126 // L panel below the diagonal
1127 for (int i = k + 1; i < full_tiles; ++i)
1128 ll_update_below_one<k_extent, tile_size, internal_piv, internal_matrix, fuse_rhs,
1129 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0, i * tile_size,
1130 y, rhs_stride);
1131 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
1132 ll_update_below_one<k_extent, last_tile_tail, internal_piv, internal_matrix, fuse_rhs,
1133 internal_rhs>(A, A_stride, piv, piv_stride, tile, k0,
1134 full_tiles * tile_size, y, rhs_stride);
1135
1136 // U row panel right of the diagonal
1137 for (int j = k + 1; j < full_tiles; ++j)
1139 A, A_stride, piv, piv_stride, tile, k0, j * tile_size);
1140 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
1142 A, A_stride, piv, piv_stride, tile, k0, full_tiles * tile_size);
1143 return true;
1144 }
1145
1146 /* =====================================================================
1147 FACTORIZE - public entry point.
1148 ===================================================================== */
1149
1177 template<bool internal_piv, bool internal_matrix, bool oot_diagnostics = true,
1178 bool fuse_rhs = false, bool internal_rhs = true>
1179 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1180 factorize(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
1181 int& oot_count, T* TDLS_RESTRICT y = nullptr, const int rhs_stride = 1) noexcept {
1182
1183 if constexpr (oot_diagnostics) oot_count = 0;
1184
1185 if constexpr (Config.unroll_inner) {
1187 for (int i = 0; i < N; ++i)
1188 TDLS_LUPP_PIV(i) = i;
1189 } else {
1190 for (int i = 0; i < N; ++i)
1191 TDLS_LUPP_PIV(i) = i;
1192 }
1193
1194 if constexpr (schedule == Schedule::RightLooking) {
1195 for (int k = 0; k < full_tiles; ++k)
1196 if (!rl_step<tile_size, internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
1197 internal_rhs>(A, A_stride, piv, piv_stride, k, oot_count, y,
1198 rhs_stride))
1199 return false;
1200 if constexpr (last_tile_tail > 0)
1201 if (!rl_step<last_tile_tail, internal_piv, internal_matrix, oot_diagnostics,
1202 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, full_tiles,
1203 oot_count, y, rhs_stride))
1204 return false;
1205 } else {
1206 for (int k = 0; k < full_tiles; ++k)
1207 if (!ll_step<tile_size, internal_piv, internal_matrix, oot_diagnostics, fuse_rhs,
1208 internal_rhs>(A, A_stride, piv, piv_stride, k, oot_count, y,
1209 rhs_stride))
1210 return false;
1211 if constexpr (last_tile_tail > 0)
1212 if (!ll_step<last_tile_tail, internal_piv, internal_matrix, oot_diagnostics,
1213 fuse_rhs, internal_rhs>(A, A_stride, piv, piv_stride, full_tiles,
1214 oot_count, y, rhs_stride))
1215 return false;
1216 }
1217
1218 return true;
1219 }
1220
1230 template<bool internal_piv, bool internal_matrix>
1231 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1232 factorize(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1233 const int piv_stride) noexcept {
1234 int unused = 0;
1235 return factorize<internal_piv, internal_matrix, false>(A, A_stride, piv, piv_stride,
1236 unused);
1237 }
1238
1239 /* =====================================================================
1240 SUBSTITUTION - schedule-independent.
1241 Forward: unit-lower solve tile by tile, each solved segment pushed
1242 into the tiles below it. Backward: each segment first pulls the
1243 trailing contributions, then upper-solves its diagonal tile.
1244 ===================================================================== */
1245
1263 template<int k_extent, int m_extent, int pass_width, bool internal_rhs, bool internal_piv,
1264 bool internal_matrix>
1265 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1266 fwd_push_one(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1267 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
1268 const int xcol_stride, const int k0, const int m0) noexcept {
1269 T Lmk[tile_size * tile_size];
1271 piv_stride, m0, k0, Lmk);
1272 if constexpr (Config.unroll_inner) {
1274 for (int i = 0; i < m_extent; ++i) {
1276 for (int w = 0; w < pass_width; ++w) {
1277 T sum = T(0);
1279 for (int j = 0; j < k_extent; ++j)
1280 sum += Lmk[i * tile_size + j] * TDLS_LUPP_XW(w, k0 + j);
1281 TDLS_LUPP_XW(w, m0 + i) -= sum;
1282 }
1283 }
1284 } else {
1285 for (int i = 0; i < m_extent; ++i) {
1286 for (int w = 0; w < pass_width; ++w) {
1287 T sum = T(0);
1288 for (int j = 0; j < k_extent; ++j)
1289 sum += Lmk[i * tile_size + j] * TDLS_LUPP_XW(w, k0 + j);
1290 TDLS_LUPP_XW(w, m0 + i) -= sum;
1291 }
1292 }
1293 }
1294 }
1295
1311 template<int k_extent, int pass_width, bool internal_rhs, bool internal_piv,
1312 bool internal_matrix>
1313 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1314 fwd_step(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1315 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride,
1316 const int k) noexcept {
1317 const int k0 = k * tile_size;
1318
1319 T Lkk[tile_size * tile_size];
1321 k0, k0, Lkk);
1322
1323 // In-tile unit-lower solve
1324 if constexpr (Config.unroll_inner) {
1326 for (int kk = 0; kk < k_extent; ++kk) {
1328 for (int i = kk + 1; i < k_extent; ++i) {
1330 for (int w = 0; w < pass_width; ++w)
1331 TDLS_LUPP_XW(w, k0 + i) -=
1332 Lkk[i * tile_size + kk] * TDLS_LUPP_XW(w, k0 + kk);
1333 }
1334 }
1335 } else {
1336 for (int kk = 0; kk < k_extent; ++kk) {
1337 for (int i = kk + 1; i < k_extent; ++i) {
1338 for (int w = 0; w < pass_width; ++w)
1339 TDLS_LUPP_XW(w, k0 + i) -=
1340 Lkk[i * tile_size + kk] * TDLS_LUPP_XW(w, k0 + kk);
1341 }
1342 }
1343 }
1344
1345 // Push into the tiles below
1346 for (int m = k + 1; m < full_tiles; ++m)
1347 fwd_push_one<k_extent, tile_size, pass_width, internal_rhs, internal_piv,
1348 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1349 k0, m * tile_size);
1350 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
1351 fwd_push_one<k_extent, last_tile_tail, pass_width, internal_rhs, internal_piv,
1352 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1353 k0, full_tiles * tile_size);
1354 }
1355
1373 template<int k_extent, int m_extent, int pass_width, bool internal_rhs, bool internal_piv,
1374 bool internal_matrix>
1375 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1376 bwd_pull_one(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1377 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
1378 const int xcol_stride, const int k0, const int m0) noexcept {
1379 T Ukm[tile_size * tile_size];
1381 piv_stride, k0, m0, Ukm);
1382 if constexpr (Config.unroll_inner) {
1384 for (int i = 0; i < k_extent; ++i) {
1386 for (int w = 0; w < pass_width; ++w) {
1387 T sum = T(0);
1389 for (int j = 0; j < m_extent; ++j)
1390 sum += Ukm[i * tile_size + j] * TDLS_LUPP_XW(w, m0 + j);
1391 TDLS_LUPP_XW(w, k0 + i) -= sum;
1392 }
1393 }
1394 } else {
1395 for (int i = 0; i < k_extent; ++i) {
1396 for (int w = 0; w < pass_width; ++w) {
1397 T sum = T(0);
1398 for (int j = 0; j < m_extent; ++j)
1399 sum += Ukm[i * tile_size + j] * TDLS_LUPP_XW(w, m0 + j);
1400 TDLS_LUPP_XW(w, k0 + i) -= sum;
1401 }
1402 }
1403 }
1404 }
1405
1421 template<int k_extent, int pass_width, bool internal_rhs, bool internal_piv,
1422 bool internal_matrix>
1423 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1424 bwd_step(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1425 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride,
1426 const int k) noexcept {
1427 const int k0 = k * tile_size;
1428
1429 // Pull the trailing contributions
1430 for (int m = k + 1; m < full_tiles; ++m)
1431 bwd_pull_one<k_extent, tile_size, pass_width, internal_rhs, internal_piv,
1432 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1433 k0, m * tile_size);
1434 if constexpr (last_tile_tail > 0 && k_extent == tile_size)
1435 bwd_pull_one<k_extent, last_tile_tail, pass_width, internal_rhs, internal_piv,
1436 internal_matrix>(A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1437 k0, full_tiles * tile_size);
1438
1439 // In-tile upper solve
1440 T Ukk[tile_size * tile_size];
1442 k0, k0, Ukk);
1443
1444 if constexpr (Config.unroll_inner) {
1446 for (int kk = k_extent - 1; kk >= 0; --kk) {
1448 for (int w = 0; w < pass_width; ++w)
1449 TDLS_LUPP_XW(w, k0 + kk) *= Ukk[kk * tile_size + kk]; // diag holds 1/pivot
1451 for (int i = 0; i < kk; ++i) {
1453 for (int w = 0; w < pass_width; ++w)
1454 TDLS_LUPP_XW(w, k0 + i) -=
1455 Ukk[i * tile_size + kk] * TDLS_LUPP_XW(w, k0 + kk);
1456 }
1457 }
1458 } else {
1459 for (int kk = k_extent - 1; kk >= 0; --kk) {
1460 for (int w = 0; w < pass_width; ++w)
1461 TDLS_LUPP_XW(w, k0 + kk) *= Ukk[kk * tile_size + kk]; // diag holds 1/pivot
1462 for (int i = 0; i < kk; ++i) {
1463 for (int w = 0; w < pass_width; ++w)
1464 TDLS_LUPP_XW(w, k0 + i) -=
1465 Ukk[i * tile_size + kk] * TDLS_LUPP_XW(w, k0 + kk);
1466 }
1467 }
1468 }
1469 }
1470
1484 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
1485 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1486 bwd_only(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1487 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
1488 const int xcol_stride) noexcept {
1489 if constexpr (last_tile_tail > 0)
1491 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, full_tiles);
1492 for (int k = full_tiles - 1; k >= 0; --k)
1494 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1495 }
1496
1512 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
1513 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1514 fwd_bwd(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1515 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
1516 const int xcol_stride) noexcept {
1517 for (int k = 0; k < full_tiles; ++k)
1519 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1520 if constexpr (last_tile_tail > 0)
1522 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, full_tiles);
1523
1525 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1526 }
1527
1528 /* =====================================================================
1529 SUBSTITUTION - public entry points.
1530 ===================================================================== */
1531
1547 template<bool internal_rhs, bool internal_piv, bool internal_matrix>
1548 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1549 substitute(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1550 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
1551 const int rhs_stride) noexcept {
1552 if constexpr (internal_rhs) {
1553 // Predicated gather: b[piv[i]] would index the internal b with
1554 // a runtime value and demote it to local memory. The equality
1555 // sweep keeps every index compile-time (measured: local memory
1556 // eliminated, identical values).
1557 if constexpr (Config.unroll_inner) {
1559 for (int i = 0; i < N; ++i) {
1560 const int p = TDLS_LUPP_PIV(i);
1561 T v = T(0);
1563 for (int j = 0; j < N; ++j)
1564 if (p == j) v = b[j];
1565 TDLS_LUPP_X(i) = v;
1566 }
1567 } else {
1568 for (int i = 0; i < N; ++i) {
1569 const int p = TDLS_LUPP_PIV(i);
1570 T v = T(0);
1571 for (int j = 0; j < N; ++j)
1572 if (p == j) v = b[j];
1573 TDLS_LUPP_X(i) = v;
1574 }
1575 }
1576 } else {
1577 if constexpr (Config.unroll_inner) {
1579 for (int i = 0; i < N; ++i)
1581 } else {
1582 for (int i = 0; i < N; ++i)
1584 }
1585 }
1586 fwd_bwd<1, internal_rhs, internal_piv, internal_matrix>(A, A_stride, piv, piv_stride, x,
1587 rhs_stride, 0);
1588 }
1589
1606 template<bool internal_rhs, bool internal_piv, bool internal_matrix>
1607 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1608 substitute_canonical(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1609 const int piv_stride, const int col, T* TDLS_RESTRICT x,
1610 const int rhs_stride) noexcept {
1612 A, A_stride, piv, piv_stride, col, x, rhs_stride, 0);
1613 }
1614
1633 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
1634 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1636 const int* TDLS_RESTRICT piv, const int piv_stride,
1637 const int col0, T* TDLS_RESTRICT x, const int rhs_stride,
1638 const int xcol_stride) noexcept {
1639 if constexpr (Config.unroll_inner) {
1641 for (int i = 0; i < N; ++i) {
1642 const int p = TDLS_LUPP_PIV(i);
1644 for (int w = 0; w < pass_width; ++w)
1645 TDLS_LUPP_XW(w, i) = (p == col0 + w) ? T(1) : T(0);
1646 }
1647 } else {
1648 for (int i = 0; i < N; ++i) {
1649 const int p = TDLS_LUPP_PIV(i);
1650 for (int w = 0; w < pass_width; ++w)
1651 TDLS_LUPP_XW(w, i) = (p == col0 + w) ? T(1) : T(0);
1652 }
1653 }
1655 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1656 }
1657
1685 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
1686 int pass_width = 0>
1687 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1688 substitute_canonical_multirhs(const T* TDLS_RESTRICT A, const int A_stride,
1689 const int* TDLS_RESTRICT piv, const int piv_stride,
1690 const int col0, T* TDLS_RESTRICT x, const int rhs_stride,
1691 const int xcol_stride) noexcept {
1692 static_assert(nrhs >= 1, "tdls: nrhs must be at least 1");
1693 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
1694 constexpr int columns_per_pass =
1695 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
1696 for (int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
1697 const unsigned off =
1698 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
1699 substitute_canonical_multirhs_pass<columns_per_pass, internal_rhs, internal_piv,
1700 internal_matrix>(
1701 A, A_stride, piv, piv_stride, col0 + c0, x + off, rhs_stride, xcol_stride);
1702 }
1703 if constexpr (nrhs % columns_per_pass > 0) {
1704 constexpr int c0 = nrhs - nrhs % columns_per_pass;
1705 const unsigned off =
1706 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
1707 substitute_canonical_multirhs_pass<nrhs % columns_per_pass, internal_rhs, internal_piv,
1708 internal_matrix>(
1709 A, A_stride, piv, piv_stride, col0 + c0, x + off, rhs_stride, xcol_stride);
1710 }
1711 }
1712
1731 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
1732 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1733 substitute_multirhs_pass(const T* TDLS_RESTRICT A, const int A_stride,
1734 const int* TDLS_RESTRICT piv, const int piv_stride,
1735 const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
1736 const int xcol_stride) noexcept {
1737 if constexpr (internal_rhs) {
1738 // Predicated gather, exactly as in substitute: b[piv[i]] would
1739 // index the internal block with a runtime value and demote it
1740 // to local memory.
1741 if constexpr (Config.unroll_inner) {
1743 for (int i = 0; i < N; ++i) {
1744 const int p = TDLS_LUPP_PIV(i);
1746 for (int w = 0; w < pass_width; ++w) {
1747 T v = T(0);
1749 for (int j = 0; j < N; ++j)
1750 if (p == j) v = TDLS_LUPP_BW(w, j);
1751 TDLS_LUPP_XW(w, i) = v;
1752 }
1753 }
1754 } else {
1755 for (int i = 0; i < N; ++i) {
1756 const int p = TDLS_LUPP_PIV(i);
1757 for (int w = 0; w < pass_width; ++w) {
1758 T v = T(0);
1759 for (int j = 0; j < N; ++j)
1760 if (p == j) v = TDLS_LUPP_BW(w, j);
1761 TDLS_LUPP_XW(w, i) = v;
1762 }
1763 }
1764 }
1765 } else {
1766 if constexpr (Config.unroll_inner) {
1768 for (int i = 0; i < N; ++i) {
1769 const int p = TDLS_LUPP_PIV(i);
1771 for (int w = 0; w < pass_width; ++w)
1772 TDLS_LUPP_XW(w, i) = TDLS_LUPP_BW(w, p);
1773 }
1774 } else {
1775 for (int i = 0; i < N; ++i) {
1776 const int p = TDLS_LUPP_PIV(i);
1777 for (int w = 0; w < pass_width; ++w)
1778 TDLS_LUPP_XW(w, i) = TDLS_LUPP_BW(w, p);
1779 }
1780 }
1781 }
1783 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1784 }
1785
1816 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
1817 int pass_width = 0>
1818 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1819 substitute_multirhs(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1820 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
1821 const int rhs_stride, const int xcol_stride) noexcept {
1822 static_assert(nrhs >= 1, "tdls: nrhs must be at least 1");
1823 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
1824 constexpr int columns_per_pass =
1825 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
1826 for (int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
1827 const unsigned off =
1828 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
1830 A, A_stride, piv, piv_stride, b + off, x + off, rhs_stride, xcol_stride);
1831 }
1832 if constexpr (nrhs % columns_per_pass > 0) {
1833 constexpr int c0 = nrhs - nrhs % columns_per_pass;
1834 const unsigned off =
1835 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
1836 substitute_multirhs_pass<nrhs % columns_per_pass, internal_rhs, internal_piv,
1837 internal_matrix>(A, A_stride, piv, piv_stride, b + off,
1838 x + off, rhs_stride, xcol_stride);
1839 }
1840 }
1841
1860 template<bool internal_rhs, bool internal_piv, bool internal_matrix>
1861 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1862 substitute_inplace(const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1863 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride) noexcept {
1864 if constexpr (N <= 64) {
1865 // Bitmask cycle decomposition: one visited bit per row, the
1866 // whole state in a single 32- or 64-bit register.
1867 using mask_t = std::conditional_t<(N <= 32), unsigned, unsigned long long>;
1868 mask_t visited = mask_t(0);
1869 if constexpr (Config.unroll_inner) {
1871 for (int s = 0; s < N; ++s) {
1872 if ((visited >> s) & mask_t(1)) continue;
1873 const T tmp = TDLS_LUPP_X(s);
1874 int cur = s;
1875 int nxt = TDLS_LUPP_PIV(cur);
1876 while (nxt != s) {
1877 TDLS_LUPP_X(cur) = TDLS_LUPP_X(nxt);
1878 visited |= mask_t(1) << cur;
1879 cur = nxt;
1880 nxt = TDLS_LUPP_PIV(cur);
1881 }
1882 TDLS_LUPP_X(cur) = tmp;
1883 visited |= mask_t(1) << cur;
1884 }
1885 } else {
1886 for (int s = 0; s < N; ++s) {
1887 if ((visited >> s) & mask_t(1)) continue;
1888 const T tmp = TDLS_LUPP_X(s);
1889 int cur = s;
1890 int nxt = TDLS_LUPP_PIV(cur);
1891 while (nxt != s) {
1892 TDLS_LUPP_X(cur) = TDLS_LUPP_X(nxt);
1893 visited |= mask_t(1) << cur;
1894 cur = nxt;
1895 nxt = TDLS_LUPP_PIV(cur);
1896 }
1897 TDLS_LUPP_X(cur) = tmp;
1898 visited |= mask_t(1) << cur;
1899 }
1900 }
1901 } else {
1902 // Cycle-leader scan (no visited storage): a cycle is rotated
1903 // only when reached from its smallest index, detected by
1904 // walking the orbit. Values and order match the mask path.
1905 for (int s = 0; s < N; ++s) {
1906 int probe = TDLS_LUPP_PIV(s);
1907 while (probe > s)
1908 probe = TDLS_LUPP_PIV(probe);
1909 if (probe != s) continue;
1910
1911 const T tmp = TDLS_LUPP_X(s);
1912 int cur = s;
1913 int nxt = TDLS_LUPP_PIV(cur);
1914 while (nxt != s) {
1915 TDLS_LUPP_X(cur) = TDLS_LUPP_X(nxt);
1916 cur = nxt;
1917 nxt = TDLS_LUPP_PIV(cur);
1918 }
1919 TDLS_LUPP_X(cur) = tmp;
1920 }
1921 }
1922 fwd_bwd<1, internal_rhs, internal_piv, internal_matrix>(A, A_stride, piv, piv_stride, x,
1923 rhs_stride, 0);
1924 }
1925
1944 template<int pass_width, bool internal_rhs, bool internal_piv, bool internal_matrix>
1945 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1946 substitute_inplace_multirhs_pass(const T* TDLS_RESTRICT A, const int A_stride,
1947 const int* TDLS_RESTRICT piv, const int piv_stride,
1948 T* TDLS_RESTRICT x, const int rhs_stride,
1949 const int xcol_stride) noexcept {
1950 if constexpr (N <= 64) {
1951 // Bitmask cycle decomposition, as in substitute_inplace; tmp
1952 // widens to one entry per column.
1953 using mask_t = std::conditional_t<(N <= 32), unsigned, unsigned long long>;
1954 mask_t visited = mask_t(0);
1955 if constexpr (Config.unroll_inner) {
1957 for (int s = 0; s < N; ++s) {
1958 if ((visited >> s) & mask_t(1)) continue;
1959 T tmp[pass_width];
1961 for (int w = 0; w < pass_width; ++w)
1962 tmp[w] = TDLS_LUPP_XW(w, s);
1963 int cur = s;
1964 int nxt = TDLS_LUPP_PIV(cur);
1965 while (nxt != s) {
1967 for (int w = 0; w < pass_width; ++w)
1968 TDLS_LUPP_XW(w, cur) = TDLS_LUPP_XW(w, nxt);
1969 visited |= mask_t(1) << cur;
1970 cur = nxt;
1971 nxt = TDLS_LUPP_PIV(cur);
1972 }
1974 for (int w = 0; w < pass_width; ++w)
1975 TDLS_LUPP_XW(w, cur) = tmp[w];
1976 visited |= mask_t(1) << cur;
1977 }
1978 } else {
1979 for (int s = 0; s < N; ++s) {
1980 if ((visited >> s) & mask_t(1)) continue;
1981 T tmp[pass_width];
1982 for (int w = 0; w < pass_width; ++w)
1983 tmp[w] = TDLS_LUPP_XW(w, s);
1984 int cur = s;
1985 int nxt = TDLS_LUPP_PIV(cur);
1986 while (nxt != s) {
1987 for (int w = 0; w < pass_width; ++w)
1988 TDLS_LUPP_XW(w, cur) = TDLS_LUPP_XW(w, nxt);
1989 visited |= mask_t(1) << cur;
1990 cur = nxt;
1991 nxt = TDLS_LUPP_PIV(cur);
1992 }
1993 for (int w = 0; w < pass_width; ++w)
1994 TDLS_LUPP_XW(w, cur) = tmp[w];
1995 visited |= mask_t(1) << cur;
1996 }
1997 }
1998 } else {
1999 // Cycle-leader scan, as in substitute_inplace.
2000 for (int s = 0; s < N; ++s) {
2001 int probe = TDLS_LUPP_PIV(s);
2002 while (probe > s)
2003 probe = TDLS_LUPP_PIV(probe);
2004 if (probe != s) continue;
2005
2006 T tmp[pass_width];
2007 for (int w = 0; w < pass_width; ++w)
2008 tmp[w] = TDLS_LUPP_XW(w, s);
2009 int cur = s;
2010 int nxt = TDLS_LUPP_PIV(cur);
2011 while (nxt != s) {
2012 for (int w = 0; w < pass_width; ++w)
2013 TDLS_LUPP_XW(w, cur) = TDLS_LUPP_XW(w, nxt);
2014 cur = nxt;
2015 nxt = TDLS_LUPP_PIV(cur);
2016 }
2017 for (int w = 0; w < pass_width; ++w)
2018 TDLS_LUPP_XW(w, cur) = tmp[w];
2019 }
2020 }
2022 A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2023 }
2024
2050 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
2051 int pass_width = 0>
2052 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
2053 substitute_inplace_multirhs(const T* TDLS_RESTRICT A, const int A_stride,
2054 const int* TDLS_RESTRICT piv, const int piv_stride,
2055 T* TDLS_RESTRICT x, const int rhs_stride,
2056 const int xcol_stride) noexcept {
2057 static_assert(nrhs >= 1, "tdls: nrhs must be at least 1");
2058 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
2059 constexpr int columns_per_pass =
2060 (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
2061 for (int c0 = 0; c0 + columns_per_pass <= nrhs; c0 += columns_per_pass) {
2062 const unsigned off =
2063 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
2064 substitute_inplace_multirhs_pass<columns_per_pass, internal_rhs, internal_piv,
2065 internal_matrix>(A, A_stride, piv, piv_stride, x + off,
2066 rhs_stride, xcol_stride);
2067 }
2068 if constexpr (nrhs % columns_per_pass > 0) {
2069 constexpr int c0 = nrhs - nrhs % columns_per_pass;
2070 const unsigned off =
2071 internal_rhs ? unsigned(c0) * unsigned(N) : unsigned(c0) * unsigned(xcol_stride);
2072 substitute_inplace_multirhs_pass<nrhs % columns_per_pass, internal_rhs, internal_piv,
2073 internal_matrix>(A, A_stride, piv, piv_stride, x + off,
2074 rhs_stride, xcol_stride);
2075 }
2076 }
2077
2078 /* =====================================================================
2079 SOLVE - convenience wrappers.
2080 ===================================================================== */
2081
2099 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
2100 bool oot_diagnostics = true>
2101 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
2102 solve(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
2103 const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
2104 int& oot_count) noexcept {
2105 if (!factorize<internal_piv, internal_matrix, oot_diagnostics>(A, A_stride, piv, piv_stride,
2106 oot_count))
2107 return false;
2108 substitute<internal_rhs, internal_piv, internal_matrix>(A, A_stride, piv, piv_stride, b, x,
2109 rhs_stride);
2110 return true;
2111 }
2112
2127 template<bool internal_rhs, bool internal_piv, bool internal_matrix>
2128 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
2129 solve(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv, const int piv_stride,
2130 const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride) noexcept {
2131 int unused = 0;
2133 A, A_stride, piv, piv_stride, b, x, rhs_stride, unused);
2134 }
2135
2161 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
2162 int pass_width = 0, bool oot_diagnostics = true>
2163 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
2164 solve_multirhs(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2165 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
2166 const int rhs_stride, const int xcol_stride, int& oot_count) noexcept {
2167 if (!factorize<internal_piv, internal_matrix, oot_diagnostics>(A, A_stride, piv, piv_stride,
2168 oot_count))
2169 return false;
2171 A, A_stride, piv, piv_stride, b, x, rhs_stride, xcol_stride);
2172 return true;
2173 }
2174
2197 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
2198 int pass_width = 0>
2199 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
2200 solve_multirhs(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2201 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
2202 const int rhs_stride, const int xcol_stride) noexcept {
2203 int unused = 0;
2205 A, A_stride, piv, piv_stride, b, x, rhs_stride, xcol_stride, unused);
2206 }
2207
2232 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
2233 bool oot_diagnostics = true>
2234 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
2235 solve_inplace(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2236 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride,
2237 int& oot_count) noexcept {
2239 A, A_stride, piv, piv_stride, oot_count, y, rhs_stride))
2240 return false;
2241
2242 // Backward pass only: the forward one happened inside factorize.
2243 bwd_only<1, internal_rhs, internal_piv, internal_matrix>(A, A_stride, piv, piv_stride, y,
2244 rhs_stride, 0);
2245 return true;
2246 }
2247
2261 template<bool internal_rhs, bool internal_piv, bool internal_matrix>
2262 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
2263 solve_inplace(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2264 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride) noexcept {
2265 int unused = 0;
2267 A, A_stride, piv, piv_stride, y, rhs_stride, unused);
2268 }
2269
2299 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
2300 int pass_width = 0, bool oot_diagnostics = true>
2301 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
2302 solve_inplace_multirhs(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2303 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride,
2304 const int xcol_stride, int& oot_count) noexcept {
2305 if (!factorize<internal_piv, internal_matrix, oot_diagnostics>(A, A_stride, piv, piv_stride,
2306 oot_count))
2307 return false;
2309 A, A_stride, piv, piv_stride, y, rhs_stride, xcol_stride);
2310 return true;
2311 }
2312
2334 template<int nrhs, bool internal_rhs, bool internal_piv, bool internal_matrix,
2335 int pass_width = 0>
2336 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
2337 solve_inplace_multirhs(T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2338 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride,
2339 const int xcol_stride) noexcept {
2340 int unused = 0;
2341 return solve_inplace_multirhs<nrhs, internal_rhs, internal_piv, internal_matrix, pass_width,
2342 false>(A, A_stride, piv, piv_stride, y, rhs_stride,
2343 xcol_stride, unused);
2344 }
2345
2346 /* =====================================================================
2347 Note on pivot ownership: the pivot storage is ALWAYS caller-provided
2348 (internal array or remote scratch), for every entry point: one
2349 uniform calling convention. A solver-internal pivot would only be
2350 expressible for the single combination {solve/solve_inplace x internal
2351 pivot}; externalizing it there costs nothing (the caller's int
2352 piv[N] inlines to the exact same codegen), so no special case is
2353 kept. For one-shot solves the contents are simply treated as
2354 scratch; for factorize + substitute* (consistent tangent), the same
2355 storage carries the permutation across the calls.
2356 ===================================================================== */
2357};
2358
2359
2360
2361#undef TDLS_LUPP_A
2362#undef TDLS_LUPP_PIV
2363#undef TDLS_LUPP_X
2364#undef TDLS_LUPP_B
2365#undef TDLS_LUPP_XW
2366#undef TDLS_LUPP_BW
2367#undef TDLS_LUPP_Y
2368
2369
2370
2371} // namespace tdls
2372
2373#if defined(__GNUC__) && !defined(__clang__)
2374#pragma GCC diagnostic pop
2375#endif
2376
2377#if defined(__clang__)
2378#pragma clang diagnostic pop
2379#endif
2380
2381
2382
2383#endif // TDLS_SOLVERS_TILED_LUPP_SOLVER_STATIC_HPP
Compile-time configuration of the TiledLUpp solver family.
#define TDLS_RESTRICT
Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC).
Definition macros.hpp:80
#define TDLS_UNROLL_FORCE
Full-unroll pragma in the local compiler dialect.
Definition macros.hpp:103
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:45
#define TDLS_FORCEINLINE
__forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS_HOST_DEVICE for why),...
Definition macros.hpp:67
Scalar math helpers usable in constant expressions.
Schedule
Elimination schedule of a factorization.
Definition options.hpp:37
#define TDLS_LUPP_XW(w, i)
Entry i of column w of a multi right-hand-side block: pass_width contiguous columns under internal re...
Definition solver_static.hpp:140
#define TDLS_LUPP_Y(i)
Entry i of the fused right-hand side of solve_inplace (y follows the matrix rows through pivoting).
Definition solver_static.hpp:152
#define TDLS_LUPP_PIV(i)
Pivot entry i: contiguous under internal residency, strided otherwise.
Definition solver_static.hpp:126
#define TDLS_LUPP_A(r, c)
Element (r, c) of the factor matrix: contiguous under internal residency, strided otherwise,...
Definition solver_static.hpp:120
#define TDLS_LUPP_BW(w, i)
Entry i of column w of a multi right-hand-side block of b, addressed exactly as TDLS_LUPP_XW (b and x...
Definition solver_static.hpp:146
#define TDLS_LUPP_X(i)
Entry i of the solution vector: contiguous under internal residency, strided otherwise.
Definition solver_static.hpp:130
#define TDLS_LUPP_B(i)
Entry i of the right-hand side: contiguous under internal residency, strided otherwise.
Definition solver_static.hpp:134
Tiled dense LU factorization with logical partial pivoting and out-of-tile pivot recovery,...
Definition solver_static.hpp:200
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical_multirhs(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int col0, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
nrhs canonical columns e_col0 .. e_{col0+nrhs-1} solved from a prior factorize: the columns of the in...
Definition solver_static.hpp:1688
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factor_diag_tile(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k0, T *TDLS_RESTRICT tile, int &oot_count, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
Factor the KExKE diagonal tile in registers, with out-of-tile pivot recovery (drives the per-column l...
Definition solver_static.hpp:686
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ll_update_below_one(T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT tile, const int k0, const int i0, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
LL: correct + TRSM one L-panel tile below the diagonal, with the optional fused forward-substitution ...
Definition solver_static.hpp:1006
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void rl_trsm_right_one(T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT pk, const T *TDLS_RESTRICT tile, const int j0) noexcept
RL: one row-panel tile update, Akj := L^-1 Akj.
Definition solver_static.hpp:727
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace_multirhs(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride, const int xcol_stride, int &oot_count) noexcept
factorize + substitute_inplace_multirhs in one call: y holds the nrhs unpermuted right-hand-side colu...
Definition solver_static.hpp:2302
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, int &oot_count) noexcept
factorize + substitute in one call.
Definition solver_static.hpp:2102
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Solve in place: x already holds the unpermuted RHS on entry and the solution on exit.
Definition solver_static.hpp:1862
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ll_correct_tile(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, const int k0, T *TDLS_RESTRICT t) noexcept
LL: t (RExCE, rows row0.., cols col0..) -= sum over prior tiles bj < k0/tile_size of L(row0....
Definition solver_static.hpp:972
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int col, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Solve with b = e_col generated on the fly: the consistent-tangent-operator path.
Definition solver_static.hpp:1608
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void store_tile(T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT prow, const int col0, const T *TDLS_RESTRICT t) noexcept
Store an RxC tile through a cached physical-row segment.
Definition solver_static.hpp:280
static constexpr int num_tiles
Tiles per dimension, the partial one included.
Definition solver_static.hpp:228
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factorize(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, int &oot_count, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
Factor A := P*L*U in place.
Definition solver_static.hpp:1180
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_multirhs(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride, int &oot_count) noexcept
factorize + substitute_multirhs in one call.
Definition solver_static.hpp:2164
static constexpr T oot_threshold
Acceptable-pivot threshold of the out-of-tile search, read once from the configuration (see TiledLUpp...
Definition solver_static.hpp:208
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical_multirhs_pass(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int col0, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
One pass of the canonical multi right-hand-side substitution: pass_width canonical columns e_col0 ....
Definition solver_static.hpp:1635
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Diagnostics-free solve overload: no out-of-tile out-parameter at all.
Definition solver_static.hpp:2129
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile_piv_upper(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, T *TDLS_RESTRICT t) noexcept
Triangular variant of the diagonal-tile load for the backward substitution: only the upper triangle i...
Definition solver_static.hpp:418
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void fwd_push_one(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride, const int k0, const int m0) noexcept
Forward push: subtract L(m,k) * x_k from the x_m segment, for pass_width columns at once.
Definition solver_static.hpp:1266
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile_piv(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, T *TDLS_RESTRICT t) noexcept
Load an RxC tile, reading the permutation inline (one read per row).
Definition solver_static.hpp:312
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool rl_step(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k, int &oot_count, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
RL: one full factorization step (diagonal tile + trailing updates).
Definition solver_static.hpp:881
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factorize(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride) noexcept
Diagnostics-free factorize overload: no out-of-tile out-parameter at all.
Definition solver_static.hpp:1232
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void bwd_only(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
Backward pass alone, used by fwd_bwd and by solve_inplace (whose forward pass happens inside the fact...
Definition solver_static.hpp:1486
static constexpr int tile_size
tile size (int)
Definition solver_static.hpp:202
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_multirhs_pass(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
One pass of the multi right-hand-side substitution: pass_width columns of b gathered in permuted orde...
Definition solver_static.hpp:1733
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void fwd_bwd(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
Triangular solves on already-permuted column(s) x. pass_width columns are processed per tile visit,...
Definition solver_static.hpp:1514
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT prow, const int col0, T *TDLS_RESTRICT t) noexcept
Load an RxC tile through a cached physical-row segment.
Definition solver_static.hpp:252
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace_multirhs_pass(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
One pass of the in-place multi right-hand-side substitution: pass_width columns of x permuted by the ...
Definition solver_static.hpp:1946
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Solve x := U^-1 L^-1 P b from a prior factorize. b and x must not alias (use substitute_inplace for t...
Definition solver_static.hpp:1549
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void fwd_step(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride, const int k) noexcept
Forward step: unit-lower solve the diagonal tile's segment, then push it into the tiles below.
Definition solver_static.hpp:1314
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void ll_update_right_one(T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT tile, const int k0, const int j0) noexcept
LL: correct + TRSM one U-panel tile right of the diagonal.
Definition solver_static.hpp:1056
static constexpr int full_tiles
Definition solver_static.hpp:224
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace_multirhs(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
nrhs columns solved in place: x holds the nrhs unpermuted right-hand-side columns on entry and the nr...
Definition solver_static.hpp:2053
static constexpr Schedule schedule
elimination schedule (RightLooking or LeftLooking)
Definition solver_static.hpp:203
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void bwd_pull_one(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride, const int k0, const int m0) noexcept
Backward pull: subtract U(k,m) * x_m from the x_k segment, for pass_width columns at once.
Definition solver_static.hpp:1376
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void rl_schur_one(T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT pk, const int *TDLS_RESTRICT pi, const T *TDLS_RESTRICT Aik, const int j0) noexcept
RL: one Schur-complement tile update, Aij -= Aik * Akj, streaming the factored Akj row by row from re...
Definition solver_static.hpp:750
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void load_tile_piv_lower(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, T *TDLS_RESTRICT t) noexcept
Triangular variant of the diagonal-tile load for the forward substitution: only the strict lower tria...
Definition solver_static.hpp:384
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride, int &oot_count) noexcept
Factorization with the forward substitution folded in.
Definition solver_static.hpp:2235
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void rl_update_row_one(T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int *TDLS_RESTRICT pk, const T *TDLS_RESTRICT tile, const int k, const int i0, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
RL: TRSM down + Schur sweep of one row block below the diagonal, with the optional fused forward-subs...
Definition solver_static.hpp:807
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void bwd_step(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride, const int k) noexcept
Backward step: pull the trailing contributions, then upper-solve the diagonal tile's segment.
Definition solver_static.hpp:1424
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factor_diag_column(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k0, T *TDLS_RESTRICT tile, int &oot_count, const int c, T *TDLS_RESTRICT y, const int rhs_stride) noexcept
One column step of the diagonal-tile factorization: pivot search (in-tile, then out-of-tile recovery)...
Definition solver_static.hpp:474
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_multirhs(const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
Solve nrhs right-hand-side columns from a prior factorize: X := U^-1 L^-1 P B, column by column....
Definition solver_static.hpp:1819
static constexpr T singular_floor
Singularity floor of the out-of-tile recovery, read once from the configuration (see TiledLUppConfig:...
Definition solver_static.hpp:211
static constexpr int last_tile_tail
Extent of the last, partial tile: 0 when N is a multiple of tile_size.
Definition solver_static.hpp:226
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void store_tile_piv(T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, const T *TDLS_RESTRICT t) noexcept
Store an RxC tile, reading the permutation inline (one read per row).
Definition solver_static.hpp:347
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace_multirhs(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride, const int xcol_stride) noexcept
Diagnostics-free solve_inplace_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_static.hpp:2337
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride) noexcept
Diagnostics-free solve_inplace overload: no out-of-tile out-parameter at all.
Definition solver_static.hpp:2263
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_multirhs(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
Diagnostics-free solve_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_static.hpp:2200
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool ll_step(T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k, int &oot_count, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
LL: one full factorization step (correct + factor the diagonal tile, then its L and U panels).
Definition solver_static.hpp:1089
tile_size x tile_size register-tile micro-kernels of the TiledLUpp solvers: the shared kernels plus t...
Definition tile_operations.hpp:46
LU register-tile micro-kernel of the TiledLUpp solver family.