TDLS 0.2.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
53
54
55
56#include <tdls/core/math.hpp>
57#include <type_traits>
58
60
61
62
63// gcc's flow analysis cannot prove that the replay buffer of the
64// out-of-tile search is written before being read (the loop structure
65// guarantees it, a property validated bitwise against the static
66// solver), and emits spurious -Wmaybe-uninitialized warnings at -O2.
67//
68// Under UBSan instrumentation (-fsanitize=undefined) at -O2 and
69// above, gcc additionally loses the value ranges of the loop indices
70// and reports impossible subscripts through -Warray-bounds: the
71// unsigned cast of a provably non-negative index read as a wrapped
72// negative, or a loop counter assumed below its own loop guard.
73// Observed from gcc 12 to gcc 16; plain builds and clang are clean,
74// and the exercised paths are certified free of out-of-bounds
75// accesses by the constexpr suite.
76//
77// Both suppressions are scoped to this header and to those warnings
78// only, exactly as in solver_static.hpp.
79#if defined(__GNUC__) && !defined(__clang__)
80#pragma GCC diagnostic push
81#pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
82#pragma GCC diagnostic ignored "-Warray-bounds"
83#endif
84
85namespace tdls {
86
87
88
89/* Addressing macros. They expand inside member functions where n, the
90 parameter names (A, A_stride, piv, piv_stride, x, b, y, rhs_stride,
91 xcol_stride) and the Config value are in scope. #undef'd at the end
92 of this header. */
93
97#define TDLS_LUPP_DYN_A(r, c) A[TDLS_LAYOUT_INDEX(r, c, n) * unsigned(A_stride)]
100#define TDLS_LUPP_DYN_PIV(i) piv[unsigned(i) * unsigned(piv_stride)]
103#define TDLS_LUPP_DYN_X(i) x[unsigned(i) * unsigned(rhs_stride)]
106#define TDLS_LUPP_DYN_B(i) b[unsigned(i) * unsigned(rhs_stride)]
109#define TDLS_LUPP_DYN_XW(w, i) \
110 x[unsigned(i) * unsigned(rhs_stride) + unsigned(w) * unsigned(xcol_stride)]
114#define TDLS_LUPP_DYN_BW(w, i) \
115 b[unsigned(i) * unsigned(rhs_stride) + unsigned(w) * unsigned(xcol_stride)]
118#define TDLS_LUPP_DYN_Y(i) y[unsigned(i) * unsigned(rhs_stride)]
119
120
121
158template<typename T, TiledLUppConfig<T> Config = TiledLUppConfig<T>{}>
160
161 static constexpr int tile_size = Config.tile_size;
162 static constexpr Schedule schedule =
163 Config.schedule;
164
167 static constexpr T oot_threshold = Config.oot_threshold;
170 static constexpr T singular_floor = Config.singular_floor;
171
172 static_assert(
173 Config.oot_threshold.is_finite() && Config.singular_floor.is_finite(),
174 "TiledLUppSolverDynamic: oot_threshold and singular_floor must be finite (and fit a "
175 "63-bit mantissa)");
176 static_assert(singular_floor > T(0), "TiledLUppSolverDynamic: singular_floor must be positive");
177 static_assert(singular_floor <= oot_threshold,
178 "TiledLUppSolverDynamic: singular_floor must not exceed oot_threshold (the "
179 "floor applies to the out-of-tile recovery path)");
180 static_assert(tile_size >= 1, "TiledLUppSolverDynamic: tile size must be >= 1");
181
184 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int num_tiles(const int n) noexcept {
185 return (n + tile_size - 1) / tile_size;
186 }
187
192 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr int tile_size_at(const int t0,
193 const int n) noexcept {
194 return (n - t0 < tile_size) ? (n - t0) : tile_size;
195 }
196
197 /* =====================================================================
198 Register-tile micro-kernels, runtime extents. A tile is a
199 tile_size x tile_size array with row stride tile_size; the runtime
200 extents bound the active part.
201 ===================================================================== */
202
208 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
209 ops_swap_rows(T* TDLS_RESTRICT t, const int k, const int r, const int ke) noexcept {
210 for (int j = 0; j < ke; ++j) {
211 const T tmp = t[k * tile_size + j];
212 t[k * tile_size + j] = t[r * tile_size + j];
213 t[r * tile_size + j] = tmp;
214 }
215 }
216
226 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
227 ops_eliminate_column(T* TDLS_RESTRICT t, const int k, const int re, const int ce) noexcept {
228 const T inv_pivot = T(1) / t[k * tile_size + k];
229 t[k * tile_size + k] = inv_pivot;
230 for (int i = k + 1; i < re; ++i) {
231 t[i * tile_size + k] *= inv_pivot;
232 for (int j = k + 1; j < ce; ++j)
233 t[i * tile_size + j] -= t[i * tile_size + k] * t[k * tile_size + j];
234 }
235 }
236
243 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
244 ops_trsm_left_unit(const T* TDLS_RESTRICT lu, T* TDLS_RESTRICT B, const int kd,
245 const int ce) noexcept {
246 for (int k = 0; k < kd; ++k) {
247 for (int i = k + 1; i < kd; ++i) {
248 const T L_ik = lu[i * tile_size + k];
249 for (int j = 0; j < ce; ++j)
250 B[i * tile_size + j] -= L_ik * B[k * tile_size + j];
251 }
252 }
253 }
254
264 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
265 ops_trsm_right(const T* TDLS_RESTRICT lu, T* TDLS_RESTRICT B, const int kd,
266 const int re) noexcept {
267 for (int k = 0; k < kd; ++k) {
268 const T U_kk_inv = lu[k * tile_size + k];
269 for (int i = 0; i < re; ++i)
270 B[i * tile_size + k] *= U_kk_inv;
271 for (int j = k + 1; j < kd; ++j) {
272 const T U_kj = lu[k * tile_size + j];
273 for (int i = 0; i < re; ++i)
274 B[i * tile_size + j] -= B[i * tile_size + k] * U_kj;
275 }
276 }
277 }
278
287 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
288 ops_gemm_sub(T* TDLS_RESTRICT Ct, const T* TDLS_RESTRICT At, const T* TDLS_RESTRICT Bt,
289 const int re, const int ce, const int kd) noexcept {
290 for (int i = 0; i < re; ++i) {
291 for (int j = 0; j < ce; ++j) {
292 T sum = T(0);
293 for (int k = 0; k < kd; ++k)
294 sum += At[i * tile_size + k] * Bt[k * tile_size + j];
295 Ct[i * tile_size + j] -= sum;
296 }
297 }
298 }
299
300 /* =====================================================================
301 Remote <-> register tile movement, runtime extents.
302 ===================================================================== */
303
313 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
314 load_tile(const int n, const T* TDLS_RESTRICT A, const int A_stride,
315 const int* TDLS_RESTRICT prow, const int col0, T* TDLS_RESTRICT t, const int re,
316 const int ce) noexcept {
317 for (int i = 0; i < re; ++i) {
318 for (int j = 0; j < ce; ++j)
319 t[i * tile_size + j] = TDLS_LUPP_DYN_A(prow[i], col0 + j);
320 }
321 }
322
332 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
333 store_tile(const int n, T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT prow,
334 const int col0, const T* TDLS_RESTRICT t, const int re, const int ce) noexcept {
335 for (int i = 0; i < re; ++i) {
336 for (int j = 0; j < ce; ++j)
337 TDLS_LUPP_DYN_A(prow[i], col0 + j) = t[i * tile_size + j];
338 }
339 }
340
353 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
354 load_tile_piv(const int n, const T* TDLS_RESTRICT A, const int A_stride,
355 const int* TDLS_RESTRICT piv, const int piv_stride, const int row0,
356 const int col0, T* TDLS_RESTRICT t, const int re, const int ce) noexcept {
357 for (int i = 0; i < re; ++i) {
358 const int phys = TDLS_LUPP_DYN_PIV(row0 + i);
359 for (int j = 0; j < ce; ++j)
360 t[i * tile_size + j] = TDLS_LUPP_DYN_A(phys, col0 + j);
361 }
362 }
363
376 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
377 store_tile_piv(const int n, T* TDLS_RESTRICT A, const int A_stride,
378 const int* TDLS_RESTRICT piv, const int piv_stride, const int row0,
379 const int col0, const T* TDLS_RESTRICT t, const int re, const int ce) noexcept {
380 for (int i = 0; i < re; ++i) {
381 const int phys = TDLS_LUPP_DYN_PIV(row0 + i);
382 for (int j = 0; j < ce; ++j)
383 TDLS_LUPP_DYN_A(phys, col0 + j) = t[i * tile_size + j];
384 }
385 }
386
398 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
399 load_tile_piv_lower(const int n, const T* TDLS_RESTRICT A, const int A_stride,
400 const int* TDLS_RESTRICT piv, const int piv_stride, const int row0,
401 const int col0, T* TDLS_RESTRICT t, const int re) noexcept {
402 for (int i = 1; i < re; ++i) {
403 const int phys = TDLS_LUPP_DYN_PIV(row0 + i);
404 for (int j = 0; j < i; ++j)
405 t[i * tile_size + j] = TDLS_LUPP_DYN_A(phys, col0 + j);
406 }
407 }
408
421 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
422 load_tile_piv_upper(const int n, const T* TDLS_RESTRICT A, const int A_stride,
423 const int* TDLS_RESTRICT piv, const int piv_stride, const int row0,
424 const int col0, T* TDLS_RESTRICT t, const int re) noexcept {
425 for (int i = 0; i < re; ++i) {
426 const int phys = TDLS_LUPP_DYN_PIV(row0 + i);
427 for (int j = i; j < re; ++j)
428 t[i * tile_size + j] = TDLS_LUPP_DYN_A(phys, col0 + j);
429 }
430 }
431
432 /* =====================================================================
433 Diagonal-tile factorization with out-of-tile pivoting. Same
434 algorithm as the static TiledLUpp solver, runtime column/extent parameters,
435 direct (unpredicated) fused-RHS swap.
436 ===================================================================== */
437
456 template<bool oot_diagnostics, bool fuse_rhs>
457 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
458 factor_diag_column(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
459 const int piv_stride, const int k0, T* TDLS_RESTRICT tile, int& oot_count,
460 const int c, const int ke, T* TDLS_RESTRICT y,
461 const int rhs_stride) noexcept {
462
463 const int gc = k0 + c; // global column
464
465 // In-tile pivot search (rows c..ke of the register tile)
466 int best_r = c;
467 T best = detail::abs(tile[c * tile_size + c]);
468 for (int r = c + 1; r < ke; ++r) {
469 const T v = detail::abs(tile[r * tile_size + c]);
470 if (v > best) {
471 best = v;
472 best_r = r;
473 }
474 }
475
476 int piv_row; // winning global (logical) row
477
478 if (best >= oot_threshold) {
479 piv_row = k0 + best_r;
480 } else if (ke < tile_size) {
481 // Trailing tile: no rows below to recover from. Diagnostic
482 // order: singularity verdict first, then count the weak pivot
483 // (full tiles count before the verdict).
484 if (best < singular_floor) return false;
485 if constexpr (oot_diagnostics) ++oot_count;
486 piv_row = k0 + best_r;
487 } else {
488 // Out-of-tile recovery: scan the rows below the tile and
489 // evaluate each candidate as if it had received the
490 // eliminations it is missing, keeping the best (or, with
491 // Config.oot_first_acceptable, the first to reach the threshold).
492 if constexpr (oot_diagnostics) ++oot_count;
493 T gbest = best;
494 int gbest_row = k0 + best_r;
495
496 for (int row = k0 + ke; row < n; ++row) {
497 const int phys = TDLS_LUPP_DYN_PIV(row);
498
499 T corrected = TDLS_LUPP_DYN_A(phys, gc);
500 if constexpr (schedule == Schedule::LeftLooking) {
501 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
502 for (int p = 0; p < tile_size; ++p)
503 corrected -= TDLS_LUPP_DYN_A(phys, bj0 + p) *
505 }
506
507 if (c > 0) {
508 T L_row[tile_size];
509 for (int t = 0; t < c; ++t) {
510 T a_t = TDLS_LUPP_DYN_A(phys, k0 + t);
511 if constexpr (schedule == Schedule::LeftLooking) {
512 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
513 for (int p = 0; p < tile_size; ++p)
514 a_t -= TDLS_LUPP_DYN_A(phys, bj0 + p) *
515 TDLS_LUPP_DYN_A(TDLS_LUPP_DYN_PIV(bj0 + p), k0 + t);
516 }
517 for (int p = 0; p < t; ++p)
518 a_t -= L_row[p] * tile[p * tile_size + t];
519 L_row[t] = a_t * tile[t * tile_size + t]; // diag holds 1/pivot
520 corrected -= L_row[t] * tile[t * tile_size + c];
521 }
522 }
523
524 const T v = detail::abs(corrected);
525 if (v > gbest) {
526 gbest = v;
527 gbest_row = row;
528 }
529
530 // First-acceptable out-of-tile pivot: a candidate that
531 // reaches the threshold already beats the sub-threshold
532 // in-tile pivot, so stop scanning.
533 if constexpr (Config.oot_first_acceptable)
534 if (v >= oot_threshold) break;
535 }
536
537 if (gbest < singular_floor) return false;
538 piv_row = gbest_row;
539 }
540
541 {
542 // Swap the permutation entries unconditionally (self-swap is a
543 // bitwise no-op, and the branch would diverge on GPU).
544 const int tmp = TDLS_LUPP_DYN_PIV(gc);
546 TDLS_LUPP_DYN_PIV(piv_row) = tmp;
547
548 if constexpr (fuse_rhs) {
549 // The fused RHS follows the rows through pivoting.
550 const T ty = TDLS_LUPP_DYN_Y(gc);
551 TDLS_LUPP_DYN_Y(gc) = TDLS_LUPP_DYN_Y(piv_row);
552 TDLS_LUPP_DYN_Y(piv_row) = ty;
553 }
554
555 if (piv_row < k0 + ke) {
556 ops_swap_rows(tile, c, piv_row - k0, ke);
557 } else {
558 // Cross-tile swap: pull the new row into the tile and
559 // replay everything it missed: prior tiles (LL), then
560 // the current tile's factored columns, on the FULL row.
561 const int phys = TDLS_LUPP_DYN_PIV(gc);
562 for (int j = 0; j < ke; ++j) {
563 tile[c * tile_size + j] = TDLS_LUPP_DYN_A(phys, k0 + j);
564 if constexpr (schedule == Schedule::LeftLooking) {
565 for (int bj0 = 0; bj0 < k0; bj0 += tile_size)
566 for (int p = 0; p < tile_size; ++p)
567 tile[c * tile_size + j] -=
568 TDLS_LUPP_DYN_A(phys, bj0 + p) *
569 TDLS_LUPP_DYN_A(TDLS_LUPP_DYN_PIV(bj0 + p), k0 + j);
570 }
571 }
572
573 if (c > 0) {
574 T L_row[tile_size];
575 for (int t = 0; t < c; ++t) {
576 T a_t = tile[c * tile_size + t];
577 for (int p = 0; p < t; ++p)
578 a_t -= L_row[p] * tile[p * tile_size + t];
579 L_row[t] = a_t * tile[t * tile_size + t]; // diag holds 1/pivot
580 tile[c * tile_size + t] = L_row[t];
581 }
582 for (int j = c; j < ke; ++j) {
583 for (int t = 0; t < c; ++t)
584 tile[c * tile_size + j] -= L_row[t] * tile[t * tile_size + j];
585 }
586 }
587 }
588 }
589
590 ops_eliminate_column(tile, c, ke, ke);
591 return true;
592 }
593
612 template<bool oot_diagnostics, bool fuse_rhs>
613 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
614 factor_diag_tile(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
615 const int piv_stride, const int k0, T* TDLS_RESTRICT tile, int& oot_count,
616 const int ke, T* TDLS_RESTRICT y = nullptr,
617 const int rhs_stride = 1) noexcept {
618 for (int c = 0; c < ke; ++c) {
620 n, A, A_stride, piv, piv_stride, k0, tile, oot_count, c, ke, y, rhs_stride))
621 return false;
622 }
623 return true;
624 }
625
626 /* =====================================================================
627 RIGHT-LOOKING schedule.
628 ===================================================================== */
629
639 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
640 rl_trsm_right_one(const int n, T* TDLS_RESTRICT A, const int A_stride,
641 const int* TDLS_RESTRICT pk, const T* TDLS_RESTRICT tile, const int j0,
642 const int ke, const int je) noexcept {
643 T Akj[tile_size * tile_size];
644 load_tile(n, A, A_stride, pk, j0, Akj, ke, je);
645 ops_trsm_left_unit(tile, Akj, ke, je);
646 store_tile(n, A, A_stride, pk, j0, Akj, ke, je);
647 }
648
661 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
662 rl_schur_one(const int n, T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT pk,
663 const int* TDLS_RESTRICT pi, const T* TDLS_RESTRICT Aik, const int j0,
664 const int ke, const int ie, const int je) noexcept {
665 T Aij[tile_size * tile_size];
666 load_tile(n, A, A_stride, pi, j0, Aij, ie, je);
667
668 for (int p = 0; p < ke; ++p) {
669 T Akj_row[tile_size];
670 for (int j = 0; j < je; ++j)
671 Akj_row[j] = TDLS_LUPP_DYN_A(pk[p], j0 + j);
672 for (int i = 0; i < ie; ++i) {
673 const T L_ip = Aik[i * tile_size + p];
674 for (int j = 0; j < je; ++j)
675 Aij[i * tile_size + j] -= L_ip * Akj_row[j];
676 }
677 }
678
679 store_tile(n, A, A_stride, pi, j0, Aij, ie, je);
680 }
681
698 template<bool fuse_rhs>
699 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
700 rl_update_row_one(const int n, T* TDLS_RESTRICT A, const int A_stride,
701 const int* TDLS_RESTRICT piv, const int piv_stride,
702 const int* TDLS_RESTRICT pk, const T* TDLS_RESTRICT tile, const int k,
703 const int i0, const int ke, const int ie, T* TDLS_RESTRICT y = nullptr,
704 const int rhs_stride = 1) noexcept {
705 const int k0 = k * tile_size;
706 const int nt = num_tiles(n);
707
708 int pi[tile_size];
709 for (int i = 0; i < ie; ++i)
710 pi[i] = TDLS_LUPP_DYN_PIV(i0 + i);
711
712 // TRSM down: Aik := Aik * U^-1
713 T Aik[tile_size * tile_size];
714 load_tile(n, A, A_stride, pi, k0, Aik, ie, ke);
715 ops_trsm_right(tile, Aik, ke, ie);
716 store_tile(n, A, A_stride, pi, k0, Aik, ie, ke);
717
718 // Fused forward substitution: push the solved y_k segment into
719 // this row block while its L panel sits in registers.
720 if constexpr (fuse_rhs) {
721 for (int r = 0; r < ie; ++r) {
722 T sum = T(0);
723 for (int j = 0; j < ke; ++j)
724 sum += Aik[r * tile_size + j] * TDLS_LUPP_DYN_Y(k0 + j);
725 TDLS_LUPP_DYN_Y(i0 + r) -= sum;
726 }
727 }
728
729 // Schur sweep over the trailing columns
730 for (int tj = k + 1; tj < nt; ++tj)
731 rl_schur_one(n, A, A_stride, pk, pi, Aik, tj * tile_size, ke, ie,
732 tile_size_at(tj * tile_size, n));
733 }
734
749 template<bool oot_diagnostics, bool fuse_rhs>
750 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
751 rl_step(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
752 const int piv_stride, const int k, int& oot_count, T* TDLS_RESTRICT y = nullptr,
753 const int rhs_stride = 1) noexcept {
754 const int k0 = k * tile_size;
755 const int nt = num_tiles(n);
756 const int ke = tile_size_at(k0, n);
757
758 T tile[tile_size * tile_size];
759 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
760
761 if (!factor_diag_tile<oot_diagnostics, fuse_rhs>(n, A, A_stride, piv, piv_stride, k0, tile,
762 oot_count, ke, y, rhs_stride))
763 return false;
764
765 // Physical rows of the tile after the swaps of this step
766 int pk[tile_size];
767 for (int i = 0; i < ke; ++i)
768 pk[i] = TDLS_LUPP_DYN_PIV(k0 + i);
769
770 store_tile(n, A, A_stride, pk, k0, tile, ke, ke);
771
772 // Fused forward substitution: this tile's y segment is final from
773 // here on: unit-lower-solve it while the tile is in registers.
774 if constexpr (fuse_rhs) {
775 for (int kk = 0; kk < ke; ++kk) {
776 for (int i = kk + 1; i < ke; ++i)
777 TDLS_LUPP_DYN_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_DYN_Y(k0 + kk);
778 }
779 }
780
781 // TRSM right over the row panel
782 for (int tj = k + 1; tj < nt; ++tj)
783 rl_trsm_right_one(n, A, A_stride, pk, tile, tj * tile_size, ke,
784 tile_size_at(tj * tile_size, n));
785
786 // TRSM down + Schur over the rows below
787 for (int ti = k + 1; ti < nt; ++ti)
788 rl_update_row_one<fuse_rhs>(n, A, A_stride, piv, piv_stride, pk, tile, k,
789 ti * tile_size, ke, tile_size_at(ti * tile_size, n), y,
790 rhs_stride);
791 return true;
792 }
793
794 /* =====================================================================
795 LEFT-LOOKING schedule.
796 ===================================================================== */
797
812 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
813 ll_correct_tile(const int n, const T* TDLS_RESTRICT A, const int A_stride,
814 const int* TDLS_RESTRICT piv, const int piv_stride, const int row0,
815 const int col0, const int k0, T* TDLS_RESTRICT t, const int re,
816 const int ce) noexcept {
817 for (int bj0 = 0; bj0 < k0; bj0 += tile_size) {
818 T Lt[tile_size * tile_size];
819 load_tile_piv(n, A, A_stride, piv, piv_stride, row0, bj0, Lt, re, tile_size);
820 T Ut[tile_size * tile_size];
821 load_tile_piv(n, A, A_stride, piv, piv_stride, bj0, col0, Ut, tile_size, ce);
822 ops_gemm_sub(t, Lt, Ut, re, ce, tile_size);
823 }
824 }
825
841 template<bool fuse_rhs>
843 const int n, T* TDLS_RESTRICT A, const int A_stride, const int* TDLS_RESTRICT piv,
844 const int piv_stride, const T* TDLS_RESTRICT tile, const int k0, const int i0, const int ke,
845 const int ie, T* TDLS_RESTRICT y = nullptr, const int rhs_stride = 1) noexcept {
846 T B[tile_size * tile_size];
847 load_tile_piv(n, A, A_stride, piv, piv_stride, i0, k0, B, ie, ke);
848 ll_correct_tile(n, A, A_stride, piv, piv_stride, i0, k0, k0, B, ie, ke);
849 ops_trsm_right(tile, B, ke, ie);
850 store_tile_piv(n, A, A_stride, piv, piv_stride, i0, k0, B, ie, ke);
851
852 // Fused forward substitution: B is the final L(i,k) panel, so push
853 // the solved y_k segment into this row block.
854 if constexpr (fuse_rhs) {
855 for (int r = 0; r < ie; ++r) {
856 T sum = T(0);
857 for (int j = 0; j < ke; ++j)
858 sum += B[r * tile_size + j] * TDLS_LUPP_DYN_Y(k0 + j);
859 TDLS_LUPP_DYN_Y(i0 + r) -= sum;
860 }
861 }
862 }
863
875 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
876 ll_update_right_one(const int n, T* TDLS_RESTRICT A, const int A_stride,
877 const int* TDLS_RESTRICT piv, const int piv_stride,
878 const T* TDLS_RESTRICT tile, const int k0, const int j0, const int ke,
879 const int je) noexcept {
880 T B[tile_size * tile_size];
881 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
882 ll_correct_tile(n, A, A_stride, piv, piv_stride, k0, j0, k0, B, ke, je);
883 ops_trsm_left_unit(tile, B, ke, je);
884 store_tile_piv(n, A, A_stride, piv, piv_stride, k0, j0, B, ke, je);
885 }
886
901 template<bool oot_diagnostics, bool fuse_rhs>
902 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
903 ll_step(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
904 const int piv_stride, const int k, int& oot_count, T* TDLS_RESTRICT y = nullptr,
905 const int rhs_stride = 1) noexcept {
906 const int k0 = k * tile_size;
907 const int nt = num_tiles(n);
908 const int ke = tile_size_at(k0, n);
909
910 T tile[tile_size * tile_size];
911 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
912 ll_correct_tile(n, A, A_stride, piv, piv_stride, k0, k0, k0, tile, ke, ke);
913
914 if (!factor_diag_tile<oot_diagnostics, fuse_rhs>(n, A, A_stride, piv, piv_stride, k0, tile,
915 oot_count, ke, y, rhs_stride))
916 return false;
917
918 store_tile_piv(n, A, A_stride, piv, piv_stride, k0, k0, tile, ke, ke);
919
920 // Fused forward substitution: this tile's y segment is final from
921 // here on: unit-lower-solve it while the tile is in registers.
922 if constexpr (fuse_rhs) {
923 for (int kk = 0; kk < ke; ++kk) {
924 for (int i = kk + 1; i < ke; ++i)
925 TDLS_LUPP_DYN_Y(k0 + i) -= tile[i * tile_size + kk] * TDLS_LUPP_DYN_Y(k0 + kk);
926 }
927 }
928
929 // L panel below the diagonal
930 for (int ti = k + 1; ti < nt; ++ti)
931 ll_update_below_one<fuse_rhs>(n, A, A_stride, piv, piv_stride, tile, k0, ti * tile_size,
932 ke, tile_size_at(ti * tile_size, n), y, rhs_stride);
933
934 // U row panel right of the diagonal
935 for (int tj = k + 1; tj < nt; ++tj)
936 ll_update_right_one(n, A, A_stride, piv, piv_stride, tile, k0, tj * tile_size, ke,
937 tile_size_at(tj * tile_size, n));
938 return true;
939 }
940
941 /* =====================================================================
942 FACTORIZE - public entry point.
943 ===================================================================== */
944
965 template<bool oot_diagnostics = true, bool fuse_rhs = false>
966 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
967 factorize(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
968 const int piv_stride, int& oot_count, T* TDLS_RESTRICT y = nullptr,
969 const int rhs_stride = 1) noexcept {
970
971 if constexpr (oot_diagnostics) oot_count = 0;
972
973 for (int i = 0; i < n; ++i)
974 TDLS_LUPP_DYN_PIV(i) = i;
975
976 const int nt = num_tiles(n);
977 for (int k = 0; k < nt; ++k) {
978 if constexpr (schedule == Schedule::RightLooking) {
979 if (!rl_step<oot_diagnostics, fuse_rhs>(n, A, A_stride, piv, piv_stride, k,
980 oot_count, y, rhs_stride))
981 return false;
982 } else {
983 if (!ll_step<oot_diagnostics, fuse_rhs>(n, A, A_stride, piv, piv_stride, k,
984 oot_count, y, rhs_stride))
985 return false;
986 }
987 }
988
989 return true;
990 }
991
1000 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1001 factorize(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1002 const int piv_stride) noexcept {
1003 int unused = 0;
1004 return factorize<false>(n, A, A_stride, piv, piv_stride, unused);
1005 }
1006
1007 /* =====================================================================
1008 SUBSTITUTION - schedule-independent.
1009 ===================================================================== */
1010
1026 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1027 fwd_push_one(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1028 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1029 const int rhs_stride, const int xcol_stride, const int k0, const int m0,
1030 const int ke, const int me) noexcept {
1031 T Lmk[tile_size * tile_size];
1032 load_tile_piv(n, A, A_stride, piv, piv_stride, m0, k0, Lmk, me, ke);
1033 for (int i = 0; i < me; ++i) {
1034 for (int w = 0; w < nrhs; ++w) {
1035 T sum = T(0);
1036 for (int j = 0; j < ke; ++j)
1037 sum += Lmk[i * tile_size + j] * TDLS_LUPP_DYN_XW(w, k0 + j);
1038 TDLS_LUPP_DYN_XW(w, m0 + i) -= sum;
1039 }
1040 }
1041 }
1042
1055 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1056 fwd_step(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1057 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1058 const int rhs_stride, const int xcol_stride, const int k) noexcept {
1059 const int k0 = k * tile_size;
1060 const int nt = num_tiles(n);
1061 const int ke = tile_size_at(k0, n);
1062
1063 T Lkk[tile_size * tile_size];
1064 load_tile_piv_lower(n, A, A_stride, piv, piv_stride, k0, k0, Lkk, ke);
1065
1066 // In-tile unit-lower solve
1067 for (int kk = 0; kk < ke; ++kk) {
1068 for (int i = kk + 1; i < ke; ++i) {
1069 for (int w = 0; w < nrhs; ++w)
1070 TDLS_LUPP_DYN_XW(w, k0 + i) -=
1071 Lkk[i * tile_size + kk] * TDLS_LUPP_DYN_XW(w, k0 + kk);
1072 }
1073 }
1074
1075 // Push into the tiles below
1076 for (int m = k + 1; m < nt; ++m)
1077 fwd_push_one(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k0,
1078 m * tile_size, ke, tile_size_at(m * tile_size, n));
1079 }
1080
1096 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1097 bwd_pull_one(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1098 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1099 const int rhs_stride, const int xcol_stride, const int k0, const int m0,
1100 const int ke, const int me) noexcept {
1101 T Ukm[tile_size * tile_size];
1102 load_tile_piv(n, A, A_stride, piv, piv_stride, k0, m0, Ukm, ke, me);
1103 for (int i = 0; i < ke; ++i) {
1104 for (int w = 0; w < nrhs; ++w) {
1105 T sum = T(0);
1106 for (int j = 0; j < me; ++j)
1107 sum += Ukm[i * tile_size + j] * TDLS_LUPP_DYN_XW(w, m0 + j);
1108 TDLS_LUPP_DYN_XW(w, k0 + i) -= sum;
1109 }
1110 }
1111 }
1112
1125 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1126 bwd_step(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1127 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1128 const int rhs_stride, const int xcol_stride, const int k) noexcept {
1129 const int k0 = k * tile_size;
1130 const int nt = num_tiles(n);
1131 const int ke = tile_size_at(k0, n);
1132
1133 // Pull the trailing contributions
1134 for (int m = k + 1; m < nt; ++m)
1135 bwd_pull_one(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k0,
1136 m * tile_size, ke, tile_size_at(m * tile_size, n));
1137
1138 // In-tile upper solve
1139 T Ukk[tile_size * tile_size];
1140 load_tile_piv_upper(n, A, A_stride, piv, piv_stride, k0, k0, Ukk, ke);
1141
1142 for (int kk = ke - 1; kk >= 0; --kk) {
1143 for (int w = 0; w < nrhs; ++w)
1144 TDLS_LUPP_DYN_XW(w, k0 + kk) *= Ukk[kk * tile_size + kk]; // diag holds 1/pivot
1145 for (int i = 0; i < kk; ++i) {
1146 for (int w = 0; w < nrhs; ++w)
1147 TDLS_LUPP_DYN_XW(w, k0 + i) -=
1148 Ukk[i * tile_size + kk] * TDLS_LUPP_DYN_XW(w, k0 + kk);
1149 }
1150 }
1151 }
1152
1164 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1165 bwd_only(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1166 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1167 const int rhs_stride, const int xcol_stride) noexcept {
1168 for (int k = num_tiles(n) - 1; k >= 0; --k)
1169 bwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1170 }
1171
1185 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1186 fwd_bwd(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1187 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1188 const int rhs_stride, const int xcol_stride) noexcept {
1189 const int nt = num_tiles(n);
1190 for (int k = 0; k < nt; ++k)
1191 fwd_step(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride, k);
1192
1193 bwd_only(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1194 }
1195
1211 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1212 permute_inplace(const int n, const int* TDLS_RESTRICT piv, const int piv_stride,
1213 T* TDLS_RESTRICT x, const int rhs_stride) noexcept {
1214 if (n <= 64) {
1215 // Bitmask cycle decomposition: one visited bit per row, the
1216 // whole state in a single 64-bit register.
1217 unsigned long long visited = 0ull;
1218 for (int s = 0; s < n; ++s) {
1219 if ((visited >> s) & 1ull) continue;
1220 const T tmp = TDLS_LUPP_DYN_X(s);
1221 int cur = s;
1222 int nxt = TDLS_LUPP_DYN_PIV(cur);
1223 while (nxt != s) {
1224 TDLS_LUPP_DYN_X(cur) = TDLS_LUPP_DYN_X(nxt);
1225 visited |= 1ull << cur;
1226 cur = nxt;
1227 nxt = TDLS_LUPP_DYN_PIV(cur);
1228 }
1229 TDLS_LUPP_DYN_X(cur) = tmp;
1230 visited |= 1ull << cur;
1231 }
1232 } else {
1233 // Cycle-leader scan (no visited storage): a cycle is rotated
1234 // only when reached from its smallest index, detected by
1235 // walking the orbit. Values and order match the mask path.
1236 for (int s = 0; s < n; ++s) {
1237 int probe = TDLS_LUPP_DYN_PIV(s);
1238 while (probe > s)
1239 probe = TDLS_LUPP_DYN_PIV(probe);
1240 if (probe != s) continue;
1241
1242 const T tmp = TDLS_LUPP_DYN_X(s);
1243 int cur = s;
1244 int nxt = TDLS_LUPP_DYN_PIV(cur);
1245 while (nxt != s) {
1246 TDLS_LUPP_DYN_X(cur) = TDLS_LUPP_DYN_X(nxt);
1247 cur = nxt;
1248 nxt = TDLS_LUPP_DYN_PIV(cur);
1249 }
1250 TDLS_LUPP_DYN_X(cur) = tmp;
1251 }
1252 }
1253 }
1254
1255 /* =====================================================================
1256 SUBSTITUTION - public entry points.
1257 ===================================================================== */
1258
1270 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1271 substitute(const int n, const T* TDLS_RESTRICT A, const int A_stride,
1272 const int* TDLS_RESTRICT piv, const int piv_stride, const T* TDLS_RESTRICT b,
1273 T* TDLS_RESTRICT x, const int rhs_stride) noexcept {
1274 for (int i = 0; i < n; ++i)
1276 fwd_bwd(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
1277 }
1278
1293 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1294 substitute_canonical(const int n, const T* TDLS_RESTRICT A, const int A_stride,
1295 const int* TDLS_RESTRICT piv, const int piv_stride, const int col,
1296 T* TDLS_RESTRICT x, const int rhs_stride) noexcept {
1297 substitute_canonical_multirhs_pass(n, 1, A, A_stride, piv, piv_stride, col, x, rhs_stride,
1298 0);
1299 }
1300
1316 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1317 substitute_canonical_multirhs_pass(const int n, const int nrhs, const T* TDLS_RESTRICT A,
1318 const int A_stride, const int* TDLS_RESTRICT piv,
1319 const int piv_stride, const int col0, T* TDLS_RESTRICT x,
1320 const int rhs_stride, const int xcol_stride) noexcept {
1321 for (int i = 0; i < n; ++i) {
1322 const int p = TDLS_LUPP_DYN_PIV(i);
1323 for (int w = 0; w < nrhs; ++w)
1324 TDLS_LUPP_DYN_XW(w, i) = (p == col0 + w) ? T(1) : T(0);
1325 }
1326 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1327 }
1328
1352 template<int pass_width = 0>
1353 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1354 substitute_canonical_multirhs(const int n, const int nrhs, const T* TDLS_RESTRICT A,
1355 const int A_stride, const int* TDLS_RESTRICT piv,
1356 const int piv_stride, const int col0, T* TDLS_RESTRICT x,
1357 const int rhs_stride, const int xcol_stride) noexcept {
1358 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
1359 if constexpr (pass_width == 0) {
1360 substitute_canonical_multirhs_pass(n, nrhs, A, A_stride, piv, piv_stride, col0, x,
1361 rhs_stride, xcol_stride);
1362 } else {
1363 // A pass_width of nrhs or more runs the loop once, the whole
1364 // block being its single pass.
1365 for (int c0 = 0; c0 < nrhs; c0 += pass_width) {
1366 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
1367 const unsigned off = unsigned(c0) * unsigned(xcol_stride);
1368 substitute_canonical_multirhs_pass(n, cw, A, A_stride, piv, piv_stride, col0 + c0,
1369 x + off, rhs_stride, xcol_stride);
1370 }
1371 }
1372 }
1373
1389 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1390 substitute_multirhs_pass(const int n, const int nrhs, const T* TDLS_RESTRICT A,
1391 const int A_stride, const int* TDLS_RESTRICT piv, const int piv_stride,
1392 const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
1393 const int xcol_stride) noexcept {
1394 for (int i = 0; i < n; ++i) {
1395 const int p = TDLS_LUPP_DYN_PIV(i);
1396 for (int w = 0; w < nrhs; ++w)
1397 TDLS_LUPP_DYN_XW(w, i) = TDLS_LUPP_DYN_BW(w, p);
1398 }
1399 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1400 }
1401
1428 template<int pass_width = 0>
1429 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1430 substitute_multirhs(const int n, const int nrhs, const T* TDLS_RESTRICT A, const int A_stride,
1431 const int* TDLS_RESTRICT piv, const int piv_stride,
1432 const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
1433 const int xcol_stride) noexcept {
1434 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
1435 if constexpr (pass_width == 0) {
1436 substitute_multirhs_pass(n, nrhs, A, A_stride, piv, piv_stride, b, x, rhs_stride,
1437 xcol_stride);
1438 } else {
1439 // A pass_width of nrhs or more runs the loop once, the whole
1440 // block being its single pass.
1441 for (int c0 = 0; c0 < nrhs; c0 += pass_width) {
1442 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
1443 const unsigned off = unsigned(c0) * unsigned(xcol_stride);
1444 substitute_multirhs_pass(n, cw, A, A_stride, piv, piv_stride, b + off, x + off,
1445 rhs_stride, xcol_stride);
1446 }
1447 }
1448 }
1449
1462 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1463 substitute_inplace(const int n, const T* TDLS_RESTRICT A, const int A_stride,
1464 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1465 const int rhs_stride) noexcept {
1466 permute_inplace(n, piv, piv_stride, x, rhs_stride);
1467 fwd_bwd(n, 1, A, A_stride, piv, piv_stride, x, rhs_stride, 0);
1468 }
1469
1485 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1486 substitute_inplace_multirhs_pass(const int n, const int nrhs, const T* TDLS_RESTRICT A,
1487 const int A_stride, const int* TDLS_RESTRICT piv,
1488 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
1489 const int xcol_stride) noexcept {
1490 for (int w = 0; w < nrhs; ++w)
1491 permute_inplace(n, piv, piv_stride, x + unsigned(w) * unsigned(xcol_stride),
1492 rhs_stride);
1493 fwd_bwd(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride, xcol_stride);
1494 }
1495
1519 template<int pass_width = 0>
1520 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
1521 substitute_inplace_multirhs(const int n, const int nrhs, const T* TDLS_RESTRICT A,
1522 const int A_stride, const int* TDLS_RESTRICT piv,
1523 const int piv_stride, T* TDLS_RESTRICT x, const int rhs_stride,
1524 const int xcol_stride) noexcept {
1525 static_assert(pass_width >= 0, "tdls: pass_width must not be negative");
1526 if constexpr (pass_width == 0) {
1527 substitute_inplace_multirhs_pass(n, nrhs, A, A_stride, piv, piv_stride, x, rhs_stride,
1528 xcol_stride);
1529 } else {
1530 // A pass_width of nrhs or more runs the loop once, the whole
1531 // block being its single pass.
1532 for (int c0 = 0; c0 < nrhs; c0 += pass_width) {
1533 const int cw = (pass_width > nrhs - c0) ? nrhs - c0 : pass_width;
1534 const unsigned off = unsigned(c0) * unsigned(xcol_stride);
1535 substitute_inplace_multirhs_pass(n, cw, A, A_stride, piv, piv_stride, x + off,
1536 rhs_stride, xcol_stride);
1537 }
1538 }
1539 }
1540
1541 /* =====================================================================
1542 SOLVE - convenience wrappers.
1543 ===================================================================== */
1544
1560 template<bool oot_diagnostics = true>
1561 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1562 solve(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1563 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
1564 int& oot_count) noexcept {
1565 if (!factorize<oot_diagnostics>(n, A, A_stride, piv, piv_stride, oot_count)) return false;
1566 substitute(n, A, A_stride, piv, piv_stride, b, x, rhs_stride);
1567 return true;
1568 }
1569
1582 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1583 solve(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1584 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x,
1585 const int rhs_stride) noexcept {
1586 int unused = 0;
1587 return solve<false>(n, A, A_stride, piv, piv_stride, b, x, rhs_stride, unused);
1588 }
1589
1612 template<int pass_width = 0, bool oot_diagnostics = true>
1613 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1614 solve_multirhs(const int n, const int nrhs, T* TDLS_RESTRICT A, const int A_stride,
1615 int* TDLS_RESTRICT piv, const int piv_stride, const T* TDLS_RESTRICT b,
1616 T* TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride,
1617 int& oot_count) noexcept {
1618 if (!factorize<oot_diagnostics>(n, A, A_stride, piv, piv_stride, oot_count)) return false;
1619 substitute_multirhs<pass_width>(n, nrhs, A, A_stride, piv, piv_stride, b, x, rhs_stride,
1620 xcol_stride);
1621 return true;
1622 }
1623
1643 template<int pass_width = 0>
1644 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1645 solve_multirhs(const int n, const int nrhs, T* TDLS_RESTRICT A, const int A_stride,
1646 int* TDLS_RESTRICT piv, const int piv_stride, const T* TDLS_RESTRICT b,
1647 T* TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept {
1648 int unused = 0;
1649 return solve_multirhs<pass_width, false>(n, nrhs, A, A_stride, piv, piv_stride, b, x,
1650 rhs_stride, xcol_stride, unused);
1651 }
1652
1674 template<bool oot_diagnostics = true>
1675 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1676 solve_inplace(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1677 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride,
1678 int& oot_count) noexcept {
1679 if (!factorize<oot_diagnostics, true>(n, A, A_stride, piv, piv_stride, oot_count, y,
1680 rhs_stride))
1681 return false;
1682
1683 // Backward pass only: the forward one happened inside factorize.
1684 bwd_only(n, 1, A, A_stride, piv, piv_stride, y, rhs_stride, 0);
1685 return true;
1686 }
1687
1699 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1700 solve_inplace(const int n, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1701 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride) noexcept {
1702 int unused = 0;
1703 return solve_inplace<false>(n, A, A_stride, piv, piv_stride, y, rhs_stride, unused);
1704 }
1705
1732 template<int pass_width = 0, bool oot_diagnostics = true>
1733 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1734 solve_inplace_multirhs(const int n, const int nrhs, T* TDLS_RESTRICT A, const int A_stride,
1735 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT y,
1736 const int rhs_stride, const int xcol_stride, int& oot_count) noexcept {
1737 if (!factorize<oot_diagnostics>(n, A, A_stride, piv, piv_stride, oot_count)) return false;
1738 substitute_inplace_multirhs<pass_width>(n, nrhs, A, A_stride, piv, piv_stride, y,
1739 rhs_stride, xcol_stride);
1740 return true;
1741 }
1742
1761 template<int pass_width = 0>
1762 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr bool
1763 solve_inplace_multirhs(const int n, const int nrhs, T* TDLS_RESTRICT A, const int A_stride,
1764 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT y,
1765 const int rhs_stride, const int xcol_stride) noexcept {
1766 int unused = 0;
1767 return solve_inplace_multirhs<pass_width, false>(n, nrhs, A, A_stride, piv, piv_stride, y,
1768 rhs_stride, xcol_stride, unused);
1769 }
1770};
1771
1772
1773
1774#undef TDLS_LUPP_DYN_A
1775#undef TDLS_LUPP_DYN_PIV
1776#undef TDLS_LUPP_DYN_X
1777#undef TDLS_LUPP_DYN_B
1778#undef TDLS_LUPP_DYN_XW
1779#undef TDLS_LUPP_DYN_BW
1780#undef TDLS_LUPP_DYN_Y
1781
1782
1783
1784} // namespace tdls
1785
1786#if defined(__GNUC__) && !defined(__clang__)
1787#pragma GCC diagnostic pop
1788#endif
1789
1790
1791
1792#endif // TDLS_SOLVERS_TILED_LUPP_SOLVER_DYNAMIC_HPP
Compile-time configuration of the TiledLUpp solver family.
#define TDLS_RESTRICT
Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC).
Definition macros.hpp:80
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:45
#define TDLS_FORCEINLINE
__forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS_HOST_DEVICE for why),...
Definition macros.hpp:67
Scalar math helpers usable in constant expressions.
Schedule
Elimination schedule of a factorization.
Definition options.hpp:37
#define TDLS_LUPP_DYN_X(i)
Strided entry i of the solution vector.
Definition solver_dynamic.hpp:103
#define TDLS_LUPP_DYN_PIV(i)
Strided pivot entry i.
Definition solver_dynamic.hpp:100
#define TDLS_LUPP_DYN_Y(i)
Strided entry i of the fused right-hand side of solve_inplace.
Definition solver_dynamic.hpp:118
#define TDLS_LUPP_DYN_A(r, c)
Strided element (r, c) of the factor matrix, flat index remapped by Config.layout.
Definition solver_dynamic.hpp:97
#define TDLS_LUPP_DYN_XW(w, i)
Strided entry i of column w of a multi right-hand-side block.
Definition solver_dynamic.hpp:109
#define TDLS_LUPP_DYN_B(i)
Strided entry i of the right-hand side.
Definition solver_dynamic.hpp:106
#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:114
Runtime-size tiled dense LU factorization with logical partial pivoting and out-of-tile pivot recover...
Definition solver_dynamic.hpp:159
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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 int 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:1165
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:903
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:354
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace(const 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:1700
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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: only the strict lower tria...
Definition solver_dynamic.hpp:399
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:662
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:876
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:614
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace_multirhs(const int n, const 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 int xcol_stride) noexcept
Diagnostics-free solve_inplace_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:1763
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factorize(const 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:967
static constexpr Schedule schedule
elimination schedule (RightLooking or LeftLooking)
Definition solver_dynamic.hpp:162
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:227
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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 int xcol_stride) noexcept
Triangular solves on already-permuted column(s) x. nrhs columns are processed per tile visit,...
Definition solver_dynamic.hpp:1186
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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 int 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:1097
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:314
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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 int xcol_stride) noexcept
One pass of the canonical multi right-hand-side substitution: nrhs canonical columns e_col0 ....
Definition solver_dynamic.hpp:1317
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:751
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:192
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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 int 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:1486
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_multirhs(const int n, const 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 int xcol_stride) noexcept
Diagnostics-free solve_multirhs overload: no out-of-tile out-parameter at all.
Definition solver_dynamic.hpp:1645
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace(const 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:1463
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:1212
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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 int 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:1027
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool factorize(const 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:1001
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace_multirhs(const int n, const 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 int xcol_stride, int &oot_count) noexcept
factorize + substitute_inplace_multirhs in one call: y holds the nrhs unpermuted right-hand-side colu...
Definition solver_dynamic.hpp:1734
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:265
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:700
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:813
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical(const 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:1294
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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 int xcol_stride, const int k) noexcept
Backward step: pull the trailing contributions, then upper-solve the diagonal tile's segment.
Definition solver_dynamic.hpp:1126
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_multirhs(const int n, const 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 int xcol_stride, int &oot_count) noexcept
factorize + substitute_multirhs in one call.
Definition solver_dynamic.hpp:1614
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr int num_tiles(const int n) noexcept
Number of tiles per dimension (last one possibly partial).
Definition solver_dynamic.hpp:184
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_inplace_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 int xcol_stride) noexcept
nrhs columns solved in place: x holds the nrhs unpermuted right-hand-side columns on entry and the nr...
Definition solver_dynamic.hpp:1521
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve(const 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:1562
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve(const 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:1583
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:842
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr bool solve_inplace(const 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:1676
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:244
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:333
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_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, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
Solve nrhs right-hand-side columns from a prior factorize: X := U^-1 L^-1 P B, column by column....
Definition solver_dynamic.hpp:1430
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:640
static constexpr T singular_floor
Singularity floor of the out-of-tile recovery, read once from the configuration (see TiledLUppConfig:...
Definition solver_dynamic.hpp:170
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:458
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute(const 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. b and x must not alias (use substitute_inplace for t...
Definition solver_dynamic.hpp:1271
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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: only the upper triangle i...
Definition solver_dynamic.hpp:422
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:377
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:209
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void substitute_canonical_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, const int col0, T *TDLS_RESTRICT x, const int rhs_stride, const int xcol_stride) noexcept
nrhs canonical columns e_col0 .. e_{col0+nrhs-1} solved from a prior factorize: the columns of the in...
Definition solver_dynamic.hpp:1354
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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:288
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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 int xcol_stride, const int k) noexcept
Forward step: unit-lower solve the diagonal tile's segment, then push it into the tiles below.
Definition solver_dynamic.hpp:1056
static constexpr int tile_size
tile size (int)
Definition solver_dynamic.hpp:161
TDLS_HOST_DEVICE static TDLS_FORCEINLINE 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 int 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:1390
static constexpr T oot_threshold
Acceptable-pivot threshold of the out-of-tile search, read once from the configuration (see TiledLUpp...
Definition solver_dynamic.hpp:167