1#ifndef TDLS_SOLVERS_COOPERATIVE_LUPP_SOLVER_DYNAMIC_HPP
2#define TDLS_SOLVERS_COOPERATIVE_LUPP_SOLVER_DYNAMIC_HPP
123#include <type_traits>
144#if defined(TDLS_DETAIL_COMPILER_GCC)
145#pragma GCC diagnostic push
146#pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
147#pragma GCC diagnostic ignored "-Warray-bounds"
164#define TDLS_COOP_LUPP_DYN_A(r, c) A[TDLS_STRIDED((c) * n + (r), A_stride)]
167#define TDLS_COOP_LUPP_DYN_PIV(r) piv[TDLS_STRIDED(r, piv_stride)]
170#define TDLS_COOP_LUPP_DYN_B(r) b[TDLS_STRIDED(r, rhs_stride)]
173#define TDLS_COOP_LUPP_DYN_X(r) x[TDLS_STRIDED(r, rhs_stride)]
176#define TDLS_COOP_LUPP_DYN_Y(r) y[TDLS_STRIDED(r, rhs_stride)]
207template<
typename T, CooperativeLUppConfig<T> Config = CooperativeLUppConfig<T>{}>
211 Config.row_interchange != RowInterchange::Auto,
212 "CooperativeLUppSolverDynamic: row_interchange must be Logical or Physical (Auto is the "
213 "default of the blocks of ArrowheadLUpp only)");
214 static_assert(Config.rows_per_thread >= 1,
215 "CooperativeLUppSolverDynamic: rows_per_thread must be >= 1");
216 static_assert(Config.relative_pivot_threshold.is_finite() && Config.singular_floor.is_finite(),
217 "CooperativeLUppSolverDynamic: relative_pivot_threshold and singular_floor must "
218 "be finite (and fit a 63-bit mantissa)");
234 static constexpr bool transposed = Config.layout == MatrixLayout::RowMajor;
237 "CooperativeLUppSolverDynamic: relative_pivot_threshold must be in (0, 1]");
239 "CooperativeLUppSolverDynamic: singular_floor must be positive");
258 return detail::barrier_elements + (Config.row_interchange == RowInterchange::Logical
284 rpiv[K] = tx + K * threads;
298 const int r = tx + K * threads;
314 const int r = tx + K * threads;
332 const int r = tx + K * threads;
349 const int r = tx + K * threads;
404 template<
bool fuse_rhs,
typename Sync>
410 Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
413 const int r = tx + K * threads;
420 T rx_abs_max = dsx[i];
422 for (
int j = i + 1; j < n; j++) {
423 if (dsx[j] > rx_abs_max) {
434 linfo = (zero_pivot && linfo == 0) ? (i + 1) : linfo;
435 const T update = zero_pivot ? T(0) : T(1);
440 const int r = tx + K * threads;
442 if (rowid[K] == max_id) {
444 for (
int j = i; j < n; j++)
446 if constexpr (fuse_rhs && !
transposed) sB[0] = rB[K];
447 }
else if (rowid[K] == i) {
453 if (r == i) sB[0] = rB[K];
458 const T reg = zero_pivot ? T(1) : T(1) / sx[i];
459 [[maybe_unused]] T z_i = T(0);
460 if constexpr (fuse_rhs &&
transposed) z_i = sB[0] * reg;
464 const int r = tx + K * threads;
468 for (
int j = i + 1; j < n; j++)
475 if (r == i) rB[K] = z_i;
476 if (r > i) rB[K] -= sx[r] * z_i;
491 T* record)
noexcept {
492 for (
int c = 0; c < n; c++)
505 const T* record)
noexcept {
506 for (
int c = 0; c < n; c++)
530 template<
bool fuse_rhs,
typename Sync>
536 Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
540 constexpr bool carries_rhs = fuse_rhs && !
transposed;
542 T*
const candidate = work;
543 T*
const where = candidate + threads;
544 T*
const in_place = where + threads;
545 T*
const leaving = in_place + 1;
546 T*
const entering = leaving + n + 2;
555 const int pos = tx + K * threads;
557 if (pos < i)
continue;
559 if (pos == i || v > best) {
563 if (pos == i) local = v;
565 candidate[tx] = best;
566 where[tx] =
static_cast<T
>(best_pos);
567 if (tx == i % threads) in_place[0] = local;
572 T rx_abs_max = candidate[i % threads];
573 for (
int t = 0; t < threads; t++)
574 rx_abs_max = candidate[t] > rx_abs_max ? candidate[t] : rx_abs_max;
575 bool keep = !(in_place[0] < rx_abs_max);
580 for (
int t = 0; t < threads; t++) {
581 const int pos =
static_cast<int>(where[t]);
582 if (candidate[t] == rx_abs_max && pos < max_id) max_id = pos;
586 linfo = (zero_pivot && linfo == 0) ? (i + 1) : linfo;
596 if (tx == i % threads)
save_row(n, A, A_stride, i, leaving);
597 if (max_id != i && tx == max_id % threads)
save_row(n, A, A_stride, max_id, entering);
599 const int pos = tx + K * threads;
600 if (pos != i && (pos != max_id || max_id == i))
continue;
601 T*
const record = pos == i ? leaving : entering;
602 if constexpr (fuse_rhs) record[n] = rB[K];
603 record[n + 1] =
static_cast<T
>(origin[K]);
608 if (tx == i % threads)
load_row(n, A, A_stride, i, entering);
609 if (tx == max_id % threads)
load_row(n, A, A_stride, max_id, leaving);
611 const int pos = tx + K * threads;
612 if (pos != max_id && pos != i)
continue;
613 const T*
const record = pos == i ? entering : leaving;
614 if constexpr (carries_rhs) rB[K] = record[n];
615 origin[K] =
static_cast<int>(record[n + 1]);
619 const T*
const prow = max_id == i ? leaving : entering;
620 const T reg = zero_pivot ? T(1) : T(1) / prow[i];
621 [[maybe_unused]] T z_i = T(0);
622 if constexpr (fuse_rhs &&
transposed) z_i = leaving[n] * reg;
628 const int pos = tx + K * threads;
632 for (
int j = i + 1; j < n; j++)
637 if (pos == i) rB[K] = z_i;
638 if (pos > i) rB[K] -= prow[pos] * z_i;
660 template<
bool fuse_rhs,
typename Sync>
664 Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
666 if constexpr (Config.row_interchange == RowInterchange::Logical) {
668 T*
const sx = work + n;
669 T*
const dsx = work + 2 * n;
670 for (
int i = 0; i < n; i++)
671 logical_elimination_step<fuse_rhs>(n, i, tx, threads, A, A_stride, rB, rpiv, sB, sx,
674 for (
int i = 0; i < n; i++)
707 template<
typename Sync>
711 T* work, Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
713 for (
int i = 0; i < n; i++) {
715 if (tx + K * threads >= n)
break;
716 if (rowid[K] == i) sB[0] = rB[K];
721 const int r = tx + K * threads;
749 template<
typename Sync>
753 T* work, Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
755 T*
const sx = work + n;
757 if (tx + K * threads >= n)
break;
758 sB[rowid[K]] = rB[K];
762 for (
int i = n - 1; i >= 0; i--) {
764 const int r = tx + K * threads;
770 const T reg = sB[i] / sx[i];
773 const int idx = tx + K * threads;
775 if (idx < i) sB[idx] = sB[idx] - reg * sx[idx];
776 rB[K] = (idx == i) ? reg : rB[K];
804 row_at([[maybe_unused]]
const T* rows,
const int k)
noexcept {
805 if constexpr (Config.row_interchange == RowInterchange::Logical)
806 return static_cast<int>(rows[k]);
835 template<
bool backward_only,
typename Sync>
839 Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
840 if constexpr (!backward_only) {
841 for (
int i = 0; i < n; i++) {
842 const int ri =
row_at(rows, i);
846 const int c = tx + K * threads;
849 if (c == i) buf[i & 1] = rB[K];
853 const T z_i = buf[i & 1] * (T(1) / u_ii);
855 const int c = tx + K * threads;
857 if (c == i) rB[K] = z_i;
858 if (c > i) rB[K] -= u[K] * z_i;
862 for (
int k = n - 1; k >= 1; k--) {
863 const int rk =
row_at(rows, k);
866 const int c = tx + K * threads;
869 if (c == k) buf[(k + 1) & 1] = rB[K];
873 const T w_k = buf[(k + 1) & 1];
875 const int c = tx + K * threads;
877 if (c < k) rB[K] -= l[K] * w_k;
906 template<
bool backward_only,
typename Sync>
910 T* work, Sync& sync)
noexcept(detail::nothrow_sync<Sync>) {
911 constexpr bool logical = Config.row_interchange == RowInterchange::Logical;
914 if constexpr (logical) {
916 const int r = tx + K * threads;
918 work[rpiv[K]] =
static_cast<T
>(r);
926 T*
const staged = work;
928 const int c = tx + K * threads;
930 staged[logical ? c : rpiv[K]] = rB[K];
934 const int c = tx + K * threads;
936 rB[K] = staged[logical ? rpiv[K] : c];
938 }
else if constexpr (logical) {
939 if constexpr (!backward_only)
940 forward_pass(n, tx, threads, A, A_stride, rB, rpiv, work, sync);
941 backward_pass(n, tx, threads, A, A_stride, rB, rpiv, work, sync);
945 if constexpr (!backward_only) {
950 T*
const staged = work + n;
952 const int r = tx + K * threads;
958 if (tx + K * threads >= n)
break;
959 rB[K] = staged[rpiv[K]];
961 forward_pass(n, tx, threads, A, A_stride, rB, position, work, sync);
963 backward_pass(n, tx, threads, A, A_stride, rB, position, work, sync);
1005 template<
typename Sync = tdls::AutoSync>
1009 Sync&& sync = Sync{})
noexcept(detail::nothrow_sync<Sync>) {
1011 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1012 T*
const exchange = work + detail::barrier_elements;
1016 const int linfo =
eliminate<false>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1017 store_piv(n, tx, threads, piv, piv_stride, rpiv);
1039 template<
typename Sync = tdls::AutoSync>
1044 Sync&& sync = Sync{})
noexcept(detail::nothrow_sync<Sync>) {
1046 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1047 T*
const exchange = work + detail::barrier_elements;
1050 load_piv(n, tx, threads, piv, piv_stride, rpiv);
1051 load_rhs(n, tx, threads, b, rhs_stride, rB);
1074 template<
typename Sync = tdls::AutoSync>
1077 const int*
TDLS_RESTRICT piv,
const int piv_stride,
const int col,
1079 Sync&& sync = Sync{})
noexcept(detail::nothrow_sync<Sync>) {
1081 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1082 T*
const exchange = work + detail::barrier_elements;
1085 load_piv(n, tx, threads, piv, piv_stride, rpiv);
1087 rB[K] = (tx + K * threads == col) ? T(1) : T(0);
1113 template<
typename Sync = tdls::AutoSync>
1117 const int rhs_stride, T* work,
1118 Sync&& sync = Sync{})
noexcept(detail::nothrow_sync<Sync>) {
1120 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1121 T*
const exchange = work + detail::barrier_elements;
1124 load_piv(n, tx, threads, piv, piv_stride, rpiv);
1125 load_rhs(n, tx, threads, x, rhs_stride, rB);
1154 template<
typename Sync = tdls::AutoSync>
1158 T* work, Sync&& sync = Sync{})
noexcept(detail::nothrow_sync<Sync>) {
1160 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1161 T*
const exchange = work + detail::barrier_elements;
1165 const int linfo =
eliminate<false>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1166 store_piv(n, tx, threads, piv, piv_stride, rpiv);
1167 load_rhs(n, tx, threads, b, rhs_stride, rB);
1198 template<
typename Sync = tdls::AutoSync>
1202 const int rhs_stride, T* work,
1203 Sync&& sync = Sync{})
noexcept(detail::nothrow_sync<Sync>) {
1205 auto&& group = detail::resolve_sync<0>(sync, work, threads);
1206 T*
const exchange = work + detail::barrier_elements;
1210 load_rhs(n, tx, threads, y, rhs_stride, rB);
1211 const int linfo =
eliminate<true>(n, tx, threads, A, A_stride, rB, rpiv, exchange, group);
1212 store_piv(n, tx, threads, piv, piv_stride, rpiv);
1215 const int r = tx + K * threads;
1226#undef TDLS_COOP_LUPP_DYN_A
1227#undef TDLS_COOP_LUPP_DYN_PIV
1228#undef TDLS_COOP_LUPP_DYN_B
1229#undef TDLS_COOP_LUPP_DYN_X
1230#undef TDLS_COOP_LUPP_DYN_Y
1236#if defined(TDLS_DETAIL_COMPILER_GCC)
1237#pragma GCC diagnostic pop
Compile-time configuration of the CooperativeLUpp solver family.
#define TDLS_COOP_LUPP_DYN_X(r)
Strided entry r of the solution.
Definition solver_dynamic.hpp:173
#define TDLS_COOP_LUPP_DYN_B(r)
Strided entry r of the right-hand side.
Definition solver_dynamic.hpp:170
#define TDLS_COOP_LUPP_DYN_PIV(r)
Strided pivot entry of row r.
Definition solver_dynamic.hpp:167
#define TDLS_COOP_LUPP_DYN_A(r, c)
Strided element (r, c) of S, the stored matrix read column-major (A under the column-major layout,...
Definition solver_dynamic.hpp:164
#define TDLS_COOP_LUPP_DYN_Y(r)
Strided entry r of the right-hand side solved in place.
Definition solver_dynamic.hpp:176
Barriers of the groups of threads that share a system: the barrier deduced from the compilation targe...
Toolchain detection and portability macros.
#define TDLS_FORCE_INLINE_GPU
Inlining forced on GPU only: __forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS...
Definition macros.hpp:77
#define TDLS_RESTRICT
Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC; nothing in the code that nvc++ gene...
Definition macros.hpp:174
#define TDLS_EXEC_CHECK_DISABLE
Lifts the execution-space check of nvcc on the next function template; empty elsewhere.
Definition macros.hpp:274
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:50
Scalar math helpers usable in constant expressions.
Barrier argument of the cooperative solvers that asks them to deduce the barrier from the compilation...
Definition group.hpp:268
Runtime-size dense LU factorization with partial pivoting shared by a group of threads,...
Definition solver_dynamic.hpp:208
static constexpr bool transposed
Whether the solve uses the transpose of the factored matrix, as in the compile-time solver: S = A and...
Definition solver_dynamic.hpp:234
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool factorize(const int n, const int tx, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
In-place LU factorization with partial pivoting.
Definition solver_dynamic.hpp:1007
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_piv(const int n, const int tx, const int threads, const int *TDLS_RESTRICT piv, const int piv_stride, int(&rpiv)[rows_per_thread]) noexcept
Load the pivot entries recorded by factorize.
Definition solver_dynamic.hpp:295
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_canonical(const int n, const int tx, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const int col, T *TDLS_RESTRICT x, const int rhs_stride, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
Solve with b = e_col generated on the fly: the consistent-tangent-operator path.
Definition solver_dynamic.hpp:1076
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute(const int n, const int tx, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
Solve from a factorization produced by factorize (b and x distinct).
Definition solver_dynamic.hpp:1041
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void forward_pass(const int n, const int tx, const int threads, const T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], const int(&rowid)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Forward pass: L y = P b on the given row positions, y overwriting the right-hand-side entries of the ...
Definition solver_dynamic.hpp:709
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr auto make_sync(const int n, T *work) noexcept
The barrier deduced for the group of the calling thread.
Definition solver_dynamic.hpp:984
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void init_piv(const int tx, const int threads, int(&rpiv)[rows_per_thread]) noexcept
Initial pivot entries: no row has moved yet, so the entry of each slot is the index of its row,...
Definition solver_dynamic.hpp:282
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void physical_elimination_step(const int n, const int i, const int tx, const int threads, T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], int(&origin)[rows_per_thread], T *work, int &linfo, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
One column of the factorization under physical row interchanges, as in the compile-time solver: the r...
Definition solver_dynamic.hpp:532
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void transposed_passes(const int n, const int tx, const int threads, const T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], const T *rows, T *buf, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Forward then backward pass of S^T x = b: b in the entries of the thread on entry, w at the entries of...
Definition solver_dynamic.hpp:837
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitute_inplace(const int n, const int tx, const T *TDLS_RESTRICT A, const int A_stride, const int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT x, const int rhs_stride, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
Solve in place from a factorization produced by factorize: x holds the right-hand side on entry and t...
Definition solver_dynamic.hpp:1115
static constexpr T relative_pivot_threshold
Relative threshold of the pivot choice, read once from the configuration (see CooperativeLUppConfig::...
Definition solver_dynamic.hpp:227
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve(const int n, const int tx, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, const T *TDLS_RESTRICT b, T *TDLS_RESTRICT x, const int rhs_stride, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
factorize + substitute in one call, the pivot entries staying with the threads in between.
Definition solver_dynamic.hpp:1156
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_rhs(const int n, const int tx, const int threads, const T *TDLS_RESTRICT b, const int rhs_stride, T(&rB)[rows_per_thread]) noexcept
Load the right-hand-side entries of the rows of the calling thread.
Definition solver_dynamic.hpp:329
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void substitution(const int n, const int tx, const int threads, const T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], const int(&rpiv)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Substitution from the factored rows, the right-hand side in the entries of the thread on entry,...
Definition solver_dynamic.hpp:908
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void save_row(const int n, const T *TDLS_RESTRICT A, const int A_stride, const int r, T *record) noexcept
Copy a row of the matrix to the first n entries of an exchange record of the workspace.
Definition solver_dynamic.hpp:490
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void logical_elimination_step(const int n, const int i, const int tx, const int threads, T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], int(&rowid)[rows_per_thread], T *sB, T *sx, T *dsx, int &linfo, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
One column of the factorization under logical row interchanges. With fuse_rhs, the forward pass of th...
Definition solver_dynamic.hpp:406
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int row_at(const T *rows, const int k) noexcept
Index of the row of S at position k: read from the published positions under logical row interchanges...
Definition solver_dynamic.hpp:804
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int eliminate(const int n, const int tx, const int threads, T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], int(&rpiv)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Factorization of the rows of the group, in place.
Definition solver_dynamic.hpp:662
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool solve_inplace(const int n, const int tx, T *TDLS_RESTRICT A, const int A_stride, int *TDLS_RESTRICT piv, const int piv_stride, T *TDLS_RESTRICT y, const int rhs_stride, T *work, Sync &&sync=Sync{}) noexcept(detail::nothrow_sync< Sync >)
Factorization with the forward substitution folded in: the original MAGMA kernel.
Definition solver_dynamic.hpp:1200
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr bool keeps_pivot(const T in_place, const T column_max) noexcept
Relative threshold of the pivot choice: whether the row in place keeps the pivot (see CooperativeLUpp...
Definition solver_dynamic.hpp:375
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void store_solution(const int n, const int tx, const int threads, T *TDLS_RESTRICT x, const int rhs_stride, const T(&rB)[rows_per_thread]) noexcept
Store the solution entries held by the calling thread.
Definition solver_dynamic.hpp:346
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void store_piv(const int n, const int tx, const int threads, int *TDLS_RESTRICT piv, const int piv_stride, const int(&rpiv)[rows_per_thread]) noexcept
Record the pivot entries of the calling thread.
Definition solver_dynamic.hpp:311
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void load_row(const int n, T *TDLS_RESTRICT A, const int A_stride, const int r, const T *record) noexcept
Copy the first n entries of an exchange record of the workspace to a row of the matrix,...
Definition solver_dynamic.hpp:504
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int workspace_size(const int n) noexcept
Elements of the workspace shared by the threads of a group: the two elements of the deduced CPU barri...
Definition solver_dynamic.hpp:257
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int threads_per_system(const int n) noexcept
Threads solving one system of dimension n: ceil(n / rows_per_thread).
Definition solver_dynamic.hpp:246
static constexpr T singular_floor
Singularity floor, read once from the configuration (see CooperativeLUppConfig::singular_floor).
Definition solver_dynamic.hpp:230
static constexpr int rows_per_thread
Row slots of each thread: Config.rows_per_thread. When n is smaller, one thread solves the system and...
Definition solver_dynamic.hpp:224
TDLS_EXEC_CHECK_DISABLE TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void backward_pass(const int n, const int tx, const int threads, const T *TDLS_RESTRICT A, const int A_stride, T(&rB)[rows_per_thread], const int(&rowid)[rows_per_thread], T *work, Sync &sync) noexcept(detail::nothrow_sync< Sync >)
Backward pass: U x = y on the given row positions, the forward-reduced right-hand side y in the entri...
Definition solver_dynamic.hpp:751
Detection of the compiler, of the language and of the side of the compilation pass: the one place whe...