TDLS 0.4.0
Tiny Device-callable Linear Solvers
Loading...
Searching...
No Matches
solver_static.hpp
Go to the documentation of this file.
1#ifndef TDLS_SOLVERS_COOPERATIVE_LUPP_SOLVER_STATIC_HPP
2#define TDLS_SOLVERS_COOPERATIVE_LUPP_SOLVER_STATIC_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
251
252
253
254#include <type_traits>
255#include <utility>
256
257#include <tdls/core/group.hpp>
258#include <tdls/core/macros.hpp>
259#include <tdls/core/math.hpp>
260#include <tdls/core/target.hpp>
262
263
264
265// Older gcc releases cannot prove that the rows of a thread are loaded
266// before being read: the load of a phantom-capable row slot is guarded by
267// the row index and its later uses by rowid, which starts equal to it and
268// stays at or above N for a phantom row. The resulting
269// -Wmaybe-uninitialized reports are spurious. Under UBSan instrumentation
270// at -O2 and above, gcc additionally loses the value ranges of the fully
271// unrolled loop indices and reports impossible subscripts through
272// -Warray-bounds, exactly as in the TiledLUpp headers; the exercised paths
273// are certified free of out-of-bounds accesses by the constexpr suite.
274// Both suppressions are scoped to this header and to those warnings only.
275#if defined(TDLS_DETAIL_COMPILER_GCC)
276#pragma GCC diagnostic push
277#pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
278#pragma GCC diagnostic ignored "-Warray-bounds"
279#endif
280
281
282
283// clang reports a forced unrolling that the optimizer could not perform
284// through -Wpass-failed. The unrolling requested by TDLS_UNROLL_FORCE is
285// a performance hint: a failed hint does not affect correctness. The
286// suppression is scoped to this header and to that warning only.
287#if defined(TDLS_DETAIL_COMPILER_CLANG)
288#pragma clang diagnostic push
289#pragma clang diagnostic ignored "-Wpass-failed"
290#endif
291
292
293
294// nvc++ reports through no_device_stack, under the managed memory of
295// -stdpar=gpu, every lambda that captures local objects by reference:
296// such a lambda handed to a parallel algorithm would read the host stack.
297// The lambdas of this header are called on the spot, on the stack of the
298// thread that runs the solver: the report is a false positive. The
299// suppression is scoped to this header and to that warning only.
300#if defined(TDLS_DETAIL_COMPILER_NVCXX)
301#pragma diag_push
302#pragma diag_suppress no_device_stack
303#endif
304
305namespace tdls {
306
307
308
309namespace detail {
310
316template<typename Fn, int... Ks>
318unroll_K_impl(Fn& fn, std::integer_sequence<int, Ks...>) noexcept {
319 (fn(std::integral_constant<int, Ks>{}), ...);
320}
321
333template<int R, typename Fn>
334TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void unroll_K(Fn&& fn) noexcept {
335 unroll_K_impl(fn, std::make_integer_sequence<int, R>{});
336}
337
338} // namespace detail
339
340
341
342/* Addressing macros. They expand inside member functions where the
343 parameter names (tx, A, A_stride, piv, piv_stride, b, x, y,
344 rhs_stride), the residency template booleans and the class constants
345 are in scope. K is a row slot of the calling thread, its row is
346 tx + K * threads_per_system. #undef'd at the end of this header. */
347
354#define TDLS_COOP_LUPP_A(K, c) \
355 A[internal_matrix ? unsigned((c) * rows_per_thread + (K)) \
356 : TDLS_STRIDED((c) * N + tx + (K) * threads_per_system, A_stride)]
360#define TDLS_COOP_LUPP_PIV(K) \
361 piv[internal_piv ? unsigned(K) : TDLS_STRIDED(tx + (K) * threads_per_system, piv_stride)]
365#define TDLS_COOP_LUPP_B(K) \
366 b[internal_rhs ? unsigned(K) : TDLS_STRIDED(tx + (K) * threads_per_system, rhs_stride)]
370#define TDLS_COOP_LUPP_X(K) \
371 x[internal_rhs ? unsigned(K) : TDLS_STRIDED(tx + (K) * threads_per_system, rhs_stride)]
375#define TDLS_COOP_LUPP_Y(K) \
376 y[internal_rhs ? unsigned(K) : TDLS_STRIDED(tx + (K) * threads_per_system, rhs_stride)]
377
378
379
411template<typename T, int N, CooperativeLUppConfig<T> Config = CooperativeLUppConfig<T>{}>
413
414 static_assert(N >= 1, "CooperativeLUppSolverStatic: N must be >= 1");
415 static_assert(
416 Config.row_interchange != RowInterchange::Auto,
417 "CooperativeLUppSolverStatic: row_interchange must be Logical or Physical (Auto is the "
418 "default of the blocks of ArrowheadLUpp only)");
419 static_assert(Config.rows_per_thread >= 1,
420 "CooperativeLUppSolverStatic: rows_per_thread must be >= 1");
421 static_assert(Config.relative_pivot_threshold.is_finite() && Config.singular_floor.is_finite(),
422 "CooperativeLUppSolverStatic: relative_pivot_threshold and singular_floor must "
423 "be finite (and fit a 63-bit mantissa)");
424
428 static constexpr int rows_per_thread = Config.rows_per_thread < 1 ? 1
429 : Config.rows_per_thread < N ? Config.rows_per_thread
430 : N;
432 static constexpr int threads_per_system = (N + rows_per_thread - 1) / rows_per_thread;
437 static constexpr int workspace_size =
438 detail::barrier_elements + (Config.row_interchange == RowInterchange::Logical
439 ? 3 * N
440 : 2 * N + 2 * threads_per_system + 5);
443 static constexpr T relative_pivot_threshold = Config.relative_pivot_threshold;
446 static constexpr T singular_floor = Config.singular_floor;
454 static constexpr bool transposed = Config.layout == MatrixLayout::RowMajor;
455
456 static_assert(relative_pivot_threshold > T(0) && relative_pivot_threshold <= T(1),
457 "CooperativeLUppSolverStatic: relative_pivot_threshold must be in (0, 1]");
458 static_assert(singular_floor > T(0),
459 "CooperativeLUppSolverStatic: singular_floor must be positive");
460
461 /* =====================================================================
462 Row slots.
463 Slot K of thread tx holds row tx + K * threads_per_system. When N is
464 not a multiple of rows_per_thread, the last slots hold phantom rows
465 (index >= N) for some or all threads. The classification is a
466 compile-time function of K: a slot without any real row is skipped
467 entirely, a slot full for every thread needs no guard, and only a
468 mixed slot tests its row at run time.
469 ===================================================================== */
470
474 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
475 slot_has_rows(const int K) noexcept {
476 return K * threads_per_system < N;
477 }
478
482 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
483 slot_is_full(const int K) noexcept {
484 return (K + 1) * threads_per_system <= N;
485 }
486
501 template<typename Fn>
502 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
503 at_position(const int tx, const int pos, Fn&& fn) noexcept {
504 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
505 constexpr int K = decltype(K_tag)::value;
506 if constexpr (slot_has_rows(K)) {
507 if (K == pos / threads_per_system && tx == pos % threads_per_system) fn(K_tag);
508 }
509 });
510 }
511
512 /* =====================================================================
513 Operand movement.
514 Each thread reads and writes only the entries of its own rows, so
515 the operands themselves never carry data between threads: the
516 workspace does.
517 ===================================================================== */
518
525 template<bool internal_matrix>
526 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
527 load_rows(const int tx, const T* TDLS_RESTRICT A, const int A_stride,
528 T (&rA)[rows_per_thread][N]) noexcept {
529 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
530 constexpr int K = decltype(K_tag)::value;
531 if constexpr (slot_has_rows(K)) {
532 if (slot_is_full(K) || tx + K * threads_per_system < N) {
533 if constexpr (Config.unroll_loops) {
535 for (int c = 0; c < N; ++c)
536 rA[K][c] = TDLS_COOP_LUPP_A(K, c);
537 } else {
538 for (int c = 0; c < N; ++c)
539 rA[K][c] = TDLS_COOP_LUPP_A(K, c);
540 }
541 }
542 }
543 });
544 }
545
553 template<bool internal_matrix>
554 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
555 store_rows(const int tx, T* TDLS_RESTRICT A, const int A_stride,
556 const T (&rA)[rows_per_thread][N]) noexcept {
557 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
558 constexpr int K = decltype(K_tag)::value;
559 if constexpr (slot_has_rows(K)) {
560 if (slot_is_full(K) || tx + K * threads_per_system < N) {
561 if constexpr (Config.unroll_loops) {
563 for (int c = 0; c < N; ++c)
564 TDLS_COOP_LUPP_A(K, c) = rA[K][c];
565 } else {
566 for (int c = 0; c < N; ++c)
567 TDLS_COOP_LUPP_A(K, c) = rA[K][c];
568 }
569 }
570 }
571 });
572 }
573
581 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
582 init_piv(const int tx, int (&rpiv)[rows_per_thread]) noexcept {
583 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
584 constexpr int K = decltype(K_tag)::value;
585 rpiv[K] = tx + K * threads_per_system;
586 });
587 }
588
596 template<bool internal_piv>
597 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
598 load_piv(const int tx, const int* TDLS_RESTRICT piv, const int piv_stride,
599 int (&rpiv)[rows_per_thread]) noexcept {
600 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
601 constexpr int K = decltype(K_tag)::value;
602 if (slot_is_full(K) || tx + K * threads_per_system < N)
603 rpiv[K] = TDLS_COOP_LUPP_PIV(K);
604 else
605 rpiv[K] = tx + K * threads_per_system;
606 });
607 }
608
616 template<bool internal_piv>
617 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
618 store_piv(const int tx, int* TDLS_RESTRICT piv, const int piv_stride,
619 const int (&rpiv)[rows_per_thread]) noexcept {
620 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
621 constexpr int K = decltype(K_tag)::value;
622 if constexpr (slot_has_rows(K)) {
623 if (slot_is_full(K) || tx + K * threads_per_system < N)
624 TDLS_COOP_LUPP_PIV(K) = rpiv[K];
625 }
626 });
627 }
628
637 template<bool internal_rhs>
638 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
639 load_rhs(const int tx, const T* TDLS_RESTRICT b, const int rhs_stride,
640 T (&rB)[rows_per_thread]) noexcept {
641 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
642 constexpr int K = decltype(K_tag)::value;
643 if constexpr (slot_has_rows(K)) {
644 if (slot_is_full(K) || tx + K * threads_per_system < N) rB[K] = TDLS_COOP_LUPP_B(K);
645 }
646 });
647 }
648
656 template<bool internal_rhs>
657 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
658 store_solution(const int tx, T* TDLS_RESTRICT x, const int rhs_stride,
659 const T (&rB)[rows_per_thread]) noexcept {
660 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
661 constexpr int K = decltype(K_tag)::value;
662 if constexpr (slot_has_rows(K)) {
663 if (slot_is_full(K) || tx + K * threads_per_system < N) TDLS_COOP_LUPP_X(K) = rB[K];
664 }
665 });
666 }
667
668 /* =====================================================================
669 FACTORIZATION, right-looking, one column per step, as in MAGMA.
670 Logical row interchanges, three barrier intervals per column: the
671 magnitudes of the column are published at the positions of their
672 rows, every thread finds the same pivot, the owner of the pivot row
673 publishes it, and every thread eliminates its own rows below the
674 pivot. Physical row interchanges, two barrier intervals per column:
675 each thread publishes its best candidate, every thread finds the
676 same pivot, the owner of position i publishes its row, the owner of
677 the pivot row as well when the pivot is not in place, and every
678 thread eliminates its rows below the pivot after the exchange. The
679 pivot, the operations and their order are the same in both schemes.
680 ===================================================================== */
681
688 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
689 keeps_pivot(const T in_place, const T column_max) noexcept {
690 const T bound = relative_pivot_threshold * column_max;
691 return in_place >= (bound > singular_floor ? bound : singular_floor);
692 }
693
722 template<bool fuse_rhs, typename Sync>
723 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
724 logical_elimination_step([[maybe_unused]] const int tx, const int i,
725 T (&rA)[rows_per_thread][N], [[maybe_unused]] T (&rB)[rows_per_thread],
726 int (&rowid)[rows_per_thread], T* sB, T* sx, T* dsx, int& linfo,
727 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
728 // izamax: the magnitudes of column i, at the positions of their rows.
729 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
730 constexpr int K = decltype(K_tag)::value;
731 if constexpr (slot_has_rows(K)) {
732 if (slot_is_full(K) || rowid[K] < N) dsx[rowid[K]] = detail::abs(rA[K][i]);
733 }
734 });
735 sync();
736
737 // Every thread scans the same magnitudes and finds the same pivot.
738 // The scan reads the workspace, not the registers: no pragma.
739 T rx_abs_max = dsx[i];
740 int max_id = i;
741 for (int j = i + 1; j < N; j++) {
742 if (dsx[j] > rx_abs_max) {
743 max_id = j;
744 rx_abs_max = dsx[j];
745 }
746 }
747 // Below the default threshold 1, the row in place keeps the pivot
748 // when it is large enough.
749 if constexpr (relative_pivot_threshold < T(1)) {
750 if (keeps_pivot(dsx[i], rx_abs_max)) max_id = i;
751 }
752 const bool zero_pivot = (rx_abs_max < singular_floor);
753 linfo = (zero_pivot && linfo == 0) ? (i + 1) : linfo;
754 const T update = zero_pivot ? T(0) : T(1);
755
756 // The owner of the pivot row publishes it; positions i and max_id
757 // are swapped. A phantom row never matches: its position is >= N.
758 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
759 constexpr int K = decltype(K_tag)::value;
760 if constexpr (slot_has_rows(K)) {
761 if (rowid[K] == max_id) {
762 rowid[K] = i;
763 if constexpr (Config.unroll_loops) {
765 for (int j = i; j < N; j++)
766 sx[j] = update * rA[K][j];
767 } else {
768 for (int j = i; j < N; j++)
769 sx[j] = update * rA[K][j];
770 }
771 if constexpr (fuse_rhs && !transposed) sB[0] = rB[K];
772 } else if (rowid[K] == i) {
773 rowid[K] = max_id;
774 }
775 }
776 });
777 if constexpr (fuse_rhs && transposed)
778 at_position(tx, i, [&](auto K_tag) { sB[0] = rB[decltype(K_tag)::value]; });
779 sync();
780
781 const T reg = zero_pivot ? T(1) : T(1) / sx[i];
782 [[maybe_unused]] T z_i = T(0);
783 if constexpr (fuse_rhs && transposed) z_i = sB[0] * reg;
784
785 // scal and ger: every thread eliminates its rows below the pivot.
786 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
787 constexpr int K = decltype(K_tag)::value;
788 if constexpr (slot_has_rows(K)) {
789 if ((slot_is_full(K) || rowid[K] < N) && rowid[K] > i) {
790 rA[K][i] *= reg;
791 if constexpr (Config.unroll_loops) {
793 for (int j = i + 1; j < N; j++)
794 rA[K][j] -= rA[K][i] * sx[j];
795 } else {
796 for (int j = i + 1; j < N; j++)
797 rA[K][j] -= rA[K][i] * sx[j];
798 }
799 if constexpr (fuse_rhs && !transposed) rB[K] -= rA[K][i] * sB[0];
800 }
801 if constexpr (fuse_rhs && transposed) {
802 const int c = tx + K * threads_per_system;
803 if (c == i) rB[K] = z_i;
804 if ((slot_is_full(K) || c < N) && c > i) rB[K] -= sx[c] * z_i;
805 }
806 }
807 });
808 sync();
809 }
810
821 template<bool fuse_rhs>
822 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
823 save_row(const T (&row)[N], [[maybe_unused]] const T& rhs, const int origin,
824 T* record) noexcept {
825 if constexpr (Config.unroll_loops) {
827 for (int c = 0; c < N; c++)
828 record[c] = row[c];
829 } else {
830 for (int c = 0; c < N; c++)
831 record[c] = row[c];
832 }
833 if constexpr (fuse_rhs) record[N] = rhs;
834 record[N + 1] = static_cast<T>(origin);
835 }
836
844 template<bool fuse_rhs>
845 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
846 load_row(T (&row)[N], [[maybe_unused]] T& rhs, int& origin, const T* record) noexcept {
847 if constexpr (Config.unroll_loops) {
849 for (int c = 0; c < N; c++)
850 row[c] = record[c];
851 } else {
852 for (int c = 0; c < N; c++)
853 row[c] = record[c];
854 }
855 if constexpr (fuse_rhs) rhs = record[N];
856 origin = static_cast<int>(record[N + 1]);
857 }
858
867 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool below(const int tx, const int K,
868 const int i) noexcept {
869 return (K + 1) * threads_per_system > i + 1 &&
870 (slot_is_full(K) || tx + K * threads_per_system < N);
871 }
872
889 template<bool fuse_rhs>
890 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
891 gather_row(const int tx, const int i, const int pos, const T (&rA)[rows_per_thread][N],
892 [[maybe_unused]] const T (&rB)[rows_per_thread],
893 const int (&origin)[rows_per_thread], T* record) noexcept {
894 const int slot = pos / threads_per_system;
895 const auto entry = [&](const int c) {
896 T v = T(0);
897 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
898 constexpr int K = decltype(K_tag)::value;
899 if constexpr (slot_has_rows(K)) {
900 if (below(tx, K, i)) v = K == slot ? rA[K][c] : v;
901 }
902 });
903 return v;
904 };
905 if constexpr (Config.unroll_loops) {
907 for (int c = 0; c < N; c++)
908 record[c] = entry(c);
909 } else {
910 for (int c = 0; c < N; c++)
911 record[c] = entry(c);
912 }
913 T rhs = T(0);
914 int index = 0;
915 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
916 constexpr int K = decltype(K_tag)::value;
917 if constexpr (slot_has_rows(K)) {
918 if (below(tx, K, i)) {
919 if constexpr (fuse_rhs) rhs = K == slot ? rB[K] : rhs;
920 index = K == slot ? origin[K] : index;
921 }
922 }
923 });
924 if constexpr (fuse_rhs) record[N] = rhs;
925 record[N + 1] = static_cast<T>(index);
926 }
927
940 template<bool fuse_rhs>
941 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
942 scatter_row(const int tx, const int i, const int pos, T (&rA)[rows_per_thread][N],
943 [[maybe_unused]] T (&rB)[rows_per_thread], int (&origin)[rows_per_thread],
944 const T* record) noexcept {
945 const int slot = pos / threads_per_system;
946 const auto entry = [&](const int c) {
947 const T v = record[c];
948 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
949 constexpr int K = decltype(K_tag)::value;
950 if constexpr (slot_has_rows(K)) {
951 if (below(tx, K, i)) rA[K][c] = K == slot ? v : rA[K][c];
952 }
953 });
954 };
955 if constexpr (Config.unroll_loops) {
957 for (int c = 0; c < N; c++)
958 entry(c);
959 } else {
960 for (int c = 0; c < N; c++)
961 entry(c);
962 }
963 [[maybe_unused]] const T rhs = fuse_rhs ? record[N] : T(0);
964 const int index = static_cast<int>(record[N + 1]);
965 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
966 constexpr int K = decltype(K_tag)::value;
967 if constexpr (slot_has_rows(K)) {
968 if (below(tx, K, i)) {
969 if constexpr (fuse_rhs) rB[K] = K == slot ? rhs : rB[K];
970 origin[K] = K == slot ? index : origin[K];
971 }
972 }
973 });
974 }
975
1010 template<bool fuse_rhs, typename Sync>
1011 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1012 physical_elimination_step(const int tx, const int i, T (&rA)[rows_per_thread][N],
1013 [[maybe_unused]] T (&rB)[rows_per_thread],
1014 int (&origin)[rows_per_thread], T* work, int& linfo,
1015 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1016 // The right-hand side moves with the rows only under S x = b.
1017 constexpr bool carries_rhs = fuse_rhs && !transposed;
1018 // Workspace: the candidates and their positions, the magnitude of
1019 // the row in place, then the exchange records of the rows leaving
1020 // and entering position i. Positions travel as values of T, as in
1021 // save_row.
1022 T* const candidate = work;
1023 T* const where = candidate + threads_per_system;
1024 T* const in_place = where + threads_per_system;
1025 T* const leaving = in_place + 1;
1026 T* const entering = leaving + N + 2;
1027
1028 // Candidate of the thread. The owner of position i starts from the
1029 // row in place, so that a NaN there stops the search, as in the
1030 // logical scan, and publishes its magnitude.
1031 T best = T(-1);
1032 int best_pos = N;
1033 T local = T(0);
1034 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1035 constexpr int K = decltype(K_tag)::value;
1036 if constexpr (slot_has_rows(K)) {
1037 const int pos = tx + K * threads_per_system;
1038 if ((slot_is_full(K) || pos < N) && pos >= i) {
1039 const T v = detail::abs(rA[K][i]);
1040 if (pos == i || v > best) {
1041 best = v;
1042 best_pos = pos;
1043 }
1044 if (pos == i) local = v;
1045 }
1046 }
1047 });
1048 candidate[tx] = best;
1049 where[tx] = static_cast<T>(best_pos);
1050 if (tx == i % threads_per_system) in_place[0] = local;
1051 sync();
1052
1053 // Every thread finds the same column maximum. The row in place keeps
1054 // the pivot when it reaches it, as in the logical scan, which keeps
1055 // the first maximum; otherwise the pivot is the first position
1056 // holding it. Only that rare case reads the positions. The loops
1057 // read the workspace, not the registers: no pragma.
1058 T rx_abs_max = candidate[i % threads_per_system];
1059 for (int t = 0; t < threads_per_system; t++)
1060 rx_abs_max = candidate[t] > rx_abs_max ? candidate[t] : rx_abs_max;
1061 bool keep = !(in_place[0] < rx_abs_max);
1062 if constexpr (relative_pivot_threshold < T(1)) keep = keeps_pivot(in_place[0], rx_abs_max);
1063 int max_id = i;
1064 if (!keep) {
1065 max_id = N;
1066 for (int t = 0; t < threads_per_system; t++) {
1067 const int pos = static_cast<int>(where[t]);
1068 if (candidate[t] == rx_abs_max && pos < max_id) max_id = pos;
1069 }
1070 }
1071 const bool zero_pivot = (rx_abs_max < singular_floor);
1072 linfo = (zero_pivot && linfo == 0) ? (i + 1) : linfo;
1073
1074 // The owner of position i publishes its row whole, from a slot
1075 // known at compile time in an unrolled sweep. When the pivot is not
1076 // in place, the owner of position max_id publishes the pivot row as
1077 // well: it enters position i, and the row in place leaves it for
1078 // position max_id. The elimination reads the pivot row and its
1079 // right-hand-side entry from the record of the row that ends at
1080 // position i, under S x = b; under S^T x = b, the entry b'(i) of
1081 // the owner of position i, from the record of the row in place. A
1082 // zero pivot is not scaled away as in the logical scheme: the
1083 // factorization of a singular matrix is unspecified, and reg = 1
1084 // keeps it finite.
1085 at_position(tx, i, [&](auto K_tag) {
1086 constexpr int K = decltype(K_tag)::value;
1087 save_row<fuse_rhs>(rA[K], rB[K], origin[K], leaving);
1088 });
1089 if (max_id != i && tx == max_id % threads_per_system)
1090 gather_row<carries_rhs>(tx, i, max_id, rA, rB, origin, entering);
1091 sync();
1092
1093 if (max_id != i) {
1094 at_position(tx, i, [&](auto K_tag) {
1095 constexpr int K = decltype(K_tag)::value;
1096 load_row<carries_rhs>(rA[K], rB[K], origin[K], entering);
1097 });
1098 if (tx == max_id % threads_per_system)
1099 scatter_row<carries_rhs>(tx, i, max_id, rA, rB, origin, leaving);
1100 }
1101
1102 const T* const prow = max_id == i ? leaving : entering;
1103 const T reg = zero_pivot ? T(1) : T(1) / prow[i];
1104 [[maybe_unused]] T z_i = T(0);
1105 if constexpr (fuse_rhs && transposed) z_i = leaving[N] * reg;
1106
1107 // scal and ger: every thread eliminates its rows below the pivot.
1108 // The next step writes only the candidates, which no thread reads
1109 // any more: no barrier closes the step.
1110 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1111 constexpr int K = decltype(K_tag)::value;
1112 if constexpr (slot_has_rows(K)) {
1113 const int pos = tx + K * threads_per_system;
1114 if ((slot_is_full(K) || pos < N) && pos > i) {
1115 rA[K][i] *= reg;
1116 if constexpr (Config.unroll_loops) {
1118 for (int j = i + 1; j < N; j++)
1119 rA[K][j] -= rA[K][i] * prow[j];
1120 } else {
1121 for (int j = i + 1; j < N; j++)
1122 rA[K][j] -= rA[K][i] * prow[j];
1123 }
1124 if constexpr (carries_rhs) rB[K] -= rA[K][i] * prow[N];
1125 }
1126 if constexpr (fuse_rhs && transposed) {
1127 if (pos == i) rB[K] = z_i;
1128 if ((slot_is_full(K) || pos < N) && pos > i) rB[K] -= prow[pos] * z_i;
1129 }
1130 }
1131 });
1132 }
1133
1148 template<bool fuse_rhs, typename Sync>
1149 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
1150 eliminate(const int tx, T (&rA)[rows_per_thread][N], T (&rB)[rows_per_thread],
1151 int (&rpiv)[rows_per_thread], T* work,
1152 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1153 int linfo = 0;
1154 if constexpr (Config.row_interchange == RowInterchange::Logical) {
1155 T* const sB = work;
1156 T* const sx = work + N;
1157 T* const dsx = work + 2 * N;
1158 if constexpr (Config.unroll_loops) {
1160 for (int i = 0; i < N; i++)
1161 logical_elimination_step<fuse_rhs>(tx, i, rA, rB, rpiv, sB, sx, dsx, linfo,
1162 sync);
1163 } else {
1164 for (int i = 0; i < N; i++)
1165 logical_elimination_step<fuse_rhs>(tx, i, rA, rB, rpiv, sB, sx, dsx, linfo,
1166 sync);
1167 }
1168 } else {
1169 if constexpr (Config.unroll_loops) {
1171 for (int i = 0; i < N; i++)
1172 physical_elimination_step<fuse_rhs>(tx, i, rA, rB, rpiv, work, linfo, sync);
1173 } else {
1174 for (int i = 0; i < N; i++)
1175 physical_elimination_step<fuse_rhs>(tx, i, rA, rB, rpiv, work, linfo, sync);
1176 }
1177 sync(); // no barrier closes the steps: this one frees the workspace
1178 }
1179 return linfo;
1180 }
1181
1182 /* =====================================================================
1183 SUBSTITUTION of S x = b, the column-major layout.
1184 Forward: the owner of the row at position i publishes its
1185 right-hand-side entry, every thread updates its rows below. This is
1186 the elimination of the right-hand side that solve_inplace folds into
1187 the factorization, replayed from the factored rows. Backward: the
1188 right-hand side is staged at the positions of its rows, then column
1189 i of U is published, every thread computes x_i, the owner of entry i
1190 keeps it and every thread updates its entries above. Under physical
1191 row interchanges, the right-hand side is first gathered in the
1192 pivoted order, and both passes run on the positions of the rows,
1193 which are static: the arithmetic is the logical one by construction,
1194 down to the contractions a compiler applies, and in an unrolled
1195 sweep the owner of each position is known at compile time.
1196 ===================================================================== */
1197
1208 template<typename Sync>
1209 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1210 forward_step(const int i, const T (&rA)[rows_per_thread][N], T (&rB)[rows_per_thread],
1211 const int (&rowid)[rows_per_thread], T* sB,
1212 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1213 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1214 constexpr int K = decltype(K_tag)::value;
1215 if constexpr (slot_has_rows(K)) {
1216 if (rowid[K] == i) sB[0] = rB[K];
1217 }
1218 });
1219 sync();
1220
1221 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1222 constexpr int K = decltype(K_tag)::value;
1223 if constexpr (slot_has_rows(K)) {
1224 if ((slot_is_full(K) || rowid[K] < N) && rowid[K] > i) rB[K] -= rA[K][i] * sB[0];
1225 }
1226 });
1227 sync();
1228 }
1229
1239 template<typename Sync>
1240 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1241 forward_pass(const T (&rA)[rows_per_thread][N], T (&rB)[rows_per_thread],
1242 const int (&rowid)[rows_per_thread], T* work,
1243 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1244 T* const sB = work;
1245 if constexpr (Config.unroll_loops) {
1247 for (int i = 0; i < N; i++)
1248 forward_step(i, rA, rB, rowid, sB, sync);
1249 } else {
1250 for (int i = 0; i < N; i++)
1251 forward_step(i, rA, rB, rowid, sB, sync);
1252 }
1253 }
1254
1273 template<typename Sync>
1274 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1275 backward_step(const int i, const int tx, const T (&rA)[rows_per_thread][N],
1276 T (&rB)[rows_per_thread], const int (&rowid)[rows_per_thread], T* sB, T* sx,
1277 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1278 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1279 constexpr int K = decltype(K_tag)::value;
1280 if constexpr (slot_has_rows(K)) {
1281 if (slot_is_full(K) || rowid[K] < N) sx[rowid[K]] = rA[K][i];
1282 }
1283 });
1284 sync();
1285
1286 const T reg = sB[i] / sx[i];
1287
1288 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1289 constexpr int K = decltype(K_tag)::value;
1290 if constexpr (slot_has_rows(K)) {
1291 const int idx = tx + K * threads_per_system;
1292 if (slot_is_full(K) || idx < N) {
1293 if (idx < i) sB[idx] = sB[idx] - reg * sx[idx];
1294 rB[K] = (idx == i) ? reg : rB[K];
1295 }
1296 }
1297 });
1298 sync();
1299 }
1300
1314 template<typename Sync>
1315 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1316 backward_pass(const int tx, const T (&rA)[rows_per_thread][N], T (&rB)[rows_per_thread],
1317 const int (&rowid)[rows_per_thread], T* work,
1318 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1319 T* const sB = work;
1320 T* const sx = work + N;
1321 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1322 constexpr int K = decltype(K_tag)::value;
1323 if constexpr (slot_has_rows(K)) {
1324 if (slot_is_full(K) || rowid[K] < N) sB[rowid[K]] = rB[K];
1325 }
1326 });
1327 sync();
1328
1329 if constexpr (Config.unroll_loops) {
1331 for (int i = N - 1; i >= 0; i--)
1332 backward_step(i, tx, rA, rB, rowid, sB, sx, sync);
1333 } else {
1334 for (int i = N - 1; i >= 0; i--)
1335 backward_step(i, tx, rA, rB, rowid, sB, sx, sync);
1336 }
1337 }
1338
1352 template<typename Sync>
1353 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1354 gather_rhs(const int tx, T (&rB)[rows_per_thread], const int (&origin)[rows_per_thread],
1355 T* staged, Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1356 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1357 constexpr int K = decltype(K_tag)::value;
1358 if constexpr (slot_has_rows(K)) {
1359 if (slot_is_full(K) || tx + K * threads_per_system < N)
1360 staged[tx + K * threads_per_system] = rB[K];
1361 }
1362 });
1363 sync();
1364
1365 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1366 constexpr int K = decltype(K_tag)::value;
1367 if constexpr (slot_has_rows(K)) {
1368 if (slot_is_full(K) || tx + K * threads_per_system < N) rB[K] = staged[origin[K]];
1369 }
1370 });
1371 }
1372
1373 /* =====================================================================
1374 SUBSTITUTION of S^T x = b, the row-major layout.
1375 With P S = L U, S^T = U^T L^T P. Both passes run on the positions of
1376 the rows, held at the entries of the same index: z_p and w_p live in
1377 the slot of entry p, which b'(p) leaves when it is consumed, so no
1378 second vector is needed. Forward, U^T z = b: the row at position i
1379 publishes its row of U, the owner of entry i its b'(i), every thread
1380 computes z_i = b'(i) / U(i, i) through the reciprocal of the pivot,
1381 the owner of entry i keeps z_i, and each entry c > i loses
1382 U(i, c) z_i. This is the forward pass that solve_inplace folds into
1383 the factorization, replayed from the factored rows. Backward,
1384 L^T w = z: the row at position k publishes its multipliers L(k, c),
1385 c < k, the owner of entry k its w_k, and each entry c < k loses
1386 L(k, c) w_k; L has a unit diagonal, so no division sits on the chain
1387 of the steps. One barrier per step: each step writes its own half of
1388 a double buffer, which no thread still reads (see segment). Then
1389 x = P^T w, one exchange through the workspace: under logical row
1390 interchanges, row r of S takes w at its position; under physical
1391 ones, w_p goes to the original index of the row at position p. These
1392 are the mirror of the staging of the right-hand side that the
1393 substitutions of S x = b make at the start.
1394 ===================================================================== */
1395
1407 const int step) noexcept {
1408 return work + (step & 1) * (N + 1);
1409 }
1410
1423 template<typename Sync>
1424 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1425 transposed_forward_step(const int tx, const int i, const T (&rA)[rows_per_thread][N],
1426 T (&rB)[rows_per_thread], const int (&rowid)[rows_per_thread], T* work,
1427 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1428 T* const sx = segment(work, i);
1429 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1430 constexpr int K = decltype(K_tag)::value;
1431 if constexpr (slot_has_rows(K)) {
1432 if (rowid[K] == i) {
1433 // Unrolled, the copy runs over all N entries under a
1434 // guard: a trip count set by the step lets the AMD
1435 // compiler unroll it before the step loop, and the
1436 // rows then fall to scratch memory.
1437 if constexpr (Config.unroll_loops) {
1439 for (int j = 0; j < N; j++)
1440 if (j >= i) sx[j] = rA[K][j];
1441 } else {
1442 for (int j = i; j < N; j++)
1443 sx[j] = rA[K][j];
1444 }
1445 }
1446 }
1447 });
1448 at_position(tx, i, [&](auto K_tag) { sx[N] = rB[decltype(K_tag)::value]; });
1449 sync();
1450
1451 const T z_i = sx[N] * (T(1) / sx[i]);
1452 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1453 constexpr int K = decltype(K_tag)::value;
1454 if constexpr (slot_has_rows(K)) {
1455 const int c = tx + K * threads_per_system;
1456 if (c == i) rB[K] = z_i;
1457 if ((slot_is_full(K) || c < N) && c > i) rB[K] -= sx[c] * z_i;
1458 }
1459 });
1460 }
1461
1478 template<typename Sync>
1479 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1480 transposed_backward_step(const int tx, const int k, const T (&rA)[rows_per_thread][N],
1481 T (&rB)[rows_per_thread], const int (&rowid)[rows_per_thread], T* work,
1482 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1483 T* const sx = segment(work, k + 1);
1484 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1485 constexpr int K = decltype(K_tag)::value;
1486 if constexpr (slot_has_rows(K)) {
1487 if (rowid[K] == k) {
1488 // Unrolled over all N entries under a guard, as in the
1489 // forward step.
1490 if constexpr (Config.unroll_loops) {
1492 for (int c = 0; c < N; c++)
1493 if (c < k) sx[c] = rA[K][c];
1494 } else {
1495 for (int c = 0; c < k; c++)
1496 sx[c] = rA[K][c];
1497 }
1498 }
1499 }
1500 });
1501 at_position(tx, k, [&](auto K_tag) { sx[N] = rB[decltype(K_tag)::value]; });
1502 sync();
1503
1504 const T w_k = sx[N];
1505 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1506 constexpr int K = decltype(K_tag)::value;
1507 if constexpr (slot_has_rows(K)) {
1508 const int c = tx + K * threads_per_system;
1509 if (c < k) rB[K] -= sx[c] * w_k;
1510 }
1511 });
1512 }
1513
1529 template<bool backward_only, typename Sync>
1530 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1531 transposed_passes(const int tx, const T (&rA)[rows_per_thread][N], T (&rB)[rows_per_thread],
1532 const int (&rowid)[rows_per_thread], T* work,
1533 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1534 if constexpr (Config.unroll_loops) {
1535 if constexpr (!backward_only) {
1537 for (int i = 0; i < N; i++)
1538 transposed_forward_step(tx, i, rA, rB, rowid, work, sync);
1539 }
1541 for (int k = N - 1; k >= 1; k--)
1542 transposed_backward_step(tx, k, rA, rB, rowid, work, sync);
1543 } else {
1544 if constexpr (!backward_only) {
1545 for (int i = 0; i < N; i++)
1546 transposed_forward_step(tx, i, rA, rB, rowid, work, sync);
1547 }
1548 for (int k = N - 1; k >= 1; k--)
1549 transposed_backward_step(tx, k, rA, rB, rowid, work, sync);
1550 }
1551 }
1552
1567 template<typename Sync>
1568 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1569 unpermute_solution(const int tx, T (&rB)[rows_per_thread], const int (&rpiv)[rows_per_thread],
1570 T* work, Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1571 constexpr bool logical = Config.row_interchange == RowInterchange::Logical;
1572 T* const staged = segment(work, 1);
1573 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1574 constexpr int K = decltype(K_tag)::value;
1575 if constexpr (slot_has_rows(K)) {
1576 if (slot_is_full(K) || tx + K * threads_per_system < N)
1577 staged[logical ? tx + K * threads_per_system : rpiv[K]] = rB[K];
1578 }
1579 });
1580 sync();
1581
1582 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1583 constexpr int K = decltype(K_tag)::value;
1584 if constexpr (slot_has_rows(K)) {
1585 if (slot_is_full(K) || tx + K * threads_per_system < N)
1586 rB[K] = staged[logical ? rpiv[K] : tx + K * threads_per_system];
1587 }
1588 });
1589 }
1590
1591 /* =====================================================================
1592 SUBSTITUTION, dispatch on the layout.
1593 ===================================================================== */
1594
1611 template<bool backward_only, typename Sync>
1612 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1613 substitution(const int tx, const T (&rA)[rows_per_thread][N], T (&rB)[rows_per_thread],
1614 const int (&rpiv)[rows_per_thread], T* work,
1615 Sync& sync) noexcept(detail::nothrow_sync<Sync>) {
1616 constexpr bool logical = Config.row_interchange == RowInterchange::Logical;
1617 if constexpr (transposed) {
1618 if constexpr (logical) {
1619 transposed_passes<backward_only>(tx, rA, rB, rpiv, work, sync);
1620 } else {
1621 int position[rows_per_thread];
1622 init_piv(tx, position);
1623 transposed_passes<backward_only>(tx, rA, rB, position, work, sync);
1624 }
1625 unpermute_solution(tx, rB, rpiv, work, sync);
1626 } else if constexpr (logical) {
1627 if constexpr (!backward_only) forward_pass(rA, rB, rpiv, work, sync);
1628 backward_pass(tx, rA, rB, rpiv, work, sync);
1629 } else {
1630 int position[rows_per_thread];
1631 init_piv(tx, position);
1632 if constexpr (!backward_only) {
1633 // Staged apart from the broadcast slot of the pass, which a
1634 // faster thread may already write.
1635 gather_rhs(tx, rB, rpiv, work + N, sync);
1636 forward_pass(rA, rB, position, work, sync);
1637 }
1638 backward_pass(tx, rA, rB, position, work, sync);
1639 }
1640 }
1641
1642 /* =====================================================================
1643 Entry points.
1644 ===================================================================== */
1645
1658 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr auto make_sync(T* work) noexcept {
1659 AutoSync deduce;
1660 return detail::resolve_sync<threads_per_system>(deduce, work, threads_per_system);
1661 }
1662
1684 template<bool internal_piv, bool internal_matrix, int A_static_stride = dynamic_stride,
1685 int piv_static_stride = dynamic_stride, typename Sync = tdls::AutoSync>
1686 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1687 factorize(const int tx, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1688 const int piv_stride, T* work,
1689 Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1690 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride>) {
1692 tx, A, detail::stride_or<A_static_stride>(A_stride), piv,
1693 detail::stride_or<piv_static_stride>(piv_stride), work, std::forward<Sync>(sync));
1694 } else {
1695 auto&& group = detail::resolve_sync<threads_per_system>(sync, work, threads_per_system);
1696 T* const exchange = work + detail::barrier_elements;
1697 T rA[rows_per_thread][N];
1698 T rB[rows_per_thread];
1699 int rpiv[rows_per_thread];
1700 load_rows<internal_matrix>(tx, A, A_stride, rA);
1701 init_piv(tx, rpiv);
1702 const int linfo = eliminate<false>(tx, rA, rB, rpiv, exchange, group);
1703 store_rows<internal_matrix>(tx, A, A_stride, rA);
1704 store_piv<internal_piv>(tx, piv, piv_stride, rpiv);
1705 group(); // every output visible to the whole group on return
1706 return linfo == 0;
1707 }
1708 }
1709
1735 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
1736 int A_static_stride = dynamic_stride, int piv_static_stride = dynamic_stride,
1737 int rhs_static_stride = dynamic_stride, typename Sync = tdls::AutoSync>
1738 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1739 substitute(const int tx, const T* TDLS_RESTRICT A, const int A_stride,
1740 const int* TDLS_RESTRICT piv, const int piv_stride, const T* TDLS_RESTRICT b,
1741 T* TDLS_RESTRICT x, const int rhs_stride, T* work,
1742 Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1743 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
1744 rhs_static_stride>) {
1746 tx, A, detail::stride_or<A_static_stride>(A_stride), piv,
1747 detail::stride_or<piv_static_stride>(piv_stride), b, x,
1748 detail::stride_or<rhs_static_stride>(rhs_stride), work, std::forward<Sync>(sync));
1749 } else {
1750 auto&& group = detail::resolve_sync<threads_per_system>(sync, work, threads_per_system);
1751 T* const exchange = work + detail::barrier_elements;
1752 T rA[rows_per_thread][N];
1753 T rB[rows_per_thread];
1754 int rpiv[rows_per_thread];
1755 load_rows<internal_matrix>(tx, A, A_stride, rA);
1756 load_piv<internal_piv>(tx, piv, piv_stride, rpiv);
1757 load_rhs<internal_rhs>(tx, b, rhs_stride, rB);
1758 substitution<false>(tx, rA, rB, rpiv, exchange, group);
1759 store_solution<internal_rhs>(tx, x, rhs_stride, rB);
1760 group(); // every output visible to the whole group on return
1761 }
1762 }
1763
1789 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
1790 int A_static_stride = dynamic_stride, int piv_static_stride = dynamic_stride,
1791 int rhs_static_stride = dynamic_stride, typename Sync = tdls::AutoSync>
1792 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1793 substitute_canonical(const int tx, const T* TDLS_RESTRICT A, const int A_stride,
1794 const int* TDLS_RESTRICT piv, const int piv_stride, const int col,
1795 T* TDLS_RESTRICT x, const int rhs_stride, T* work,
1796 Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1797 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
1798 rhs_static_stride>) {
1800 tx, A, detail::stride_or<A_static_stride>(A_stride), piv,
1801 detail::stride_or<piv_static_stride>(piv_stride), col, x,
1802 detail::stride_or<rhs_static_stride>(rhs_stride), work, std::forward<Sync>(sync));
1803 } else {
1804 auto&& group = detail::resolve_sync<threads_per_system>(sync, work, threads_per_system);
1805 T* const exchange = work + detail::barrier_elements;
1806 T rA[rows_per_thread][N];
1807 T rB[rows_per_thread];
1808 int rpiv[rows_per_thread];
1809 load_rows<internal_matrix>(tx, A, A_stride, rA);
1810 load_piv<internal_piv>(tx, piv, piv_stride, rpiv);
1811 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
1812 constexpr int K = decltype(K_tag)::value;
1813 rB[K] = (tx + K * threads_per_system == col) ? T(1) : T(0);
1814 });
1815 substitution<false>(tx, rA, rB, rpiv, exchange, group);
1816 store_solution<internal_rhs>(tx, x, rhs_stride, rB);
1817 group(); // every output visible to the whole group on return
1818 }
1819 }
1820
1849 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
1850 int A_static_stride = dynamic_stride, int piv_static_stride = dynamic_stride,
1851 int rhs_static_stride = dynamic_stride, typename Sync = tdls::AutoSync>
1852 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
1853 substitute_inplace(const int tx, const T* TDLS_RESTRICT A, const int A_stride,
1854 const int* TDLS_RESTRICT piv, const int piv_stride, T* TDLS_RESTRICT x,
1855 const int rhs_stride, T* work,
1856 Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1857 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
1858 rhs_static_stride>) {
1860 tx, A, detail::stride_or<A_static_stride>(A_stride), piv,
1861 detail::stride_or<piv_static_stride>(piv_stride), x,
1862 detail::stride_or<rhs_static_stride>(rhs_stride), work, std::forward<Sync>(sync));
1863 } else {
1864 auto&& group = detail::resolve_sync<threads_per_system>(sync, work, threads_per_system);
1865 T* const exchange = work + detail::barrier_elements;
1866 T rA[rows_per_thread][N];
1867 T rB[rows_per_thread];
1868 int rpiv[rows_per_thread];
1869 load_rows<internal_matrix>(tx, A, A_stride, rA);
1870 load_piv<internal_piv>(tx, piv, piv_stride, rpiv);
1871 load_rhs<internal_rhs>(tx, x, rhs_stride, rB);
1872 substitution<false>(tx, rA, rB, rpiv, exchange, group);
1873 store_solution<internal_rhs>(tx, x, rhs_stride, rB);
1874 group(); // every output visible to the whole group on return
1875 }
1876 }
1877
1911 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
1912 int A_static_stride = dynamic_stride, int piv_static_stride = dynamic_stride,
1913 int rhs_static_stride = dynamic_stride, typename Sync = tdls::AutoSync>
1914 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1915 solve(const int tx, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1916 const int piv_stride, const T* TDLS_RESTRICT b, T* TDLS_RESTRICT x, const int rhs_stride,
1917 T* work, Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1918 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
1919 rhs_static_stride>) {
1921 tx, A, detail::stride_or<A_static_stride>(A_stride), piv,
1922 detail::stride_or<piv_static_stride>(piv_stride), b, x,
1923 detail::stride_or<rhs_static_stride>(rhs_stride), work, std::forward<Sync>(sync));
1924 } else {
1925 auto&& group = detail::resolve_sync<threads_per_system>(sync, work, threads_per_system);
1926 T* const exchange = work + detail::barrier_elements;
1927 T rA[rows_per_thread][N];
1928 T rB[rows_per_thread];
1929 int rpiv[rows_per_thread];
1930 load_rows<internal_matrix>(tx, A, A_stride, rA);
1931 init_piv(tx, rpiv);
1932 const int linfo = eliminate<false>(tx, rA, rB, rpiv, exchange, group);
1933 store_rows<internal_matrix>(tx, A, A_stride, rA);
1934 store_piv<internal_piv>(tx, piv, piv_stride, rpiv);
1935 load_rhs<internal_rhs>(tx, b, rhs_stride, rB);
1936 substitution<false>(tx, rA, rB, rpiv, exchange, group);
1937 store_solution<internal_rhs>(tx, x, rhs_stride, rB);
1938 group(); // every output visible to the whole group on return
1939 return linfo == 0;
1940 }
1941 }
1942
1979 template<bool internal_rhs, bool internal_piv, bool internal_matrix,
1980 int A_static_stride = dynamic_stride, int piv_static_stride = dynamic_stride,
1981 int rhs_static_stride = dynamic_stride, typename Sync = tdls::AutoSync>
1982 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr bool
1983 solve_inplace(const int tx, T* TDLS_RESTRICT A, const int A_stride, int* TDLS_RESTRICT piv,
1984 const int piv_stride, T* TDLS_RESTRICT y, const int rhs_stride, T* work,
1985 Sync&& sync = Sync{}) noexcept(detail::nothrow_sync<Sync>) {
1986 if constexpr (detail::any_static_stride<A_static_stride, piv_static_stride,
1987 rhs_static_stride>) {
1989 tx, A, detail::stride_or<A_static_stride>(A_stride), piv,
1990 detail::stride_or<piv_static_stride>(piv_stride), y,
1991 detail::stride_or<rhs_static_stride>(rhs_stride), work, std::forward<Sync>(sync));
1992 } else {
1993 auto&& group = detail::resolve_sync<threads_per_system>(sync, work, threads_per_system);
1994 T* const exchange = work + detail::barrier_elements;
1995 T rA[rows_per_thread][N];
1996 T rB[rows_per_thread];
1997 int rpiv[rows_per_thread];
1998 load_rows<internal_matrix>(tx, A, A_stride, rA);
1999 init_piv(tx, rpiv);
2000 load_rhs<internal_rhs>(tx, y, rhs_stride, rB);
2001 const int linfo = eliminate<true>(tx, rA, rB, rpiv, exchange, group);
2002 store_rows<internal_matrix>(tx, A, A_stride, rA);
2003 store_piv<internal_piv>(tx, piv, piv_stride, rpiv);
2004 substitution<true>(tx, rA, rB, rpiv, exchange, group);
2005 detail::unroll_K<rows_per_thread>([&](auto K_tag) {
2006 constexpr int K = decltype(K_tag)::value;
2007 if constexpr (slot_has_rows(K)) {
2008 if (slot_is_full(K) || tx + K * threads_per_system < N)
2009 TDLS_COOP_LUPP_Y(K) = rB[K];
2010 }
2011 });
2012 group(); // every output visible to the whole group on return
2013 return linfo == 0;
2014 }
2015 }
2016};
2017
2018
2019
2020#undef TDLS_COOP_LUPP_A
2021#undef TDLS_COOP_LUPP_PIV
2022#undef TDLS_COOP_LUPP_B
2023#undef TDLS_COOP_LUPP_X
2024#undef TDLS_COOP_LUPP_Y
2025
2026
2027
2028} // namespace tdls
2029
2030#if defined(TDLS_DETAIL_COMPILER_NVCXX)
2031#pragma diag_pop
2032#endif
2033
2034#if defined(TDLS_DETAIL_COMPILER_GCC)
2035#pragma GCC diagnostic pop
2036#endif
2037
2038#if defined(TDLS_DETAIL_COMPILER_CLANG)
2039#pragma clang diagnostic pop
2040#endif
2041
2042
2043
2044#endif // TDLS_SOLVERS_COOPERATIVE_LUPP_SOLVER_STATIC_HPP
Compile-time configuration of the CooperativeLUpp solver family.
#define TDLS_COOP_LUPP_Y(K)
Entry of the row of slot K of the right-hand side solved in place, addressed as TDLS_COOP_LUPP_B.
Definition solver_static.hpp:375
#define TDLS_COOP_LUPP_PIV(K)
Pivot entry of the row of slot K: slice of the calling thread under internal residency,...
Definition solver_static.hpp:360
#define TDLS_COOP_LUPP_X(K)
Solution entry of the row of slot K, addressed as TDLS_COOP_LUPP_B.
Definition solver_static.hpp:370
#define TDLS_COOP_LUPP_A(K, c)
Element (row of slot K, c) of S, the stored matrix read column-major (A under the column-major layout...
Definition solver_static.hpp:354
#define TDLS_COOP_LUPP_B(K)
Right-hand-side entry of the row of slot K: slice of the calling thread under internal residency,...
Definition solver_static.hpp:365
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void unroll_K(Fn &&fn) noexcept
Compile-time iteration over the row slots K = 0 .. R-1 of a thread.
Definition solver_static.hpp:334
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void unroll_K_impl(Fn &fn, std::integer_sequence< int, Ks... >) noexcept
Calls fn once per index of the sequence, in increasing order, with the index as a std::integral_const...
Definition solver_static.hpp:318
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_UNROLL_FORCE
Full-unroll pragma in the local compiler dialect.
Definition macros.hpp:250
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:50
Scalar math helpers usable in constant expressions.
constexpr int dynamic_stride
The compile-time stride of an operand whose stride is known at run time only: the default of the stri...
Definition options.hpp:181
Barrier argument of the cooperative solvers that asks them to deduce the barrier from the compilation...
Definition group.hpp:268
Dense LU factorization with partial pivoting held in the registers of a group of threads,...
Definition solver_static.hpp:412
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void gather_rhs(const int tx, T(&rB)[rows_per_thread], const int(&origin)[rows_per_thread], T *staged, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Gather of the right-hand side in the pivoted order, under physical row interchanges: each slot receiv...
Definition solver_static.hpp:1354
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_backward_step(const int tx, const int k, const T(&rA)[rows_per_thread][N], T(&rB)[rows_per_thread], const int(&rowid)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
One step of the backward substitution L^T w = z.
Definition solver_static.hpp:1480
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void backward_pass(const int tx, const T(&rA)[rows_per_thread][N], 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 regis...
Definition solver_static.hpp:1316
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void init_piv(const int tx, 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_static.hpp:582
static constexpr int threads_per_system
Threads solving one system: ceil(N / rows_per_thread).
Definition solver_static.hpp:432
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void store_solution(const int tx, 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_static.hpp:658
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool slot_is_full(const int K) noexcept
Whether slot K holds a real row for every thread.
Definition solver_static.hpp:483
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_static.hpp:689
static constexpr T relative_pivot_threshold
Relative threshold of the pivot choice, read once from the configuration (see CooperativeLUppConfig::...
Definition solver_static.hpp:443
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_rows(const int tx, const T *TDLS_RESTRICT A, const int A_stride, T(&rA)[rows_per_thread][N]) noexcept
Load the rows of the calling thread into registers.
Definition solver_static.hpp:527
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void at_position(const int tx, const int pos, Fn &&fn) noexcept
Calls fn on the slot of static index pos when the calling thread holds it: the position of a row unde...
Definition solver_static.hpp:503
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace(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_static.hpp:1983
static constexpr bool transposed
Whether the solve uses the transpose of the factored matrix. The solver factors S,...
Definition solver_static.hpp:454
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void gather_row(const int tx, const int i, const int pos, const T(&rA)[rows_per_thread][N], const T(&rB)[rows_per_thread], const int(&origin)[rows_per_thread], T *record) noexcept
Copy the register row at position pos below column i, held by the calling thread in a slot known at r...
Definition solver_static.hpp:891
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitution(const int tx, const T(&rA)[rows_per_thread][N], 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 registers on entry,...
Definition solver_static.hpp:1613
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void save_row(const T(&row)[N], const T &rhs, const int origin, T *record) noexcept
Copy a register row to an exchange record of the workspace: its N entries, its right-hand-side entry ...
Definition solver_static.hpp:823
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void store_piv(const int tx, 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_static.hpp:618
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr auto make_sync(T *work) noexcept
The barrier deduced for the group of the calling thread.
Definition solver_static.hpp:1658
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool factorize(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_static.hpp:1687
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_forward_step(const int tx, const int i, const T(&rA)[rows_per_thread][N], T(&rB)[rows_per_thread], const int(&rowid)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
One step of the forward substitution U^T z = b.
Definition solver_static.hpp:1425
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void backward_step(const int i, const int tx, const T(&rA)[rows_per_thread][N], T(&rB)[rows_per_thread], const int(&rowid)[rows_per_thread], T *sB, T *sx, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
One step of the backward substitution.
Definition solver_static.hpp:1275
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_passes(const int tx, const T(&rA)[rows_per_thread][N], T(&rB)[rows_per_thread], const int(&rowid)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Forward then backward pass of S^T x = b on the given row positions: b in the entries of the thread on...
Definition solver_static.hpp:1531
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve(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 rows stay in registers between the two, so the factors are st...
Definition solver_static.hpp:1915
static constexpr int rows_per_thread
Rows held by each thread: Config.rows_per_thread, or N when it is larger. A value below 1,...
Definition solver_static.hpp:428
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_rhs(const int tx, 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_static.hpp:639
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_row(T(&row)[N], T &rhs, int &origin, const T *record) noexcept
Copy an exchange record of the workspace to a register row, the converse of save_row.
Definition solver_static.hpp:846
static constexpr int workspace_size
Elements of the workspace shared by the threads of a group: the two elements of the deduced CPU barri...
Definition solver_static.hpp:437
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void forward_step(const int i, const T(&rA)[rows_per_thread][N], T(&rB)[rows_per_thread], const int(&rowid)[rows_per_thread], T *sB, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
One step of the forward substitution.
Definition solver_static.hpp:1210
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_inplace(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_static.hpp:1853
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_piv(const int tx, const int *TDLS_RESTRICT piv, const int piv_stride, int(&rpiv)[rows_per_thread]) noexcept
Load the pivot entries recorded by factorize.
Definition solver_static.hpp:598
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool below(const int tx, const int K, const int i) noexcept
Whether slot K of the calling thread holds a real row that may sit below column i: the slots a pivot ...
Definition solver_static.hpp:867
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void forward_pass(const T(&rA)[rows_per_thread][N], 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 registers.
Definition solver_static.hpp:1241
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void unpermute_solution(const int tx, T(&rB)[rows_per_thread], const int(&rpiv)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
x = P^T w, through the workspace. Logical row interchanges: the slot of row r of S takes w at the pos...
Definition solver_static.hpp:1569
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool slot_has_rows(const int K) noexcept
Whether slot K holds a real row for at least one thread.
Definition solver_static.hpp:475
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr T * segment(T *work, const int step) noexcept
Half of the double buffer of the transposed substitutions used by a step: N + 1 elements,...
Definition solver_static.hpp:1406
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void logical_elimination_step(const int tx, const int i, T(&rA)[rows_per_thread][N], 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.
Definition solver_static.hpp:724
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_canonical(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_static.hpp:1793
static constexpr T singular_floor
Singularity floor, read once from the configuration (see CooperativeLUppConfig::singular_floor).
Definition solver_static.hpp:446
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void store_rows(const int tx, T *TDLS_RESTRICT A, const int A_stride, const T(&rA)[rows_per_thread][N]) noexcept
Store the register rows of the calling thread.
Definition solver_static.hpp:555
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void scatter_row(const int tx, const int i, const int pos, T(&rA)[rows_per_thread][N], T(&rB)[rows_per_thread], int(&origin)[rows_per_thread], const T *record) noexcept
Copy an exchange record to the register row at position pos below column i, held by the calling threa...
Definition solver_static.hpp:942
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute(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_static.hpp:1739
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int eliminate(const int tx, T(&rA)[rows_per_thread][N], T(&rB)[rows_per_thread], int(&rpiv)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Factorization of the register rows of the group.
Definition solver_static.hpp:1150
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_elimination_step(const int tx, const int i, T(&rA)[rows_per_thread][N], 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.
Definition solver_static.hpp:1012
Detection of the compiler, of the language and of the side of the compilation pass: the one place whe...