TDLS 0.4.0
Tiny Device-callable Linear Solvers
Loading...
Searching...
No Matches
solver_dynamic.hpp
Go to the documentation of this file.
1#ifndef TDLS_SOLVERS_TILED_LUPP_SOLVER_DYNAMIC_HPP
2#define TDLS_SOLVERS_TILED_LUPP_SOLVER_DYNAMIC_HPP
3
4
5
66
67
68
69#include <tdls/core/math.hpp>
70#include <tdls/core/target.hpp>
71#include <type_traits>
72
74
75
76
77// gcc's flow analysis cannot prove that the replay buffer of the
78// out-of-tile search is written before being read (the loop structure
79// guarantees it, a property validated bitwise against the static
80// solver), and emits spurious -Wmaybe-uninitialized warnings at -O2.
81//
82// Under UBSan instrumentation (-fsanitize=undefined) at -O2 and
83// above, gcc additionally loses the value ranges of the loop indices
84// and reports impossible subscripts through -Warray-bounds: the
85// unsigned cast of a provably non-negative index read as a wrapped
86// negative, or a loop counter assumed below its own loop guard.
87// Observed from gcc 12 to gcc 16; plain builds and clang are clean,
88// and the exercised paths are certified free of out-of-bounds
89// accesses by the constexpr suite.
90//
91// Both suppressions are scoped to this header and to those warnings
92// only, exactly as in solver_static.hpp.
93#if defined(TDLS_DETAIL_COMPILER_GCC)
94#pragma GCC diagnostic push
95#pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
96#pragma GCC diagnostic ignored "-Warray-bounds"
97#endif
98
99namespace tdls {
100
101
102
103/* Addressing macros. They expand inside member functions where n, the
104 parameter names (A, A_stride, piv, piv_stride, x, b, y, rhs_stride,
105 xcol_stride) and the Config value are in scope. #undef'd at the end
106 of this header. */
107
111#define TDLS_LUPP_DYN_A(r, c) A[TDLS_STRIDED((r) * n + (c), A_stride)]
114#define TDLS_LUPP_DYN_PIV(i) piv[TDLS_STRIDED(i, piv_stride)]
117#define TDLS_LUPP_DYN_X(i) x[TDLS_STRIDED(i, rhs_stride)]
120#define TDLS_LUPP_DYN_B(i) b[TDLS_STRIDED(i, rhs_stride)]
123#define TDLS_LUPP_DYN_XW(w, i) x[TDLS_STRIDED(i, rhs_stride) + TDLS_STRIDED_WIDE(w, xcol_stride)]
127#define TDLS_LUPP_DYN_BW(w, i) b[TDLS_STRIDED(i, rhs_stride) + TDLS_STRIDED_WIDE(w, xcol_stride)]
130#define TDLS_LUPP_DYN_Y(i) y[TDLS_STRIDED(i, rhs_stride)]
131
132
133
175template<typename T, TiledLUppConfig<T> Config = TiledLUppConfig<T>{}>
177
178 static constexpr int tile_size = Config.tile_size;
179 static constexpr Schedule schedule =
180 Config.schedule;
181
184 static constexpr T oot_pivot_threshold = Config.oot_pivot_threshold;
187 static constexpr T relative_pivot_threshold = Config.relative_pivot_threshold;
190 static constexpr T singular_floor = Config.singular_floor;
194 static constexpr bool transposed = Config.layout == MatrixLayout::ColMajor;
197 static constexpr bool physical = Config.row_interchange == RowInterchange::Physical;
198
199 static_assert(Config.row_interchange != RowInterchange::Auto,
200 "TiledLUppSolverDynamic: row_interchange must be Logical or Physical (Auto is "
201 "the default of the blocks of ArrowheadLUpp only)");
202 static_assert(
203 Config.oot_pivot_threshold.is_finite() && Config.singular_floor.is_finite(),
204 "TiledLUppSolverDynamic: oot_pivot_threshold and singular_floor must be finite (and fit a "
205 "63-bit mantissa)");
206 static_assert(Config.relative_pivot_threshold.is_finite(),
207 "TiledLUppSolverDynamic: relative_pivot_threshold must be finite (and fit a "
208 "63-bit mantissa)");
209 static_assert(relative_pivot_threshold > T(0) && relative_pivot_threshold <= T(1),
210 "TiledLUppSolverDynamic: relative_pivot_threshold must be in (0, 1]");
211 static_assert(singular_floor > T(0), "TiledLUppSolverDynamic: singular_floor must be positive");
212 static_assert(singular_floor <= oot_pivot_threshold,
213 "TiledLUppSolverDynamic: singular_floor must not exceed oot_pivot_threshold (the "
214 "floor applies to the out-of-tile recovery path)");
215 static_assert(tile_size >= 1, "TiledLUppSolverDynamic: tile size must be >= 1");
216
220 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int num_tiles(const int n) noexcept {
221 return (n + tile_size - 1) / tile_size;
222 }
223
229 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int tile_size_at(const int t0,
230 const int n) noexcept {
231 return (n - t0 < tile_size) ? (n - t0) : tile_size;
232 }
233
234 /* =====================================================================
235 Register-tile micro-kernels, runtime extents. A tile is a
236 tile_size x tile_size array with row stride tile_size; the runtime
237 extents bound the active part.
238 ===================================================================== */
239
245 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
246 ops_swap_rows(T* TDLS_RESTRICT t, const int k, const int r, const int ke) noexcept {
247 for (int j = 0; j < ke; ++j) {
248 const T tmp = t[k * tile_size + j];
249 t[k * tile_size + j] = t[r * tile_size + j];
250 t[r * tile_size + j] = tmp;
251 }
252 }
253
263 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
264 ops_eliminate_column(T* TDLS_RESTRICT t, const int k, const int re, const int ce) noexcept {
265 const T inv_pivot = T(1) / t[k * tile_size + k];
266 t[k * tile_size + k] = inv_pivot;
267 for (int i = k + 1; i < re; ++i) {
268 t[i * tile_size + k] *= inv_pivot;
269 for (int j = k + 1; j < ce; ++j)
270 t[i * tile_size + j] -= t[i * tile_size + k] * t[k * tile_size + j];
271 }
272 }
273
280 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
281 ops_trsm_left_unit(const T* TDLS_RESTRICT lu, T* TDLS_RESTRICT B, const int kd,
282 const int ce) noexcept {
283 for (int k = 0; k < kd; ++k) {
284 for (int i = k + 1; i < kd; ++i) {
285 const T L_ik = lu[i * tile_size + k];
286 for (int j = 0; j < ce; ++j)
287 B[i * tile_size + j] -= L_ik * B[k * tile_size + j];
288 }
289 }
290 }
291
301 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
302 ops_trsm_right(const T* TDLS_RESTRICT lu, T* TDLS_RESTRICT B, const int kd,
303 const int re) noexcept {
304 for (int k = 0; k < kd; ++k) {
305 const T U_kk_inv = lu[k * tile_size + k];
306 for (int i = 0; i < re; ++i)
307 B[i * tile_size + k] *= U_kk_inv;
308 for (int j = k + 1; j < kd; ++j) {
309 const T U_kj = lu[k * tile_size + j];
310 for (int i = 0; i < re; ++i)
311 B[i * tile_size + j] -= B[i * tile_size + k] * U_kj;
312 }
313 }
314 }
315
324 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
325 ops_gemm_sub(T* TDLS_RESTRICT Ct, const T* TDLS_RESTRICT At, const T* TDLS_RESTRICT Bt,
326 const int re, const int ce, const int kd) noexcept {
327 for (int i = 0; i < re; ++i) {
328 for (int j = 0; j < ce; ++j) {
329 T sum = T(0);
330 for (int k = 0; k < kd; ++k)
331 sum += At[i * tile_size + k] * Bt[k * tile_size + j];
332 Ct[i * tile_size + j] -= sum;
333 }
334 }
335 }
336
337 /* =====================================================================
338 Remote <-> register tile movement, runtime extents.
339 ===================================================================== */
340
350 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
351 load_tile(const int n, const T* TDLS_RESTRICT A, const int A_stride,
352 const int* TDLS_RESTRICT prow, const int col0, T* TDLS_RESTRICT t, const int re,
353 const int ce) noexcept {
354 for (int i = 0; i < re; ++i) {
355 for (int j = 0; j < ce; ++j)
356 t[i * tile_size + j] = TDLS_LUPP_DYN_A(prow[i], col0 + j);
357 }
358 }
359
369 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
370 store_tile(const int n, T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT prow,
371 const int col0, const T* TDLS_RESTRICT t, const int re, const int ce) noexcept {
372 for (int i = 0; i < re; ++i) {
373 for (int j = 0; j < ce; ++j)
374 TDLS_LUPP_DYN_A(prow[i], col0 + j) = t[i * tile_size + j];
375 }
376 }
377
390 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
391 load_tile_piv(const int n, const T* TDLS_RESTRICT A, const int A_stride,
392 const int* TDLS_RESTRICT piv, const int piv_stride, const int row0,
393 const int col0, T* TDLS_RESTRICT t, const int re, const int ce) noexcept {
394 for (int i = 0; i < re; ++i) {
395 const int phys = TDLS_LUPP_DYN_PIV(row0 + i);
396 for (int j = 0; j < ce; ++j)
397 t[i * tile_size + j] = TDLS_LUPP_DYN_A(phys, col0 + j);
398 }
399 }
400
413 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
414 store_tile_piv(const int n, T* TDLS_RESTRICT A, const int A_stride,
415 const int* TDLS_RESTRICT piv, const int piv_stride, const int row0,
416 const int col0, const T* TDLS_RESTRICT t, const int re, const int ce) noexcept {
417 for (int i = 0; i < re; ++i) {
418 const int phys = TDLS_LUPP_DYN_PIV(row0 + i);
419 for (int j = 0; j < ce; ++j)
420 TDLS_LUPP_DYN_A(phys, col0 + j) = t[i * tile_size + j];
421 }
422 }
423
436 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
437 load_tile_piv_lower(const int n, const T* TDLS_RESTRICT A, const int A_stride,
438 const int* TDLS_RESTRICT piv, const int piv_stride, const int row0,
439 const int col0, T* TDLS_RESTRICT t, const int re) noexcept {
440 for (int i = 1; i < re; ++i) {
441 const int phys = TDLS_LUPP_DYN_PIV(row0 + i);
442 for (int j = 0; j < i; ++j)
443 t[i * tile_size + j] = TDLS_LUPP_DYN_A(phys, col0 + j);
444 }
445 }
446
459 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
460 load_tile_piv_upper(const int n, const T* TDLS_RESTRICT A, const int A_stride,
461 const int* TDLS_RESTRICT piv, const int piv_stride, const int row0,
462 const int col0, T* TDLS_RESTRICT t, const int re) noexcept {
463 for (int i = 0; i < re; ++i) {
464 const int phys = TDLS_LUPP_DYN_PIV(row0 + i);
465 for (int j = i; j < re; ++j)
466 t[i * tile_size + j] = TDLS_LUPP_DYN_A(phys, col0 + j);
467 }
468 }
469
470 /* =====================================================================
471 Diagonal-tile factorization with out-of-tile pivoting. Same
472 algorithm as the static TiledLUpp solver, runtime column/extent parameters,
473 direct (unpredicated) fused-RHS swap.
474 ===================================================================== */
475
495 template<bool oot_diagnostics, bool fuse_rhs>
496 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
497 factor_diag_column(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
498 const int piv_stride, const int k0, T* TDLS_RESTRICT tile, int& oot_count,
499 const int c, const int ke, T* TDLS_RESTRICT y,
500 const int rhs_stride) noexcept {
501
502 const int gc = k0 + c; // global column
503
504 // In-tile pivot search (rows c..ke of the register tile)
505 int best_r = c;
506 T best = detail::abs(tile[c * tile_size + c]);
507 for (int r = c + 1; r < ke; ++r) {
508 const T v = detail::abs(tile[r * tile_size + c]);
509 if (v > best) {
510 best = v;
511 best_r = r;
512 }
513 }
514
515 // Threshold pivoting inside the tile: the row in place keeps the
516 // pivot when it reaches both relative_pivot_threshold times the
517 // best in-tile candidate and oot_pivot_threshold. The out-of-tile
518 // recovery below keeps its own absolute trigger. A threshold of 1
519 // compiles the test out: the best in-tile candidate wins.
520 if constexpr (relative_pivot_threshold < T(1)) {
521 const T in_place = detail::abs(tile[c * tile_size + c]);
522 if (in_place >= relative_pivot_threshold * best && in_place >= oot_pivot_threshold)
523 best_r = c;
524 }
525
526 int piv_row; // winning global (logical) row
527
528 if (best >= oot_pivot_threshold) {
529 piv_row = k0 + best_r;
530 } else if (ke < tile_size) {
531 // Trailing tile: no rows below to recover from. Diagnostic
532 // order: singularity verdict first, then count the weak pivot
533 // (full tiles count before the verdict).
534 if (best < singular_floor) return false;
535 if constexpr (oot_diagnostics) ++oot_count;
536 piv_row = k0 + best_r;
537 } else {
538 // Out-of-tile recovery: scan the rows below the tile and
539 // evaluate each candidate as if it had received the
540 // eliminations it is missing, keeping the best (or, with
541 // Config.oot_first_acceptable, the first to reach the threshold).
542 if constexpr (oot_diagnostics) ++oot_count;
543 T gbest = best;
544 int gbest_row = k0 + best_r;
545
546 for (int row = k0 + ke; row < n; ++row) {
547 const int phys = TDLS_LUPP_DYN_PIV(row);
548
549 T corrected = TDLS_LUPP_DYN_A(phys, gc);
550 if constexpr (schedule == Schedule::LeftLooking) {
551 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
552 for (int p = 0; p < tile_size; ++p)
553 corrected -= TDLS_LUPP_DYN_A(phys, bj0 + p) *
555 }
556
557 if (c > 0) {
558 T L_row[tile_size];
559 for (int t = 0; t < c; ++t) {
560 T a_t = TDLS_LUPP_DYN_A(phys, k0 + t);
561 if constexpr (schedule == Schedule::LeftLooking) {
562 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
563 for (int p = 0; p < tile_size; ++p)
564 a_t -= TDLS_LUPP_DYN_A(phys, bj0 + p) *
565 TDLS_LUPP_DYN_A(TDLS_LUPP_DYN_PIV(bj0 + p), k0 + t);
566 }
567 for (int p = 0; p < t; ++p)
568 a_t -= L_row[p] * tile[p * tile_size + t];
569 L_row[t] = a_t * tile[t * tile_size + t]; // diag holds 1/pivot
570 corrected -= L_row[t] * tile[t * tile_size + c];
571 }
572 }
573
574 const T v = detail::abs(corrected);
575 if (v > gbest) {
576 gbest = v;
577 gbest_row = row;
578 }
579
580 // First-acceptable out-of-tile pivot: a candidate that
581 // reaches the threshold already beats the sub-threshold
582 // in-tile pivot, so stop scanning.
583 if constexpr (Config.oot_first_acceptable)
584 if (v >= oot_pivot_threshold) break;
585 }
586
587 if (gbest < singular_floor) return false;
588 piv_row = gbest_row;
589 }
590
591 {
592 // Swap the permutation entries unconditionally (self-swap is a
593 // bitwise no-op, and the branch would diverge on GPU).
594 const int tmp = TDLS_LUPP_DYN_PIV(gc);
596 TDLS_LUPP_DYN_PIV(piv_row) = tmp;
597
598 if constexpr (fuse_rhs && !transposed) {
599 // The fused RHS follows the rows through pivoting (row-major
600 // layout; under the column-major one y is indexed by the
601 // columns of S, which pivoting never moves).
602 const T ty = TDLS_LUPP_DYN_Y(gc);
603 TDLS_LUPP_DYN_Y(gc) = TDLS_LUPP_DYN_Y(piv_row);
604 TDLS_LUPP_DYN_Y(piv_row) = ty;
605 }
606
607 if (piv_row < k0 + ke) {
608 ops_swap_rows(tile, c, piv_row - k0, ke);
609 } else {
610 // Cross-tile swap: pull the new row into the tile and
611 // replay everything it missed: prior tiles (LL), then
612 // the current tile's factored columns, on the FULL row.
613 const int phys = TDLS_LUPP_DYN_PIV(gc);
614 for (int j = 0; j < ke; ++j) {
615 tile[c * tile_size + j] = TDLS_LUPP_DYN_A(phys, k0 + j);
616 if constexpr (schedule == Schedule::LeftLooking) {
617 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
618 for (int p = 0; p < tile_size; ++p)
619 tile[c * tile_size + j] -=
620 TDLS_LUPP_DYN_A(phys, bj0 + p) *
621 TDLS_LUPP_DYN_A(TDLS_LUPP_DYN_PIV(bj0 + p), k0 + j);
622 }
623 }
624
625 if (c > 0) {
626 T L_row[tile_size];
627 for (int t = 0; t < c; ++t) {
628 T a_t = tile[c * tile_size + t];
629 for (int p = 0; p < t; ++p)
630 a_t -= L_row[p] * tile[p * tile_size + t];
631 L_row[t] = a_t * tile[t * tile_size + t]; // diag holds 1/pivot
632 tile[c * tile_size + t] = L_row[t];
633 }
634 for (int j = c; j < ke; ++j) {
635 for (int t = 0; t < c; ++t)
636 tile[c * tile_size + j] -= L_row[t] * tile[t * tile_size + j];
637 }
638 }
639 }
640 }
641
642 ops_eliminate_column(tile, c, ke, ke);
643 return true;
644 }
645
665 template<bool oot_diagnostics, bool fuse_rhs>
666 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
667 factor_diag_tile(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
668 const int piv_stride, const int k0, T* TDLS_RESTRICT tile, int& oot_count,
669 const int ke, T* TDLS_RESTRICT y = nullptr,
670 const int rhs_stride = 1) noexcept {
671 for (int c = 0; c < ke; ++c) {
673 n, A, A_stride, piv, piv_stride, k0, tile, oot_count, c, ke, y, rhs_stride))
674 return false;
675 }
676 return true;
677 }
678
679 /* =====================================================================
680 RIGHT-LOOKING schedule.
681 ===================================================================== */
682
692 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
693 rl_trsm_right_one(const int n, T* TDLS_RESTRICT A, const int A_stride,
694 const int* TDLS_RESTRICT pk, const T* TDLS_RESTRICT tile, const int j0,
695 const int ke, const int je) noexcept {
696 T Akj[tile_size * tile_size];
697 load_tile(n, A, A_stride, pk, j0, Akj, ke, je);
698 ops_trsm_left_unit(tile, Akj, ke, je);
699 store_tile(n, A, A_stride, pk, j0, Akj, ke, je);
700 }
701
717 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
718 rl_trsm_right_push_one(const int n, T* TDLS_RESTRICT A, const int A_stride,
719 const int* TDLS_RESTRICT pk, const T* TDLS_RESTRICT tile, const int j0,
720 const int ke, const int je, const int k0, T* TDLS_RESTRICT y,
721 const int rhs_stride) noexcept {
722 T Akj[tile_size * tile_size];
723 load_tile(n, A, A_stride, pk, j0, Akj, ke, je);
724 ops_trsm_left_unit(tile, Akj, ke, je);
725 store_tile(n, A, A_stride, pk, j0, Akj, ke, je);
726 transposed_push_one(Akj, y, rhs_stride, 0, 1, k0, j0, ke, je);
727 }
728
741 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
742 rl_schur_one(const int n, T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT pk,
743 const int* TDLS_RESTRICT pi, const T* TDLS_RESTRICT Aik, const int j0,
744 const int ke, const int ie, const int je) noexcept {
745 T Aij[tile_size * tile_size];
746 load_tile(n, A, A_stride, pi, j0, Aij, ie, je);
747
748 for (int p = 0; p < ke; ++p) {
749 T Akj_row[tile_size];
750 for (int j = 0; j < je; ++j)
751 Akj_row[j] = TDLS_LUPP_DYN_A(pk[p], j0 + j);
752 for (int i = 0; i < ie; ++i) {
753 const T L_ip = Aik[i * tile_size + p];
754 for (int j = 0; j < je; ++j)
755 Aij[i * tile_size + j] -= L_ip * Akj_row[j];
756 }
757 }
758
759 store_tile(n, A, A_stride, pi, j0, Aij, ie, je);
760 }
761
778 template<bool fuse_rhs>
779 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
780 rl_update_row_one(const int n, T* TDLS_RESTRICT A, const int A_stride,
781 const int* TDLS_RESTRICT piv, const int piv_stride,
782 const int* TDLS_RESTRICT pk, const T* TDLS_RESTRICT tile, const int k,
783 const int i0, const int ke, const int ie, T* TDLS_RESTRICT y = nullptr,
784 const int rhs_stride = 1) noexcept {
785 const int k0 = k * tile_size;
786 const int nt = num_tiles(n);
787
788 int pi[tile_size];
789 for (int i = 0; i < ie; ++i)
790 pi[i] = TDLS_LUPP_DYN_PIV(i0 + i);
791
792 // TRSM down: Aik := Aik * U^-1
793 T Aik[tile_size * tile_size];
794 load_tile(n, A, A_stride, pi, k0, Aik, ie, ke);
795 ops_trsm_right(tile, Aik, ke, ie);
796 store_tile(n, A, A_stride, pi, k0, Aik, ie, ke);
797
798 // Fused forward substitution: push the solved y_k segment into
799 // this row block while its L panel sits in registers (row-major
800 // layout; the column-major one pushes from the U row panel).
801 if constexpr (fuse_rhs && !transposed) {
802 for (int r = 0; r < ie; ++r) {
803 T sum = T(0);
804 for (int j = 0; j < ke; ++j)
805 sum += Aik[r * tile_size + j] * TDLS_LUPP_DYN_Y(k0 + j);
806 TDLS_LUPP_DYN_Y(i0 + r) -= sum;
807 }
808 }
809
810 // Schur sweep over the trailing columns
811 for (int tj = k + 1; tj < nt; ++tj)
812 rl_schur_one(n, A, A_stride, pk, pi, Aik, tj * tile_size, ke, ie,
813 tile_size_at(tj * tile_size, n));
814 }
815
830 template<bool oot_diagnostics, bool fuse_rhs>
831 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
832 rl_step(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
833 const int piv_stride, const int k, int& oot_count, T* TDLS_RESTRICT y = nullptr,
834 const int rhs_stride = 1) noexcept {
835 const int k0 = k * tile_size;
836 const int nt = num_tiles(n);
837 const int ke = tile_size_at(k0, n);
838
839 T tile[tile_size * tile_size];
840 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
841
842 if (!factor_diag_tile<oot_diagnostics, fuse_rhs>(n, A, A_stride, piv, piv_stride, k0, tile,
843 oot_count, ke, y, rhs_stride))
844 return false;
845
846 // Physical rows of the tile after the swaps of this step
847 int pk[tile_size];
848 for (int i = 0; i < ke; ++i)
849 pk[i] = TDLS_LUPP_DYN_PIV(k0 + i);
850
851 store_tile(n, A, A_stride, pk, k0, tile, ke, ke);
852
853 // Fused forward substitution: this tile's y segment is final from
854 // here on: solve it while the tile is in registers, unit-lower
855 // (row-major layout) or with U^T (column-major layout).
856 if constexpr (fuse_rhs && transposed) {
857 transposed_solve_tile_upper(tile, y, rhs_stride, 0, 1, k0, ke);
858 } else if constexpr (fuse_rhs) {
859 for (int kk = 0; kk < ke; ++kk) {
860 for (int i = kk + 1; i < ke; ++i)
861 TDLS_LUPP_DYN_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_DYN_Y(k0 + kk);
862 }
863 }
864
865 // TRSM right over the row panel; under the column-major layout,
866 // solve_inplace pushes the solved z_k segment along
867 for (int tj = k + 1; tj < nt; ++tj) {
868 if constexpr (fuse_rhs && transposed)
869 rl_trsm_right_push_one(n, A, A_stride, pk, tile, tj * tile_size, ke,
870 tile_size_at(tj * tile_size, n), k0, y, rhs_stride);
871 else
872 rl_trsm_right_one(n, A, A_stride, pk, tile, tj * tile_size, ke,
873 tile_size_at(tj * tile_size, n));
874 }
875
876 // TRSM down + Schur over the rows below
877 for (int ti = k + 1; ti < nt; ++ti)
878 rl_update_row_one<fuse_rhs>(n, A, A_stride, piv, piv_stride, pk, tile, k,
879 ti * tile_size, ke, tile_size_at(ti * tile_size, n), y,
880 rhs_stride);
881 return true;
882 }
883
884 /* =====================================================================
885 LEFT-LOOKING schedule.
886 ===================================================================== */
887
902 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
903 ll_correct_tile(const int n, const T* TDLS_RESTRICT A, const int A_stride,
904 const int* TDLS_RESTRICT piv, const int piv_stride, const int row0,
905 const int col0, const int k0, T* TDLS_RESTRICT t, const int re,
906 const int ce) noexcept {
907 for (int bj0 = 0; bj0 < k0; bj0 += tile_size) {
908 T Lt[tile_size * tile_size];
909 load_tile_piv(n, A, A_stride, piv, piv_stride, row0, bj0, Lt, re, tile_size);
910 T Ut[tile_size * tile_size];
911 load_tile_piv(n, A, A_stride, piv, piv_stride, bj0, col0, Ut, tile_size, ce);
912 ops_gemm_sub(t, Lt, Ut, re, ce, tile_size);
913 }
914 }
915
931 template<bool fuse_rhs>
933 const int n, T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
934 const int piv_stride, const T* TDLS_RESTRICT tile, const int k0, const int i0, const int ke,
935 const int ie, T* TDLS_RESTRICT y = nullptr, const int rhs_stride = 1) noexcept {
936 T B[tile_size * tile_size];
937 load_tile_piv(n, A, A_stride, piv, piv_stride, i0, k0, B, ie, ke);
938 ll_correct_tile(n, A, A_stride, piv, piv_stride, i0, k0, k0, B, ie, ke);
939 ops_trsm_right(tile, B, ke, ie);
940 store_tile_piv(n, A, A_stride, piv, piv_stride, i0, k0, B, ie, ke);
941
942 // Fused forward substitution: B is the final L(i,k) panel, so push
943 // the solved y_k segment into this row block (row-major layout; the
944 // column-major one pushes from the U row panel).
945 if constexpr (fuse_rhs && !transposed) {
946 for (int r = 0; r < ie; ++r) {
947 T sum = T(0);
948 for (int j = 0; j < ke; ++j)
949 sum += B[r * tile_size + j] * TDLS_LUPP_DYN_Y(k0 + j);
950 TDLS_LUPP_DYN_Y(i0 + r) -= sum;
951 }
952 }
953 }
954
966 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
967 ll_update_right_one(const int n, T* TDLS_RESTRICT A, const int A_stride,
968 const int* TDLS_RESTRICT piv, const int piv_stride,
969 const T* TDLS_RESTRICT tile, const int k0, const int j0, const int ke,
970 const int je) noexcept {
971 T B[tile_size * tile_size];
972 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
973 ll_correct_tile(n, A, A_stride, piv, piv_stride, k0, j0, k0, B, ke, je);
974 ops_trsm_left_unit(tile, B, ke, je);
975 store_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
976 }
977
994 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
995 ll_update_right_push_one(const int n, T* TDLS_RESTRICT A, const int A_stride,
996 const int* TDLS_RESTRICT piv, const int piv_stride,
997 const T* TDLS_RESTRICT tile, const int k0, const int j0, const int ke,
998 const int je, T* TDLS_RESTRICT y, const int rhs_stride) noexcept {
999 T B[tile_size * tile_size];
1000 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
1001 ll_correct_tile(n, A, A_stride, piv, piv_stride, k0, j0, k0, B, ke, je);
1002 ops_trsm_left_unit(tile, B, ke, je);
1003 store_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
1004 transposed_push_one(B, y, rhs_stride, 0, 1, k0, j0, ke, je);
1005 }
1006
1021 template<bool oot_diagnostics, bool fuse_rhs>
1022 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1023 ll_step(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1024 const int piv_stride, const int k, int& oot_count, T* TDLS_RESTRICT y = nullptr,
1025 const int rhs_stride = 1) noexcept {
1026 const int k0 = k * tile_size;
1027 const int nt = num_tiles(n);
1028 const int ke = tile_size_at(k0, n);
1029
1030 T tile[tile_size * tile_size];
1031 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
1032 ll_correct_tile(n, A, A_stride, piv, piv_stride, k0, k0, k0, tile, ke, ke);
1033
1034 if (!factor_diag_tile<oot_diagnostics, fuse_rhs>(n, A, A_stride, piv, piv_stride, k0, tile,
1035 oot_count, ke, y, rhs_stride))
1036 return false;
1037
1038 store_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
1039
1040 // Fused forward substitution: this tile's y segment is final from
1041 // here on: solve it while the tile is in registers, unit-lower
1042 // (row-major layout) or with U^T (column-major layout).
1043 if constexpr (fuse_rhs && transposed) {
1044 transposed_solve_tile_upper(tile, y, rhs_stride, 0, 1, k0, ke);
1045 } else if constexpr (fuse_rhs) {
1046 for (int kk = 0; kk < ke; ++kk) {
1047 for (int i = kk + 1; i < ke; ++i)
1048 TDLS_LUPP_DYN_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_DYN_Y(k0 + kk);
1049 }
1050 }
1051
1052 // L panel below the diagonal
1053 for (int ti = k + 1; ti < nt; ++ti)
1054 ll_update_below_one<fuse_rhs>(n, A, A_stride, piv, piv_stride, tile, k0, ti * tile_size,
1055 ke, tile_size_at(ti * tile_size, n), y, rhs_stride);
1056
1057 // U row panel right of the diagonal; under the column-major layout,
1058 // solve_inplace pushes the solved z_k segment along
1059 for (int tj = k + 1; tj < nt; ++tj) {
1060 if constexpr (fuse_rhs && transposed)
1061 ll_update_right_push_one(n, A, A_stride, piv, piv_stride, tile, k0, tj * tile_size,
1062 ke, tile_size_at(tj * tile_size, n), y, rhs_stride);
1063 else
1064 ll_update_right_one(n, A, A_stride, piv, piv_stride, tile, k0, tj * tile_size, ke,
1065 tile_size_at(tj * tile_size, n));
1066 }
1067 return true;
1068 }
1069
1070 /* =====================================================================
1071 Physical row interchanges (Config.row_interchange == Physical): the
1072 whole matrix is one tile, the rows exchanged as in LAPACK; the
1073 operations of TiledLUppSolverStatic in the same order, so that both
1074 solvers agree bit for bit.
1075 ===================================================================== */
1076
1093 template<bool check>
1094 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
1095 physical_factorize_impl(const int n, T* TDLS_RESTRICT A, const int A_stride,
1096 int* TDLS_RESTRICT piv, const int piv_stride,
1097 [[maybe_unused]] const T u, int& weak) noexcept {
1098 bool ok = true;
1099 bool singular = false;
1100 weak = 0;
1101 for (int j = 0; j < n; ++j) {
1102 T best = detail::abs(TDLS_LUPP_DYN_A(j, j));
1103 int r = j;
1104 for (int i = j + 1; i < n; ++i) {
1105 const T a = detail::abs(TDLS_LUPP_DYN_A(i, j));
1106 if (a > best) {
1107 best = a;
1108 r = i;
1109 }
1110 }
1111 if constexpr (relative_pivot_threshold < T(1)) {
1112 const T in_place = detail::abs(TDLS_LUPP_DYN_A(j, j));
1113 if (in_place >= relative_pivot_threshold * best && in_place >= oot_pivot_threshold)
1114 r = j;
1115 }
1116 const bool pivot_ok = best >= singular_floor;
1117 singular = singular || !pivot_ok;
1118 weak += best < oot_pivot_threshold ? 1 : 0;
1119 TDLS_LUPP_DYN_PIV(j) = r;
1120 if (r != j) {
1121 for (int c = 0; c < n; ++c) {
1122 const T t = TDLS_LUPP_DYN_A(j, c);
1123 TDLS_LUPP_DYN_A(j, c) = TDLS_LUPP_DYN_A(r, c);
1124 TDLS_LUPP_DYN_A(r, c) = t;
1125 }
1126 }
1127 const T inv = pivot_ok ? T(1) / TDLS_LUPP_DYN_A(j, j) : T(0);
1128 TDLS_LUPP_DYN_A(j, j) = inv;
1129 for (int i = j + 1; i < n; ++i) {
1130 const T l = TDLS_LUPP_DYN_A(i, j) * inv;
1131 TDLS_LUPP_DYN_A(i, j) = l;
1132 if constexpr (check) ok = ok && u * detail::abs(l) <= T(1);
1133 for (int c = j + 1; c < n; ++c)
1134 TDLS_LUPP_DYN_A(i, c) -= l * TDLS_LUPP_DYN_A(j, c);
1135 }
1136 }
1137 return singular ? 2 : ok ? 0 : 1;
1138 }
1139
1149 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1150 physical_factorize(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1151 const int piv_stride, int& weak) noexcept {
1152 return physical_factorize_impl<false>(n, A, A_stride, piv, piv_stride, T(0), weak) != 2;
1153 }
1154
1166 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
1167 physical_factorize_checked(const int n, T* TDLS_RESTRICT A, const int A_stride,
1168 int* TDLS_RESTRICT piv, const int piv_stride, const T u) noexcept {
1169 int weak = 0;
1170 return physical_factorize_impl<true>(n, A, A_stride, piv, piv_stride, u, weak);
1171 }
1172
1183 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1184 physical_permute(const int n, const int w, const int* TDLS_RESTRICT piv, const int piv_stride,
1185 T* TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride,
1186 const bool reverse) noexcept {
1187 for (int t = 0; t < n; ++t) {
1188 const int j = reverse ? n - 1 - t : t;
1189 const int r = TDLS_LUPP_DYN_PIV(j);
1190 if (r != j)
1191 for (int c = 0; c < w; ++c) {
1192 const T v = TDLS_LUPP_DYN_XW(c, j);
1193 TDLS_LUPP_DYN_XW(c, j) = TDLS_LUPP_DYN_XW(c, r);
1194 TDLS_LUPP_DYN_XW(c, r) = v;
1195 }
1196 }
1197 }
1198
1211 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1212 physical_forward(const int n, const int w, const T* TDLS_RESTRICT A, const int A_stride,
1213 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1214 const int rhs_stride, const std::size_t xcol_stride) noexcept {
1215 if constexpr (!transposed) {
1216 physical_permute(n, w, piv, piv_stride, x, rhs_stride, xcol_stride, false);
1217 for (int j = 0; j < n; ++j)
1218 for (int i = j + 1; i < n; ++i) {
1219 const T l = TDLS_LUPP_DYN_A(i, j);
1220 for (int c = 0; c < w; ++c)
1221 TDLS_LUPP_DYN_XW(c, i) -= l * TDLS_LUPP_DYN_XW(c, j);
1222 }
1223 } else {
1224 for (int j = 0; j < n; ++j) {
1225 for (int c = 0; c < w; ++c)
1226 TDLS_LUPP_DYN_XW(c, j) *= TDLS_LUPP_DYN_A(j, j);
1227 for (int q = j + 1; q < n; ++q) {
1228 const T uq = TDLS_LUPP_DYN_A(j, q);
1229 for (int c = 0; c < w; ++c)
1230 TDLS_LUPP_DYN_XW(c, q) -= uq * TDLS_LUPP_DYN_XW(c, j);
1231 }
1232 }
1233 }
1234 }
1235
1248 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1249 physical_backward(const int n, const int w, const T* TDLS_RESTRICT A, const int A_stride,
1250 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1251 const int rhs_stride, const std::size_t xcol_stride) noexcept {
1252 if constexpr (!transposed) {
1253 for (int j = n - 1; j >= 0; --j)
1254 for (int c = 0; c < w; ++c) {
1255 T v = TDLS_LUPP_DYN_XW(c, j);
1256 for (int q = j + 1; q < n; ++q)
1257 v -= TDLS_LUPP_DYN_A(j, q) * TDLS_LUPP_DYN_XW(c, q);
1258 TDLS_LUPP_DYN_XW(c, j) = v * TDLS_LUPP_DYN_A(j, j);
1259 }
1260 } else {
1261 for (int j = n - 1; j >= 0; --j)
1262 for (int q = j + 1; q < n; ++q) {
1263 const T l = TDLS_LUPP_DYN_A(q, j);
1264 for (int c = 0; c < w; ++c)
1265 TDLS_LUPP_DYN_XW(c, j) -= l * TDLS_LUPP_DYN_XW(c, q);
1266 }
1267 physical_permute(n, w, piv, piv_stride, x, rhs_stride, xcol_stride, true);
1268 }
1269 }
1270
1284 template<int pass_width, int stages>
1285 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1286 physical_multirhs(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1287 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1288 const int rhs_stride, const std::size_t xcol_stride) noexcept {
1289 const int cpp = (pass_width <= 0 || pass_width >= nrhs) ? nrhs : pass_width;
1290 for (int c0 = 0; c0 < nrhs; c0 += cpp) {
1291 const int w = nrhs - c0 < cpp ? nrhs - c0 : cpp;
1292 T* TDLS_RESTRICT xp = x + TDLS_STRIDED_WIDE(c0, xcol_stride);
1293 if constexpr ((stages & 1) != 0)
1294 physical_forward(n, w, A, A_stride, piv, piv_stride, xp, rhs_stride, xcol_stride);
1295 if constexpr ((stages & 2) != 0)
1296 physical_backward(n, w, A, A_stride, piv, piv_stride, xp, rhs_stride, xcol_stride);
1297 }
1298 }
1299
1300 /* =====================================================================
1301 FACTORIZE - public entry point.
1302 ===================================================================== */
1303
1327 template<bool oot_diagnostics = true, bool fuse_rhs = false>
1328 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1329 factorize(int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1330 const int piv_stride, int& oot_count, T* TDLS_RESTRICT y = nullptr,
1331 const int rhs_stride = 1) noexcept {
1332 n = detail::runtime_dimension(n);
1333 if constexpr (physical) {
1334 int weak = 0;
1335 const bool ok = physical_factorize(n, A, A_stride, piv, piv_stride, weak);
1336 if constexpr (oot_diagnostics) oot_count = weak;
1337 if (!ok) return false;
1338 if constexpr (fuse_rhs)
1339 physical_forward(n, 1, A, A_stride, piv, piv_stride, y, rhs_stride, 0);
1340 return true;
1341 }
1342
1343 if constexpr (oot_diagnostics) oot_count = 0;
1344
1345 for (int i = 0; i < n; ++i)
1346 TDLS_LUPP_DYN_PIV(i) = i;
1347
1348 const int nt = num_tiles(n);
1349 for (int k = 0; k < nt; ++k) {
1350 if constexpr (schedule == Schedule::RightLooking) {
1351 if (!rl_step<oot_diagnostics, fuse_rhs>(n, A, A_stride, piv, piv_stride, k,
1352 oot_count, y, rhs_stride))
1353 return false;
1354 } else {
1355 if (!ll_step<oot_diagnostics, fuse_rhs>(n, A, A_stride, piv, piv_stride, k,
1356 oot_count, y, rhs_stride))
1357 return false;
1358 }
1359 }
1360
1361 return true;
1362 }
1363
1372 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1373 factorize(int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1374 const int piv_stride) noexcept {
1375 n = detail::runtime_dimension(n);
1376 int unused = 0;
1377 return factorize<false>(n, A, A_stride, piv, piv_stride, unused);
1378 }
1379
1380 /* =====================================================================
1381 SUBSTITUTION - schedule-independent.
1382 ===================================================================== */
1383
1399 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1400 fwd_push_one(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1401 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1402 const int rhs_stride, const std::size_t xcol_stride, const int k0, const int m0,
1403 const int ke, const int me) noexcept {
1404 T Lmk[tile_size * tile_size];
1405 load_tile_piv(n, A, A_stride, piv, piv_stride, m0, k0, Lmk, me, ke);
1406 for (int i = 0; i < me; ++i) {
1407 for (int w = 0; w < nrhs; ++w) {
1408 T sum = T(0);
1409 for (int j = 0; j < ke; ++j)
1410 sum += Lmk[i * tile_size + j] * TDLS_LUPP_DYN_XW(w, k0 + j);
1411 TDLS_LUPP_DYN_XW(w, m0 + i) -= sum;
1412 }
1413 }
1414 }
1415
1428 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1429 fwd_step(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1430 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1431 const int rhs_stride, const std::size_t xcol_stride, const int k) noexcept {
1432 const int k0 = k * tile_size;
1433 const int nt = num_tiles(n);
1434 const int ke = tile_size_at(k0, n);
1435
1436 T Lkk[tile_size * tile_size];
1437 load_tile_piv_lower(n, A, A_stride, piv, piv_stride, k0, k0, Lkk, ke);
1438
1439 // In-tile unit-lower solve
1440 for (int kk = 0; kk < ke; ++kk) {
1441 for (int i = kk + 1; i < ke; ++i) {
1442 for (int w = 0; w < nrhs; ++w)
1443 TDLS_LUPP_DYN_XW(w, k0 + i) -=
1444 Lkk[i * tile_size + kk] * TDLS_LUPP_DYN_XW(w, k0 + kk);
1445 }
1446 }
1447
1448 // Push into the tiles below
1449 for (int m = k + 1; m < nt; ++m)
1450 fwd_push_one(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k0,
1451 m * tile_size, ke, tile_size_at(m * tile_size, n));
1452 }
1453
1469 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1470 bwd_pull_one(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1471 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1472 const int rhs_stride, const std::size_t xcol_stride, const int k0, const int m0,
1473 const int ke, const int me) noexcept {
1474 T Ukm[tile_size * tile_size];
1475 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, m0, Ukm, ke, me);
1476 for (int i = 0; i < ke; ++i) {
1477 for (int w = 0; w < nrhs; ++w) {
1478 T sum = T(0);
1479 for (int j = 0; j < me; ++j)
1480 sum += Ukm[i * tile_size + j] * TDLS_LUPP_DYN_XW(w, m0 + j);
1481 TDLS_LUPP_DYN_XW(w, k0 + i) -= sum;
1482 }
1483 }
1484 }
1485
1498 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1499 bwd_step(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1500 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1501 const int rhs_stride, const std::size_t xcol_stride, const int k) noexcept {
1502 const int k0 = k * tile_size;
1503 const int nt = num_tiles(n);
1504 const int ke = tile_size_at(k0, n);
1505
1506 // Pull the trailing contributions
1507 for (int m = k + 1; m < nt; ++m)
1508 bwd_pull_one(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k0,
1509 m * tile_size, ke, tile_size_at(m * tile_size, n));
1510
1511 // In-tile upper solve
1512 T Ukk[tile_size * tile_size];
1513 load_tile_piv_upper(n, A, A_stride, piv, piv_stride, k0, k0, Ukk, ke);
1514
1515 for (int kk = ke - 1; kk >= 0; --kk) {
1516 for (int w = 0; w < nrhs; ++w)
1517 TDLS_LUPP_DYN_XW(w, k0 + kk) *= Ukk[kk * tile_size + kk]; // diag holds 1/pivot
1518 for (int i = 0; i < kk; ++i) {
1519 for (int w = 0; w < nrhs; ++w)
1520 TDLS_LUPP_DYN_XW(w, k0 + i) -=
1521 Ukk[i * tile_size + kk] * TDLS_LUPP_DYN_XW(w, k0 + kk);
1522 }
1523 }
1524 }
1525
1537 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1538 bwd_only(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1539 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1540 const int rhs_stride, const std::size_t xcol_stride) noexcept {
1541 for (int k = num_tiles(n) - 1; k >= 0; --k)
1542 bwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1543 }
1544
1558 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1559 fwd_bwd(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1560 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1561 const int rhs_stride, const std::size_t xcol_stride) noexcept {
1562 const int nt = num_tiles(n);
1563 for (int k = 0; k < nt; ++k)
1564 fwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1565
1566 bwd_only(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1567 }
1568
1584 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1585 permute_inplace(const int n, const int* TDLS_RESTRICT piv, const int piv_stride,
1586 T* TDLS_RESTRICT x, const int rhs_stride) noexcept {
1587 if (n <= 64) {
1588 // Bitmask cycle decomposition: one visited bit per row, the
1589 // whole state in a single 64-bit register.
1590 unsigned long long visited = 0ull;
1591 for (int s = 0; s < n; ++s) {
1592 if ((visited >> s) & 1ull) continue;
1593 const T tmp = TDLS_LUPP_DYN_X(s);
1594 int cur = s;
1595 int nxt = TDLS_LUPP_DYN_PIV(cur);
1596 while (nxt != s) {
1597 TDLS_LUPP_DYN_X(cur) = TDLS_LUPP_DYN_X(nxt);
1598 visited |= 1ull << cur;
1599 cur = nxt;
1600 nxt = TDLS_LUPP_DYN_PIV(cur);
1601 }
1602 TDLS_LUPP_DYN_X(cur) = tmp;
1603 visited |= 1ull << cur;
1604 }
1605 } else {
1606 // Cycle-leader scan (no visited storage): a cycle is rotated
1607 // only when reached from its smallest index, detected by
1608 // walking the orbit. Values and order match the mask path.
1609 for (int s = 0; s < n; ++s) {
1610 int probe = TDLS_LUPP_DYN_PIV(s);
1611 while (probe > s)
1612 probe = TDLS_LUPP_DYN_PIV(probe);
1613 if (probe != s) continue;
1614
1615 const T tmp = TDLS_LUPP_DYN_X(s);
1616 int cur = s;
1617 int nxt = TDLS_LUPP_DYN_PIV(cur);
1618 while (nxt != s) {
1619 TDLS_LUPP_DYN_X(cur) = TDLS_LUPP_DYN_X(nxt);
1620 cur = nxt;
1621 nxt = TDLS_LUPP_DYN_PIV(cur);
1622 }
1623 TDLS_LUPP_DYN_X(cur) = tmp;
1624 }
1625 }
1626 }
1627
1628 /* =====================================================================
1629 TRANSPOSED SUBSTITUTION - column-major layout, as in the
1630 compile-time solver: A = U^T L^T P, so A x = b is U^T z = b, then
1631 L^T w = z, then x = P^T w, with the same tile kernels for the fused
1632 forward pass of solve_inplace and for the separate passes.
1633 ===================================================================== */
1634
1644 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1645 transposed_solve_tile_upper(const T* TDLS_RESTRICT t, T* TDLS_RESTRICT x, const int rhs_stride,
1646 const std::size_t xcol_stride, const int nrhs, const int k0,
1647 const int ke) noexcept {
1648 for (int kk = 0; kk < ke; ++kk) {
1649 for (int w = 0; w < nrhs; ++w)
1650 TDLS_LUPP_DYN_XW(w, k0 + kk) *= t[kk * tile_size + kk]; // diag holds 1/pivot
1651 for (int i = kk + 1; i < ke; ++i) {
1652 for (int w = 0; w < nrhs; ++w)
1653 TDLS_LUPP_DYN_XW(w, k0 + i) -=
1654 t[kk * tile_size + i] * TDLS_LUPP_DYN_XW(w, k0 + kk);
1655 }
1656 }
1657 }
1658
1669 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1670 transposed_push_one(const T* TDLS_RESTRICT u, T* TDLS_RESTRICT x, const int rhs_stride,
1671 const std::size_t xcol_stride, const int nrhs, const int k0, const int m0,
1672 const int ke, const int me) noexcept {
1673 for (int i = 0; i < me; ++i) {
1674 for (int w = 0; w < nrhs; ++w) {
1675 T sum = T(0);
1676 for (int j = 0; j < ke; ++j)
1677 sum += u[j * tile_size + i] * TDLS_LUPP_DYN_XW(w, k0 + j);
1678 TDLS_LUPP_DYN_XW(w, m0 + i) -= sum;
1679 }
1680 }
1681 }
1682
1693 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1694 transposed_pull_one(const T* TDLS_RESTRICT l, T* TDLS_RESTRICT x, const int rhs_stride,
1695 const std::size_t xcol_stride, const int nrhs, const int k0, const int m0,
1696 const int ke, const int me) noexcept {
1697 for (int i = 0; i < ke; ++i) {
1698 for (int w = 0; w < nrhs; ++w) {
1699 T sum = T(0);
1700 for (int j = 0; j < me; ++j)
1701 sum += l[j * tile_size + i] * TDLS_LUPP_DYN_XW(w, m0 + j);
1702 TDLS_LUPP_DYN_XW(w, k0 + i) -= sum;
1703 }
1704 }
1705 }
1706
1715 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1717 const int rhs_stride, const std::size_t xcol_stride,
1718 const int nrhs, const int k0, const int ke) noexcept {
1719 for (int kk = ke - 1; kk > 0; --kk) {
1720 for (int i = 0; i < kk; ++i) {
1721 for (int w = 0; w < nrhs; ++w)
1722 TDLS_LUPP_DYN_XW(w, k0 + i) -=
1723 l[kk * tile_size + i] * TDLS_LUPP_DYN_XW(w, k0 + kk);
1724 }
1725 }
1726 }
1727
1740 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1741 transposed_fwd_step(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1742 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1743 const int rhs_stride, const std::size_t xcol_stride, const int k) noexcept {
1744 const int k0 = k * tile_size;
1745 const int nt = num_tiles(n);
1746 const int ke = tile_size_at(k0, n);
1747 T Ukk[tile_size * tile_size];
1748 load_tile_piv_upper(n, A, A_stride, piv, piv_stride, k0, k0, Ukk, ke);
1749 transposed_solve_tile_upper(Ukk, x, rhs_stride, xcol_stride, nrhs, k0, ke);
1750 for (int m = k + 1; m < nt; ++m) {
1751 const int me = tile_size_at(m * tile_size, n);
1752 T Ukm[tile_size * tile_size];
1753 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, m * tile_size, Ukm, ke, me);
1754 transposed_push_one(Ukm, x, rhs_stride, xcol_stride, nrhs, k0, m * tile_size, ke, me);
1755 }
1756 }
1757
1770 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1771 transposed_bwd_step(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1772 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1773 const int rhs_stride, const std::size_t xcol_stride, const int k) noexcept {
1774 const int k0 = k * tile_size;
1775 const int nt = num_tiles(n);
1776 const int ke = tile_size_at(k0, n);
1777 for (int m = k + 1; m < nt; ++m) {
1778 const int me = tile_size_at(m * tile_size, n);
1779 T Lmk[tile_size * tile_size];
1780 load_tile_piv(n, A, A_stride, piv, piv_stride, m * tile_size, k0, Lmk, me, ke);
1781 transposed_pull_one(Lmk, x, rhs_stride, xcol_stride, nrhs, k0, m * tile_size, ke, me);
1782 }
1783 T Lkk[tile_size * tile_size];
1784 load_tile_piv_lower(n, A, A_stride, piv, piv_stride, k0, k0, Lkk, ke);
1785 transposed_solve_tile_lower_unit(Lkk, x, rhs_stride, xcol_stride, nrhs, k0, ke);
1786 }
1787
1799 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1800 transposed_bwd_only(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1801 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1802 const int rhs_stride, const std::size_t xcol_stride) noexcept {
1803 for (int k = num_tiles(n) - 1; k >= 0; --k)
1804 transposed_bwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1805 k);
1806 }
1807
1820 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1821 transposed_fwd_bwd(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1822 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1823 const int rhs_stride, const std::size_t xcol_stride) noexcept {
1824 const int nt = num_tiles(n);
1825 for (int k = 0; k < nt; ++k)
1826 transposed_fwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride,
1827 k);
1828 transposed_bwd_only(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1829 }
1830
1840 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1841 unpermute_inplace(const int n, const int* TDLS_RESTRICT piv, const int piv_stride,
1842 T* TDLS_RESTRICT x, const int rhs_stride) noexcept {
1843 if (n <= 64) {
1844 unsigned long long visited = 0ull;
1845 for (int s = 0; s < n; ++s) {
1846 if ((visited >> s) & 1ull) continue;
1847 T carry = TDLS_LUPP_DYN_X(s);
1848 int dst = TDLS_LUPP_DYN_PIV(s);
1849 while (dst != s) {
1850 const T displaced = TDLS_LUPP_DYN_X(dst);
1851 TDLS_LUPP_DYN_X(dst) = carry;
1852 carry = displaced;
1853 visited |= 1ull << dst;
1854 dst = TDLS_LUPP_DYN_PIV(dst);
1855 }
1856 TDLS_LUPP_DYN_X(s) = carry;
1857 visited |= 1ull << s;
1858 }
1859 } else {
1860 for (int s = 0; s < n; ++s) {
1861 int probe = TDLS_LUPP_DYN_PIV(s);
1862 while (probe > s)
1863 probe = TDLS_LUPP_DYN_PIV(probe);
1864 if (probe != s) continue;
1865
1866 T carry = TDLS_LUPP_DYN_X(s);
1867 int dst = TDLS_LUPP_DYN_PIV(s);
1868 while (dst != s) {
1869 const T displaced = TDLS_LUPP_DYN_X(dst);
1870 TDLS_LUPP_DYN_X(dst) = carry;
1871 carry = displaced;
1872 dst = TDLS_LUPP_DYN_PIV(dst);
1873 }
1874 TDLS_LUPP_DYN_X(s) = carry;
1875 }
1876 }
1877 }
1878
1879 /* =====================================================================
1880 SUBSTITUTION - public entry points.
1881 ===================================================================== */
1882
1894 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1895 substitute(int n, const T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
1896 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
1897 const int rhs_stride) noexcept {
1898 n = detail::runtime_dimension(n);
1899 if constexpr (physical) {
1900 for (int i = 0; i < n; ++i)
1902 physical_multirhs<0, 3>(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
1903 return;
1904 }
1905 if constexpr (transposed) {
1906 // Column-major layout: b consumed unpermuted, x = P^T w at the end.
1907 for (int i = 0; i < n; ++i)
1909 transposed_fwd_bwd(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
1910 unpermute_inplace(n, piv, piv_stride, x, rhs_stride);
1911 } else {
1912 for (int i = 0; i < n; ++i)
1914 fwd_bwd(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
1915 }
1916 }
1917
1932 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1933 substitute_canonical(int n, const T* TDLS_RESTRICT A, const int A_stride,
1934 const int* TDLS_RESTRICT piv, const int piv_stride, const int col,
1935 T* TDLS_RESTRICT x, const int rhs_stride) noexcept {
1936 n = detail::runtime_dimension(n);
1937 if constexpr (physical) {
1938 for (int i = 0; i < n; ++i)
1939 TDLS_LUPP_DYN_X(i) = i == col ? T(1) : T(0);
1940 physical_multirhs<0, 3>(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
1941 return;
1942 }
1943 substitute_canonical_multirhs_pass(n, 1, A, A_stride, piv, piv_stride, col, x, rhs_stride,
1944 0);
1945 }
1946
1965 const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1966 const int* TDLS_RESTRICT piv, const int piv_stride, const int col0, T* TDLS_RESTRICT x,
1967 const int rhs_stride, const std::size_t xcol_stride) noexcept {
1968 if constexpr (transposed) {
1969 // Column-major layout: e_col unpermuted, x = P^T w at the end.
1970 for (int i = 0; i < n; ++i)
1971 for (int w = 0; w < nrhs; ++w)
1972 TDLS_LUPP_DYN_XW(w, i) = (i == col0 + w) ? T(1) : T(0);
1973 transposed_fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1974 for (int w = 0; w < nrhs; ++w)
1975 unpermute_inplace(n, piv, piv_stride, x + TDLS_STRIDED_WIDE(w, xcol_stride),
1976 rhs_stride);
1977 } else {
1978 for (int i = 0; i < n; ++i) {
1979 const int p = TDLS_LUPP_DYN_PIV(i);
1980 for (int w = 0; w < nrhs; ++w)
1981 TDLS_LUPP_DYN_XW(w, i) = (p == col0 + w) ? T(1) : T(0);
1982 }
1983 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1984 }
1985 }
1986
2011 template<int pass_width = 0>
2012 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2013 substitute_canonical_multirhs(int n, int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
2014 const int* TDLS_RESTRICT piv, const int piv_stride,
2015 const int col0, T* TDLS_RESTRICT x, const int rhs_stride,
2016 const std::size_t xcol_stride) noexcept {
2017 n = detail::runtime_dimension(n);
2018 nrhs = detail::runtime_dimension(nrhs);
2019 if constexpr (physical) {
2020 for (int c = 0; c < nrhs; ++c)
2021 for (int i = 0; i < n; ++i)
2022 TDLS_LUPP_DYN_XW(c, i) = i == col0 + c ? T(1) : T(0);
2023 physical_multirhs<pass_width, 3>(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride,
2024 xcol_stride);
2025 return;
2026 }
2027 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
2028 if constexpr (pass_width == 0) {
2029 substitute_canonical_multirhs_pass(n, nrhs, A, A_stride, piv, piv_stride, col0, x,
2030 rhs_stride, xcol_stride);
2031 } else {
2032 // A pass_width of nrhs or more runs the loop once, the whole
2033 // block being its single pass.
2034 for (int c0 = 0; c0 < nrhs; c0 += pass_width) {
2035 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
2036 const std::size_t off = TDLS_STRIDED_WIDE(c0, xcol_stride);
2037 substitute_canonical_multirhs_pass(n, cw, A, A_stride, piv, piv_stride, col0 + c0,
2038 x + off, rhs_stride, xcol_stride);
2039 }
2040 }
2041 }
2042
2060 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2061 substitute_multirhs_pass(const int n, const int nrhs, const T* TDLS_RESTRICT A,
2062 const int A_stride, const int* TDLS_RESTRICT piv, const int piv_stride,
2063 const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
2064 const std::size_t xcol_stride) noexcept {
2065 if constexpr (transposed) {
2066 // Column-major layout, as in substitute.
2067 for (int i = 0; i < n; ++i)
2068 for (int w = 0; w < nrhs; ++w)
2069 TDLS_LUPP_DYN_XW(w, i) = TDLS_LUPP_DYN_BW(w, i);
2070 transposed_fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2071 for (int w = 0; w < nrhs; ++w)
2072 unpermute_inplace(n, piv, piv_stride, x + TDLS_STRIDED_WIDE(w, xcol_stride),
2073 rhs_stride);
2074 } else {
2075 for (int i = 0; i < n; ++i) {
2076 const int p = TDLS_LUPP_DYN_PIV(i);
2077 for (int w = 0; w < nrhs; ++w)
2078 TDLS_LUPP_DYN_XW(w, i) = TDLS_LUPP_DYN_BW(w, p);
2079 }
2080 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2081 }
2082 }
2083
2112 template<int pass_width = 0>
2113 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2114 substitute_multirhs(int n, int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
2115 const int* TDLS_RESTRICT piv, const int piv_stride,
2116 const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
2117 const std::size_t xcol_stride) noexcept {
2118 n = detail::runtime_dimension(n);
2119 nrhs = detail::runtime_dimension(nrhs);
2120 if constexpr (physical) {
2121 for (int c = 0; c < nrhs; ++c)
2122 for (int i = 0; i < n; ++i)
2123 TDLS_LUPP_DYN_XW(c, i) = TDLS_LUPP_DYN_BW(c, i);
2124 physical_multirhs<pass_width, 3>(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride,
2125 xcol_stride);
2126 return;
2127 }
2128 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
2129 if constexpr (pass_width == 0) {
2130 substitute_multirhs_pass(n, nrhs, A, A_stride, piv, piv_stride, b, x, rhs_stride,
2131 xcol_stride);
2132 } else {
2133 // A pass_width of nrhs or more runs the loop once, the whole
2134 // block being its single pass.
2135 for (int c0 = 0; c0 < nrhs; c0 += pass_width) {
2136 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
2137 const std::size_t off = TDLS_STRIDED_WIDE(c0, xcol_stride);
2138 substitute_multirhs_pass(n, cw, A, A_stride, piv, piv_stride, b + off, x + off,
2139 rhs_stride, xcol_stride);
2140 }
2141 }
2142 }
2143
2158 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2159 substitute_inplace(int n, const T* TDLS_RESTRICT A, const int A_stride,
2160 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
2161 const int rhs_stride) noexcept {
2162 n = detail::runtime_dimension(n);
2163 if constexpr (physical) {
2164 physical_multirhs<0, 3>(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2165 return;
2166 }
2167 if constexpr (transposed) {
2168 // Column-major layout: the sweeps on the unpermuted x, then the
2169 // inverse cycle rotation (x = P^T w).
2170 transposed_fwd_bwd(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2171 unpermute_inplace(n, piv, piv_stride, x, rhs_stride);
2172 } else {
2173 permute_inplace(n, piv, piv_stride, x, rhs_stride);
2174 fwd_bwd(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
2175 }
2176 }
2177
2195 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2196 substitute_inplace_multirhs_pass(const int n, const int nrhs, const T* TDLS_RESTRICT A,
2197 const int A_stride, const int* TDLS_RESTRICT piv,
2198 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
2199 const std::size_t xcol_stride) noexcept {
2200 if constexpr (transposed) {
2201 // Column-major layout, as in substitute_inplace.
2202 transposed_fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2203 for (int w = 0; w < nrhs; ++w)
2204 unpermute_inplace(n, piv, piv_stride, x + TDLS_STRIDED_WIDE(w, xcol_stride),
2205 rhs_stride);
2206 } else {
2207 for (int w = 0; w < nrhs; ++w)
2208 permute_inplace(n, piv, piv_stride, x + TDLS_STRIDED_WIDE(w, xcol_stride),
2209 rhs_stride);
2210 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2211 }
2212 }
2213
2227 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2228 forward_inplace_multirhs_pass(const int n, const int nrhs, const T* TDLS_RESTRICT A,
2229 const int A_stride, const int* TDLS_RESTRICT piv,
2230 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
2231 const std::size_t xcol_stride) noexcept {
2232 const int nt = num_tiles(n);
2233 if constexpr (transposed) {
2234 for (int k = 0; k < nt; ++k)
2235 transposed_fwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride,
2236 xcol_stride, k);
2237 } else {
2238 for (int w = 0; w < nrhs; ++w)
2239 permute_inplace(n, piv, piv_stride, x + TDLS_STRIDED_WIDE(w, xcol_stride),
2240 rhs_stride);
2241 for (int k = 0; k < nt; ++k)
2242 fwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
2243 }
2244 }
2245
2259 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2260 backward_inplace_multirhs_pass(const int n, const int nrhs, const T* TDLS_RESTRICT A,
2261 const int A_stride, const int* TDLS_RESTRICT piv,
2262 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
2263 const std::size_t xcol_stride) noexcept {
2264 if constexpr (transposed) {
2265 transposed_bwd_only(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2266 for (int w = 0; w < nrhs; ++w)
2267 unpermute_inplace(n, piv, piv_stride, x + TDLS_STRIDED_WIDE(w, xcol_stride),
2268 rhs_stride);
2269 } else {
2270 bwd_only(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
2271 }
2272 }
2273
2299 template<int pass_width = 0>
2300 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2301 substitute_inplace_multirhs(int n, int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
2302 const int* TDLS_RESTRICT piv, const int piv_stride,
2303 T* TDLS_RESTRICT x, const int rhs_stride,
2304 const std::size_t xcol_stride) noexcept {
2305 n = detail::runtime_dimension(n);
2306 nrhs = detail::runtime_dimension(nrhs);
2307 if constexpr (physical) {
2308 physical_multirhs<pass_width, 3>(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride,
2309 xcol_stride);
2310 return;
2311 }
2312 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
2313 if constexpr (pass_width == 0) {
2314 substitute_inplace_multirhs_pass(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride,
2315 xcol_stride);
2316 } else {
2317 // A pass_width of nrhs or more runs the loop once, the whole
2318 // block being its single pass.
2319 for (int c0 = 0; c0 < nrhs; c0 += pass_width) {
2320 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
2321 const std::size_t off = TDLS_STRIDED_WIDE(c0, xcol_stride);
2322 substitute_inplace_multirhs_pass(n, cw, A, A_stride, piv, piv_stride, x + off,
2323 rhs_stride, xcol_stride);
2324 }
2325 }
2326 }
2327
2345 template<int pass_width = 0>
2346 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2347 forward_inplace_multirhs(int n, int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
2348 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
2349 const int rhs_stride, const std::size_t xcol_stride) noexcept {
2350 n = detail::runtime_dimension(n);
2351 nrhs = detail::runtime_dimension(nrhs);
2352 if constexpr (physical) {
2353 physical_multirhs<pass_width, 1>(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride,
2354 xcol_stride);
2355 return;
2356 }
2357 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
2358 if constexpr (pass_width == 0) {
2359 forward_inplace_multirhs_pass(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride,
2360 xcol_stride);
2361 } else {
2362 for (int c0 = 0; c0 < nrhs; c0 += pass_width) {
2363 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
2364 const std::size_t off = TDLS_STRIDED_WIDE(c0, xcol_stride);
2365 forward_inplace_multirhs_pass(n, cw, A, A_stride, piv, piv_stride, x + off,
2366 rhs_stride, xcol_stride);
2367 }
2368 }
2369 }
2370
2385 template<int pass_width = 0>
2386 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
2387 backward_inplace_multirhs(int n, int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
2388 const int* TDLS_RESTRICT piv, const int piv_stride,
2389 T* TDLS_RESTRICT x, const int rhs_stride,
2390 const std::size_t xcol_stride) noexcept {
2391 n = detail::runtime_dimension(n);
2392 nrhs = detail::runtime_dimension(nrhs);
2393 if constexpr (physical) {
2394 physical_multirhs<pass_width, 2>(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride,
2395 xcol_stride);
2396 return;
2397 }
2398 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
2399 if constexpr (pass_width == 0) {
2400 backward_inplace_multirhs_pass(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride,
2401 xcol_stride);
2402 } else {
2403 for (int c0 = 0; c0 < nrhs; c0 += pass_width) {
2404 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
2405 const std::size_t off = TDLS_STRIDED_WIDE(c0, xcol_stride);
2406 backward_inplace_multirhs_pass(n, cw, A, A_stride, piv, piv_stride, x + off,
2407 rhs_stride, xcol_stride);
2408 }
2409 }
2410 }
2411
2412 /* =====================================================================
2413 SOLVE - convenience wrappers.
2414 ===================================================================== */
2415
2431 template<bool oot_diagnostics = true>
2432 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
2433 solve(int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2434 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
2435 int& oot_count) noexcept {
2436 n = detail::runtime_dimension(n);
2437 if (!factorize<oot_diagnostics>(n, A, A_stride, piv, piv_stride, oot_count)) return false;
2438 substitute(n, A, A_stride, piv, piv_stride, b, x, rhs_stride);
2439 return true;
2440 }
2441
2454 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
2455 solve(int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2456 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
2457 const int rhs_stride) noexcept {
2458 n = detail::runtime_dimension(n);
2459 int unused = 0;
2460 return solve<false>(n, A, A_stride, piv, piv_stride, b, x, rhs_stride, unused);
2461 }
2462
2485 template<int pass_width = 0, bool oot_diagnostics = true>
2486 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
2487 solve_multirhs(int n, int nrhs, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2488 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
2489 const int rhs_stride, const std::size_t xcol_stride, int& oot_count) noexcept {
2490 n = detail::runtime_dimension(n);
2491 nrhs = detail::runtime_dimension(nrhs);
2492 if (!factorize<oot_diagnostics>(n, A, A_stride, piv, piv_stride, oot_count)) return false;
2493 substitute_multirhs<pass_width>(n, nrhs, A, A_stride, piv, piv_stride, b, x, rhs_stride,
2494 xcol_stride);
2495 return true;
2496 }
2497
2517 template<int pass_width = 0>
2518 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
2519 solve_multirhs(int n, int nrhs, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2520 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
2521 const int rhs_stride, const std::size_t xcol_stride) noexcept {
2522 n = detail::runtime_dimension(n);
2523 nrhs = detail::runtime_dimension(nrhs);
2524 int unused = 0;
2525 return solve_multirhs<pass_width, false>(n, nrhs, A, A_stride, piv, piv_stride, b, x,
2526 rhs_stride, xcol_stride, unused);
2527 }
2528
2556 template<bool oot_diagnostics = true>
2557 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
2558 solve_inplace(int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2559 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride,
2560 int& oot_count) noexcept {
2561 n = detail::runtime_dimension(n);
2562 if constexpr (physical) {
2563 int weak = 0;
2564 const bool ok = physical_factorize(n, A, A_stride, piv, piv_stride, weak);
2565 if constexpr (oot_diagnostics) oot_count = weak;
2566 if (!ok) return false;
2567 physical_multirhs<0, 3>(n, 1, A, A_stride, piv, piv_stride, y, rhs_stride, 0);
2568 return true;
2569 }
2570 if (!factorize<oot_diagnostics, true>(n, A, A_stride, piv, piv_stride, oot_count, y,
2571 rhs_stride))
2572 return false;
2573
2574 // Backward pass only: the forward one happened inside factorize.
2575 if constexpr (transposed) {
2576 // Column-major layout: L^T w = z, then x = P^T w.
2577 transposed_bwd_only(n, 1, A, A_stride, piv, piv_stride, y, rhs_stride, 0);
2578 unpermute_inplace(n, piv, piv_stride, y, rhs_stride);
2579 } else {
2580 bwd_only(n, 1, A, A_stride, piv, piv_stride, y, rhs_stride, 0);
2581 }
2582 return true;
2583 }
2584
2596 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
2597 solve_inplace(int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
2598 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride) noexcept {
2599 n = detail::runtime_dimension(n);
2600 int unused = 0;
2601 return solve_inplace<false>(n, A, A_stride, piv, piv_stride, y, rhs_stride, unused);
2602 }
2603
2630 template<int pass_width = 0, bool oot_diagnostics = true>
2631 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
2632 solve_inplace_multirhs(int n, int nrhs, T* TDLS_RESTRICT A, const int A_stride,
2633 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT y,
2634 const int rhs_stride, const std::size_t xcol_stride,
2635 int& oot_count) noexcept {
2636 n = detail::runtime_dimension(n);
2637 nrhs = detail::runtime_dimension(nrhs);
2638 if (!factorize<oot_diagnostics>(n, A, A_stride, piv, piv_stride, oot_count)) return false;
2639 substitute_inplace_multirhs<pass_width>(n, nrhs, A, A_stride, piv, piv_stride, y,
2640 rhs_stride, xcol_stride);
2641 return true;
2642 }
2643
2662 template<int pass_width = 0>
2663 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
2664 solve_inplace_multirhs(int n, int nrhs, T* TDLS_RESTRICT A, const int A_stride,
2665 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT y,
2666 const int rhs_stride, const std::size_t xcol_stride) noexcept {
2667 n = detail::runtime_dimension(n);
2668 nrhs = detail::runtime_dimension(nrhs);
2669 int unused = 0;
2670 return solve_inplace_multirhs<pass_width, false>(n, nrhs, A, A_stride, piv, piv_stride, y,
2671 rhs_stride, xcol_stride, unused);
2672 }
2673};
2674
2675
2676
2677#undef TDLS_LUPP_DYN_A
2678#undef TDLS_LUPP_DYN_PIV
2679#undef TDLS_LUPP_DYN_X
2680#undef TDLS_LUPP_DYN_B
2681#undef TDLS_LUPP_DYN_XW
2682#undef TDLS_LUPP_DYN_BW
2683#undef TDLS_LUPP_DYN_Y
2684
2685
2686
2687} // namespace tdls
2688
2689#if defined(TDLS_DETAIL_COMPILER_GCC)
2690#pragma GCC diagnostic pop
2691#endif
2692
2693
2694
2695#endif // TDLS_SOLVERS_TILED_LUPP_SOLVER_DYNAMIC_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_RESTRICT
Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC; nothing in the code that nvc++ gene...
Definition macros.hpp:174
#define TDLS_STRIDED_WIDE(i, s)
Offset of entry i of a strided object, i times stride s, for a 64-bit stride s (std::size_t); i is a ...
Definition macros.hpp:227
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:50
Scalar math helpers usable in constant expressions.
Schedule
Elimination schedule of a factorization.
Definition options.hpp:51
Runtime-size tiled dense LU factorization with logical partial pivoting and out-of-tile pivot recover...
Definition solver_dynamic.hpp:176
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool rl_step(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k, int &oot_count, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
RL: one full factorization step (diagonal tile + trailing updates).
Definition solver_dynamic.hpp:832
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_multirhs(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Physical substitution of nrhs columns in passes of pass_width columns: the first stage,...
Definition solver_dynamic.hpp:1286
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_tile_piv_lower(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, T *TDLS_RESTRICT t, const int re) noexcept
Triangular variant of the diagonal-tile load for the forward substitution (the backward one under the...
Definition solver_dynamic.hpp:437
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool factorize(int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, int &oot_count, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
Factor A := P*L*U in place.
Definition solver_dynamic.hpp:1329
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void permute_inplace(const int n, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
In-place application of the permutation to one strided column x, by cycle decomposition: the permute ...
Definition solver_dynamic.hpp:1585
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void bwd_pull_one(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int k0, const int m0, const int ke, const int me) noexcept
Backward pull: subtract U(k,m) * x_m from the x_k segment, for nrhs columns at once.
Definition solver_dynamic.hpp:1470
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void fwd_step(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int k) noexcept
Forward step: unit-lower solve the diagonal tile's segment, then push it into the tiles below.
Definition solver_dynamic.hpp:1429
static constexpr T relative_pivot_threshold
Relative threshold of the in-tile pivot choice, read once from the configuration (see TiledLUppConfig...
Definition solver_dynamic.hpp:187
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ll_correct_tile(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, const int k0, T *TDLS_RESTRICT t, const int re, const int ce) noexcept
LL: t (re x ce, rows row0.., cols col0..) -= sum over prior tiles bj < k0/tile_size of L(row0....
Definition solver_dynamic.hpp:903
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ll_update_right_push_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT tile, const int k0, const int j0, const int ke, const int je, T *TDLS_RESTRICT y, const int rhs_stride) noexcept
LL, column-major layout, solve_inplace: ll_update_right_one, then the fused U^T forward push....
Definition solver_dynamic.hpp:995
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ops_trsm_left_unit(const T *TDLS_RESTRICT lu, T *TDLS_RESTRICT B, const int kd, const int ce) noexcept
B := L^-1 B, with L the unit lower part of the factored diagonal tile. L is kd x kd,...
Definition solver_dynamic.hpp:281
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_multirhs(int n, int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Solve nrhs right-hand-side columns from a prior factorize: X := U^-1 L^-1 P B (column-major layout: X...
Definition solver_dynamic.hpp:2114
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_multirhs(int n, int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, int &oot_count) noexcept
factorize + substitute_multirhs in one call.
Definition solver_dynamic.hpp:2487
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int physical_factorize_checked(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T u) noexcept
Physical factorization with its multipliers bounded, the check fused into the elimination (the conden...
Definition solver_dynamic.hpp:1167
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace(int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride) noexcept
Diagnostics-free solve_inplace overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:2597
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_push_one(const T *TDLS_RESTRICT u, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int nrhs, const int k0, const int m0, const int ke, const int me) noexcept
Forward push with U^T: x_m segment -= U(k,m)^T * x_k.
Definition solver_dynamic.hpp:1670
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_inplace_multirhs_pass(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
One pass of the in-place multi right-hand-side substitution: nrhs columns of x permuted by the cycle ...
Definition solver_dynamic.hpp:2196
static constexpr Schedule schedule
elimination schedule (RightLooking or LeftLooking)
Definition solver_dynamic.hpp:179
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_pull_one(const T *TDLS_RESTRICT l, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int nrhs, const int k0, const int m0, const int ke, const int me) noexcept
Backward pull with L^T: x_k segment -= L(m,k)^T * x_m.
Definition solver_dynamic.hpp:1694
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_tile_piv_upper(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, T *TDLS_RESTRICT t, const int re) noexcept
Triangular variant of the diagonal-tile load for the backward substitution (the forward one under the...
Definition solver_dynamic.hpp:460
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void bwd_step(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int k) noexcept
Backward step: pull the trailing contributions, then upper-solve the diagonal tile's segment.
Definition solver_dynamic.hpp:1499
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_multirhs(int n, int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Diagnostics-free solve_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:2519
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_fwd_step(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int k) noexcept
Transposed forward step: U^T-solve the diagonal tile's segment, then push it into the following segme...
Definition solver_dynamic.hpp:1741
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace_multirhs(int n, int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride, const std::size_t xcol_stride) noexcept
Diagnostics-free solve_inplace_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:2664
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void unpermute_inplace(const int n, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
x := P^T x in place for one strided column (x[piv[i]] receives x[i]): the inverse rotation of permute...
Definition solver_dynamic.hpp:1841
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int num_tiles(const int n) noexcept
Number of tiles per dimension (last one possibly partial).
Definition solver_dynamic.hpp:220
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void forward_inplace_multirhs(int n, int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
First of the two stages of substitute_inplace_multirhs, in passes of pass_width columns: x := L^-1 P ...
Definition solver_dynamic.hpp:2347
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool factor_diag_tile(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k0, T *TDLS_RESTRICT tile, int &oot_count, const int ke, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
Factor the ke x ke diagonal tile in registers, with out-of-tile pivot recovery (drives the per-column...
Definition solver_dynamic.hpp:667
static constexpr bool transposed
Whether the solve uses the transpose of the factored matrix: S = A^T under the column-major layout (s...
Definition solver_dynamic.hpp:194
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool factorize(int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride) noexcept
Diagnostics-free factorize overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:1373
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void rl_trsm_right_push_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT pk, const T *TDLS_RESTRICT tile, const int j0, const int ke, const int je, const int k0, T *TDLS_RESTRICT y, const int rhs_stride) noexcept
RL, column-major layout, solve_inplace: rl_trsm_right_one, then the fused U^T forward push....
Definition solver_dynamic.hpp:718
static constexpr bool physical
Whether the rows are interchanged physically, the whole matrix a single tile (see TiledLUppConfig::ro...
Definition solver_dynamic.hpp:197
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void rl_trsm_right_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT pk, const T *TDLS_RESTRICT tile, const int j0, const int ke, const int je) noexcept
RL: one row-panel tile update, Akj := L^-1 Akj.
Definition solver_dynamic.hpp:693
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_forward(const int n, const int w, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
First stage of the physical substitution of w columns: x := L^-1 P x under the row-major layout,...
Definition solver_dynamic.hpp:1212
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_solve_tile_lower_unit(const T *TDLS_RESTRICT l, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int nrhs, const int k0, const int ke) noexcept
In-tile backward solve with L^T (unit diagonal).
Definition solver_dynamic.hpp:1716
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_fwd_bwd(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Transposed triangular sweeps on unpermuted column(s) x: U^T z = x, then L^T w = z,...
Definition solver_dynamic.hpp:1821
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ll_update_right_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT tile, const int k0, const int j0, const int ke, const int je) noexcept
LL: correct + TRSM one U-panel tile right of the diagonal.
Definition solver_dynamic.hpp:967
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void forward_inplace_multirhs_pass(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Forward stage of one pass of substitute_inplace_multirhs: the permutation and the sweep with L under ...
Definition solver_dynamic.hpp:2228
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace(int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride, int &oot_count) noexcept
Factorization with the forward substitution folded in.
Definition solver_dynamic.hpp:2558
static constexpr T oot_pivot_threshold
Acceptable-pivot threshold of the out-of-tile search, read once from the configuration (see TiledLUpp...
Definition solver_dynamic.hpp:184
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void bwd_only(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Backward pass alone, used by fwd_bwd and by solve_inplace (whose forward pass happens inside the fact...
Definition solver_dynamic.hpp:1538
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ops_eliminate_column(T *TDLS_RESTRICT t, const int k, const int re, const int ce) noexcept
Gaussian elimination of column k inside the diagonal tile.
Definition solver_dynamic.hpp:264
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void rl_update_row_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int *TDLS_RESTRICT pk, const T *TDLS_RESTRICT tile, const int k, const int i0, const int ke, const int ie, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
RL: TRSM down + Schur sweep of one row block below the diagonal, with the optional fused forward-subs...
Definition solver_dynamic.hpp:780
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_tile_piv(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, T *TDLS_RESTRICT t, const int re, const int ce) noexcept
Load a re x ce tile, reading the permutation inline (one read per row).
Definition solver_dynamic.hpp:391
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace_multirhs(int n, int nrhs, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride, const std::size_t xcol_stride, int &oot_count) noexcept
factorize + substitute_inplace_multirhs in one call: y holds the nrhs unpermuted right-hand-side colu...
Definition solver_dynamic.hpp:2632
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_canonical_multirhs(int n, int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int col0, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
nrhs canonical columns e_col0 .. e_{col0+nrhs-1} solved from a prior factorize: the columns of the in...
Definition solver_dynamic.hpp:2013
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_tile(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT prow, const int col0, T *TDLS_RESTRICT t, const int re, const int ce) noexcept
Load a re x ce tile through a cached physical-row segment.
Definition solver_dynamic.hpp:351
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void rl_schur_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT pk, const int *TDLS_RESTRICT pi, const T *TDLS_RESTRICT Aik, const int j0, const int ke, const int ie, const int je) noexcept
RL: one Schur-complement tile update, Aij -= Aik * Akj, streaming the factored Akj row by row from re...
Definition solver_dynamic.hpp:742
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void fwd_push_one(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int k0, const int m0, const int ke, const int me) noexcept
Forward push: subtract L(m,k) * x_k from the x_m segment, for nrhs columns at once.
Definition solver_dynamic.hpp:1400
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_permute(const int n, const int w, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const bool reverse) noexcept
Applies the interchanges to w columns of x, in their order (P x) or in reverse order (P^T x).
Definition solver_dynamic.hpp:1184
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void store_tile_piv(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int row0, const int col0, const T *TDLS_RESTRICT t, const int re, const int ce) noexcept
Store a re x ce tile, reading the permutation inline (one read per row).
Definition solver_dynamic.hpp:414
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ops_trsm_right(const T *TDLS_RESTRICT lu, T *TDLS_RESTRICT B, const int kd, const int re) noexcept
B := B U^-1, with U the upper part of the factored diagonal tile. U is kd x kd, B is re x kd.
Definition solver_dynamic.hpp:302
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_backward(const int n, const int w, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Second stage of the physical substitution of w columns: x := U^-1 x under the row-major layout,...
Definition solver_dynamic.hpp:1249
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool physical_factorize(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, int &weak) noexcept
Physical factorization (see physical_factorize_impl).
Definition solver_dynamic.hpp:1150
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void backward_inplace_multirhs_pass(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Backward stage of one pass of substitute_inplace_multirhs: the sweep with U under the row-major layou...
Definition solver_dynamic.hpp:2260
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_canonical(int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int col, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Solve with b = e_col generated on the fly: the consistent-tangent-operator path.
Definition solver_dynamic.hpp:1933
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ops_gemm_sub(T *TDLS_RESTRICT Ct, const T *TDLS_RESTRICT At, const T *TDLS_RESTRICT Bt, const int re, const int ce, const int kd) noexcept
Ct -= At*Bt with per-element dot-product accumulation. At is re x kd, Bt is kd x ce,...
Definition solver_dynamic.hpp:325
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_solve_tile_upper(const T *TDLS_RESTRICT t, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int nrhs, const int k0, const int ke) noexcept
In-tile forward solve with U^T (the diagonal of the tile holds the pivot reciprocals).
Definition solver_dynamic.hpp:1645
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_bwd_only(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Transposed backward pass alone, used by transposed_fwd_bwd and by solve_inplace (whose forward pass h...
Definition solver_dynamic.hpp:1800
static constexpr T singular_floor
Singularity floor of the out-of-tile recovery, read once from the configuration (see TiledLUppConfig:...
Definition solver_dynamic.hpp:190
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_inplace_multirhs(int n, int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
nrhs columns solved in place: x holds the nrhs unpermuted right-hand-side columns on entry and the nr...
Definition solver_dynamic.hpp:2301
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void fwd_bwd(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Triangular solves on already-permuted column(s) x. nrhs columns are processed per tile visit,...
Definition solver_dynamic.hpp:1559
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute(int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Solve x := U^-1 L^-1 P b from a prior factorize (column-major layout: x := P^T L^-T U^-T b)....
Definition solver_dynamic.hpp:1895
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_canonical_multirhs_pass(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int col0, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
One pass of the canonical multi right-hand-side substitution: nrhs canonical columns e_col0 ....
Definition solver_dynamic.hpp:1964
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool factor_diag_column(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k0, T *TDLS_RESTRICT tile, int &oot_count, const int c, const int ke, T *TDLS_RESTRICT y, const int rhs_stride) noexcept
One column step of the diagonal-tile factorization: pivot search (in-tile, then out-of-tile recovery)...
Definition solver_dynamic.hpp:497
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_multirhs_pass(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
One pass of the multi right-hand-side substitution: nrhs columns of b gathered in permuted order into...
Definition solver_dynamic.hpp:2061
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ops_swap_rows(T *TDLS_RESTRICT t, const int k, const int r, const int ke) noexcept
Row swap k <-> r inside the tile (direct, first ke columns).
Definition solver_dynamic.hpp:246
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void store_tile(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT prow, const int col0, const T *TDLS_RESTRICT t, const int re, const int ce) noexcept
Store a re x ce tile through a cached physical-row segment.
Definition solver_dynamic.hpp:370
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve(int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Diagnostics-free solve overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:2455
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void backward_inplace_multirhs(int n, int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride) noexcept
Second of the two stages of substitute_inplace_multirhs, in passes of pass_width columns: x := U^-1 x...
Definition solver_dynamic.hpp:2387
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_bwd_step(const int n, const int nrhs, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, const std::size_t xcol_stride, const int k) noexcept
Transposed backward step: pull the trailing contributions with L^T, then L^T-solve the diagonal tile'...
Definition solver_dynamic.hpp:1771
static constexpr int tile_size
tile size (int)
Definition solver_dynamic.hpp:178
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve(int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, int &oot_count) noexcept
factorize + substitute in one call.
Definition solver_dynamic.hpp:2433
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void ll_update_below_one(const int n, T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT tile, const int k0, const int i0, const int ke, const int ie, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
LL: correct + TRSM one L-panel tile below the diagonal, with the optional fused forward-substitution ...
Definition solver_dynamic.hpp:932
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool ll_step(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const int k, int &oot_count, T *TDLS_RESTRICT y=nullptr, const int rhs_stride=1) noexcept
LL: one full factorization step (correct + factor the diagonal tile, then its L and U panels).
Definition solver_dynamic.hpp:1023
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_inplace(int n, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride) noexcept
Solve in place: x already holds the unpermuted RHS on entry and the solution on exit.
Definition solver_dynamic.hpp:2159
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int physical_factorize_impl(const int n, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T u, int &weak) noexcept
Physical factorization: P S = L U in place, L below the diagonal and U above it in the pivoted row or...
Definition solver_dynamic.hpp:1095
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int tile_size_at(const int t0, const int n) noexcept
Extent of the tile starting at row/column t0: tile_size, or less for the last tile.
Definition solver_dynamic.hpp:229
Detection of the compiler, of the language and of the side of the compilation pass: the one place whe...
Compile-time configuration of the TiledLUpp solver family.
#define TDLS_LUPP_DYN_X(i)
Strided entry i of the solution vector.
Definition solver_dynamic.hpp:117
#define TDLS_LUPP_DYN_PIV(i)
Strided pivot entry i.
Definition solver_dynamic.hpp:114
#define TDLS_LUPP_DYN_Y(i)
Strided entry i of the fused right-hand side of solve_inplace.
Definition solver_dynamic.hpp:130
#define TDLS_LUPP_DYN_A(r, c)
Strided element (r, c) of S, the stored matrix read row-major (A or A^T, see the file documentation).
Definition solver_dynamic.hpp:111
#define TDLS_LUPP_DYN_XW(w, i)
Strided entry i of column w of a multi right-hand-side block.
Definition solver_dynamic.hpp:123
#define TDLS_LUPP_DYN_B(i)
Strided entry i of the right-hand side.
Definition solver_dynamic.hpp:120
#define TDLS_LUPP_DYN_BW(w, i)
Strided entry i of column w of a multi right-hand-side block of b, addressed exactly as TDLS_LUPP_DYN...
Definition solver_dynamic.hpp:127