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_COOPERATIVE_LUPP_SOLVER_DYNAMIC_HPP
2#define TDLS_SOLVERS_COOPERATIVE_LUPP_SOLVER_DYNAMIC_HPP
3
4
5
6/* The elimination and substitution kernels of this file are derived from
7 MAGMA 2.10.0, magmablas/zgesv_batched_small.cu, function
8 zgesv_batched_small_device (authors Azzam Haidar and Ahmad Abdelfattah).
9 Those portions are distributed under the MAGMA license, reproduced
10 below as required by it:
11
12 Copyright (c) 2009-2023, The University of Tennessee
13 All rights reserved.
14
15 Redistribution and use in source and binary forms, with or without
16 modification, are permitted provided that the following conditions
17 are met:
18
19 * Redistributions of source code must retain the above copyright
20 notice, this list of conditions and the following disclaimer.
21 * Redistributions in binary form must reproduce the above copyright
22 notice, this list of conditions and the following disclaimer in the
23 documentation and/or other materials provided with the distribution.
24 * Neither the name of the University of Tennessee, Knoxville nor the
25 names of its contributors may be used to endorse or promote products
26 derived from this software without specific prior written permission.
27
28 This software is provided by the copyright holders and contributors
29 ``as is'' and any express or implied warranties, including, but not
30 limited to, the implied warranties of merchantability and fitness for
31 a particular purpose are disclaimed. In no event shall the copyright
32 holders or contributors be liable for any direct, indirect, incidental,
33 special, exemplary, or consequential damages (including, but not
34 limited to, procurement of substitute goods or services; loss of use,
35 data, or profits; or business interruption) however caused and on any
36 theory of liability, whether in contract, strict liability, or tort
37 (including negligence or otherwise) arising in any way out of the use
38 of this software, even if advised of the possibility of such damage.
39*/
40
41
42
120
121
122
123#include <type_traits>
124
125#include <tdls/core/group.hpp>
126#include <tdls/core/macros.hpp>
127#include <tdls/core/math.hpp>
128#include <tdls/core/target.hpp>
130
131
132
133// gcc's flow analysis cannot prove that the right-hand-side entries of a
134// thread are loaded before being read: the load stops at the first
135// phantom row and every later use stops at the same row, a property
136// validated bitwise against the compile-time solver. The resulting
137// -Wmaybe-uninitialized reports are spurious. Under UBSan instrumentation
138// at -O2 and above, gcc additionally loses the value ranges of the loop
139// indices and reports impossible subscripts through -Warray-bounds,
140// exactly as in the TiledLUpp headers; the exercised paths are certified
141// free of out-of-bounds accesses by the constexpr suite. Both
142// suppressions are scoped to this header and to those warnings only,
143// exactly as in solver_static.hpp.
144#if defined(TDLS_DETAIL_COMPILER_GCC)
145#pragma GCC diagnostic push
146#pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
147#pragma GCC diagnostic ignored "-Warray-bounds"
148#endif
149
150namespace tdls {
151
152
153
154/* Addressing macros. They expand inside member functions where n, the
155 parameter names (A, A_stride, piv, piv_stride, b, x, y, rhs_stride)
156 and the Config value are in scope. #undef'd at the end of this
157 header. */
158
164#define TDLS_COOP_LUPP_DYN_A(r, c) A[TDLS_STRIDED((c) * n + (r), A_stride)]
167#define TDLS_COOP_LUPP_DYN_PIV(r) piv[TDLS_STRIDED(r, piv_stride)]
170#define TDLS_COOP_LUPP_DYN_B(r) b[TDLS_STRIDED(r, rhs_stride)]
173#define TDLS_COOP_LUPP_DYN_X(r) x[TDLS_STRIDED(r, rhs_stride)]
176#define TDLS_COOP_LUPP_DYN_Y(r) y[TDLS_STRIDED(r, rhs_stride)]
177
178
179
207template<typename T, CooperativeLUppConfig<T> Config = CooperativeLUppConfig<T>{}>
209
210 static_assert(
211 Config.row_interchange != RowInterchange::Auto,
212 "CooperativeLUppSolverDynamic: row_interchange must be Logical or Physical (Auto is the "
213 "default of the blocks of ArrowheadLUpp only)");
214 static_assert(Config.rows_per_thread >= 1,
215 "CooperativeLUppSolverDynamic: rows_per_thread must be >= 1");
216 static_assert(Config.relative_pivot_threshold.is_finite() && Config.singular_floor.is_finite(),
217 "CooperativeLUppSolverDynamic: relative_pivot_threshold and singular_floor must "
218 "be finite (and fit a 63-bit mantissa)");
219
224 static constexpr int rows_per_thread = Config.rows_per_thread < 1 ? 1 : Config.rows_per_thread;
227 static constexpr T relative_pivot_threshold = Config.relative_pivot_threshold;
230 static constexpr T singular_floor = Config.singular_floor;
234 static constexpr bool transposed = Config.layout == MatrixLayout::RowMajor;
235
236 static_assert(relative_pivot_threshold > T(0) && relative_pivot_threshold <= T(1),
237 "CooperativeLUppSolverDynamic: relative_pivot_threshold must be in (0, 1]");
238 static_assert(singular_floor > T(0),
239 "CooperativeLUppSolverDynamic: singular_floor must be positive");
240
245 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
246 threads_per_system(const int n) noexcept {
247 return (n + rows_per_thread - 1) / rows_per_thread;
248 }
249
256 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
257 workspace_size(const int n) noexcept {
258 return detail::barrier_elements + (Config.row_interchange == RowInterchange::Logical
259 ? 3 * n
260 : 2 * n + 2 * threads_per_system(n) + 5);
261 }
262
263 /* =====================================================================
264 Pivot entries and right-hand-side entries.
265 Slot K of thread tx holds row tx + K * threads, a phantom row when
266 it reaches n. The rows themselves stay in the matrix; only their
267 pivot entries and their right-hand-side entries are kept by the
268 thread. Each thread reads and writes only the entries of its own
269 rows, so the operands never carry data between threads: the
270 workspace does.
271 ===================================================================== */
272
281 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
282 init_piv(const int tx, const int threads, int (&rpiv)[rows_per_thread]) noexcept {
283 for (int K = 0; K < rows_per_thread; ++K)
284 rpiv[K] = tx + K * threads;
285 }
286
294 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
295 load_piv(const int n, const int tx, const int threads, const int* TDLS_RESTRICT piv,
296 const int piv_stride, int (&rpiv)[rows_per_thread]) noexcept {
297 for (int K = 0; K < rows_per_thread; ++K) {
298 const int r = tx + K * threads;
299 rpiv[K] = r < n ? TDLS_COOP_LUPP_DYN_PIV(r) : r;
300 }
301 }
302
310 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
311 store_piv(const int n, const int tx, const int threads, int* TDLS_RESTRICT piv,
312 const int piv_stride, const int (&rpiv)[rows_per_thread]) noexcept {
313 for (int K = 0; K < rows_per_thread; ++K) {
314 const int r = tx + K * threads;
315 if (r >= n) break;
316 TDLS_COOP_LUPP_DYN_PIV(r) = rpiv[K];
317 }
318 }
319
328 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
329 load_rhs(const int n, const int tx, const int threads, const T* TDLS_RESTRICT b,
330 const int rhs_stride, T (&rB)[rows_per_thread]) noexcept {
331 for (int K = 0; K < rows_per_thread; ++K) {
332 const int r = tx + K * threads;
333 if (r >= n) break;
334 rB[K] = TDLS_COOP_LUPP_DYN_B(r);
335 }
336 }
337
345 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
346 store_solution(const int n, const int tx, const int threads, T* TDLS_RESTRICT x,
347 const int rhs_stride, const T (&rB)[rows_per_thread]) noexcept {
348 for (int K = 0; K < rows_per_thread; ++K) {
349 const int r = tx + K * threads;
350 if (r >= n) break;
351 TDLS_COOP_LUPP_DYN_X(r) = rB[K];
352 }
353 }
354
355 /* =====================================================================
356 FACTORIZATION, right-looking, one column per step, as in the
357 compile-time solver, in place in the matrix. Logical row
358 interchanges: the magnitudes of the column are published at the
359 positions of their rows, every thread finds the same pivot, the
360 owner of the pivot row publishes it, and every thread eliminates its
361 own rows below the pivot. Physical row interchanges: each thread
362 publishes its best candidate, every thread finds the same pivot, the
363 owner of position i publishes its row, the owner of the pivot row as
364 well when the pivot is not in place, and every thread eliminates its
365 rows below the pivot after the exchange.
366 ===================================================================== */
367
374 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
375 keeps_pivot(const T in_place, const T column_max) noexcept {
376 const T bound = relative_pivot_threshold * column_max;
377 return in_place >= (bound > singular_floor ? bound : singular_floor);
378 }
379
404 template<bool fuse_rhs, typename Sync>
405 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
406 logical_elimination_step(const int n, const int i, const int tx, const int threads,
407 T* TDLS_RESTRICT A, const int A_stride,
408 [[maybe_unused]] T (&rB)[rows_per_thread],
409 int (&rowid)[rows_per_thread], T* sB, T* sx, T* dsx, int& linfo,
410 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
411 // izamax: the magnitudes of column i, at the positions of their rows.
412 for (int K = 0; K < rows_per_thread; ++K) {
413 const int r = tx + K * threads;
414 if (r >= n) break;
415 dsx[rowid[K]] = detail::abs(TDLS_COOP_LUPP_DYN_A(r, i));
416 }
417 sync();
418
419 // Every thread scans the same magnitudes and finds the same pivot.
420 T rx_abs_max = dsx[i];
421 int max_id = i;
422 for (int j = i + 1; j < n; j++) {
423 if (dsx[j] > rx_abs_max) {
424 max_id = j;
425 rx_abs_max = dsx[j];
426 }
427 }
428 // Below the default threshold 1, the row in place keeps the pivot
429 // when it is large enough.
430 if constexpr (relative_pivot_threshold < T(1)) {
431 if (keeps_pivot(dsx[i], rx_abs_max)) max_id = i;
432 }
433 const bool zero_pivot = (rx_abs_max < singular_floor);
434 linfo = (zero_pivot && linfo == 0) ? (i + 1) : linfo;
435 const T update = zero_pivot ? T(0) : T(1);
436
437 // The owner of the pivot row publishes it; positions i and max_id
438 // are swapped.
439 for (int K = 0; K < rows_per_thread; ++K) {
440 const int r = tx + K * threads;
441 if (r >= n) break;
442 if (rowid[K] == max_id) {
443 rowid[K] = i;
444 for (int j = i; j < n; j++)
445 sx[j] = update * TDLS_COOP_LUPP_DYN_A(r, j);
446 if constexpr (fuse_rhs && !transposed) sB[0] = rB[K];
447 } else if (rowid[K] == i) {
448 rowid[K] = max_id;
449 }
450 // S^T x = b: the owner of entry i publishes b'(i), final at
451 // this step.
452 if constexpr (fuse_rhs && transposed) {
453 if (r == i) sB[0] = rB[K];
454 }
455 }
456 sync();
457
458 const T reg = zero_pivot ? T(1) : T(1) / sx[i];
459 [[maybe_unused]] T z_i = T(0);
460 if constexpr (fuse_rhs && transposed) z_i = sB[0] * reg;
461
462 // scal and ger: every thread eliminates its rows below the pivot.
463 for (int K = 0; K < rows_per_thread; ++K) {
464 const int r = tx + K * threads;
465 if (r >= n) break;
466 if (rowid[K] > i) {
467 TDLS_COOP_LUPP_DYN_A(r, i) *= reg;
468 for (int j = i + 1; j < n; j++)
469 TDLS_COOP_LUPP_DYN_A(r, j) -= TDLS_COOP_LUPP_DYN_A(r, i) * sx[j];
470 if constexpr (fuse_rhs && !transposed) rB[K] -= TDLS_COOP_LUPP_DYN_A(r, i) * sB[0];
471 }
472 // the owner of entry i keeps z_i in its place; each entry c > i
473 // loses U(i, c) z_i
474 if constexpr (fuse_rhs && transposed) {
475 if (r == i) rB[K] = z_i;
476 if (r > i) rB[K] -= sx[r] * z_i;
477 }
478 }
479 sync();
480 }
481
489 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
490 save_row(const int n, const T* TDLS_RESTRICT A, const int A_stride, const int r,
491 T* record) noexcept {
492 for (int c = 0; c < n; c++)
493 record[c] = TDLS_COOP_LUPP_DYN_A(r, c);
494 }
495
503 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
504 load_row(const int n, T* TDLS_RESTRICT A, const int A_stride, const int r,
505 const T* record) noexcept {
506 for (int c = 0; c < n; c++)
507 TDLS_COOP_LUPP_DYN_A(r, c) = record[c];
508 }
509
530 template<bool fuse_rhs, typename Sync>
531 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
532 physical_elimination_step(const int n, const int i, const int tx, const int threads,
533 T* TDLS_RESTRICT A, const int A_stride,
534 [[maybe_unused]] T (&rB)[rows_per_thread],
535 int (&origin)[rows_per_thread], T* work, int& linfo,
536 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
537 // The right-hand side moves with the rows only under S x = b; under
538 // S^T x = b, the owner of position i is the owner of entry i: its
539 // record carries b'(i), and it keeps z_i in its place.
540 constexpr bool carries_rhs = fuse_rhs && !transposed;
541 // Workspace layout of the compile-time solver.
542 T* const candidate = work;
543 T* const where = candidate + threads;
544 T* const in_place = where + threads;
545 T* const leaving = in_place + 1;
546 T* const entering = leaving + n + 2;
547
548 // Candidate of the thread. The owner of position i starts from the
549 // row in place and publishes its magnitude, as in the compile-time
550 // solver.
551 T best = T(-1);
552 int best_pos = n;
553 T local = T(0);
554 for (int K = 0; K < rows_per_thread; ++K) {
555 const int pos = tx + K * threads;
556 if (pos >= n) break;
557 if (pos < i) continue;
558 const T v = detail::abs(TDLS_COOP_LUPP_DYN_A(pos, i));
559 if (pos == i || v > best) {
560 best = v;
561 best_pos = pos;
562 }
563 if (pos == i) local = v;
564 }
565 candidate[tx] = best;
566 where[tx] = static_cast<T>(best_pos);
567 if (tx == i % threads) in_place[0] = local;
568 sync();
569
570 // Every thread finds the same column maximum and the same pivot, as
571 // in the compile-time solver.
572 T rx_abs_max = candidate[i % threads];
573 for (int t = 0; t < threads; t++)
574 rx_abs_max = candidate[t] > rx_abs_max ? candidate[t] : rx_abs_max;
575 bool keep = !(in_place[0] < rx_abs_max);
576 if constexpr (relative_pivot_threshold < T(1)) keep = keeps_pivot(in_place[0], rx_abs_max);
577 int max_id = i;
578 if (!keep) {
579 max_id = n;
580 for (int t = 0; t < threads; t++) {
581 const int pos = static_cast<int>(where[t]);
582 if (candidate[t] == rx_abs_max && pos < max_id) max_id = pos;
583 }
584 }
585 const bool zero_pivot = (rx_abs_max < singular_floor);
586 linfo = (zero_pivot && linfo == 0) ? (i + 1) : linfo;
587
588 // As in the compile-time solver: the owner of position i publishes
589 // its row whole and, when the pivot is not in place, the owner of
590 // position max_id publishes the pivot row. Each record holds the n
591 // entries of its row, then its right-hand-side entry (fuse_rhs
592 // only) and its original index. The rows are addressed by position,
593 // and the entries of the slots move through loops kept small, which
594 // the compiler unrolls: a slot indexed at run time would demote rB
595 // and origin to the stack.
596 if (tx == i % threads) save_row(n, A, A_stride, i, leaving);
597 if (max_id != i && tx == max_id % threads) save_row(n, A, A_stride, max_id, entering);
598 for (int K = 0; K < rows_per_thread; ++K) {
599 const int pos = tx + K * threads;
600 if (pos != i && (pos != max_id || max_id == i)) continue;
601 T* const record = pos == i ? leaving : entering;
602 if constexpr (fuse_rhs) record[n] = rB[K];
603 record[n + 1] = static_cast<T>(origin[K]);
604 }
605 sync();
606
607 if (max_id != i) {
608 if (tx == i % threads) load_row(n, A, A_stride, i, entering);
609 if (tx == max_id % threads) load_row(n, A, A_stride, max_id, leaving);
610 for (int K = 0; K < rows_per_thread; ++K) {
611 const int pos = tx + K * threads;
612 if (pos != max_id && pos != i) continue;
613 const T* const record = pos == i ? entering : leaving;
614 if constexpr (carries_rhs) rB[K] = record[n];
615 origin[K] = static_cast<int>(record[n + 1]);
616 }
617 }
618
619 const T* const prow = max_id == i ? leaving : entering;
620 const T reg = zero_pivot ? T(1) : T(1) / prow[i];
621 [[maybe_unused]] T z_i = T(0);
622 if constexpr (fuse_rhs && transposed) z_i = leaving[n] * reg;
623
624 // scal and ger: every thread eliminates its rows below the pivot.
625 // The next step writes only the candidates, which no thread reads
626 // any more: no barrier closes the step.
627 for (int K = 0; K < rows_per_thread; ++K) {
628 const int pos = tx + K * threads;
629 if (pos >= n) break;
630 if (pos > i) {
631 TDLS_COOP_LUPP_DYN_A(pos, i) *= reg;
632 for (int j = i + 1; j < n; j++)
633 TDLS_COOP_LUPP_DYN_A(pos, j) -= TDLS_COOP_LUPP_DYN_A(pos, i) * prow[j];
634 if constexpr (carries_rhs) rB[K] -= TDLS_COOP_LUPP_DYN_A(pos, i) * prow[n];
635 }
636 if constexpr (fuse_rhs && transposed) {
637 if (pos == i) rB[K] = z_i;
638 if (pos > i) rB[K] -= prow[pos] * z_i;
639 }
640 }
641 }
642
660 template<bool fuse_rhs, typename Sync>
661 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
662 eliminate(const int n, const int tx, const int threads, T* TDLS_RESTRICT A, const int A_stride,
663 T (&rB)[rows_per_thread], int (&rpiv)[rows_per_thread], T* work,
664 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
665 int linfo = 0;
666 if constexpr (Config.row_interchange == RowInterchange::Logical) {
667 T* const sB = work;
668 T* const sx = work + n;
669 T* const dsx = work + 2 * n;
670 for (int i = 0; i < n; i++)
671 logical_elimination_step<fuse_rhs>(n, i, tx, threads, A, A_stride, rB, rpiv, sB, sx,
672 dsx, linfo, sync);
673 } else {
674 for (int i = 0; i < n; i++)
675 physical_elimination_step<fuse_rhs>(n, i, tx, threads, A, A_stride, rB, rpiv, work,
676 linfo, sync);
677 sync(); // no barrier closes the steps: this one frees the workspace
678 }
679 return linfo;
680 }
681
682 /* =====================================================================
683 SUBSTITUTION of S x = b, the column-major layout, as in the
684 compile-time solver: the forward pass
685 replays the elimination of the right-hand side from the factored
686 rows, the backward pass publishes column i of U, every thread
687 computes x_i, the owner of entry i keeps it and every thread updates
688 its entries above. Under physical row interchanges, the right-hand
689 side is first gathered in the pivoted order, and both passes run on
690 the positions of the rows.
691 ===================================================================== */
692
707 template<typename Sync>
708 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
709 forward_pass(const int n, const int tx, const int threads, const T* TDLS_RESTRICT A,
710 const int A_stride, T (&rB)[rows_per_thread], const int (&rowid)[rows_per_thread],
711 T* work, Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
712 T* const sB = work;
713 for (int i = 0; i < n; i++) {
714 for (int K = 0; K < rows_per_thread; ++K) {
715 if (tx + K * threads >= n) break;
716 if (rowid[K] == i) sB[0] = rB[K];
717 }
718 sync();
719
720 for (int K = 0; K < rows_per_thread; ++K) {
721 const int r = tx + K * threads;
722 if (r >= n) break;
723 if (rowid[K] > i) rB[K] -= TDLS_COOP_LUPP_DYN_A(r, i) * sB[0];
724 }
725 sync();
726 }
727 }
728
749 template<typename Sync>
750 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
751 backward_pass(const int n, const int tx, const int threads, const T* TDLS_RESTRICT A,
752 const int A_stride, T (&rB)[rows_per_thread], const int (&rowid)[rows_per_thread],
753 T* work, Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
754 T* const sB = work;
755 T* const sx = work + n;
756 for (int K = 0; K < rows_per_thread; ++K) {
757 if (tx + K * threads >= n) break;
758 sB[rowid[K]] = rB[K];
759 }
760 sync();
761
762 for (int i = n - 1; i >= 0; i--) {
763 for (int K = 0; K < rows_per_thread; ++K) {
764 const int r = tx + K * threads;
765 if (r >= n) break;
766 sx[rowid[K]] = TDLS_COOP_LUPP_DYN_A(r, i);
767 }
768 sync();
769
770 const T reg = sB[i] / sx[i];
771
772 for (int K = 0; K < rows_per_thread; ++K) {
773 const int idx = tx + K * threads;
774 if (idx >= n) break;
775 if (idx < i) sB[idx] = sB[idx] - reg * sx[idx];
776 rB[K] = (idx == i) ? reg : rB[K];
777 }
778 sync();
779 }
780 }
781
782 /* =====================================================================
783 SUBSTITUTION of S^T x = b, the row-major layout: the passes of the
784 compile-time solver, U^T z = b then L^T w = z on the positions held
785 at the entries of the same index, one barrier per step, then
786 x = P^T w, with the rows of S in memory. The positions of the rows
787 are published once; then each thread loads the entries of S it needs
788 itself, before the barrier of the step, so that the latency of the
789 memory sits off the chain of the steps. The arithmetic is that of the
790 compile-time solver. Workspace: the row of S at each position
791 (logical row interchanges), then a double-buffered broadcast value,
792 n + 2 elements; the first n elements stage x = P^T w at the end,
793 once no step reads them any more.
794 ===================================================================== */
795
803 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
804 row_at([[maybe_unused]] const T* rows, const int k) noexcept {
805 if constexpr (Config.row_interchange == RowInterchange::Logical)
806 return static_cast<int>(rows[k]);
807 else
808 return k;
809 }
810
835 template<bool backward_only, typename Sync>
836 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
837 transposed_passes(const int n, const int tx, const int threads, const T* TDLS_RESTRICT A,
838 const int A_stride, T (&rB)[rows_per_thread], const T* rows, T* buf,
839 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
840 if constexpr (!backward_only) {
841 for (int i = 0; i < n; i++) {
842 const int ri = row_at(rows, i);
843 const T u_ii = TDLS_COOP_LUPP_DYN_A(ri, i);
844 T u[rows_per_thread];
845 for (int K = 0; K < rows_per_thread; ++K) {
846 const int c = tx + K * threads;
847 if (c >= n) break;
848 u[K] = c > i ? TDLS_COOP_LUPP_DYN_A(ri, c) : T(0);
849 if (c == i) buf[i & 1] = rB[K];
850 }
851 sync();
852
853 const T z_i = buf[i & 1] * (T(1) / u_ii);
854 for (int K = 0; K < rows_per_thread; ++K) {
855 const int c = tx + K * threads;
856 if (c >= n) break;
857 if (c == i) rB[K] = z_i;
858 if (c > i) rB[K] -= u[K] * z_i;
859 }
860 }
861 }
862 for (int k = n - 1; k >= 1; k--) {
863 const int rk = row_at(rows, k);
864 T l[rows_per_thread];
865 for (int K = 0; K < rows_per_thread; ++K) {
866 const int c = tx + K * threads;
867 if (c >= n) break;
868 l[K] = c < k ? TDLS_COOP_LUPP_DYN_A(rk, c) : T(0);
869 if (c == k) buf[(k + 1) & 1] = rB[K];
870 }
871 sync();
872
873 const T w_k = buf[(k + 1) & 1];
874 for (int K = 0; K < rows_per_thread; ++K) {
875 const int c = tx + K * threads;
876 if (c >= n) break;
877 if (c < k) rB[K] -= l[K] * w_k;
878 }
879 }
880 }
881
882 /* =====================================================================
883 SUBSTITUTION, dispatch on the layout.
884 ===================================================================== */
885
906 template<bool backward_only, typename Sync>
907 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
908 substitution(const int n, const int tx, const int threads, const T* TDLS_RESTRICT A,
909 const int A_stride, T (&rB)[rows_per_thread], const int (&rpiv)[rows_per_thread],
910 T* work, Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
911 constexpr bool logical = Config.row_interchange == RowInterchange::Logical;
912 if constexpr (transposed) {
913 // the row of S at each position
914 if constexpr (logical) {
915 for (int K = 0; K < rows_per_thread; ++K) {
916 const int r = tx + K * threads;
917 if (r >= n) break;
918 work[rpiv[K]] = static_cast<T>(r);
919 }
920 sync();
921 }
922 transposed_passes<backward_only>(n, tx, threads, A, A_stride, rB, work, work + n, sync);
923 // x = P^T w: under logical row interchanges, row r of S takes w
924 // at its position; under physical ones, w_p goes to the original
925 // index of the row at position p.
926 T* const staged = work;
927 for (int K = 0; K < rows_per_thread; ++K) {
928 const int c = tx + K * threads;
929 if (c >= n) break;
930 staged[logical ? c : rpiv[K]] = rB[K];
931 }
932 sync();
933 for (int K = 0; K < rows_per_thread; ++K) {
934 const int c = tx + K * threads;
935 if (c >= n) break;
936 rB[K] = staged[logical ? rpiv[K] : c];
937 }
938 } else if constexpr (logical) {
939 if constexpr (!backward_only)
940 forward_pass(n, tx, threads, A, A_stride, rB, rpiv, work, sync);
941 backward_pass(n, tx, threads, A, A_stride, rB, rpiv, work, sync);
942 } else {
943 int position[rows_per_thread];
944 init_piv(tx, threads, position);
945 if constexpr (!backward_only) {
946 // Gather in the pivoted order: each slot receives the entry
947 // of the original row it holds. The entries are staged apart
948 // from the broadcast slot of the pass, which a faster thread
949 // may already write.
950 T* const staged = work + n;
951 for (int K = 0; K < rows_per_thread; ++K) {
952 const int r = tx + K * threads;
953 if (r >= n) break;
954 staged[r] = rB[K];
955 }
956 sync();
957 for (int K = 0; K < rows_per_thread; ++K) {
958 if (tx + K * threads >= n) break;
959 rB[K] = staged[rpiv[K]];
960 }
961 forward_pass(n, tx, threads, A, A_stride, rB, position, work, sync);
962 }
963 backward_pass(n, tx, threads, A, A_stride, rB, position, work, sync);
964 }
965 }
966
967 /* =====================================================================
968 Entry points.
969 ===================================================================== */
970
984 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr auto make_sync(const int n,
985 T* work) noexcept {
986 AutoSync deduce;
987 return detail::resolve_sync<0>(deduce, work, threads_per_system(n));
988 }
989
1005 template<typename Sync = tdls::AutoSync>
1006 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1007 factorize(const int n, const int tx, T* TDLS_RESTRICT A, const int A_stride,
1008 int* TDLS_RESTRICT piv, const int piv_stride, T* work,
1009 Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1010 const int threads = threads_per_system(n);
1011 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1012 T* const exchange = work + detail::barrier_elements;
1013 T rB[rows_per_thread];
1014 int rpiv[rows_per_thread];
1015 init_piv(tx, threads, rpiv);
1016 const int linfo = eliminate<false>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1017 store_piv(n, tx, threads, piv, piv_stride, rpiv);
1018 group(); // every output visible to the whole group on return
1019 return linfo == 0;
1020 }
1021
1039 template<typename Sync = tdls::AutoSync>
1040 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1041 substitute(const int n, const int tx, const T* TDLS_RESTRICT A, const int A_stride,
1042 const int* TDLS_RESTRICT piv, const int piv_stride, const T* TDLS_RESTRICT b,
1043 T* TDLS_RESTRICT x, const int rhs_stride, T* work,
1044 Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1045 const int threads = threads_per_system(n);
1046 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1047 T* const exchange = work + detail::barrier_elements;
1048 T rB[rows_per_thread];
1049 int rpiv[rows_per_thread];
1050 load_piv(n, tx, threads, piv, piv_stride, rpiv);
1051 load_rhs(n, tx, threads, b, rhs_stride, rB);
1052 substitution<false>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1053 store_solution(n, tx, threads, x, rhs_stride, rB);
1054 group(); // every output visible to the whole group on return
1055 }
1056
1074 template<typename Sync = tdls::AutoSync>
1075 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1076 substitute_canonical(const int n, const int tx, const T* TDLS_RESTRICT A, const int A_stride,
1077 const int* TDLS_RESTRICT piv, const int piv_stride, const int col,
1078 T* TDLS_RESTRICT x, const int rhs_stride, T* work,
1079 Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1080 const int threads = threads_per_system(n);
1081 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1082 T* const exchange = work + detail::barrier_elements;
1083 T rB[rows_per_thread];
1084 int rpiv[rows_per_thread];
1085 load_piv(n, tx, threads, piv, piv_stride, rpiv);
1086 for (int K = 0; K < rows_per_thread; ++K)
1087 rB[K] = (tx + K * threads == col) ? T(1) : T(0);
1088 substitution<false>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1089 store_solution(n, tx, threads, x, rhs_stride, rB);
1090 group(); // every output visible to the whole group on return
1091 }
1092
1113 template<typename Sync = tdls::AutoSync>
1114 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1115 substitute_inplace(const int n, const int tx, const T* TDLS_RESTRICT A, const int A_stride,
1116 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1117 const int rhs_stride, T* work,
1118 Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1119 const int threads = threads_per_system(n);
1120 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1121 T* const exchange = work + detail::barrier_elements;
1122 T rB[rows_per_thread];
1123 int rpiv[rows_per_thread];
1124 load_piv(n, tx, threads, piv, piv_stride, rpiv);
1125 load_rhs(n, tx, threads, x, rhs_stride, rB);
1126 substitution<false>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1127 store_solution(n, tx, threads, x, rhs_stride, rB);
1128 group(); // every output visible to the whole group on return
1129 }
1130
1154 template<typename Sync = tdls::AutoSync>
1155 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1156 solve(const int n, const int tx, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1157 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
1158 T* work, Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1159 const int threads = threads_per_system(n);
1160 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1161 T* const exchange = work + detail::barrier_elements;
1162 T rB[rows_per_thread];
1163 int rpiv[rows_per_thread];
1164 init_piv(tx, threads, rpiv);
1165 const int linfo = eliminate<false>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1166 store_piv(n, tx, threads, piv, piv_stride, rpiv);
1167 load_rhs(n, tx, threads, b, rhs_stride, rB);
1168 substitution<false>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1169 store_solution(n, tx, threads, x, rhs_stride, rB);
1170 group(); // every output visible to the whole group on return
1171 return linfo == 0;
1172 }
1173
1198 template<typename Sync = tdls::AutoSync>
1199 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1200 solve_inplace(const int n, const int tx, T* TDLS_RESTRICT A, const int A_stride,
1201 int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT y,
1202 const int rhs_stride, T* work,
1203 Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1204 const int threads = threads_per_system(n);
1205 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1206 T* const exchange = work + detail::barrier_elements;
1207 T rB[rows_per_thread];
1208 int rpiv[rows_per_thread];
1209 init_piv(tx, threads, rpiv);
1210 load_rhs(n, tx, threads, y, rhs_stride, rB);
1211 const int linfo = eliminate<true>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1212 store_piv(n, tx, threads, piv, piv_stride, rpiv);
1213 substitution<true>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1214 for (int K = 0; K < rows_per_thread; ++K) {
1215 const int r = tx + K * threads;
1216 if (r >= n) break;
1217 TDLS_COOP_LUPP_DYN_Y(r) = rB[K];
1218 }
1219 group(); // every output visible to the whole group on return
1220 return linfo == 0;
1221 }
1222};
1223
1224
1225
1226#undef TDLS_COOP_LUPP_DYN_A
1227#undef TDLS_COOP_LUPP_DYN_PIV
1228#undef TDLS_COOP_LUPP_DYN_B
1229#undef TDLS_COOP_LUPP_DYN_X
1230#undef TDLS_COOP_LUPP_DYN_Y
1231
1232
1233
1234} // namespace tdls
1235
1236#if defined(TDLS_DETAIL_COMPILER_GCC)
1237#pragma GCC diagnostic pop
1238#endif
1239
1240
1241
1242#endif // TDLS_SOLVERS_COOPERATIVE_LUPP_SOLVER_DYNAMIC_HPP
Compile-time configuration of the CooperativeLUpp solver family.
#define TDLS_COOP_LUPP_DYN_X(r)
Strided entry r of the solution.
Definition solver_dynamic.hpp:173
#define TDLS_COOP_LUPP_DYN_B(r)
Strided entry r of the right-hand side.
Definition solver_dynamic.hpp:170
#define TDLS_COOP_LUPP_DYN_PIV(r)
Strided pivot entry of row r.
Definition solver_dynamic.hpp:167
#define TDLS_COOP_LUPP_DYN_A(r, c)
Strided element (r, c) of S, the stored matrix read column-major (A under the column-major layout,...
Definition solver_dynamic.hpp:164
#define TDLS_COOP_LUPP_DYN_Y(r)
Strided entry r of the right-hand side solved in place.
Definition solver_dynamic.hpp:176
Barriers of the groups of threads that share a system: the barrier deduced from the compilation targe...
Toolchain detection and portability macros.
#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_EXEC_CHECK_DISABLE
Lifts the execution-space check of nvcc on the next function template; empty elsewhere.
Definition macros.hpp:274
#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.
Barrier argument of the cooperative solvers that asks them to deduce the barrier from the compilation...
Definition group.hpp:268
Runtime-size dense LU factorization with partial pivoting shared by a group of threads,...
Definition solver_dynamic.hpp:208
static constexpr bool transposed
Whether the solve uses the transpose of the factored matrix, as in the compile-time solver: S = A and...
Definition solver_dynamic.hpp:234
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool factorize(const int n, const int tx, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
In-place LU factorization with partial pivoting.
Definition solver_dynamic.hpp:1007
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_piv(const int n, const int tx, const int threads, const int *TDLS_RESTRICT piv, const int piv_stride, int(&rpiv)[rows_per_thread]) noexcept
Load the pivot entries recorded by factorize.
Definition solver_dynamic.hpp:295
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_canonical(const int n, const int tx, 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, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
Solve with b = e_col generated on the fly: the consistent-tangent-operator path.
Definition solver_dynamic.hpp:1076
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute(const int n, const int tx, 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, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
Solve from a factorization produced by factorize (b and x distinct).
Definition solver_dynamic.hpp:1041
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void forward_pass(const int n, const int tx, const int threads, const T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], const int(&rowid)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Forward pass: L y = P b on the given row positions, y overwriting the right-hand-side entries of the ...
Definition solver_dynamic.hpp:709
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr auto make_sync(const int n, T *work) noexcept
The barrier deduced for the group of the calling thread.
Definition solver_dynamic.hpp:984
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void init_piv(const int tx, const int threads, int(&rpiv)[rows_per_thread]) noexcept
Initial pivot entries: no row has moved yet, so the entry of each slot is the index of its row,...
Definition solver_dynamic.hpp:282
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_elimination_step(const int n, const int i, const int tx, const int threads, T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], int(&origin)[rows_per_thread], T *work, int &linfo, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
One column of the factorization under physical row interchanges, as in the compile-time solver: the r...
Definition solver_dynamic.hpp:532
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_passes(const int n, const int tx, const int threads, const T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], const T *rows, T *buf, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Forward then backward pass of S^T x = b: b in the entries of the thread on entry, w at the entries of...
Definition solver_dynamic.hpp:837
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_inplace(const int n, const int tx, 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, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
Solve in place from a factorization produced by factorize: x holds the right-hand side on entry and t...
Definition solver_dynamic.hpp:1115
static constexpr T relative_pivot_threshold
Relative threshold of the pivot choice, read once from the configuration (see CooperativeLUppConfig::...
Definition solver_dynamic.hpp:227
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve(const int n, const int tx, 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, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
factorize + substitute in one call, the pivot entries staying with the threads in between.
Definition solver_dynamic.hpp:1156
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_rhs(const int n, const int tx, const int threads, const T *TDLS_RESTRICT b, const int rhs_stride, T(&rB)[rows_per_thread]) noexcept
Load the right-hand-side entries of the rows of the calling thread.
Definition solver_dynamic.hpp:329
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitution(const int n, const int tx, const int threads, const T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], const int(&rpiv)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Substitution from the factored rows, the right-hand side in the entries of the thread on entry,...
Definition solver_dynamic.hpp:908
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void save_row(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int r, T *record) noexcept
Copy a row of the matrix to the first n entries of an exchange record of the workspace.
Definition solver_dynamic.hpp:490
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void logical_elimination_step(const int n, const int i, const int tx, const int threads, T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], int(&rowid)[rows_per_thread], T *sB, T *sx, T *dsx, int &linfo, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
One column of the factorization under logical row interchanges. With fuse_rhs, the forward pass of th...
Definition solver_dynamic.hpp:406
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int row_at(const T *rows, const int k) noexcept
Index of the row of S at position k: read from the published positions under logical row interchanges...
Definition solver_dynamic.hpp:804
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int eliminate(const int n, const int tx, const int threads, T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], int(&rpiv)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Factorization of the rows of the group, in place.
Definition solver_dynamic.hpp:662
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace(const int n, const int tx, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
Factorization with the forward substitution folded in: the original MAGMA kernel.
Definition solver_dynamic.hpp:1200
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool keeps_pivot(const T in_place, const T column_max) noexcept
Relative threshold of the pivot choice: whether the row in place keeps the pivot (see CooperativeLUpp...
Definition solver_dynamic.hpp:375
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void store_solution(const int n, const int tx, const int threads, T *TDLS_RESTRICT x, const int rhs_stride, const T(&rB)[rows_per_thread]) noexcept
Store the solution entries held by the calling thread.
Definition solver_dynamic.hpp:346
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void store_piv(const int n, const int tx, const int threads, int *TDLS_RESTRICT piv, const int piv_stride, const int(&rpiv)[rows_per_thread]) noexcept
Record the pivot entries of the calling thread.
Definition solver_dynamic.hpp:311
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_row(const int n, T *TDLS_RESTRICT A, const int A_stride, const int r, const T *record) noexcept
Copy the first n entries of an exchange record of the workspace to a row of the matrix,...
Definition solver_dynamic.hpp:504
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int workspace_size(const int n) noexcept
Elements of the workspace shared by the threads of a group: the two elements of the deduced CPU barri...
Definition solver_dynamic.hpp:257
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int threads_per_system(const int n) noexcept
Threads solving one system of dimension n: ceil(n / rows_per_thread).
Definition solver_dynamic.hpp:246
static constexpr T singular_floor
Singularity floor, read once from the configuration (see CooperativeLUppConfig::singular_floor).
Definition solver_dynamic.hpp:230
static constexpr int rows_per_thread
Row slots of each thread: Config.rows_per_thread. When n is smaller, one thread solves the system and...
Definition solver_dynamic.hpp:224
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void backward_pass(const int n, const int tx, const int threads, const T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], const int(&rowid)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Backward pass: U x = y on the given row positions, the forward-reduced right-hand side y in the entri...
Definition solver_dynamic.hpp:751
Detection of the compiler, of the language and of the side of the compilation pass: the one place whe...