1#ifndef TDLS_SOLVERS_COOPERATIVE_LUPP_SOLVER_STATIC_HPP
2#define TDLS_SOLVERS_COOPERATIVE_LUPP_SOLVER_STATIC_HPP
254#include <type_traits>
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"
287#if defined(TDLS_DETAIL_COMPILER_CLANG)
288#pragma clang diagnostic push
289#pragma clang diagnostic ignored "-Wpass-failed"
300#if defined(TDLS_DETAIL_COMPILER_NVCXX)
302#pragma diag_suppress no_device_stack
316template<
typename Fn,
int... Ks>
319 (fn(std::integral_constant<int, Ks>{}), ...);
333template<
int R,
typename Fn>
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)]
411template<
typename T,
int N, CooperativeLUppConfig<T> Config = CooperativeLUppConfig<T>{}>
414 static_assert(N >= 1,
"CooperativeLUppSolverStatic: N must be >= 1");
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)");
429 : Config.rows_per_thread < N ? Config.rows_per_thread
438 detail::barrier_elements + (Config.row_interchange == RowInterchange::Logical
454 static constexpr bool transposed = Config.layout == MatrixLayout::RowMajor;
457 "CooperativeLUppSolverStatic: relative_pivot_threshold must be in (0, 1]");
459 "CooperativeLUppSolverStatic: singular_floor must be positive");
501 template<
typename Fn>
504 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
505 constexpr int K =
decltype(K_tag)::value;
525 template<
bool internal_matrix>
529 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
530 constexpr int K =
decltype(K_tag)::value;
533 if constexpr (Config.unroll_loops) {
535 for (
int c = 0; c < N; ++c)
538 for (
int c = 0; c < N; ++c)
553 template<
bool internal_matrix>
557 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
558 constexpr int K =
decltype(K_tag)::value;
561 if constexpr (Config.unroll_loops) {
563 for (
int c = 0; c < N; ++c)
566 for (
int c = 0; c < N; ++c)
583 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
584 constexpr int K =
decltype(K_tag)::value;
596 template<
bool internal_piv>
600 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
601 constexpr int K =
decltype(K_tag)::value;
616 template<
bool internal_piv>
620 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
621 constexpr int K =
decltype(K_tag)::value;
637 template<
bool internal_rhs>
641 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
642 constexpr int K =
decltype(K_tag)::value;
656 template<
bool internal_rhs>
660 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
661 constexpr int K =
decltype(K_tag)::value;
722 template<
bool fuse_rhs,
typename Sync>
727 Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
729 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
730 constexpr int K =
decltype(K_tag)::value;
732 if (
slot_is_full(K) || rowid[K] < N) dsx[rowid[K]] = detail::abs(rA[K][i]);
739 T rx_abs_max = dsx[i];
741 for (
int j = i + 1; j < N; j++) {
742 if (dsx[j] > rx_abs_max) {
753 linfo = (zero_pivot && linfo == 0) ? (i + 1) : linfo;
754 const T update = zero_pivot ? T(0) : T(1);
758 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
759 constexpr int K =
decltype(K_tag)::value;
761 if (rowid[K] == max_id) {
763 if constexpr (Config.unroll_loops) {
765 for (
int j = i; j < N; j++)
766 sx[j] = update * rA[K][j];
768 for (
int j = i; j < N; j++)
769 sx[j] = update * rA[K][j];
771 if constexpr (fuse_rhs && !
transposed) sB[0] = rB[K];
772 }
else if (rowid[K] == i) {
778 at_position(tx, i, [&](
auto K_tag) { sB[0] = rB[
decltype(K_tag)::value]; });
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;
786 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
787 constexpr int K =
decltype(K_tag)::value;
789 if ((
slot_is_full(K) || rowid[K] < N) && rowid[K] > i) {
791 if constexpr (Config.unroll_loops) {
793 for (
int j = i + 1; j < N; j++)
794 rA[K][j] -= rA[K][i] * sx[j];
796 for (
int j = i + 1; j < N; j++)
797 rA[K][j] -= rA[K][i] * sx[j];
799 if constexpr (fuse_rhs && !
transposed) rB[K] -= rA[K][i] * sB[0];
803 if (c == i) rB[K] = z_i;
804 if ((
slot_is_full(K) || c < N) && c > i) rB[K] -= sx[c] * z_i;
821 template<
bool fuse_rhs>
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++)
830 for (
int c = 0; c < N; c++)
833 if constexpr (fuse_rhs) record[N] = rhs;
834 record[N + 1] =
static_cast<T
>(origin);
844 template<
bool fuse_rhs>
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++)
852 for (
int c = 0; c < N; c++)
855 if constexpr (fuse_rhs) rhs = record[N];
856 origin =
static_cast<int>(record[N + 1]);
868 const int i)
noexcept {
889 template<
bool fuse_rhs>
895 const auto entry = [&](
const int c) {
897 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
898 constexpr int K =
decltype(K_tag)::value;
900 if (
below(tx, K, i)) v = K == slot ? rA[K][c] : v;
905 if constexpr (Config.unroll_loops) {
907 for (
int c = 0; c < N; c++)
908 record[c] = entry(c);
910 for (
int c = 0; c < N; c++)
911 record[c] = entry(c);
915 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
916 constexpr int K =
decltype(K_tag)::value;
918 if (
below(tx, K, i)) {
919 if constexpr (fuse_rhs) rhs = K == slot ? rB[K] : rhs;
920 index = K == slot ? origin[K] : index;
924 if constexpr (fuse_rhs) record[N] = rhs;
925 record[N + 1] =
static_cast<T
>(index);
940 template<
bool fuse_rhs>
944 const T* record)
noexcept {
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;
951 if (
below(tx, K, i)) rA[K][c] = K == slot ? v : rA[K][c];
955 if constexpr (Config.unroll_loops) {
957 for (
int c = 0; c < N; c++)
960 for (
int c = 0; c < N; c++)
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;
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];
1010 template<
bool fuse_rhs,
typename Sync>
1015 Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
1017 constexpr bool carries_rhs = fuse_rhs && !
transposed;
1022 T*
const candidate = work;
1025 T*
const leaving = in_place + 1;
1026 T*
const entering = leaving + N + 2;
1034 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
1035 constexpr int K =
decltype(K_tag)::value;
1039 const T v = detail::abs(rA[K][i]);
1040 if (pos == i || v > best) {
1044 if (pos == i) local = v;
1048 candidate[tx] = best;
1049 where[tx] =
static_cast<T
>(best_pos);
1060 rx_abs_max = candidate[t] > rx_abs_max ? candidate[t] : rx_abs_max;
1061 bool keep = !(in_place[0] < rx_abs_max);
1067 const int pos =
static_cast<int>(where[t]);
1068 if (candidate[t] == rx_abs_max && pos < max_id) max_id = pos;
1072 linfo = (zero_pivot && linfo == 0) ? (i + 1) : linfo;
1086 constexpr int K =
decltype(K_tag)::value;
1095 constexpr int K =
decltype(K_tag)::value;
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;
1110 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
1111 constexpr int K =
decltype(K_tag)::value;
1116 if constexpr (Config.unroll_loops) {
1118 for (
int j = i + 1; j < N; j++)
1119 rA[K][j] -= rA[K][i] * prow[j];
1121 for (
int j = i + 1; j < N; j++)
1122 rA[K][j] -= rA[K][i] * prow[j];
1124 if constexpr (carries_rhs) rB[K] -= rA[K][i] * prow[N];
1127 if (pos == i) rB[K] = z_i;
1128 if ((
slot_is_full(K) || pos < N) && pos > i) rB[K] -= prow[pos] * z_i;
1148 template<
bool fuse_rhs,
typename Sync>
1152 Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
1154 if constexpr (Config.row_interchange == RowInterchange::Logical) {
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++)
1164 for (
int i = 0; i < N; i++)
1169 if constexpr (Config.unroll_loops) {
1171 for (
int i = 0; i < N; i++)
1174 for (
int i = 0; i < N; i++)
1208 template<
typename Sync>
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;
1216 if (rowid[K] == i) sB[0] = rB[K];
1221 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
1222 constexpr int K =
decltype(K_tag)::value;
1224 if ((
slot_is_full(K) || rowid[K] < N) && rowid[K] > i) rB[K] -= rA[K][i] * sB[0];
1239 template<
typename Sync>
1243 Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
1245 if constexpr (Config.unroll_loops) {
1247 for (
int i = 0; i < N; i++)
1250 for (
int i = 0; i < N; i++)
1273 template<
typename Sync>
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;
1281 if (
slot_is_full(K) || rowid[K] < N) sx[rowid[K]] = rA[K][i];
1286 const T reg = sB[i] / sx[i];
1288 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
1289 constexpr int K =
decltype(K_tag)::value;
1293 if (idx < i) sB[idx] = sB[idx] - reg * sx[idx];
1294 rB[K] = (idx == i) ? reg : rB[K];
1314 template<
typename Sync>
1318 Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
1320 T*
const sx = work + N;
1321 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
1322 constexpr int K =
decltype(K_tag)::value;
1324 if (
slot_is_full(K) || rowid[K] < N) sB[rowid[K]] = rB[K];
1329 if constexpr (Config.unroll_loops) {
1331 for (
int i = N - 1; i >= 0; i--)
1334 for (
int i = N - 1; i >= 0; i--)
1352 template<
typename Sync>
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;
1365 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
1366 constexpr int K =
decltype(K_tag)::value;
1407 const int step)
noexcept {
1408 return work + (step & 1) * (N + 1);
1423 template<
typename Sync>
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;
1432 if (rowid[K] == i) {
1437 if constexpr (Config.unroll_loops) {
1439 for (
int j = 0; j < N; j++)
1440 if (j >= i) sx[j] = rA[K][j];
1442 for (
int j = i; j < N; j++)
1448 at_position(tx, i, [&](
auto K_tag) { sx[N] = rB[
decltype(K_tag)::value]; });
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;
1456 if (c == i) rB[K] = z_i;
1457 if ((
slot_is_full(K) || c < N) && c > i) rB[K] -= sx[c] * z_i;
1478 template<
typename Sync>
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;
1487 if (rowid[K] == k) {
1490 if constexpr (Config.unroll_loops) {
1492 for (
int c = 0; c < N; c++)
1493 if (c < k) sx[c] = rA[K][c];
1495 for (
int c = 0; c < k; c++)
1501 at_position(tx, k, [&](
auto K_tag) { sx[N] = rB[
decltype(K_tag)::value]; });
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;
1509 if (c < k) rB[K] -= sx[c] * w_k;
1529 template<
bool backward_only,
typename Sync>
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++)
1541 for (
int k = N - 1; k >= 1; k--)
1544 if constexpr (!backward_only) {
1545 for (
int i = 0; i < N; i++)
1548 for (
int k = N - 1; k >= 1; k--)
1567 template<
typename Sync>
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;
1582 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
1583 constexpr int K =
decltype(K_tag)::value;
1611 template<
bool backward_only,
typename Sync>
1615 Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
1616 constexpr bool logical = Config.row_interchange == RowInterchange::Logical;
1618 if constexpr (logical) {
1626 }
else if constexpr (logical) {
1627 if constexpr (!backward_only)
forward_pass(rA, rB, rpiv, work, sync);
1632 if constexpr (!backward_only) {
1684 template<
bool internal_piv,
bool internal_matrix,
int A_static_stride =
dynamic_stride,
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));
1695 auto&& group = detail::resolve_sync<threads_per_system>(sync, work,
threads_per_system);
1696 T*
const exchange = work + detail::barrier_elements;
1735 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
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));
1750 auto&& group = detail::resolve_sync<threads_per_system>(sync, work,
threads_per_system);
1751 T*
const exchange = work + detail::barrier_elements;
1789 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
1794 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int col,
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));
1804 auto&& group = detail::resolve_sync<threads_per_system>(sync, work,
threads_per_system);
1805 T*
const exchange = work + detail::barrier_elements;
1811 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
1812 constexpr int K =
decltype(K_tag)::value;
1849 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
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));
1864 auto&& group = detail::resolve_sync<threads_per_system>(sync, work,
threads_per_system);
1865 T*
const exchange = work + detail::barrier_elements;
1911 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
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));
1925 auto&& group = detail::resolve_sync<threads_per_system>(sync, work,
threads_per_system);
1926 T*
const exchange = work + detail::barrier_elements;
1979 template<
bool internal_rhs,
bool internal_piv,
bool internal_matrix,
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));
1993 auto&& group = detail::resolve_sync<threads_per_system>(sync, work,
threads_per_system);
1994 T*
const exchange = work + detail::barrier_elements;
2005 detail::unroll_K<rows_per_thread>([&](
auto K_tag) {
2006 constexpr int K =
decltype(K_tag)::value;
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
2030#if defined(TDLS_DETAIL_COMPILER_NVCXX)
2034#if defined(TDLS_DETAIL_COMPILER_GCC)
2035#pragma GCC diagnostic pop
2038#if defined(TDLS_DETAIL_COMPILER_CLANG)
2039#pragma clang diagnostic pop
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...