iluk.hh
Go to the documentation of this file.
1#pragma once
2
3#include <batchlas/export.hh>
9
10#include <cstddef>
11#include <cstdint>
13
14namespace batchlas {
15
16/// @brief Parameters of an incomplete LU factorization with level-of-fill k, ILU(k).
17///
18/// The symbolic phase keeps every entry whose level of fill is at most
19/// `levels_of_fill`. The numeric phase then prunes each row: entries at or below
20/// `drop_tolerance` times the row's scale (floored at 1) are dropped, and at most
21/// ceil(`fill_factor` x the row's original non-zero count) - 1 off-diagonal
22/// entries survive, largest magnitude first (ties by column). A pivot smaller
23/// than max(|`diagonal_shift`|, `diag_pivot_threshold` x max(row scale, 1)) is
24/// shifted by `diagonal_shift`, replaced by it, or raised to that threshold.
25/// @tparam T scalar type; real-valued fields use its real counterpart
26/// @ingroup api_sparse
27template <typename T>
28struct ILUKParams {
29 int levels_of_fill = 0; ///< k; 0 keeps exactly A's pattern. Must be >= 0.
30 T diagonal_shift = T(1e-8); ///< added to (or substituted for) a pivot that is too small
31 typename base_type<T>::type drop_tolerance = typename base_type<T>::type(1e-4); ///< relative drop threshold; >= 0
32 typename base_type<T>::type fill_factor = typename base_type<T>::type(10); ///< per-row cap on kept entries relative to A's row; >= 1
33 typename base_type<T>::type diag_pivot_threshold = typename base_type<T>::type(0.1); ///< relative pivot threshold; >= 0
34 bool modified_ilu = true; ///< add dropped entries to the diagonal (MILU) instead of discarding them
35 /// Check that every batch item shares one CSR pattern. Must stay true for
36 /// batch > 1: disabling the check is rejected.
38};
39
40/// @brief Non-owning description of a factored ILU(k) preconditioner.
41///
42/// This is what iluk_apply consumes, so the factor can live in an
43/// ILUKPreconditioner or in a caller-supplied workspace (the `Span` overload of
44/// iluk_factorize). Valid only while that storage is.
45/// @tparam T scalar type
46/// @ingroup api_sparse
47template <typename T>
48struct ILUKView {
49 MatrixView<T, MatrixFormat::CSR> lu; ///< L (strict lower, unit diagonal implied) and U (upper, with diagonal)
50 Span<int> diag_positions; ///< size n * batch_size, absolute indices into lu's values
51
52 Span<int> l_rows; ///< row indices ordered by forward-solve level
53 Span<int> l_level_ptr; ///< size l_levels + 1
54 Span<int> u_rows; ///< row indices ordered by backward-solve level
55 Span<int> u_level_ptr; ///< size u_levels + 1
56 int l_levels = 0; ///< number of forward-solve levels
57 int u_levels = 0; ///< number of backward-solve levels
58
59 bool u_diagonals_usable = false; ///< every U diagonal is non-zero after the shift
60
61 int n = 0; ///< order of each matrix
62 int batch_size = 0; ///< number of matrices
63 T diagonal_shift = T(1e-8); ///< shift the apply uses on U's diagonal
64};
65
66/// @brief Owning ILU(k) factor of a batch of CSR matrices sharing one pattern.
67///
68/// Holds L (unit diagonal, strict lower triangle) and U (upper triangle with its
69/// diagonal) in one CSR matrix, plus a level schedule for the two triangular
70/// solves: rows inside one level depend only on earlier levels, so iluk_apply
71/// walks `l_levels` / `u_levels` steps instead of n. One schedule serves the
72/// whole batch. Build it with iluk_factorize; `view()` gives the ILUKView that
73/// iluk_apply takes.
74/// @tparam T scalar type
75/// @ingroup api_sparse
76template <typename T>
78 /// 1x1 placeholder with a single stored non-zero; iluk_factorize replaces it.
79 ILUKPreconditioner() : lu(1, 1, NonZeros{1}, 1) {}
80
81 Matrix<T, MatrixFormat::CSR> lu; ///< L (strict lower, unit diagonal implied) and U (upper, with diagonal)
82 UnifiedVector<int> diag_positions; ///< size n * batch_size
83
84 UnifiedVector<int> l_rows; ///< row indices ordered by forward-solve level
85 UnifiedVector<int> l_level_ptr; ///< size l_levels + 1
86 UnifiedVector<int> u_rows; ///< row indices ordered by backward-solve level
87 UnifiedVector<int> u_level_ptr; ///< size u_levels + 1
88 int l_levels = 0; ///< number of forward-solve levels
89 int u_levels = 0; ///< number of backward-solve levels
90
91 /// Whether every U diagonal can be made non-singular with `diagonal_shift`.
92 /// Decided once when the factor is built, so iluk_apply never syncs with the
93 /// host to check it.
94 bool u_diagonals_usable = false;
95
96 int n = 0; ///< order of each matrix
97 int batch_size = 0; ///< number of matrices
98 int levels_of_fill = 0; ///< the ILUKParams value it was built with
99 T diagonal_shift = T(1e-8); ///< the ILUKParams value it was built with
100 typename base_type<T>::type drop_tolerance = typename base_type<T>::type(1e-4); ///< the ILUKParams value it was built with
101 typename base_type<T>::type fill_factor = typename base_type<T>::type(10); ///< the ILUKParams value it was built with
102 typename base_type<T>::type diag_pivot_threshold = typename base_type<T>::type(0.1); ///< the ILUKParams value it was built with
103 bool modified_ilu = true; ///< the ILUKParams value it was built with
104
105 /// @brief Non-owning view of this factor; valid while `*this` is alive and unmodified.
107 ILUKView<T> v;
108 v.lu = lu.view();
110 v.l_rows = l_rows;
112 v.u_rows = u_rows;
114 v.l_levels = l_levels;
115 v.u_levels = u_levels;
117 v.n = n;
120 return v;
121 }
122};
123
124/// @brief Builds the triangular-solve level schedule from `M.lu`'s sparsity pattern.
125///
126/// iluk_factorize already does this; call it only for an ILUKPreconditioner
127/// assembled by hand, because iluk_apply requires a valid schedule. Runs on the
128/// host and reads `M.lu`'s row offsets and column indices.
129/// @param M factor whose `lu`, `n` are set; its schedule fields are overwritten
130/// @throws batchlas::invalid_argument if `M.n <= 0`
131/// @ingroup api_sparse_lowlevel
132template <typename T>
134
135/// @brief Computes the ILU(k) factor of a batch of CSR matrices into owned storage.
136///
137/// Every item of `A` must share one sparsity pattern; the factor's pattern is the
138/// union of what each item kept, so the batch keeps sharing one pattern. Blocking:
139/// returns once the factor is complete. The numeric phase runs on the device for
140/// large batches and on the host otherwise (`BATCHLAS_ILUK_DEVICE` forces either).
141/// @tparam B backend the call is compiled for
142/// @tparam T scalar type
143/// @param ctx queue used for the device phase and allocations
144/// @param A batch of square CSR matrices; not modified
145/// @param params fill, drop and pivot controls
146/// @return the owning factor, ready for iluk_apply
147/// @throws batchlas::invalid_argument if `A` is not square, a parameter is out of
148/// range, or the batch items do not share one pattern
149/// @throws batchlas::convergence_error if a pivot is zero and `diagonal_shift`
150/// cannot stabilise it
151/// @ingroup api_sparse
152template <Backend B, typename T>
155 const ILUKParams<T>& params = ILUKParams<T>());
156
157/// @brief Bytes of workspace the `Span` overload of iluk_factorize needs for `A`.
158///
159/// An upper bound: it sizes against the symbolic ILU(k) pattern, and the numeric
160/// phase only prunes. Costs one host-side symbolic factorization of item 0.
161/// @param ctx queue the factorization will run on
162/// @param A the matrix batch that will be factored
163/// @param params the parameters that will be used
164/// @return required workspace size in bytes
165/// @throws batchlas::invalid_argument for a non-square `A`, out-of-range params,
166/// or `validate_batch_sparsity == false` with a batch of more than one
167/// @ingroup api_sparse_lowlevel
168template <Backend B, typename T>
169BATCHLAS_API size_t iluk_buffer_size(Queue& ctx,
171 const ILUKParams<T>& params = ILUKParams<T>());
172
173/// @brief Computes the ILU(k) factor into caller-supplied memory instead of allocating.
174///
175/// Same factorization as the owning overload. Lets an iterative solver carve the
176/// preconditioner out of a workspace it already holds. The exact footprint is
177/// only known after the numeric phase, so it is reported through `bytes_used`
178/// rather than predicted.
179/// @param ctx queue used for the device phase
180/// @param A batch of square CSR matrices sharing one pattern; not modified
181/// @param workspace at least iluk_buffer_size bytes of USM the host can write
182/// (the level schedule is copied in from the host)
183/// @param params fill, drop and pivot controls
184/// @param bytes_used if non-null, receives the bytes of `workspace` the factor occupies
185/// @return a view that aliases `workspace` and is valid only while it is
186/// @throws batchlas::invalid_argument, batchlas::convergence_error as the owning overload
187/// @ingroup api_sparse
188template <Backend B, typename T>
191 Span<std::byte> workspace,
192 const ILUKParams<T>& params,
193 size_t* bytes_used = nullptr);
194
195/// @brief Applies the ILU(k) preconditioner: `out` := \f$(LU)^{-1}\f$ `rhs`.
196///
197/// For every batch item and every column of `rhs`, solves \f$L y = r\f$ (unit
198/// diagonal) and then \f$U x = y\f$ by level-scheduled sparse triangular solves,
199/// and writes \f$x\f$ to `out`. Asynchronous: nothing is waited on, so repeated
200/// applications inside an iterative solver stay pipelined.
201/// @param ctx queue the kernel is enqueued on
202/// @param M factor from iluk_factorize
203/// @param rhs batch of n x nrhs right-hand sides; not modified
204/// @param out batch of n x nrhs results
205/// @param workspace unused; iluk_apply_buffer_size is 0
206/// @return event of the solve kernel
207/// @throws batchlas::invalid_argument if `rhs` / `out` rows differ from `M.n`,
208/// their column counts differ, their batch size differs from
209/// `M.batch_size`, or `M` has no level schedule
210/// @throws batchlas::convergence_error if `M.u_diagonals_usable` is false
211/// @ingroup api_sparse
212template <Backend B, typename T>
213BATCHLAS_API Event iluk_apply(Queue& ctx,
214 const ILUKView<T>& M,
217 Span<std::byte> workspace = Span<std::byte>());
218
219/// @brief iluk_apply on an owning factor; forwards `M.view()`.
220/// @ingroup api_sparse
221template <Backend B, typename T>
223 const ILUKPreconditioner<T>& M,
226 Span<std::byte> workspace = Span<std::byte>()) {
227 return iluk_apply<B, T>(ctx, M.view(), rhs, out, workspace);
228}
229
230/// @brief Bytes of workspace iluk_apply needs; currently always 0.
231/// @ingroup api_sparse_lowlevel
232template <Backend B, typename T>
233BATCHLAS_API size_t iluk_apply_buffer_size(Queue& ctx,
234 const ILUKView<T>& M,
237
238/// @brief iluk_apply_buffer_size on an owning factor; forwards `M.view()`.
239/// @ingroup api_sparse_lowlevel
240template <Backend B, typename T>
242 const ILUKPreconditioner<T>& M,
245 return iluk_apply_buffer_size<B, T>(ctx, M.view(), rhs, out);
246}
247
248} // namespace batchlas
249
250namespace batchlas {
251
252// Backend-deducing overloads: `f(ctx, ...)` uses ctx.backend().
253// See BATCHLAS_DISPATCH_ON_QUEUE in blas/queue-dispatch.hh.
254
259
260} // namespace batchlas
Non-owning view of a batch of matrices: the argument type of every entry point.
Definition matrix.hh:485
Owning batch of batch_size() matrices of equal (capacity) shape.
Definition matrix.hh:1399
Enumerations, scalar traits and scalar concepts shared by every public header.
size_t iluk_apply_buffer_size(Queue &ctx, const ILUKView< T > &M, const MatrixView< T, MatrixFormat::Dense > &rhs, const MatrixView< T, MatrixFormat::Dense > &out)
Bytes of workspace iluk_apply needs; currently always 0.
void iluk_build_level_schedule(ILUKPreconditioner< T > &M)
Builds the triangular-solve level schedule from M.lu's sparsity pattern.
size_t iluk_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::CSR > &A, const ILUKParams< T > &params=ILUKParams< T >())
Bytes of workspace the Span overload of iluk_factorize needs for A.
Event iluk_apply(Queue &ctx, const ILUKView< T > &M, const MatrixView< T, MatrixFormat::Dense > &rhs, const MatrixView< T, MatrixFormat::Dense > &out, Span< std::byte > workspace=Span< std::byte >())
Applies the ILU(k) preconditioner: out := rhs.
ILUKPreconditioner< T > iluk_factorize(Queue &ctx, const MatrixView< T, MatrixFormat::CSR > &A, const ILUKParams< T > &params=ILUKParams< T >())
Computes the ILU(k) factor of a batch of CSR matrices into owned storage.
The batched containers: Matrix (owning), MatrixView (non-owning), KernelMatrixView (device-side),...
Backend deduction from the Queue and the USM pointer check.
#define BATCHLAS_DISPATCH_ON_QUEUE(NAME)
Completion handle of enqueued work, returned by every entry point.
Parameters of an incomplete LU factorization with level-of-fill k, ILU(k).
Definition iluk.hh:28
int levels_of_fill
k; 0 keeps exactly A's pattern. Must be >= 0.
Definition iluk.hh:29
base_type< T >::type drop_tolerance
relative drop threshold; >= 0
Definition iluk.hh:31
base_type< T >::type fill_factor
per-row cap on kept entries relative to A's row; >= 1
Definition iluk.hh:32
base_type< T >::type diag_pivot_threshold
relative pivot threshold; >= 0
Definition iluk.hh:33
bool validate_batch_sparsity
Check that every batch item shares one CSR pattern.
Definition iluk.hh:37
T diagonal_shift
added to (or substituted for) a pivot that is too small
Definition iluk.hh:30
bool modified_ilu
add dropped entries to the diagonal (MILU) instead of discarding them
Definition iluk.hh:34
Owning ILU(k) factor of a batch of CSR matrices sharing one pattern.
Definition iluk.hh:77
UnifiedVector< int > u_rows
row indices ordered by backward-solve level
Definition iluk.hh:86
T diagonal_shift
the ILUKParams value it was built with
Definition iluk.hh:99
bool modified_ilu
the ILUKParams value it was built with
Definition iluk.hh:103
base_type< T >::type drop_tolerance
the ILUKParams value it was built with
Definition iluk.hh:100
int n
order of each matrix
Definition iluk.hh:96
UnifiedVector< int > l_rows
row indices ordered by forward-solve level
Definition iluk.hh:84
UnifiedVector< int > u_level_ptr
size u_levels + 1
Definition iluk.hh:87
UnifiedVector< int > diag_positions
size n * batch_size
Definition iluk.hh:82
int u_levels
number of backward-solve levels
Definition iluk.hh:89
bool u_diagonals_usable
Whether every U diagonal can be made non-singular with diagonal_shift.
Definition iluk.hh:94
int batch_size
number of matrices
Definition iluk.hh:97
int levels_of_fill
the ILUKParams value it was built with
Definition iluk.hh:98
ILUKView< T > view() const
Non-owning view of this factor; valid while *this is alive and unmodified.
Definition iluk.hh:106
ILUKPreconditioner()
1x1 placeholder with a single stored non-zero; iluk_factorize replaces it.
Definition iluk.hh:79
UnifiedVector< int > l_level_ptr
size l_levels + 1
Definition iluk.hh:85
base_type< T >::type fill_factor
the ILUKParams value it was built with
Definition iluk.hh:101
Matrix< T, MatrixFormat::CSR > lu
L (strict lower, unit diagonal implied) and U (upper, with diagonal).
Definition iluk.hh:81
base_type< T >::type diag_pivot_threshold
the ILUKParams value it was built with
Definition iluk.hh:102
int l_levels
number of forward-solve levels
Definition iluk.hh:88
Non-owning description of a factored ILU(k) preconditioner.
Definition iluk.hh:48
int u_levels
number of backward-solve levels
Definition iluk.hh:57
Span< int > l_level_ptr
size l_levels + 1
Definition iluk.hh:53
Span< int > u_level_ptr
size u_levels + 1
Definition iluk.hh:55
int n
order of each matrix
Definition iluk.hh:61
Span< int > diag_positions
size n * batch_size, absolute indices into lu's values
Definition iluk.hh:50
Span< int > l_rows
row indices ordered by forward-solve level
Definition iluk.hh:52
Span< int > u_rows
row indices ordered by backward-solve level
Definition iluk.hh:54
T diagonal_shift
shift the apply uses on U's diagonal
Definition iluk.hh:63
MatrixView< T, MatrixFormat::CSR > lu
L (strict lower, unit diagonal implied) and U (upper, with diagonal).
Definition iluk.hh:49
int batch_size
number of matrices
Definition iluk.hh:62
bool u_diagonals_usable
every U diagonal is non-zero after the shift
Definition iluk.hh:59
int l_levels
number of forward-solve levels
Definition iluk.hh:56
Non-zero count (per-item capacity) of a CSR matrix, as a distinct type.
Definition matrix.hh:58
Where BatchLAS work runs: a device, a backend, an ordered SYCL queue and a workspace arena.
Non-owning view of size() contiguous elements of type T.
Definition sycl-span.hh:39
Owning, growable array of T in USM shared memory (sycl::malloc_shared).
T type
The real type.
Definition enums.hh:30
batchlas::Device, batchlas::Event and batchlas::Queue: where work runs and how it is ordered.
batchlas::Span, the non-owning pointer + length every entry point takes for 1-D data.
batchlas::UnifiedVector, an owning growable array in USM shared memory.