extensions.hh
Go to the documentation of this file.
1#pragma once
2#include <batchlas/export.hh>
3#include <complex>
11#include <numeric>
12#include <limits>
13#include <cstddef>
14#include <cstdint>
15#include <vector>
16
17
18namespace batchlas {
19
20 template <typename T>
22
23 template <typename T>
24 struct StedcParams;
25
26 /// @addtogroup api_eigen
27 /// @{
28
29 /**
30 * @brief Options for syevx(): algorithm family, convergence controls, preconditioning
31 * and which part of the spectrum to compute.
32 *
33 * Every field has a default, so designated initialisers name only what changes. The
34 * range fields (`select`, `il`, `iu`, `vl`, `vu`, `abstol`, `order`) mirror LAPACK
35 * `?syevx`'s RANGE, IL/IU, VL/VU and ABSTOL arguments.
36 *
37 * @tparam T scalar type of the matrix (real or complex); tolerances on the real type
38 * @see @ref design_syevx_range for the range semantics, @ref perf_syevx for routing
39 */
40 template <typename T>
41 struct SyevxParams {
43 SyevxAlgorithm method = SyevxAlgorithm::Auto; ///< Algorithm family; `Auto` defers to syevx_select_algorithm().
44 OrthoAlgorithm algorithm = OrthoAlgorithm::Chol2; ///< Orthogonalisation used inside the iterative solvers.
45 size_t ortho_iterations = 2; ///< Orthogonalisation passes per use.
46 size_t iterations = 100; ///< Cap on outer iterations of LOBPCG / Filtered.
47 size_t extra_directions = 0; ///< Extra search directions carried beyond `neigs`.
48 bool find_largest = true; ///< Extremal selection: largest (`true`) or smallest (`false`) eigenpairs.
49 T absolute_tolerance = T(std::numeric_limits<float_type>::epsilon()); ///< Absolute residual tolerance of the iterative paths.
50 T relative_tolerance = T(std::numeric_limits<float_type>::epsilon()); ///< Relative residual tolerance of the iterative paths.
51 /// Caller-owned ILU(k) factor used as preconditioner, or nullptr.
52 /// ILU(k) approximates \f$ A^{-1} \f$, so it is valid only for the SMALLEST
53 /// eigenpairs: `find_largest = true` with a preconditioner is rejected, not ignored.
55 /// Preconditioner family. `Auto`: ILU(k) when a factor is supplied or requested,
56 /// otherwise none (unless `BATCHLAS_SYEVX_PRECONDITIONER` names a default).
58 /// Build the ILU(k) factor inside syevx from the caller's workspace. Requires a CSR
59 /// `A` and `find_largest = false`; mutually exclusive with #preconditioner.
61 ILUKParams<T> iluk_params{}; ///< Parameters of the factor built under #build_preconditioner.
62 size_t filter_degree = 0; ///< Chebyshev degree for SyevxAlgorithm::Filtered; 0 selects a default.
63 /// LOBPCG only: power-iteration steps on the random start block; -1 selects the
64 /// default, 0 disables. Ignored unless #find_largest is true.
66 const SyevxInstrumentation<T>* instrumentation = nullptr; ///< Optional convergence-history sink; nullptr records nothing.
67
68 SyevxSelect select = SyevxSelect::Extremal; ///< Range kind: `Extremal` (with #find_largest), `Index` or `Value`.
69
70 /// `select == Index`: first wanted index, 0-based and inclusive, into the
71 /// ASCENDING spectrum. `il > iu` is an empty request and is rejected.
72 int64_t il = 0;
73 int64_t iu = -1; ///< `select == Index`: last wanted index, inclusive; `iu < 0` means n-1.
74 /// `select == Value`: lower end of the half-open interval (vl, vu], as in LAPACK.
75 /// The count is data-dependent per item and is reported through syevx's `m` output.
77 float_type vu = float_type(0); ///< `select == Value`: upper end (inclusive).
78
79 /// Absolute tolerance per eigenvalue; non-positive means \f$ \varepsilon \|T\| \f$.
80 /// Forwarded to StebzParams::abstol; ignored by paths that use a full decomposition.
82
83 SortOrder order = SortOrder::Ascending; ///< Output order for `Index` and `Value`; for `Extremal` it follows #find_largest.
84 };
85
86 /**
87 * @brief Optional per-iteration convergence histories recorded by the iterative syevx paths.
88 *
89 * Every history span is laid out `[iter][batch][eig]`. With `iteration_stride` and
90 * `batch_stride` zero the strides default to `batch_size * neigs` and `neigs`. Empty
91 * spans and null pointers record nothing. Passed through SyevxParams::instrumentation.
92 */
93 template <typename T>
96
97 Span<float_type> best_residual_history{}; ///< Best residual norm seen so far, per eigenpair.
98 Span<float_type> current_residual_history{}; ///< Residual norm of the current iterate (needs #store_current_residual).
99 Span<float_type> convergence_rate_history{}; ///< Per-iteration residual reduction (needs #store_convergence_rate).
100 Span<float_type> ritz_value_history{}; ///< Ritz values per iteration (needs #store_ritz_values).
101
102 int32_t* iterations_done = nullptr; ///< Optional per-batch-item iteration count.
103
104 size_t max_iterations = 0; ///< Rows the history spans can hold.
105 size_t store_every = 1; ///< Record every k-th iteration.
106 size_t iteration_stride = 0; ///< Stride between iterations; 0 means `batch_size * neigs`.
107 size_t batch_stride = 0; ///< Stride between batch items; 0 means `neigs`.
108
109 bool store_current_residual = false; ///< Record #current_residual_history.
110 bool store_convergence_rate = true; ///< Record #convergence_rate_history.
111 bool store_ritz_values = false; ///< Record #ritz_value_history.
112 };
113
114 /** @brief Parameters for lanczos(). */
115 template <typename T>
118 OrthoAlgorithm ortho_algorithm = OrthoAlgorithm::CGS2; ///< Orthogonalisation of each new Lanczos vector.
119 size_t ortho_iterations = 2; ///< Orthogonalisation passes per vector.
120 size_t reorthogonalization_iterations = 2; ///< Iterations between full reorthogonalisations.
121 bool sort_enabled = true; ///< Sort eigenvalues (and eigenvectors) on output.
122 SortOrder sort_order = SortOrder::Ascending; ///< Order used when #sort_enabled.
123 };
124 /// @}
125
126 /// @addtogroup api_qr
127 /// @{
128
129 /**
130 * @brief Orthonormalises the columns (or rows) of each batch item in place.
131 *
132 * With `transA == NoTrans` each n-column block of `A` is replaced by an orthonormal
133 * basis of its column space, \f$ A \leftarrow Q \f$ with \f$ Q^H Q = I \f$; with
134 * `Trans` the rows are orthonormalised instead. Asynchronous.
135 *
136 * @param ctx queue the work is enqueued on
137 * @param A batch of m x n matrices; overwritten with the orthonormal vectors
138 * @param transA `NoTrans` for columns, `Trans` for rows
139 * @param workspace at least ortho_buffer_size() bytes for the same `algo`
140 * @param algo orthogonalisation algorithm (see OrthoAlgorithm)
141 * @return event of the last enqueued kernel
142 * @note The workspace depends on `algo`: size it with the algorithm the call runs.
143 */
144 template <Backend B, typename T>
145 BATCHLAS_API Event ortho(Queue& ctx,
146 const MatrixView<T, MatrixFormat::Dense>& A, //A is overwritten with orthogonal vectors
147 Transpose transA,
148 Span<std::byte> workspace,
150
151 /**
152 * @brief Orthonormalises `A` against the columns of an external basis `M`, in place.
153 *
154 * Each of `iterations` passes projects out the span of `M`,
155 * \f$ A \leftarrow A - M (M^H A) \f$, and then orthonormalises `A` with `algo`.
156 *
157 * @param ctx queue the work is enqueued on
158 * @param A batch of vectors to orthonormalise; overwritten
159 * @param M external basis, read only
160 * @param transA `NoTrans`: the vectors are the columns of `A`; transposed: its rows
161 * @param transM `NoTrans`: the basis vectors are the columns of `M`; transposed: its rows
162 * @param workspace at least ortho_buffer_size() bytes for the same arguments
163 * @param algo orthogonalisation algorithm run after each projection
164 * @param iterations number of project-then-orthonormalise passes
165 * @return event of the last enqueued kernel
166 * @pre `M`'s vectors are orthonormal, and the vector counts of `A` and `M` sum to at
167 * most the vector length.
168 * @throws batchlas::invalid_argument if the vector counts exceed the vector length
169 */
170 template <Backend B, typename T>
171 BATCHLAS_API Event ortho(Queue& ctx,
172 const MatrixView<T, MatrixFormat::Dense>& A, //A is overwritten with orthogonal vectors
173 const MatrixView<T, MatrixFormat::Dense>& M, //External metric
174 Transpose transA,
175 Transpose transM,
176 Span<std::byte> workspace,
178 size_t iterations = 2);
179
180 /** @brief Required workspace, in bytes, for ortho(); arguments as for the call itself. @ingroup api_qr_lowlevel */
181 template <Backend B, typename T>
182 BATCHLAS_API size_t ortho_buffer_size(Queue& ctx,
184 Transpose transA,
186
187 /** @brief Required workspace, in bytes, for the external-basis ortho(); same arguments. @ingroup api_qr_lowlevel */
188 template <Backend B, typename T>
189 BATCHLAS_API size_t ortho_buffer_size(Queue& ctx,
192 Transpose transA,
193 Transpose transM,
195 size_t iterations = 2);
196
197 /// @}
198
199 /// @addtogroup api_eigen
200 /// @{
201
202 /**
203 * @brief Selected eigenpairs of a batch of symmetric/Hermitian matrices, dense or CSR.
204 *
205 * Computes \f$ A v = \lambda v \f$ for the part of the spectrum that
206 * `params.select` names: the `neigs` largest or smallest eigenpairs (`Extremal`, with
207 * SyevxParams::find_largest) or an index block (`Index`). A value interval
208 * (`Value`) has a data-dependent count and needs the overload that reports `m`.
209 *
210 * The algorithm comes from syevx_select_algorithm(): `Direct` (full syev() on a private
211 * copy, then selection), `DirectSubset` (two-stage reduction plus stebz()/stein()),
212 * `LOBPCG` or `Filtered` (Chebyshev subspace iteration). `A` is never modified; the
213 * direct paths read its lower triangle. Asynchronous: returns once the work is enqueued.
214 *
215 * @tparam MFormat `MatrixFormat::Dense` or `MatrixFormat::CSR` (CSR always runs LOBPCG
216 * or Filtered)
217 * @param ctx queue the work is enqueued on
218 * @param A batch of n x n symmetric/Hermitian matrices, dense or CSR
219 * @param W eigenvalue output (real type), `neigs` entries per batch item
220 * @param neigs CAPACITY of `W` and of `V`'s columns per item, not the number produced
221 * @param workspace at least syevx_buffer_size() bytes
222 * @param jobz whether eigenvectors are written to `V`
223 * @param V n x `neigs` eigenvector output per item; used only for `EigenVectors`
224 * @param params selection range and algorithm parameters
225 * @param info per-item convergence status (0 = converged, > 0 LAPACK-like), or an
226 * empty span to not request it; see @ref md_docs_2cpp-api
227 * @return event of the last enqueued kernel
228 * @throws batchlas::invalid_argument for `SyevxSelect::Value` (use the `m`-taking
229 * overload); an `Index` block outside [0, n) or with `neigs != iu - il + 1`;
230 * a non-extremal range with CSR input or an explicit LOBPCG/Filtered `method`;
231 * an ILU(k) or `Jacobi` preconditioner with `find_largest`, or conflicting
232 * preconditioner fields.
233 * @see @ref perf_syevx, @ref design_syevx_range
234 */
235 template <Backend B, typename T, MatrixFormat MFormat>
236 BATCHLAS_API Event syevx(Queue& ctx,
237 const MatrixView<T, MFormat>& A,
238 Span<typename base_type<T>::type> W,
239 size_t neigs,
240 Span<std::byte> workspace,
243 const SyevxParams<T>& params = SyevxParams<T>(),
245
246 /** @brief syevx() on an owning `A` without eigenvectors: no `V` argument. */
247 template <Backend B, typename T, MatrixFormat MFormat>
248 inline Event syevx(Queue& ctx,
249 const Matrix<T, MFormat>& A,
250 Span<typename base_type<T>::type> W,
251 size_t neigs,
252 Span<std::byte> workspace,
254 const SyevxParams<T>& params = SyevxParams<T>(),
255 Span<int32_t> info = Span<int32_t>()) {
256 return syevx<B,T,MFormat>(ctx, MatrixView<T, MFormat>(A), W, neigs, workspace, jobz, MatrixView<T, MatrixFormat::Dense>(), params, info);
257 }
258
259 /**
260 * @brief syevx() with a per-batch-item count of eigenpairs actually found.
261 * Required for `SyevxSelect::Value`.
262 *
263 * Item `b` gets `min(m[b], neigs)` eigenpairs (the lowest ones, when truncating), the
264 * rest of its `W` is left untouched, the remaining columns of its `V` are written as
265 * EXACTLY ZERO so the back-transform can run over a uniform column count, and `m[b]`
266 * reports the TRUE count, so `m[b] > neigs` is the caller's overflow signal.
267 * Other parameters as for the `m`-less form.
268 *
269 * @param m Per-item count, at least `A.batch_size()` entries. Device-writable.
270 *
271 * @trap Unambiguous against the `m`-less forms only because `Span`'s scalar
272 * constructor is `explicit`; a bare `{}` in positions 4-6 matches both and is
273 * ambiguous.
274 */
275 template <Backend B, typename T, MatrixFormat MFormat>
276 BATCHLAS_API Event syevx(Queue& ctx,
277 const MatrixView<T, MFormat>& A,
278 Span<typename base_type<T>::type> W,
280 size_t neigs,
281 Span<std::byte> workspace,
284 const SyevxParams<T>& params = SyevxParams<T>(),
286
287 /** @brief The `m`-reporting syevx() on an owning `A` without eigenvectors. */
288 template <Backend B, typename T, MatrixFormat MFormat>
289 inline Event syevx(Queue& ctx,
290 const Matrix<T, MFormat>& A,
291 Span<typename base_type<T>::type> W,
293 size_t neigs,
294 Span<std::byte> workspace,
296 const SyevxParams<T>& params = SyevxParams<T>(),
297 Span<int32_t> info = Span<int32_t>()) {
298 return syevx<B,T,MFormat>(ctx, MatrixView<T, MFormat>(A), W, m, neigs, workspace, jobz, MatrixView<T, MatrixFormat::Dense>(), params, info);
299 }
300
301 /**
302 * @brief Required workspace, in bytes, for syevx(); arguments as for the call minus
303 * the workspace. Unlike the solve, this accepts `SyevxSelect::Value`: sizing
304 * writes no counts.
305 * @ingroup api_eigen_lowlevel
306 */
307 template <Backend B, typename T, MatrixFormat MFormat>
308 BATCHLAS_API size_t syevx_buffer_size(Queue& ctx,
309 const MatrixView<T, MFormat>& A,
310 Span<typename base_type<T>::type> W,
311 size_t neigs,
314 const SyevxParams<T>& params = SyevxParams<T>());
315
316 /** @brief syevx_buffer_size() on an owning `A` without eigenvectors. @ingroup api_eigen_lowlevel */
317 template <Backend B, typename T, MatrixFormat MFormat>
318 inline size_t syevx_buffer_size(Queue& ctx,
319 const Matrix<T, MFormat>& A,
320 Span<typename base_type<T>::type> W,
321 size_t neigs,
323 const SyevxParams<T>& params = SyevxParams<T>()) {
325 }
326
327 /**
328 * @brief syevx_buffer_size() with the solve's `m` argument, which is ACCEPTED AND
329 * IGNORED, so a value-range caller writes sizing and solve with one argument list.
330 * @ingroup api_eigen_lowlevel
331 */
332 template <Backend B, typename T, MatrixFormat MFormat>
333 inline size_t syevx_buffer_size(Queue& ctx,
334 const MatrixView<T, MFormat>& A,
335 Span<typename base_type<T>::type> W,
337 size_t neigs,
340 const SyevxParams<T>& params = SyevxParams<T>()) {
341 (void)m;
342 return syevx_buffer_size<B,T,MFormat>(ctx, A, W, neigs, jobz, V, params);
343 }
344
345 /** @brief The `m`-taking syevx_buffer_size() on an owning `A`; `m` is ignored. @ingroup api_eigen_lowlevel */
346 template <Backend B, typename T, MatrixFormat MFormat>
347 inline size_t syevx_buffer_size(Queue& ctx,
348 const Matrix<T, MFormat>& A,
349 Span<typename base_type<T>::type> W,
351 size_t neigs,
353 const SyevxParams<T>& params = SyevxParams<T>()) {
354 (void)m;
356 }
357
358 /**
359 * @brief Resolved, algorithm-independent description of the caller's request.
360 *
361 * Solvers and `*_buffer_size` must derive their behaviour from this struct, never from
362 * SyevxParams::select / il / iu / find_largest, so solve and sizing cannot disagree.
363 * `vl`/`vu` are deliberately absent: they are forwarded verbatim from `params`.
364 */
366 bool value_range; ///< true: (vl, vu]; false: the index block [il, iu]
367 int64_t il; ///< valid iff !value_range; 0-based, inclusive
368 int64_t iu; ///< valid iff !value_range; 0-based, inclusive
369 /// Upper bound on eigenpairs per item, clamped to [0, n] so a consumer may index
370 /// [il, il + max_count) unguarded. Exact for an index block; a value range may exceed it.
371 int64_t max_count;
372 bool reverse; ///< write the selected block in descending order
373 };
374
375 /**
376 * @brief Normalizes a range request into a SyevxResolvedRange. Legality is NOT checked
377 * here; it clamps rather than throws so it stays usable when sizing.
378 *
379 * @param n Matrix dimension
380 * @param neigs Capacity of W and V per batch item (see `syevx`)
381 * @param select SyevxParams::select
382 * @param find_largest SyevxParams::find_largest (Extremal only)
383 * @param il,iu SyevxParams::il / iu (Index only; iu < 0 means n-1)
384 * @param order SyevxParams::order (Index and Value only)
385 * @return the normalised request
386 * @ingroup api_eigen_lowlevel
387 */
389 size_t neigs,
391 bool find_largest,
392 int64_t il,
393 int64_t iu,
394 SortOrder order);
395
396 /// @brief syevx_resolve_range() reading the range fields from `params`; distinguished
397 /// from the 7-argument form by arity.
398 /// @ingroup api_eigen_lowlevel
399 template <typename T>
401 size_t neigs,
402 const SyevxParams<T>& params) {
403 return syevx_resolve_range(n, neigs, params.select, params.find_largest,
404 params.il, params.iu, params.order);
405 }
406
407 /**
408 * @brief Resolves SyevxParams::method (and the BATCHLAS_SYEVX_ALGORITHM override) to a
409 * concrete, implemented algorithm. Never returns `Auto`, and is deterministic
410 * so that `syevx` and `syevx_buffer_size` always agree on the choice.
411 *
412 * @param format Matrix format of A (sparse formats run LOBPCG unless `Filtered` is
413 * requested explicitly)
414 * @param n Matrix dimension
415 * @param neigs Number of requested eigenpairs
416 * @param requested Algorithm requested via SyevxParams::method
417 * @param subset_supported Whether DirectSubset is available for this T/format
418 * @param jobz Whether eigenvectors are wanted; load-bearing for the choice
419 * @param batch_size Load-bearing: the subset solver starves at small batch
420 * @param select Excludes the algorithms that cannot answer the requested range
421 * @return the algorithm syevx() will run
422 * @throws batchlas::invalid_argument for sparse input, or an explicit LOBPCG/Filtered
423 * `method`, with a non-extremal range. The same request made through
424 * `BATCHLAS_SYEVX_ALGORITHM` degrades to `Direct` with a one-time warning.
425 * @see @ref perf_syevx for the measured thresholds behind the choice
426 * @ingroup api_eigen_lowlevel
427 */
429 int64_t n,
430 size_t neigs,
431 SyevxAlgorithm requested,
432 bool subset_supported,
434 int64_t batch_size = 1,
436
437 /**
438 * @brief Partial eigensolve by full decomposition followed by selection. Runs `syev`
439 * on a private copy of A. Dense input only, but every scalar type and every
440 * `SyevxSelect` range: the universal fallback.
441 *
442 * @param W Eigenvalue output, stride `neigs`; entries past `min(m[b], neigs)` untouched
443 * @param V Eigenvector output. Columns past `min(m[b], neigs)` are written as EXACTLY
444 * ZERO -- unlike `W` -- so this path and `syevx_direct_subset` agree.
445 * @param m Per-item count in the requested range, or an empty span to not report it.
446 * `m[b] > neigs` is the truncation signal; the LOWEST `neigs` are kept.
447 *
448 * Other parameters as for syevx().
449 * @ingroup api_eigen_lowlevel
450 */
451 template <Backend B, typename T, MatrixFormat MFormat>
452 BATCHLAS_API Event syevx_direct(Queue& ctx,
453 const MatrixView<T, MFormat>& A,
454 Span<typename base_type<T>::type> W,
456 size_t neigs,
457 Span<std::byte> workspace,
458 JobType jobz,
460 const SyevxParams<T>& params,
462
463 /**
464 * @brief `syevx_direct` without the `m` output; `Extremal` and `Index` only.
465 * Distinguished from the form above by ARITY, so no `{}` trap can fire.
466 * @ingroup api_eigen_lowlevel
467 */
468 template <Backend B, typename T, MatrixFormat MFormat>
470 const MatrixView<T, MFormat>& A,
471 Span<typename base_type<T>::type> W,
472 size_t neigs,
473 Span<std::byte> workspace,
474 JobType jobz,
476 const SyevxParams<T>& params,
477 Span<int32_t> info = Span<int32_t>()) {
478 return syevx_direct<B, T, MFormat>(ctx, A, W, Span<int32_t>(), neigs, workspace,
479 jobz, V, params, info);
480 }
481
482 /**
483 * @brief Required workspace, in bytes, for syevx_direct(). Takes no `m`: sizing writes
484 * no counts, and this path's workspace is range-independent.
485 * @ingroup api_eigen_lowlevel
486 */
487 template <Backend B, typename T, MatrixFormat MFormat>
488 BATCHLAS_API size_t syevx_direct_buffer_size(Queue& ctx,
489 const MatrixView<T, MFormat>& A,
490 Span<typename base_type<T>::type> W,
491 size_t neigs,
492 JobType jobz,
494 const SyevxParams<T>& params);
495
496 /**
497 * @brief Partial eigensolve by two-stage reduction plus a subset tridiagonal solve
498 * (`stebz` + `stein`) and a narrowed back-transform. Real scalars and dense
499 * input only; zero capacity is rejected (stein requires k >= 1).
500 *
501 * `W`, `V` and `m` follow the same contract as `syevx_direct`, with `stein` doing the
502 * zeroing. The descending reversal is applied last, in the finalize kernel: `stein`'s
503 * cluster detection walks consecutive eigenvalues and requires ascending input.
504 * Other parameters as for syevx().
505 *
506 * @pre in-order Queue; `A` square
507 * @throws batchlas::unsupported for complex or sparse input
508 * @throws batchlas::invalid_argument for `neigs == 0`, an out-of-range index block,
509 * `vl >= vu`, a too-short `m`, or an out-of-order Queue
510 * @see syevx_direct_subset_supported()
511 * @ingroup api_eigen_lowlevel
512 */
513 template <Backend B, typename T, MatrixFormat MFormat>
514 BATCHLAS_API Event syevx_direct_subset(Queue& ctx,
515 const MatrixView<T, MFormat>& A,
516 Span<typename base_type<T>::type> W,
518 size_t neigs,
519 Span<std::byte> workspace,
520 JobType jobz,
522 const SyevxParams<T>& params,
524
525 /** @brief `syevx_direct_subset` without the `m` output; arity-disambiguated as above. @ingroup api_eigen_lowlevel */
526 template <Backend B, typename T, MatrixFormat MFormat>
528 const MatrixView<T, MFormat>& A,
529 Span<typename base_type<T>::type> W,
530 size_t neigs,
531 Span<std::byte> workspace,
532 JobType jobz,
534 const SyevxParams<T>& params,
535 Span<int32_t> info = Span<int32_t>()) {
536 return syevx_direct_subset<B, T, MFormat>(ctx, A, W, Span<int32_t>(), neigs, workspace,
537 jobz, V, params, info);
538 }
539
540 /**
541 * @brief Required workspace, in bytes, for syevx_direct_subset().
542 *
543 * Sizing writes no counts, but it is NOT range-independent: a `Value` range needs
544 * room for up to n eigenvalues per item, so the sizes come from syevx_resolve_range().
545 * @ingroup api_eigen_lowlevel
546 */
547 template <Backend B, typename T, MatrixFormat MFormat>
548 BATCHLAS_API size_t syevx_direct_subset_buffer_size(Queue& ctx,
549 const MatrixView<T, MFormat>& A,
550 Span<typename base_type<T>::type> W,
551 size_t neigs,
552 JobType jobz,
554 const SyevxParams<T>& params);
555
556 /** @brief Whether `syevx_direct_subset` supports this scalar type and format. @ingroup api_eigen_lowlevel */
557 template <typename T, MatrixFormat MFormat>
558 inline constexpr bool syevx_direct_subset_supported() {
559 return MFormat == MatrixFormat::Dense && std::is_same_v<T, typename base_type<T>::type>;
560 }
561
562 /**
563 * @brief Partial eigensolve by LOBPCG (locally optimal block preconditioned conjugate
564 * gradient). Dense and CSR input; extremal ranges only.
565 *
566 * Iterates a block of `neigs` (+ SyevxParams::extra_directions) vectors until the
567 * residuals meet the tolerances or SyevxParams::iterations is reached; optionally
568 * preconditioned by ILU(k) for the smallest eigenpairs. Parameters as for syevx().
569 * @ingroup api_eigen_lowlevel
570 */
571 template <Backend B, typename T, MatrixFormat MFormat>
572 BATCHLAS_API Event syevx_lobpcg(Queue& ctx,
573 const MatrixView<T, MFormat>& A,
574 Span<typename base_type<T>::type> W,
575 size_t neigs,
576 Span<std::byte> workspace,
577 JobType jobz,
579 const SyevxParams<T>& params,
581
582 /** @brief Required workspace, in bytes, for syevx_lobpcg(). @ingroup api_eigen_lowlevel */
583 template <Backend B, typename T, MatrixFormat MFormat>
584 BATCHLAS_API size_t syevx_lobpcg_buffer_size(Queue& ctx,
585 const MatrixView<T, MFormat>& A,
586 Span<typename base_type<T>::type> W,
587 size_t neigs,
588 JobType jobz,
590 const SyevxParams<T>& params);
591
592 /**
593 * @brief Chebyshev-filtered subspace iteration: amplifies the wanted end of the
594 * spectrum, then Rayleigh-Ritz. Needs only matvecs; dense and CSR input,
595 * extremal ranges only.
596 *
597 * The polynomial degree is SyevxParams::filter_degree (0 picks a default).
598 * Parameters as for syevx().
599 * @ingroup api_eigen_lowlevel
600 */
601 template <Backend B, typename T, MatrixFormat MFormat>
602 BATCHLAS_API Event syevx_filtered(Queue& ctx,
603 const MatrixView<T, MFormat>& A,
604 Span<typename base_type<T>::type> W,
605 size_t neigs,
606 Span<std::byte> workspace,
607 JobType jobz,
609 const SyevxParams<T>& params,
611
612 /** @brief Required workspace, in bytes, for syevx_filtered(). @ingroup api_eigen_lowlevel */
613 template <Backend B, typename T, MatrixFormat MFormat>
614 BATCHLAS_API size_t syevx_filtered_buffer_size(Queue& ctx,
615 const MatrixView<T, MFormat>& A,
616 Span<typename base_type<T>::type> W,
617 size_t neigs,
618 JobType jobz,
620 const SyevxParams<T>& params);
621
622 /**
623 * @brief All eigenvalues, and optionally eigenvectors, of a batch of symmetric/Hermitian
624 * matrices by a full-length Lanczos tridiagonalisation.
625 *
626 * Builds an n-step Krylov basis with reorthogonalisation, then solves the resulting
627 * tridiagonal problem. Intended for sparse (CSR) input; dense works too.
628 *
629 * @param ctx queue the work is enqueued on
630 * @param A batch of n x n symmetric/Hermitian matrices (CSR or dense); read only
631 * @param W eigenvalue output (real type), n entries per batch item
632 * @param workspace at least lanczos_buffer_size() bytes
633 * @param jobz whether eigenvectors are written to `V`
634 * @param V n x n eigenvector output per item; used only for `EigenVectors`
635 * @param params orthogonalisation and sorting options
636 * @return event of the last enqueued kernel
637 */
638 template <Backend B, typename T, MatrixFormat MFormat>
639 BATCHLAS_API Event lanczos(Queue& ctx,
640 const MatrixView<T, MFormat>& A,
641 Span<typename base_type<T>::type> W,
642 Span<std::byte> workspace,
645 const LanczosParams<T>& params = LanczosParams<T>());
646
647 /** @brief lanczos() on an owning `A` without eigenvectors: no `V` argument. */
648 template <Backend B, typename T, MatrixFormat MFormat>
649 inline Event lanczos(Queue& ctx,
650 const Matrix<T, MFormat>& A,
651 Span<typename base_type<T>::type> W,
652 Span<std::byte> workspace,
654 const LanczosParams<T>& params = LanczosParams<T>()) {
655 return lanczos<B,T,MFormat>(ctx, MatrixView<T,MFormat>(A), W, workspace, jobz, MatrixView<T, MatrixFormat::Dense>(), params);
656 }
657
658 /** @brief Required workspace, in bytes, for lanczos(); arguments as for the call itself. @ingroup api_eigen_lowlevel */
659 template <Backend B, typename T, MatrixFormat MFormat>
660 BATCHLAS_API size_t lanczos_buffer_size(Queue& ctx,
661 const MatrixView<T, MFormat>& A,
662 Span<typename base_type<T>::type> W,
665 const LanczosParams<T>& params = LanczosParams<T>());
666
667 /// @}
668
669 /// @addtogroup api_tridiag
670 /// @{
671
672 /**
673 * @brief Eigenvalues, and optionally eigenvectors, of a batch of small symmetric
674 * tridiagonal matrices by shifted QR with 2x2 reflections.
675 *
676 * One work-group of 32 per batch item; the working arrays live in local memory when
677 * they fit and in the workspace otherwise. With `EigenVectors`, `Q` is first set to
678 * the identity.
679 *
680 * @param ctx queue the work is enqueued on
681 * @param alphas diagonals, packed: item `b` at `alphas[b*n .. b*n+n)`
682 * @param betas off-diagonals, packed with the same stride `n`
683 * @param W eigenvalue output (real type), packed with stride `n`
684 * @param workspace at least tridiagonal_solver_buffer_size() bytes
685 * @param jobz whether `Q` receives the eigenvectors
686 * @param Q n x n eigenvector output per item
687 * @param n order of each tridiagonal matrix
688 * @param batch_size number of batch items
689 * @return event of the kernel; the call waits for it before returning
690 * @throws batchlas::convergence_error if an eigenvalue of any item does not deflate
691 * within 30 QR steps
692 * @note steqr() and stedc() are the maintained tridiagonal solvers.
693 */
694 template <Backend B, typename T>
695 BATCHLAS_API Event tridiagonal_solver(Queue& ctx,
696 Span<T> alphas,
697 Span<T> betas,
698 Span<typename base_type<T>::type> W,
699 Span<std::byte> workspace,
700 JobType jobz,
702 size_t n,
703 size_t batch_size);
704
705 /** @brief Required workspace, in bytes, for tridiagonal_solver(). @ingroup api_tridiag_lowlevel */
706 template <Backend B, typename T>
707 BATCHLAS_API size_t tridiagonal_solver_buffer_size(Queue& ctx, size_t n, size_t batch_size, JobType jobz);
708
709 // Convention in the tridiagonal group below: `VectorView<T>` for one vector per batch
710 // item (it carries inc/stride/batch_size), `Span<...>` for flat per-item arrays and byte
711 // workspaces. Demoting a `VectorView` to a `Span` silently drops the stride.
712
713 /** @brief How stebz() selects a subset of the spectrum. */
714 enum class EigenRangeType {
715 All, ///< Every eigenvalue
716 Index, ///< Eigenvalues il..iu inclusive, 0-based, in ascending order
717 Value ///< Eigenvalues in the half-open interval (vl, vu]
718 };
719
720 /** @brief Parameters for stebz() (bisection on a symmetric tridiagonal matrix). */
721 template <typename T>
722 struct StebzParams {
723 EigenRangeType range = EigenRangeType::All; ///< Which eigenvalues to compute.
724 int64_t il = 0; ///< First wanted index (0-based, inclusive), range == Index
725 int64_t iu = -1; ///< Last wanted index (0-based, inclusive), range == Index; < 0 means n-1
726 T vl = T(0); ///< Lower bound (exclusive), range == Value
727 T vu = T(0); ///< Upper bound (inclusive), range == Value
728 /// Absolute tolerance on each eigenvalue. Non-positive means
729 /// \f$ \varepsilon \|T\| \f$, i.e. full working precision.
730 T abstol = T(0);
731 SortOrder order = SortOrder::Ascending; ///< Output order of `w`.
732 /// Safety cap on bisection steps per eigenvalue; the loop also exits on
733 /// interval convergence.
734 int32_t max_iterations = 128;
735 };
736
737 /**
738 * @brief Computes selected eigenvalues of a batch of symmetric tridiagonal matrices by
739 * bisection on Sturm sequence sign counts. Eigenvalues only; one work-item per
740 * eigenvalue, so a subset costs proportionally less. Use stein() for vectors.
741 *
742 * @param ctx Execution context/device queue
743 * @param d Diagonal, n entries per batch item
744 * @param e Off-diagonal, n-1 entries per batch item
745 * @param w Output eigenvalues; must hold at least the number selected (n for `Value`)
746 * @param m Output count of eigenvalues found, per batch item (device-written)
747 * @param ws Pre-allocated workspace buffer, at least stebz_buffer_size() bytes
748 * @param params Selection range, tolerance and ordering
749 * @return Event Event to track operation completion
750 * @throws batchlas::invalid_argument for n < 1, a short `e`, `m` or `w`, or an invalid
751 * index range
752 */
753 template <Backend B, typename T>
754 BATCHLAS_API Event stebz(Queue& ctx,
755 const VectorView<T>& d,
756 const VectorView<T>& e,
757 const VectorView<T>& w,
759 const Span<std::byte>& ws,
760 StebzParams<T> params = StebzParams<T>());
761
762 /**
763 * @brief Required workspace size, in bytes, for stebz().
764 *
765 * @trap `params` is REQUIRED, not defaulted: it is the only argument carrying `T`, and
766 * defaulting it silently drops the queue-deducing overload. Pass `StebzParams<T>{}`.
767 * @ingroup api_tridiag_lowlevel
768 */
769 template <Backend B, typename T>
770 BATCHLAS_API size_t stebz_buffer_size(Queue& ctx,
771 size_t n,
772 size_t batch_size,
773 StebzParams<T> params);
774
775 /** @brief Parameters for stein() (inverse iteration on a symmetric tridiagonal). */
776 template <typename T>
777 struct SteinParams {
778 /// Fixed number of inverse-iteration steps; two or three suffice for eigenvalues
779 /// accurate to working precision. No convergence test is made.
780 int32_t max_iterations = 3;
781 /// Eigenvalues closer than `ortho_threshold * ||T||` form one cluster and have their
782 /// vectors explicitly reorthogonalised (LAPACK dstein uses 1e-3).
783 T ortho_threshold = T(1e-3);
784 uint32_t seed = 0x5eed1234u; ///< Seed of the random starting vectors.
785 };
786
787 /**
788 * @brief Computes eigenvectors of a batch of symmetric tridiagonal matrices by inverse
789 * iteration, given previously computed eigenvalues. Pairs with stebz().
790 *
791 * @param ctx Execution context/device queue
792 * @param d Diagonal, n entries per batch item
793 * @param e Off-diagonal, n-1 entries per batch item
794 * @param w Eigenvalues, k per batch item, in ascending order (cluster detection walks
795 * consecutive values)
796 * @param k Number of eigenvectors to compute, at least 1
797 * @param Z Output eigenvectors, n x k per batch item, columns matching w
798 * @param ws Pre-allocated workspace buffer, at least stein_buffer_size() bytes
799 * @param params Iteration count and clustering threshold
800 * @return Event Event to track operation completion
801 * @throws batchlas::invalid_argument for n < 1, k < 1, or a short `e`, `w` or `Z`
802 */
803 template <Backend B, typename T>
804 BATCHLAS_API Event stein(Queue& ctx,
805 const VectorView<T>& d,
806 const VectorView<T>& e,
807 const VectorView<T>& w,
808 size_t k,
810 const Span<std::byte>& ws,
811 SteinParams<T> params = SteinParams<T>());
812
813 /**
814 * @brief Sentinel meaning: every batch item has all `k` eigenvalues valid.
815 *
816 * @trap Spell this rather than a bare `{}` at `counts`: `{}` in a position two
817 * overloads both accept has silently selected the wrong one in this codebase before.
818 */
820
821 /**
822 * @brief stein() with a per-batch-item count of valid eigenvalues.
823 *
824 * `k` is a capacity (columns of `Z`, entries of `w` per item); `counts[b]` is how many
825 * leading entries of item `b`'s `w` are real eigenvalues -- the rest hold stale
826 * workspace, and are neither iterated on nor joined to a cluster. Columns
827 * `[counts[b], k)` of `Z` are written as EXACTLY ZERO, not left untouched, so callers
828 * may back-transform a uniform `k` columns. `counts` is read on the device, so it may
829 * be the `m` span `stebz` just wrote; empty (or `stein_all_counts`) means all `k` are valid,
830 * and it is clamped to `[0, k]`. Other parameters as for the `counts`-less form.
831 *
832 * @throws batchlas::invalid_argument additionally if a non-empty `counts` is shorter
833 * than the batch
834 */
835 template <Backend B, typename T>
836 BATCHLAS_API Event stein(Queue& ctx,
837 const VectorView<T>& d,
838 const VectorView<T>& e,
839 const VectorView<T>& w,
840 size_t k,
841 Span<const int32_t> counts,
843 const Span<std::byte>& ws,
844 SteinParams<T> params = SteinParams<T>());
845
846 /**
847 * @brief Required workspace size, in bytes, for stein().
848 *
849 * Sizes on the capacity `k`, so the `counts` overload needs no sizing entry of its own.
850 * @trap `params` is REQUIRED, as for stebz_buffer_size(): it is the only argument
851 * carrying `T`.
852 * @ingroup api_tridiag_lowlevel
853 */
854 template <Backend B, typename T>
855 BATCHLAS_API size_t stein_buffer_size(Queue& ctx,
856 size_t n,
857 size_t k,
858 size_t batch_size,
859 SteinParams<T> params);
860
861 /** @brief Shift used by the CTA STEQR's implicit QR/QL steps. @see @ref algo_steqr */
863 Lapack = 0, ///< LAPACK-style implicit shift (the stable dsteqr formulation).
864 Wilkinson = 1, ///< Wilkinson shift computed from the trailing 2x2 block.
865 };
866
867 /**
868 * @brief Update scheme of the CTA STEQR's implicit QR/QL steps.
869 *
870 * The default (SteqrParams::cta_update_scheme) is `EXP`.
871 * @see @ref algo_steqr, section "steqr: the CTA update schemes EXP and PG"
872 */
873 enum class SteqrUpdateScheme {
874 PG = 0, ///< Parlett-Gray style scalar recurrence.
875 EXP = 1, ///< Explicit similarity update mirroring steqr.cc's bulge-chasing math.
876 };
877
878 /** @brief Parameters for steqr() and steqr_cta(), and for stedc()'s leaf solves. */
879 template <typename T>
880 struct SteqrParams {
881 /// Rotations are applied in blocks of this size; larger means excess FLOPs but
882 /// better memory reuse. 1 fully serialises them.
883 size_t block_size = 32;
884 size_t max_sweeps = 50; ///< Cap on Francis QR sweeps; 2-3 typically suffice per eigenvalue.
885 /// If false, the eigenvector matrix is set to the identity and the rotations are
886 /// applied to it; if true, they are applied to the caller's matrix (a back-transform).
887 bool back_transform = false;
888 bool block_rotations = false; ///< Apply the rotations in #block_size blocks.
889 bool sort = true; ///< Sort eigenpairs on output.
890 bool transpose_working_vectors = true; ///< Interleave the working copies of d and e across the batch.
891 SortOrder sort_order = SortOrder::Ascending; ///< Order used when #sort is set.
892
893 /// CTA STEQR only: multiplies the baseline work-group size, LCM(N, sub_group_size).
894 /// 0 picks a tuned value per scalar type and partition width.
896
898 SteqrUpdateScheme cta_update_scheme = SteqrUpdateScheme::EXP; ///< CTA STEQR only: update scheme.
899 };
900
901
902 // `info` on steqr/stedc/syev/syevx/gesvd and their tiers: per-item convergence status,
903 // an ACCUMULATOR zeroed once by the entry point the caller invoked and only raised
904 // below (contract and helpers: src/kernels/common/info_span.hh).
905 // evidence: docs/cpp-api.md#convergence-status-syev-syevx-gesvd-steqr-stedc
906 // A DEFAULTED trailing parameter, not an old-arity forwarder as in potrf: these are
907 // instantiated from macros that spell the list out, and with `jobz`, `params` and
908 // `eigvects` already defaulted a forwarder would be AMBIGUOUS with the primary.
909 // evidence: docs/design/vendor-independence.md#info-spans-on-syev-gesvd-and-steqr-forwarder-or-default
910
911 /**
912 * @brief Eigenvalues, and optionally eigenvectors, of a batch of symmetric tridiagonal
913 * matrices by implicit QR/QL iteration (LAPACK `?steqr`).
914 *
915 * Real `T` only. Runs the CTA kernel (steqr_cta()) when n is at most the device's
916 * largest sub-group size (never on ROCm) and the work-group kernel otherwise;
917 * steqr_buffer_size() returns the larger of the two needs.
918 *
919 * @param ctx queue the work is enqueued on
920 * @param d diagonal, n entries per batch item
921 * @param e off-diagonal, n-1 entries per batch item
922 * @param eigenvalues output, n per batch item; ascending when SteqrParams::sort is set
923 * @param ws at least steqr_buffer_size() bytes
924 * @param jobz whether `eigvects` receives the eigenvectors
925 * @param params sweep cap, deflation threshold, sorting and CTA tuning
926 * @param eigvects n x n per item; with SteqrParams::back_transform the rotations
927 * are applied to its contents, otherwise it is set to the identity
928 * first
929 * @param info per-item convergence status (0 = converged, > 0 LAPACK-like), or
930 * empty to not request it; see @ref md_docs_2cpp-api
931 * @throws batchlas::invalid_argument if `eigvects` is not n x n with the same batch
932 * @throws batchlas::convergence_error as steqr_cta() does, when the CTA kernel runs
933 * @return event of the last enqueued kernel
934 * @see @ref algo_steqr
935 */
936 template <Backend B, typename T>
937 BATCHLAS_API Event steqr(Queue& ctx, const VectorView<T>& d, const VectorView<T>& e,
938 const VectorView<T>& eigenvalues, const Span<std::byte>& ws, JobType jobz = JobType::NoEigenVectors, SteqrParams<T> params = SteqrParams<T>(),
941
942 /**
943 * @brief steqr() pinned to the CTA kernel: one sub-group partition per matrix, for small
944 * n (runtime-dispatched to compile-time specialised kernels). Parameters as for
945 * steqr().
946 * @throws batchlas::invalid_argument unless 1 <= n <= 32
947 * @throws batchlas::convergence_error if `BATCHLAS_STEQR_CTA_CHECK` is set and an item
948 * ran out of sweeps (the check waits on the queue; unset, nothing is thrown)
949 * @ingroup api_tridiag_lowlevel
950 */
951 template <Backend B, typename T>
952 BATCHLAS_API Event steqr_cta(Queue& ctx, const VectorView<T>& d, const VectorView<T>& e,
953 const VectorView<T>& eigenvalues, const Span<std::byte>& ws,
958
959 /** @brief Required workspace, in bytes, for steqr(); arguments as for the call. @ingroup api_tridiag_lowlevel */
960 template <typename T>
961 BATCHLAS_API size_t steqr_buffer_size(Queue& ctx, const VectorView<T>& d, const VectorView<T>& e,
962 const VectorView<T>& eigenvalues, JobType jobz = JobType::NoEigenVectors, SteqrParams<T> params = SteqrParams<T>());
963
964 /** @brief Required workspace, in bytes, for steqr_cta(); arguments as for the call. @ingroup api_tridiag_lowlevel */
965 template <typename T>
966 BATCHLAS_API size_t steqr_cta_buffer_size(Queue& ctx, const VectorView<T>& d, const VectorView<T>& e,
967 const VectorView<T>& eigenvalues, JobType jobz = JobType::NoEigenVectors, SteqrParams<T> params = SteqrParams<T>());
968
969
970 /// @}
971
972 // === CTA small-matrix extensions (n <= 32) ===
973
974 /// @addtogroup api_tridiag
975 /// @{
976
977 /**
978 * @brief CTA-optimized symmetric tridiagonal reduction (SYTD2-style), n <= 32.
979 * Overwrites A with the tridiagonal and reflector storage and returns (d,e)
980 * plus reflector scalars in tau. `ws` is unused, kept for API compatibility.
981 * @ingroup api_tridiag_lowlevel
982 */
983 template <Backend B, typename T>
984 BATCHLAS_API Event sytrd_cta(Queue& ctx,
986 const VectorView<T>& d_out,
987 const VectorView<T>& e_out,
988 const VectorView<T>& tau_out,
989 Uplo uplo,
990 const Span<std::byte>& ws,
991 size_t cta_wg_size_multiplier = 1);
992
993 /**
994 * @brief LATRD-like panel factorization used by blocked SYTRD, `Uplo::Lower` only.
995 * Householder vectors go into A, in sytrd's SYTD2-style reflector layout;
996 * only the first `ib` columns of W are written (W is treated as n x nb).
997 *
998 * @param j0 first column of the panel
999 * @param ib panel width
1000 * @param wg_hint work-group size hint; 0 lets the implementation choose
1001 * @param fuse_trailing_update fold the panel's share of the trailing update into the
1002 * kernel
1003 * @ingroup api_tridiag_lowlevel
1004 */
1005 template <Backend B, typename T>
1006 BATCHLAS_API Event latrd_lower_panel(Queue& ctx,
1008 const VectorView<T>& e_out,
1009 const VectorView<T>& tau_out,
1011 int32_t j0,
1012 int32_t ib,
1013 int32_t wg_hint = 0,
1014 bool fuse_trailing_update = false);
1015
1016 /**
1017 * @brief LATRD-like panel factorization (Lower only), view-based overload.
1018 *
1019 * Pass pre-sliced views instead of (j0, ib):
1020 * - a_panel = A({j0, SliceEnd()}, {j0, SliceEnd()}) (must be square)
1021 * - e_panel = E(Slice(j0, j0 + ib))
1022 * - tau_panel = TAU(Slice(j0, j0 + ib))
1023 * - w_panel = Wmat({j0, SliceEnd()}, {0, ib})
1024 * e_panel/tau_panel size must match w_panel.cols().
1025 * @ingroup api_tridiag_lowlevel
1026 */
1027 template <Backend B, typename T>
1028 BATCHLAS_API Event latrd_lower_panel(Queue& ctx,
1029 const MatrixView<T, MatrixFormat::Dense>& a_panel_in,
1030 const VectorView<T>& e_panel_out,
1031 const VectorView<T>& tau_panel_out,
1032 const MatrixView<T, MatrixFormat::Dense>& w_panel_in,
1033 int32_t wg_hint = 0,
1034 bool fuse_trailing_update = false);
1035
1036 /**
1037 * @brief Blocked symmetric/Hermitian tridiagonal reduction, n > 32: LATRD-style panel
1038 * plus a BLAS-3 trailing update. Same SYTD2-style output as `sytrd_cta`.
1039 * In-order queue only.
1040 *
1041 * A is overwritten with the reflector storage (LAPACK `sytd2` layout, with the tridiagonal
1042 * entries restored on the first off-diagonal), the diagonal goes to `d_out`, the
1043 * off-diagonal to `e_out` and the reflector scalars to `tau_out`.
1044 *
1045 * @pre `ctx` is in-order (SYCL kernels and backend BLAS calls are not separately ordered)
1046 * @pre `ws` holds at least sytrd_blocked_buffer_size() bytes (the W panel)
1047 * @throws batchlas::invalid_argument on a non-square A, short d/e/tau, a batch mismatch or
1048 * an out-of-order `ctx`
1049 * @throws batchlas::unsupported for `Uplo::Upper` (only Lower is implemented)
1050 */
1051 template <Backend B, typename T>
1052 BATCHLAS_API Event sytrd_blocked(Queue& ctx,
1054 const VectorView<T>& d_out,
1055 const VectorView<T>& e_out,
1056 const VectorView<T>& tau_out,
1057 Uplo uplo,
1058 const Span<std::byte>& ws,
1059 int32_t block_size = tuning::SYTRD_BLOCK_SIZE_MEDIUM);
1060
1061 /** @brief Required workspace, in bytes, for sytrd_blocked(); arguments as for the call. @ingroup api_tridiag_lowlevel */
1062 template <Backend B, typename T>
1063 BATCHLAS_API size_t sytrd_blocked_buffer_size(Queue& ctx,
1065 const VectorView<T>& d,
1066 const VectorView<T>& e,
1067 const VectorView<T>& tau,
1068 Uplo uplo,
1069 int32_t block_size = tuning::SYTRD_BLOCK_SIZE_MEDIUM);
1070
1071 /**
1072 * @brief First stage of two-stage reduction: dense -> band (LAPACK xSYTRD_SY2SB).
1073 * Overwrites `A` with reflector storage and writes the band into `AB`;
1074 * `tau_out` has size (n-kd). In-order queue, `Uplo::Lower` only.
1075 *
1076 * Band storage (AB), shape (kd+1) x n:
1077 * - `Uplo::Lower`: AB(1+i-j,j) = A(i,j) for j<=i<=min(n,j+kd).
1078 * - `Uplo::Upper`: AB(kd+1+i-j,j) = A(i,j) for max(1,j-kd)<=i<=j.
1079 * @ingroup api_tridiag_lowlevel
1080 */
1081 template <Backend B, typename T>
1082 BATCHLAS_API Event sytrd_sy2sb(Queue& ctx,
1085 const VectorView<T>& tau_out,
1086 Uplo uplo,
1087 int32_t kd,
1088 const Span<std::byte>& ws);
1089
1090 /** @brief Required workspace, in bytes, for sytrd_sy2sb(); arguments as for the call. @ingroup api_tridiag_lowlevel */
1091 template <Backend B, typename T>
1092 BATCHLAS_API size_t sytrd_sy2sb_buffer_size(Queue& ctx,
1095 const VectorView<T>& tau_out,
1096 Uplo uplo,
1097 int32_t kd);
1098
1099 /**
1100 * @brief Second stage of two-stage reduction: band -> tridiagonal (LAPACK
1101 * xSB2ST/xHB2ST bulge chasing, VECT='N'). In-order queue, `Uplo::Lower` only.
1102 *
1103 * Band storage (AB) is as for `sytrd_sy2sb`. `d_out` (n) and `e_out` (n-1) are always
1104 * real-valued; `tau_out` (n-1) is unused downstream and is set to 0.
1105 * @ingroup api_tridiag_lowlevel
1106 */
1107 template <Backend B, typename T>
1108 BATCHLAS_API Event sytrd_sb2st(Queue& ctx,
1110 const VectorView<typename base_type<T>::type>& d_out,
1111 const VectorView<typename base_type<T>::type>& e_out,
1112 const VectorView<T>& tau_out,
1113 Uplo uplo,
1114 int32_t kd,
1115 const Span<std::byte>& ws,
1116 int32_t block_size);
1117
1118 /** @brief Required workspace, in bytes, for sytrd_sb2st(); arguments as for the call. @ingroup api_tridiag_lowlevel */
1119 template <Backend B, typename T>
1120 BATCHLAS_API size_t sytrd_sb2st_buffer_size(Queue& ctx,
1122 const VectorView<typename base_type<T>::type>& d_out,
1123 const VectorView<typename base_type<T>::type>& e_out,
1124 const VectorView<T>& tau_out,
1125 Uplo uplo,
1126 int32_t kd,
1127 int32_t block_size);
1128
1129 /** @brief Sweep schedule for sytrd_band_reduction() (successive band reduction). */
1131 /// Diagonals eliminated per sweep (Algorithm 2: d^(i)); 0 means the default
1132 /// schedule. If sweeps exceed the sequence length, the last value is reused.
1133 std::vector<int32_t> d_seq{0};
1134
1135 /// Block size per sweep (Algorithm 2: nb^(i)).
1136 /// If sweeps exceed sequence length, the last value is reused.
1137 std::vector<int32_t> block_size_seq{32};
1138
1139 int32_t max_sweeps = -1; ///< Maximum number of sweeps; < 0 means the implementation default.
1140
1141 /// Debug/testing: chase steps for sytrd_band_reduction_single_step();
1142 /// <= 0 means exactly one step.
1143 int32_t max_steps = 1;
1144
1145 int32_t kd_work = 0; ///< Working band semibandwidth; <= 0 means the implementation default.
1146 };
1147
1148 /**
1149 * @brief Symmetric/Hermitian band -> tridiagonal reduction (BANDR1-style). In-order
1150 * queue, `Uplo::Lower` only. (d,e) are real; for complex input `e_out` is the
1151 * magnitude of the (possibly phased) subdiagonal, and `tau_out` is set to 0.
1152 *
1153 * Band storage is as for sytrd_sy2sb(). This overload uses one block size for every
1154 * sweep.
1155 */
1156 template <Backend B, typename T>
1159 const VectorView<typename base_type<T>::type>& d_out,
1160 const VectorView<typename base_type<T>::type>& e_out,
1161 const VectorView<T>& tau_out,
1162 Uplo uplo,
1163 int32_t kd,
1164 const Span<std::byte>& ws,
1165 int32_t block_size);
1166
1167 /** @brief sytrd_band_reduction() with an explicit per-sweep schedule. */
1168 template <Backend B, typename T>
1171 const VectorView<typename base_type<T>::type>& d_out,
1172 const VectorView<typename base_type<T>::type>& e_out,
1173 const VectorView<T>& tau_out,
1174 Uplo uplo,
1175 int32_t kd,
1176 const Span<std::byte>& ws,
1178
1179 template <Backend B, typename T>
1183 Uplo uplo,
1184 int32_t kd,
1185 const Span<std::byte>& ws,
1187
1188 /** @brief Required workspace, in bytes, for sytrd_band_reduction() with one block size. @ingroup api_tridiag_lowlevel */
1189 template <Backend B, typename T>
1190 BATCHLAS_API size_t sytrd_band_reduction_buffer_size(Queue& ctx,
1192 const VectorView<typename base_type<T>::type>& d_out,
1193 const VectorView<typename base_type<T>::type>& e_out,
1194 const VectorView<T>& tau_out,
1195 Uplo uplo,
1196 int32_t kd,
1197 int32_t block_size);
1198
1199 /** @brief Required workspace, in bytes, for the schedule-parameter sytrd_band_reduction(). @ingroup api_tridiag_lowlevel */
1200 template <Backend B, typename T>
1201 BATCHLAS_API size_t sytrd_band_reduction_buffer_size(Queue& ctx,
1203 const VectorView<typename base_type<T>::type>& d_out,
1204 const VectorView<typename base_type<T>::type>& e_out,
1205 const VectorView<T>& tau_out,
1206 Uplo uplo,
1207 int32_t kd,
1209
1210 template <Backend B, typename T>
1214 Uplo uplo,
1215 int32_t kd,
1217
1218 /**
1219 * @brief Debug/testing hook: execute exactly one BANDR1 “chase step”. Not a stable
1220 * public API. `ab_in` is lower-band with rows == kd+1, `abw_out` lower-band
1221 * with rows == kd_work+1 (kd_work from params). In-order queue, Lower only.
1222 * @ingroup api_tridiag_lowlevel
1223 */
1224 template <Backend B, typename T>
1228 Uplo uplo,
1229 int32_t kd,
1230 const Span<std::byte>& ws,
1232
1233 /** @brief Required workspace, in bytes, for sytrd_band_reduction_single_step(). @ingroup api_tridiag_lowlevel */
1234 template <Backend B, typename T>
1235 BATCHLAS_API size_t sytrd_band_reduction_single_step_buffer_size(Queue& ctx,
1238 Uplo uplo,
1239 int32_t kd,
1241
1242 /// @}
1243
1244 /// @addtogroup api_eigen
1245 /// @{
1246
1247 /**
1248 * @brief CTA-optimized symmetric eigen-solver (SYEV-like), n <= 32; real symmetric and
1249 * complex Hermitian. Overwrites A with eigenvectors when jobz == EigenVectors.
1250 * Eigenvalues ascend when SteqrParams::sort is set (default).
1251 * cta_wg_size_multiplier == 0 lets the tridiagonal solve pick its tuned value
1252 * and runs the reduction and back-transform at 1.
1253 *
1254 * This and the other `syev_*` tiers are what syev() routes between; call one directly
1255 * only to pin a tier. As in syev(), `a_in`'s `uplo` triangle is read, `eigenvalues`
1256 * gets n real values per item, and `info` is the per-item convergence status or empty.
1257 * @see @ref perf_syev (small-n kernel choice)
1258 * @ingroup api_eigen_lowlevel
1259 */
1260 template <Backend B, typename T>
1261 BATCHLAS_API Event syev_cta(Queue& ctx,
1263 Span<typename base_type<T>::type> eigenvalues,
1264 JobType jobz,
1265 Uplo uplo,
1266 const Span<std::byte>& ws,
1267 SteqrParams<T> steqr_params = SteqrParams<T>(),
1268 size_t cta_wg_size_multiplier = 0,
1269 Span<int32_t> info = Span<int32_t>());
1270
1271 /** @brief Required workspace, in bytes, for syev_cta(). @ingroup api_eigen_lowlevel */
1272 template <Backend B, typename T>
1273 BATCHLAS_API size_t syev_cta_buffer_size(Queue& ctx,
1275 JobType jobz,
1276 SteqrParams<T> steqr_params = SteqrParams<T>());
1277
1278 /**
1279 * @brief Fused (single-kernel) variant of syev_cta.
1280 *
1281 * Runs end to end inside one sub-group partition, so results track syev_cta only to
1282 * within the reassociation that fusing implies. Unlike syev_cta, A is left untouched
1283 * when jobz == NoEigenVectors. `ws` is accepted for API symmetry and ignored.
1284 * cta_wg_size_multiplier == 0 picks a tuned value per scalar type and width.
1285 * @ingroup api_eigen_lowlevel
1286 */
1287 template <Backend B, typename T>
1288 BATCHLAS_API Event syev_cta_fused(Queue& ctx,
1290 Span<typename base_type<T>::type> eigenvalues,
1291 JobType jobz,
1292 Uplo uplo,
1293 const Span<std::byte>& ws = Span<std::byte>(),
1294 SteqrParams<T> steqr_params = SteqrParams<T>(),
1295 size_t cta_wg_size_multiplier = 0,
1296 Span<int32_t> info = Span<int32_t>());
1297
1298 /** @brief Required workspace, in bytes, for syev_cta_fused(). @ingroup api_eigen_lowlevel */
1299 template <Backend B, typename T>
1300 BATCHLAS_API size_t syev_cta_fused_buffer_size(Queue& ctx,
1302 JobType jobz,
1303 SteqrParams<T> steqr_params = SteqrParams<T>());
1304
1305 /** @brief Parameters for syev_jacobi_cta() (cyclic two-sided Jacobi). */
1306 template <typename T>
1308 using Real = typename base_type<T>::type;
1309
1310 /// A rotation is applied to pivot pair (p,q) only when
1311 /// \f$ |a_{pq}| > \mathrm{tol\_multiplier} \cdot n \varepsilon \sqrt{|a_{pp}| |a_{qq}|} \f$.
1312 /// This relative test, not the classical absolute one, is what yields the high
1313 /// relative accuracy; raising it trades accuracy for sweeps.
1315
1316 size_t max_sweeps = 30; ///< Cap on cyclic sweeps; a safety net for pathological inputs.
1317
1318 bool sort = true; ///< Sort eigenpairs on output.
1319 SortOrder sort_order = SortOrder::Ascending; ///< Order used when #sort is set.
1320
1321 /// Multiplies the baseline work-group size, LCM(P, sub_group_size) for the
1322 /// compile-time partition width P chosen from n. Clamped by device limits.
1324 };
1325
1326 /**
1327 * @brief CTA-optimized Jacobi symmetric/Hermitian eigen-solver, n <= 32; real
1328 * symmetric and complex Hermitian.
1329 *
1330 * Trades throughput for high *relative* accuracy on graded or badly scaled input. The
1331 * underlying theorem (Demmel & Veselic, SIMAX 13(4), 1992) is proved for symmetric
1332 * positive definite input; indefinite matrices are solved correctly but do not inherit
1333 * the bound. Overwrites A with eigenvectors when jobz == EigenVectors and leaves it
1334 * untouched otherwise; `ws` is accepted for API symmetry and ignored.
1335 * @ingroup api_eigen_lowlevel
1336 */
1337 template <Backend B, typename T>
1338 BATCHLAS_API Event syev_jacobi_cta(Queue& ctx,
1340 Span<typename base_type<T>::type> eigenvalues,
1341 JobType jobz,
1342 Uplo uplo,
1343 const Span<std::byte>& ws = Span<std::byte>(),
1345 Span<int32_t> info = Span<int32_t>());
1346
1347 /** @brief Required workspace, in bytes, for syev_jacobi_cta(). @ingroup api_eigen_lowlevel */
1348 template <Backend B, typename T>
1349 BATCHLAS_API size_t syev_jacobi_cta_buffer_size(Queue& ctx,
1351 JobType jobz,
1352 JacobiParams<T> params = JacobiParams<T>());
1353
1354 /**
1355 * @brief Blocked symmetric/Hermitian eigen-solver (SYEV-like) for medium/large n:
1356 * sytrd_blocked -> stedc -> ormqr_blocked. Overwrites A with eigenvectors when
1357 * jobz == EigenVectors. `Uplo::Upper` input is first mirrored into the lower
1358 * triangle, in place, and then takes the Lower path.
1359 * @ingroup api_eigen_lowlevel
1360 */
1361 template <Backend B, typename T>
1362 BATCHLAS_API Event syev_blocked(Queue& ctx,
1364 Span<typename base_type<T>::type> eigenvalues,
1365 JobType jobz,
1366 Uplo uplo,
1367 const Span<std::byte>& ws,
1368 StedcParams<typename base_type<T>::type> stedc_params = StedcParams<typename base_type<T>::type>(),
1369 Span<int32_t> info = Span<int32_t>());
1370
1371 /** @brief Required workspace, in bytes, for syev_blocked(); the same for either `uplo`. @ingroup api_eigen_lowlevel */
1372 template <Backend B, typename T>
1373 BATCHLAS_API size_t syev_blocked_buffer_size(Queue& ctx,
1375 JobType jobz,
1376 Uplo uplo,
1377 StedcParams<typename base_type<T>::type> stedc_params = StedcParams<typename base_type<T>::type>());
1378
1379 /**
1380 * @brief Two-stage symmetric/Hermitian eigen-solver (SYEV-like) for large n:
1381 * sytrd_sy2sb -> sytrd_sb2st -> stedc, then phase/sign recovery and a reflector
1382 * back-transform. Band width from choose_two_stage_kd (override with
1383 * BATCHLAS_SYEV_TWO_STAGE_KD). `Uplo::Upper` input is mirrored into the lower
1384 * triangle first, as in syev_blocked().
1385 * @see @ref perf_syev for where syev() routes to it
1386 * @ingroup api_eigen_lowlevel
1387 */
1388 template <Backend B, typename T>
1389 BATCHLAS_API Event syev_two_stage(Queue& ctx,
1391 Span<typename base_type<T>::type> eigenvalues,
1392 JobType jobz,
1393 Uplo uplo,
1394 const Span<std::byte>& ws,
1395 StedcParams<typename base_type<T>::type> stedc_params = StedcParams<typename base_type<T>::type>(),
1396 Span<int32_t> info = Span<int32_t>());
1397
1398 /** @brief Required workspace, in bytes, for syev_two_stage(). @ingroup api_eigen_lowlevel */
1399 template <Backend B, typename T>
1400 BATCHLAS_API size_t syev_two_stage_buffer_size(Queue& ctx,
1402 JobType jobz,
1403 Uplo uplo,
1404 StedcParams<typename base_type<T>::type> stedc_params = StedcParams<typename base_type<T>::type>());
1405
1406 /// @}
1407
1408 /// @addtogroup api_svd
1409 /// @{
1410
1411 /**
1412 * @brief Unblocked GEBRD-like reduction to real bidiagonal form: reduces each square A
1413 * in place, returning (d,e) and the Householder scalars (tauq,taup).
1414 *
1415 * Computes \f$ A = Q B P^H \f$ with `B` upper bidiagonal: `d` gets its n diagonal and
1416 * `e` its n-1 superdiagonal entries (real), and the reflectors defining `Q` and `P`
1417 * are left in `A` with their scalars in `tauq` and `taup`, as in LAPACK `?gebrd`.
1418 * @see @ref perf_gesvd (where the unblocked form stops paying)
1419 * @ingroup api_svd_lowlevel
1420 */
1421 template <Backend B, typename T>
1422 BATCHLAS_API Event gebrd_unblocked(Queue& ctx,
1424 const VectorView<typename base_type<T>::type>& d,
1425 const VectorView<typename base_type<T>::type>& e,
1426 const VectorView<T>& tauq,
1427 const VectorView<T>& taup);
1428
1429 /**
1430 * @brief CTA-parallel small-matrix GEBRD for real square matrices, `1 <= n <= 32`,
1431 * where one CTA cooperatively reduces one matrix to upper bidiagonal form.
1432 * @ingroup api_svd_lowlevel
1433 */
1434 template <Backend B, typename T>
1435 BATCHLAS_API Event gebrd_cta(Queue& ctx,
1437 const VectorView<typename base_type<T>::type>& d,
1438 const VectorView<typename base_type<T>::type>& e,
1439 const VectorView<T>& tauq,
1440 const VectorView<T>& taup,
1441 size_t cta_wg_size_multiplier = 1);
1442
1443 /**
1444 * @brief Blocked GEBRD for real square dense matrices: DLABRD-style panel + GEMM.
1445 * Output layout as for gebrd_unblocked().
1446 * @ingroup api_svd_lowlevel
1447 */
1448 template <Backend B, typename T>
1449 BATCHLAS_API Event gebrd_blocked(Queue& ctx,
1451 const VectorView<typename base_type<T>::type>& d,
1452 const VectorView<typename base_type<T>::type>& e,
1453 const VectorView<T>& tauq,
1454 const VectorView<T>& taup,
1455 const Span<std::byte>& ws,
1456 int32_t block_size = 16);
1457
1458 /** @brief Required workspace, in bytes, for gebrd_blocked(); arguments as for the call. @ingroup api_svd_lowlevel */
1459 template <Backend B, typename T>
1460 BATCHLAS_API size_t gebrd_blocked_buffer_size(Queue& ctx,
1462 const VectorView<typename base_type<T>::type>& d,
1463 const VectorView<typename base_type<T>::type>& e,
1464 const VectorView<T>& tauq,
1465 const VectorView<T>& taup,
1466 int32_t block_size = 16);
1467
1468 /**
1469 * @brief Bidiagonal QR iteration for a real upper bidiagonal matrix. The matrix
1470 * overload also accumulates the alternating Givens rotations into `u` and `vh`,
1471 * matching LAPACK `BDSQR`: it returns `u * Q` and `P^T * vh`.
1472 *
1473 * @param d,e diagonal (n) and superdiagonal (n-1) per batch item
1474 * @param singular_values_out n singular values per item, packed
1475 * @param ws at least bdsqr_buffer_size() bytes
1476 * @param sort_desc sort the singular values (and vectors) in descending order
1477 * @param info per-item convergence status, or empty; see @ref md_docs_2cpp-api
1478 * @ingroup api_svd_lowlevel
1479 */
1480 template <Backend B, typename T>
1481 BATCHLAS_API Event bdsqr(Queue& ctx,
1482 const VectorView<T>& d,
1483 const VectorView<T>& e,
1484 Span<T> singular_values_out,
1485 const Span<std::byte>& ws,
1486 bool sort_desc = true,
1487 Span<int32_t> info = Span<int32_t>());
1488
1489 /**
1490 * @brief bdsqr() that also accumulates the rotations: `u <- u * Q`, `vh <- P^T * vh`.
1491 * Other parameters as for the values-only form.
1492 * @ingroup api_svd_lowlevel
1493 */
1494 template <Backend B, typename T>
1495 BATCHLAS_API Event bdsqr(Queue& ctx,
1496 const VectorView<T>& d,
1497 const VectorView<T>& e,
1498 Span<T> singular_values_out,
1499 const Span<std::byte>& ws,
1502 bool sort_desc = true,
1503 Span<int32_t> info = Span<int32_t>());
1504
1505 /** @brief Required workspace, in bytes, for bdsqr(). @ingroup api_svd_lowlevel */
1506 template <typename T>
1507 BATCHLAS_API size_t bdsqr_buffer_size(Queue& ctx,
1508 const VectorView<T>& d,
1509 const VectorView<T>& e,
1510 Span<T> singular_values_out);
1511
1512 /** @brief bdsqr_buffer_size() for the vector-accumulating form; `u` and `vh` do not change it. @ingroup api_svd_lowlevel */
1513 template <typename T>
1514 inline size_t bdsqr_buffer_size(Queue& ctx,
1515 const VectorView<T>& d,
1516 const VectorView<T>& e,
1517 Span<T> singular_values_out,
1520 static_cast<void>(u);
1521 static_cast<void>(vh);
1522 return bdsqr_buffer_size(ctx, d, e, singular_values_out);
1523 }
1524
1525 /**
1526 * @brief Bidiagonal divide-and-conquer SVD for a real upper bidiagonal matrix. Same
1527 * problem as `bdsqr`, via the symmetric tridiagonal eigenproblem of the
1528 * interleaved Golub-Kahan form (order `2n`) handed to `stedc`.
1529 *
1530 * Unlike `bdsqr`, which accumulates into whatever it is handed (`u <- u*Q`), `bdsdc`
1531 * WRITES the leading `n x n` block of `u` and of `vh` and leaves the rest untouched --
1532 * seed them with the identity if the trailing columns matter. Workspace is dominated
1533 * by a `2n x 2n` eigenvector matrix per batch item. Parameters as for bdsqr().
1534 * @see @ref design_gesvd, section "gesvd design: why bdsdc goes through Golub-Kahan"
1535 * @ingroup api_svd_lowlevel
1536 */
1537 template <Backend B, typename T>
1538 BATCHLAS_API Event bdsdc(Queue& ctx,
1539 const VectorView<T>& d,
1540 const VectorView<T>& e,
1541 Span<T> singular_values_out,
1542 const Span<std::byte>& ws,
1543 bool sort_desc = true,
1544 Span<int32_t> info = Span<int32_t>());
1545
1546 /** @brief bdsdc() that also writes the leading n x n blocks of `u` and `vh`. @ingroup api_svd_lowlevel */
1547 template <Backend B, typename T>
1548 BATCHLAS_API Event bdsdc(Queue& ctx,
1549 const VectorView<T>& d,
1550 const VectorView<T>& e,
1551 Span<T> singular_values_out,
1552 const Span<std::byte>& ws,
1555 bool sort_desc = true,
1556 Span<int32_t> info = Span<int32_t>());
1557
1558 /** @brief Required workspace, in bytes, for bdsdc(); `want_vectors` selects the overload. @ingroup api_svd_lowlevel */
1559 template <Backend B, typename T>
1560 BATCHLAS_API size_t bdsdc_buffer_size(Queue& ctx,
1561 const VectorView<T>& d,
1562 const VectorView<T>& e,
1563 Span<T> singular_values_out,
1564 bool want_vectors);
1565
1566
1567 /**
1568 * @brief ORMBR/UNMBR-style application of bidiagonal reduction reflectors. Supports
1569 * `vect='Q'` (tauq, CTA/blocked ORMQR path) and `vect='P'` (taup, blocked
1570 * compact-WY path over the right reflectors).
1571 *
1572 * Overwrites `c` with `op(Q) * c`, `c * op(Q)`, `op(P) * c` or `c * op(P)` according to
1573 * `vect`, `side` and `trans`, where `Q`/`P` are the reflectors that gebrd_blocked() (or
1574 * gebrd_unblocked()) left in `a` and `tau`, as LAPACK `?ormbr` / `?unmbr`: `'Q'` is of
1575 * order `a.rows()`, `'P'` of order `a.cols()`.
1576 *
1577 * @pre `tau` is unit-stride and packed by batch
1578 * @throws batchlas::invalid_argument on mismatched batch sizes or orders, a `vect`
1579 * other than `'Q'`/`'P'`, or a short or strided `tau`
1580 * @throws batchlas::unsupported for `Transpose::Trans` with complex `T` and `'P'`
1581 * (use `ConjTrans`)
1582 * @ingroup api_svd_lowlevel
1583 */
1584 template <Backend B, typename T>
1585 BATCHLAS_API Event ormbr(Queue& ctx,
1587 const VectorView<T>& tau,
1589 char vect,
1590 Side side,
1591 Transpose trans,
1592 const Span<std::byte>& ws,
1593 int32_t block_size = 32);
1594
1595 /** @brief Required workspace, in bytes, for ormbr(); arguments as for the call. @ingroup api_svd_lowlevel */
1596 template <Backend B, typename T>
1597 BATCHLAS_API size_t ormbr_buffer_size(Queue& ctx,
1599 const VectorView<T>& tau,
1601 char vect,
1602 Side side,
1603 Transpose trans,
1604 int32_t block_size = 32);
1605
1606 /**
1607 * @brief Blocked native SVD for real dense matrices: GEBRD-style dense -> bidiagonal,
1608 * then BDSQR (or BDSDC when explicitly selected), then ORMBR-style
1609 * back-transforms for full U and V^H.
1610 * @pre in-order Queue
1611 * @throws batchlas::invalid_argument for an out-of-order Queue or mis-shaped outputs
1612 * @throws batchlas::unsupported for complex input (use the Hermitian overload)
1613 * @ingroup api_svd_lowlevel
1614 */
1615 template <Backend B, typename T>
1616 BATCHLAS_API Event gesvd_blocked(Queue& ctx,
1618 Span<typename base_type<T>::type> singular_values,
1621 SvdVectors jobu,
1622 SvdVectors jobvh,
1623 const Span<std::byte>& ws,
1624 Span<int32_t> info = Span<int32_t>());
1625
1626 /**
1627 * @brief gesvd_blocked() for Hermitian input: only the `hermitian_uplo` triangle of
1628 * `a_in` is read.
1629 * @ingroup api_svd_lowlevel
1630 */
1631 template <Backend B, typename T>
1632 BATCHLAS_API Event gesvd_blocked(Queue& ctx,
1634 Span<typename base_type<T>::type> singular_values,
1637 SvdVectors jobu,
1638 SvdVectors jobvh,
1639 Uplo hermitian_uplo,
1640 const Span<std::byte>& ws,
1641 Span<int32_t> info = Span<int32_t>());
1642
1643 /** @brief Required workspace, in bytes, for gesvd_blocked(). @ingroup api_svd_lowlevel */
1644 template <Backend B, typename T>
1645 BATCHLAS_API size_t gesvd_blocked_buffer_size(Queue& ctx,
1647 Span<typename base_type<T>::type> singular_values,
1650 SvdVectors jobu,
1651 SvdVectors jobvh);
1652
1653 /** @brief Required workspace, in bytes, for the Hermitian gesvd_blocked(). @ingroup api_svd_lowlevel */
1654 template <Backend B, typename T>
1655 BATCHLAS_API size_t gesvd_blocked_buffer_size(Queue& ctx,
1657 Span<typename base_type<T>::type> singular_values,
1660 SvdVectors jobu,
1661 SvdVectors jobvh,
1662 Uplo hermitian_uplo);
1663
1664 /**
1665 * @brief CTA-oriented native SVD for small matrices, `max(m, n) <= 32`: bidiagonal
1666 * reduction, then the eigenproblem of \f$ B^T B \f$ (the normal equations).
1667 * Real input, rectangular either way.
1668 *
1669 * The normal equations square the condition number; use gesvdj_cta() where small
1670 * singular values matter.
1671 * @pre in-order Queue
1672 * @throws batchlas::invalid_argument for `max(m, n) > 32`, a genuinely `Thin` U or
1673 * V^H, or an out-of-order Queue
1674 * @throws batchlas::unsupported for complex input (use the Hermitian overload)
1675 * @see @ref perf_gesvd
1676 * @ingroup api_svd_lowlevel
1677 */
1678 template <Backend B, typename T>
1679 BATCHLAS_API Event gesvd_cta(Queue& ctx,
1681 Span<typename base_type<T>::type> singular_values,
1684 SvdVectors jobu,
1685 SvdVectors jobvh,
1686 const Span<std::byte>& ws,
1687 Span<int32_t> info = Span<int32_t>());
1688
1689 /** @brief gesvd_cta() for Hermitian input: only the `hermitian_uplo` triangle is read. @ingroup api_svd_lowlevel */
1690 template <Backend B, typename T>
1691 BATCHLAS_API Event gesvd_cta(Queue& ctx,
1693 Span<typename base_type<T>::type> singular_values,
1696 SvdVectors jobu,
1697 SvdVectors jobvh,
1698 Uplo hermitian_uplo,
1699 const Span<std::byte>& ws,
1700 Span<int32_t> info = Span<int32_t>());
1701
1702 /** @brief Required workspace, in bytes, for gesvd_cta(). @ingroup api_svd_lowlevel */
1703 template <Backend B, typename T>
1704 BATCHLAS_API size_t gesvd_cta_buffer_size(Queue& ctx,
1706 Span<typename base_type<T>::type> singular_values,
1709 SvdVectors jobu,
1710 SvdVectors jobvh);
1711
1712 /** @brief Required workspace, in bytes, for the Hermitian gesvd_cta(). @ingroup api_svd_lowlevel */
1713 template <Backend B, typename T>
1714 BATCHLAS_API size_t gesvd_cta_buffer_size(Queue& ctx,
1716 Span<typename base_type<T>::type> singular_values,
1719 SvdVectors jobu,
1720 SvdVectors jobvh,
1721 Uplo hermitian_uplo);
1722
1723 /**
1724 * @brief Parameters for the one-sided Jacobi SVD (`gesvdj_cta`).
1725 *
1726 * @trap Deliberately NOT a reuse of JacobiParams: its `sort_order` defaults to
1727 * Ascending while gesvd admits exactly one order (descending), so reuse invites
1728 * a silent flip.
1729 */
1730 template <typename T>
1732 using Real = typename base_type<T>::type;
1733
1734 /// A rotation is applied to pivot pair (p,q) only when
1735 /// \f$ |a_{pq}| > \mathrm{tol\_multiplier} \cdot n \varepsilon \sqrt{|a_{pp}| |a_{qq}|} \f$
1736 /// over the 2x2 Gram entries of the current columns; again the relative test.
1738
1739 size_t max_sweeps = 30; ///< Cap on cyclic sweeps; convergence normally takes well under 10.
1740
1741 /// Multiplies the baseline problems-per-work-group. Baseline is
1742 /// 32 / P problems, clamped by local memory and max work-group size.
1744
1745 /// \f$ \sigma_j \le \mathrm{zero\_sigma\_multiplier} \cdot \varepsilon \sigma_{\max} \f$
1746 /// means \f$ U_j \f$ is not determined by A and is filled from the orthogonal complement.
1748
1749 /// Optional per-problem diagnostic, the analogue of cusolverDnXgesvdjGetSweeps:
1750 /// when non-empty (size >= batch_size) the kernel writes each problem's sweeps.
1752 };
1753
1754 /**
1755 * @brief One-sided (Hestenes) Jacobi SVD for batches of small matrices.
1756 *
1757 * Computes \f$ A = U \operatorname{diag}(s) V^H \f$ with high RELATIVE accuracy: the
1758 * singular-value error is governed by the condition number of the column-equilibrated
1759 * matrix rather than of A, so graded input keeps its small values;
1760 * gesvd_cta() / gesvd_blocked() do not.
1761 *
1762 * max(m, n) <= 64 (32 for `complex<double>` with vectors), real and complex,
1763 * rectangular either way; needs sub-group size 32. A is DESTROYED, singular values are
1764 * returned descending, and no workspace is required.
1765 *
1766 * @throws batchlas::invalid_argument for a size above the cap or undersized outputs
1767 * @throws batchlas::unsupported if the device has no sub-group size 32
1768 * @see @ref design_gesvd, @ref perf_gesvd
1769 * @ingroup api_svd_lowlevel
1770 */
1771 template <Backend B, typename T>
1772 BATCHLAS_API Event gesvdj_cta(Queue& ctx,
1774 Span<typename base_type<T>::type> singular_values,
1777 SvdVectors jobu,
1778 SvdVectors jobvh,
1779 const Span<std::byte>& ws = Span<std::byte>(),
1781 Span<int32_t> info = Span<int32_t>());
1782
1783 /** @brief Required workspace, in bytes, for gesvdj_cta() (currently 0: all state is in local memory). @ingroup api_svd_lowlevel */
1784 template <Backend B, typename T>
1785 BATCHLAS_API size_t gesvdj_cta_buffer_size(Queue& ctx,
1787 Span<typename base_type<T>::type> singular_values,
1790 SvdVectors jobu,
1791 SvdVectors jobvh,
1792 GesvdjParams<T> params = GesvdjParams<T>());
1793
1794 /// @}
1795
1796 /**
1797 * @brief CTA-optimized application of Q from a QR/QL factorization (ORMQx/UNMQx
1798 * semantics) for very small matrices: applies the implicit Q of the Householder
1799 * reflectors (A, TAU) from GEQRF (Upper) or GEQLF (Lower) to C.
1800 *
1801 * Overwrites `c_in` with `op(Q) * C` (`Side::Left`) or `C * op(Q)` (`Side::Right`),
1802 * using the first `k` reflectors.
1803 * @ingroup api_qr_lowlevel
1804 */
1805 template <Backend B, typename T>
1806 BATCHLAS_API Event ormqx_cta(Queue& ctx,
1808 const VectorView<T>& tau_in,
1810 Uplo factorization,
1811 Side side,
1812 Transpose trans,
1813 int32_t k,
1814 const Span<std::byte>& ws,
1815 size_t cta_wg_size_multiplier = 1);
1816
1817 /// @addtogroup api_tridiag
1818 /// @{
1819
1820 /** @brief Secular-equation root finder used by stedc()'s merge step. */
1822 Rocm, ///< rocSOLVER-style solver (default).
1823 Legacy, ///< Original BatchLAS solver.
1824 };
1825
1826 /** @brief How the merge step (secular solve + eigenvector formation) is dispatched. */
1828 Auto = -1, ///< Use BatchLAS tuning tables for the current problem size
1829 Baseline, ///< Serial-per-root path (3 separate kernels)
1830 Fused, ///< One kernel: warp-parallel root solve + build/normalize Qprime columns
1831 FusedCta, ///< CTA-partitioned root solve using tunable threads per root
1832 };
1833
1834 /** @brief Which divide-and-conquer driver runs the merge tree. */
1835 enum class StedcAlgorithm {
1836 Auto = -1, ///< Currently: Levels
1837 Levels, ///< Level-synchronous: every node at a tree level merges in one launch
1838 Recursive, ///< Depth-first: one node per launch (kept for A/B comparison)
1839 };
1840
1841 /**
1842 * @brief Parameters for stedc(). Every tuning knob at 0 (or `Auto`) takes the value
1843 * from the checked-in tuning tables.
1844 * @see @ref perf_stedc, section "stedc: current tuning values"
1845 */
1846 template <typename T>
1848 int64_t recursion_threshold = 0; ///< Leaf size; <= 0 uses BatchLAS tuning, otherwise this exact threshold
1851 SteqrParams<T> leaf_steqr_params = SteqrParams<T>(); ///< Parameters of the leaf steqr() solves.
1852
1854 int merge_threads = 128; ///< work-group size for fused kernel
1855 int max_sec_iter = 50; ///< iteration cap for secular root solver
1856 bool enable_rescale = true; ///< keep ROCm-style v rescale (disable for perf experiments)
1857 int secular_threads_per_root = 0; ///< <= 0 uses BatchLAS tuning; otherwise this exact partition width
1858 int secular_cta_wg_size_multiplier = 0; ///< <= 0 uses BatchLAS tuning; otherwise this exact multiplier
1859 };
1860
1861 /**
1862 * @brief Eigenvalues and eigenvectors of a batch of symmetric tridiagonal matrices by
1863 * divide and conquer (LAPACK `?stedc`).
1864 *
1865 * Splits each matrix recursively down to leaves of `recursion_threshold`, solves the
1866 * leaves with steqr(), and merges pairs by solving the secular equation and forming
1867 * the merged eigenvectors. Eigenvalues are returned ascending.
1868 *
1869 * @param ctx queue the work is enqueued on
1870 * @param d diagonal, n entries per batch item
1871 * @param e off-diagonal, n-1 entries per batch item
1872 * @param eigenvalues output, n per batch item
1873 * @param ws at least stedc_buffer_size() bytes
1874 * @param jobz whether `eigvects` receives the eigenvectors
1875 * @param params leaf size, merge driver and tuning overrides
1876 * @param eigvects n x n eigenvector output per item
1877 * @param info per-item convergence status, or empty; see @ref md_docs_2cpp-api.
1878 * Raised by a merge whose secular solve hits its iteration cap and
1879 * by a leaf steqr() that runs out of sweeps.
1880 * @return event of the last enqueued kernel
1881 * @see @ref perf_stedc
1882 */
1883 template <Backend B, typename T>
1884 BATCHLAS_API Event stedc(Queue& ctx, const VectorView<T>& d, const VectorView<T>& e, const VectorView<T>& eigenvalues, const Span<std::byte>& ws,
1885 JobType jobz, StedcParams<T> params, const MatrixView<T, MatrixFormat::Dense>& eigvects,
1886 Span<int32_t> info = Span<int32_t>());
1887
1888 /** @brief Required workspace, in bytes, for stedc(). @ingroup api_tridiag_lowlevel */
1889 template <Backend B, typename T>
1890 BATCHLAS_API size_t stedc_buffer_size(Queue& ctx, size_t n, size_t batch_size, JobType jobz, StedcParams<T> params);
1891
1892 /// @}
1893
1894 /// @addtogroup api_eigen
1895 /// @{
1896
1897 /**
1898 * @brief Computes the Ritz values (Rayleigh quotients) of a set of trial vectors.
1899 *
1900 * \f$ \theta_j = \dfrac{v_j^H A v_j}{v_j^H v_j} \f$ for each column \f$ v_j \f$ of `V`,
1901 * per batch item.
1902 *
1903 * @param ctx Execution context/device queue
1904 * @param A Matrix (can be sparse or dense)
1905 * @param V Trial vectors (dense matrix, columns are trial eigenvectors)
1906 * @param ritz_vals Output vector for Ritz values, `V.cols()` real entries per item
1907 * @param workspace Pre-allocated workspace buffer, at least ritz_values_buffer_size() bytes
1908 * @return Event Event to track operation completion
1909 * @see @ref guide_ritz_values
1910 */
1911 template <Backend B, typename T, MatrixFormat MFormat>
1912 BATCHLAS_API Event ritz_values(Queue& ctx,
1913 const MatrixView<T, MFormat>& A,
1915 const VectorView<typename base_type<T>::type>& ritz_vals,
1916 Span<std::byte> workspace);
1917
1918 /**
1919 * @brief ritz_values() returning a new `Vector` of `V.cols()` Ritz values per item.
1920 *
1921 * Allocates its own workspace and waits on `ctx` before returning, so the result is
1922 * ready to read.
1923 * @see @ref guide_ritz_values
1924 */
1925 template <Backend B, typename T, MatrixFormat MFormat>
1926 inline auto ritz_values(Queue& ctx,
1927 const MatrixView<T, MFormat>& A,
1929 using float_type = typename base_type<T>::type;
1930 size_t nRitz = V.cols();
1931 Vector<float_type> ritz_vals(nRitz, V.batch_size());
1932 size_t workspace_size = ritz_values_buffer_size<B,T,MFormat>(ctx, A, V, static_cast<VectorView<float_type>>(ritz_vals));
1933 UnifiedVector<std::byte> workspace(workspace_size);
1934 ctx.wait();
1935 // (void) on an Event: deliberate. This Queue is in-order, so the next submission
1936 // is already ordered after this one and the Event carries nothing the caller needs.
1937 (void)ritz_values<B,T,MFormat>(ctx, A, V, static_cast<VectorView<float_type>>(ritz_vals), workspace);
1938 ctx.wait();
1939 return ritz_vals;
1940 }
1941
1942
1943
1944 /** @brief Required workspace size, in bytes, for ritz_values(). @ingroup api_eigen_lowlevel */
1945 template <Backend B, typename T, MatrixFormat MFormat>
1946 BATCHLAS_API size_t ritz_values_buffer_size(Queue& ctx,
1947 const MatrixView<T, MFormat>& A,
1949 const VectorView<typename base_type<T>::type>& ritz_vals);
1950
1951 /// @}
1952
1953 /// @addtogroup api_extra
1954 /// @{
1955
1956 /**
1957 * @brief Computes the explicit inverse of each dense square matrix in the batch,
1958 * \f$ A_{inv} := A^{-1} \f$, by LU factorisation (getrf()) and getri().
1959 *
1960 * `A` is copied into the workspace first and is not modified. Singular items are not
1961 * reported: check the result where the input may be ill-conditioned.
1962 *
1963 * @param ctx Execution context/device queue
1964 * @param A Input matrices to invert, n x n per item; read only
1965 * @param Ainv Output matrices storing the inverses, n x n per item
1966 * @param workspace Pre-allocated workspace buffer, at least inv_buffer_size() bytes
1967 * @return Event Event to track operation completion
1968 */
1969 template <Backend B, typename T>
1970 BATCHLAS_API Event inv(Queue& ctx,
1973 Span<std::byte> workspace);
1974
1975 /** @brief Required workspace size, in bytes, for inv(). @ingroup api_extra_lowlevel */
1976 template <Backend B, typename T>
1977 BATCHLAS_API size_t inv_buffer_size(Queue& ctx,
1979
1980 /**
1981 * @brief inv() returning a newly allocated result; workspace is leased from the
1982 * Queue's arena. Returns without waiting.
1983 */
1984 template <Backend B, typename T>
1987
1988 /**
1989 * @brief Scales `mat` by `cto / cfrom` without overflow or underflow (LAPACK `?lascl`).
1990 *
1991 * The factor is applied in safe steps when forming `cto / cfrom` directly would
1992 * over- or underflow. Instantiated for real `float` and `double`, dense format.
1993 * @return event of the last scaling step
1994 */
1995 template <MatrixFormat MType, typename T>
1996 BATCHLAS_API Event lascl(Queue& ctx, const MatrixView<T, MType>& mat, T cfrom, T cto);
1997
1998 /// @}
1999}
2000
2001namespace batchlas {
2002
2003// Backend-deducing overloads (BATCHLAS_DISPATCH_ON_QUEUE); see blas/queue-dispatch.hh.
2004// evidence: docs/cpp-api.md#which-type-each-parameter-takes
2005
2006
2071
2074
2075// ---- ortho: option-struct and arena-backed spellings -----------------------
2076//
2077// Same layer as blas/options.hh, but it has to live HERE: blas/linalg.hh includes
2078// blas/functions.hh (which ends by including options.hh) BEFORE this header, so `ortho` is
2079// not yet declared when options.hh is parsed, and none of its helpers are reachable here.
2080
2081/// @addtogroup api_qr
2082/// @{
2083
2084/** @brief Options for the option-struct ortho(); fields as the positional arguments. */
2086 Transpose transA = Transpose::NoTrans; ///< Columns (`NoTrans`) or rows.
2087 OrthoAlgorithm algorithm = OrthoAlgorithm::Chol2; ///< Orthogonalisation algorithm.
2088};
2089
2090/** @brief Options for the option-struct ortho() against an external basis `M`. */
2092 Transpose transA = Transpose::NoTrans; ///< Vectors of `A` are its columns (`NoTrans`) or rows.
2093 Transpose transM = Transpose::NoTrans; ///< Vectors of `M` are its columns (`NoTrans`) or rows.
2094 OrthoAlgorithm algorithm = OrthoAlgorithm::Chol2; ///< Orthogonalisation algorithm.
2095 size_t iterations = 2; ///< Project-then-orthonormalise passes.
2096};
2097
2098// ---- in-place orthogonalisation --------------------------------------------
2099
2100/** @brief ortho() with an options struct and a caller-provided workspace. */
2101template <Backend B, typename T>
2102inline Event ortho(Queue& ctx,
2104 const OrthoOptions& opts,
2105 Span<std::byte> workspace) {
2106 return ortho<B, T>(ctx, A, opts.transA, workspace, opts.algorithm);
2107}
2108
2109/**
2110 * @brief ortho() with an options struct; workspace is leased from the Queue's arena.
2111 *
2112 * The lease is released when the call returns, which on an out-of-order Queue drains
2113 * the device first: this spelling is synchronous there, the span-taking one is not.
2114 */
2115template <Backend B, typename T>
2116inline Event ortho(Queue& ctx,
2118 const OrthoOptions& opts) {
2119 // Sized with the algorithm the call actually runs: Chol2, SVQB and the Householder
2120 // path need different scratch, so a default-algorithm query under-sizes the rest.
2121 auto lease = ctx.workspace(ortho_buffer_size<B, T>(ctx, A, opts.transA, opts.algorithm));
2122 return ortho<B, T>(ctx, A, opts.transA, lease.span(), opts.algorithm);
2123}
2124
2125template <typename T>
2126inline Event ortho(Queue& ctx,
2128 const OrthoOptions& opts,
2129 Span<std::byte> workspace) {
2131 detail::require_arg_accessible(ctx, A, "ortho: A");
2132 detail::require_arg_accessible(ctx, workspace, "ortho: workspace");
2133 }
2134 return with_backend_for<T>(ctx, "ortho",
2135 [&](auto Back) { return ortho<Back.value, T>(ctx, A, opts, workspace); });
2136}
2137
2138template <typename T>
2139inline Event ortho(Queue& ctx,
2141 const OrthoOptions& opts = {}) {
2143 return with_backend_for<T>(ctx, "ortho",
2144 [&](auto Back) { return ortho<Back.value, T>(ctx, A, opts); });
2145}
2146
2147/// @}
2148
2149namespace detail {
2150// The bare-`{}` trap, in the one place ortho has it: `ortho(ctx, A, {}, ws)` matches both
2151// the positional `(ctx, A, Transpose, ws, algo)` and the option `(ctx, A, OrthoOptions, ws)`
2152// overload, and the positional one would win silently, handing the caller `Transpose{}`. A
2153// third candidate at the same exact-match rank makes the bare-`{}` call ambiguous instead.
2154enum class OrthoEmptyBracesAreAmbiguous {};
2155} // namespace detail
2156
2157/// @addtogroup api_qr
2158/// @{
2159
2160template <Backend B, typename T>
2162 detail::OrthoEmptyBracesAreAmbiguous, Span<std::byte>) = delete;
2163
2164template <typename T>
2166 detail::OrthoEmptyBracesAreAmbiguous, Span<std::byte>) = delete;
2167
2168// ---- orthogonalisation against an external metric ---------------------------
2169// No `{}` guard is needed here: the positional metric form needs at least six arguments
2170// while the option forms take four and five, so no argument list can reach both.
2171
2172/** @brief ortho() against `M` with an options struct and a caller-provided workspace. */
2173template <Backend B, typename T>
2174inline Event ortho(Queue& ctx,
2177 const OrthoAgainstOptions& opts,
2178 Span<std::byte> workspace) {
2179 return ortho<B, T>(ctx, A, M, opts.transA, opts.transM, workspace, opts.algorithm,
2180 opts.iterations);
2181}
2182
2183/** @brief ortho() against `M` with an options struct; workspace leased from the arena. */
2184template <Backend B, typename T>
2185inline Event ortho(Queue& ctx,
2188 const OrthoAgainstOptions& opts) {
2189 auto lease = ctx.workspace(ortho_buffer_size<B, T>(ctx, A, M, opts.transA, opts.transM,
2190 opts.algorithm, opts.iterations));
2191 return ortho<B, T>(ctx, A, M, opts.transA, opts.transM, lease.span(), opts.algorithm,
2192 opts.iterations);
2193}
2194
2195template <typename T>
2196inline Event ortho(Queue& ctx,
2199 const OrthoAgainstOptions& opts,
2200 Span<std::byte> workspace) {
2202 detail::require_arg_accessible(ctx, A, "ortho: A");
2203 detail::require_arg_accessible(ctx, M, "ortho: M");
2204 detail::require_arg_accessible(ctx, workspace, "ortho: workspace");
2205 }
2206 return with_backend_for<T>(ctx, "ortho",
2207 [&](auto Back) { return ortho<Back.value, T>(ctx, A, M, opts, workspace); });
2208}
2209
2210template <typename T>
2211inline Event ortho(Queue& ctx,
2214 const OrthoAgainstOptions& opts = {}) {
2216 detail::require_arg_accessible(ctx, A, "ortho: A");
2217 detail::require_arg_accessible(ctx, M, "ortho: M");
2218 }
2219 return with_backend_for<T>(ctx, "ortho",
2220 [&](auto Back) { return ortho<Back.value, T>(ctx, A, M, opts); });
2221}
2222
2223/// @}
2224
2225} // namespace batchlas
Non-owning view of a batch of matrices: the argument type of every entry point.
Definition matrix.hh:485
int cols() const
Columns per item (the capacity on a heterogeneous batch).
Definition matrix.hh:611
int batch_size() const
Number of batch items.
Definition matrix.hh:603
Owning batch of batch_size() matrices of equal (capacity) shape.
Definition matrix.hh:1399
Non-owning view of a batch of strided vectors: the vector argument type of the entry points.
Definition matrix.hh:1754
Enumerations, scalar traits and scalar concepts shared by every public header.
auto with_backend_for(Queue &ctx, const char *op, F &&f)
with_backend() for a call on element type T: throws batchlas::unsupported before dispatch when T is d...
size_t syev_cta_fused_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, JobType jobz, SteqrParams< T > steqr_params=SteqrParams< T >())
Required workspace, in bytes, for syev_cta_fused().
size_t syevx_lobpcg_buffer_size(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, size_t neigs, JobType jobz, const MatrixView< T, MatrixFormat::Dense > &V, const SyevxParams< T > &params)
Required workspace, in bytes, for syevx_lobpcg().
size_t syevx_buffer_size(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, size_t neigs, JobType jobz=JobType::NoEigenVectors, const MatrixView< T, MatrixFormat::Dense > &V=MatrixView< T, MatrixFormat::Dense >(), const SyevxParams< T > &params=SyevxParams< T >())
Required workspace, in bytes, for syevx(); arguments as for the call minus the workspace.
Event syev_two_stage(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, Span< typename base_type< T >::type > eigenvalues, JobType jobz, Uplo uplo, const Span< std::byte > &ws, StedcParams< typename base_type< T >::type > stedc_params=StedcParams< typename base_type< T >::type >(), Span< int32_t > info=Span< int32_t >())
Two-stage symmetric/Hermitian eigen-solver (SYEV-like) for large n: sytrd_sy2sb -> sytrd_sb2st -> ste...
Event syevx_direct_subset(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, Span< int32_t > m, size_t neigs, Span< std::byte > workspace, JobType jobz, const MatrixView< T, MatrixFormat::Dense > &V, const SyevxParams< T > &params, Span< int32_t > info=Span< int32_t >())
Partial eigensolve by two-stage reduction plus a subset tridiagonal solve (stebz + stein) and a narro...
Event syevx_direct(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, Span< int32_t > m, size_t neigs, Span< std::byte > workspace, JobType jobz, const MatrixView< T, MatrixFormat::Dense > &V, const SyevxParams< T > &params, Span< int32_t > info=Span< int32_t >())
Partial eigensolve by full decomposition followed by selection.
size_t syev_two_stage_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, JobType jobz, Uplo uplo, StedcParams< typename base_type< T >::type > stedc_params=StedcParams< typename base_type< T >::type >())
Required workspace, in bytes, for syev_two_stage().
size_t lanczos_buffer_size(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, JobType jobz=JobType::NoEigenVectors, const MatrixView< T, MatrixFormat::Dense > &V=MatrixView< T, MatrixFormat::Dense >(), const LanczosParams< T > &params=LanczosParams< T >())
Required workspace, in bytes, for lanczos(); arguments as for the call itself.
Event syevx_lobpcg(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, size_t neigs, Span< std::byte > workspace, JobType jobz, const MatrixView< T, MatrixFormat::Dense > &V, const SyevxParams< T > &params, Span< int32_t > info=Span< int32_t >())
Partial eigensolve by LOBPCG (locally optimal block preconditioned conjugate gradient).
size_t syev_blocked_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, JobType jobz, Uplo uplo, StedcParams< typename base_type< T >::type > stedc_params=StedcParams< typename base_type< T >::type >())
Required workspace, in bytes, for syev_blocked(); the same for either uplo.
size_t syevx_filtered_buffer_size(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, size_t neigs, JobType jobz, const MatrixView< T, MatrixFormat::Dense > &V, const SyevxParams< T > &params)
Required workspace, in bytes, for syevx_filtered().
size_t ritz_values_buffer_size(Queue &ctx, const MatrixView< T, MFormat > &A, const MatrixView< T, MatrixFormat::Dense > &V, const VectorView< typename base_type< T >::type > &ritz_vals)
Required workspace size, in bytes, for ritz_values().
Event syev_cta_fused(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, Span< typename base_type< T >::type > eigenvalues, JobType jobz, Uplo uplo, const Span< std::byte > &ws=Span< std::byte >(), SteqrParams< T > steqr_params=SteqrParams< T >(), size_t cta_wg_size_multiplier=0, Span< int32_t > info=Span< int32_t >())
Fused (single-kernel) variant of syev_cta.
size_t syev_cta_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, JobType jobz, SteqrParams< T > steqr_params=SteqrParams< T >())
Required workspace, in bytes, for syev_cta().
Event syev_blocked(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, Span< typename base_type< T >::type > eigenvalues, JobType jobz, Uplo uplo, const Span< std::byte > &ws, StedcParams< typename base_type< T >::type > stedc_params=StedcParams< typename base_type< T >::type >(), Span< int32_t > info=Span< int32_t >())
Blocked symmetric/Hermitian eigen-solver (SYEV-like) for medium/large n: sytrd_blocked -> stedc -> or...
SyevxResolvedRange syevx_resolve_range(int64_t n, size_t neigs, SyevxSelect select, bool find_largest, int64_t il, int64_t iu, SortOrder order)
Normalizes a range request into a SyevxResolvedRange.
size_t syev_jacobi_cta_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, JobType jobz, JacobiParams< T > params=JacobiParams< T >())
Required workspace, in bytes, for syev_jacobi_cta().
SyevxAlgorithm syevx_select_algorithm(MatrixFormat format, int64_t n, size_t neigs, SyevxAlgorithm requested, bool subset_supported, JobType jobz=JobType::EigenVectors, int64_t batch_size=1, SyevxSelect select=SyevxSelect::Extremal)
Resolves SyevxParams::method (and the BATCHLAS_SYEVX_ALGORITHM override) to a concrete,...
Event syevx_filtered(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, size_t neigs, Span< std::byte > workspace, JobType jobz, const MatrixView< T, MatrixFormat::Dense > &V, const SyevxParams< T > &params, Span< int32_t > info=Span< int32_t >())
Chebyshev-filtered subspace iteration: amplifies the wanted end of the spectrum, then Rayleigh-Ritz.
size_t syevx_direct_subset_buffer_size(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, size_t neigs, JobType jobz, const MatrixView< T, MatrixFormat::Dense > &V, const SyevxParams< T > &params)
Required workspace, in bytes, for syevx_direct_subset().
Event syev_jacobi_cta(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, Span< typename base_type< T >::type > eigenvalues, JobType jobz, Uplo uplo, const Span< std::byte > &ws=Span< std::byte >(), JacobiParams< T > params=JacobiParams< T >(), Span< int32_t > info=Span< int32_t >())
CTA-optimized Jacobi symmetric/Hermitian eigen-solver, n <= 32; real symmetric and complex Hermitian.
constexpr bool syevx_direct_subset_supported()
Whether syevx_direct_subset supports this scalar type and format.
Event syev_cta(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, Span< typename base_type< T >::type > eigenvalues, JobType jobz, Uplo uplo, const Span< std::byte > &ws, SteqrParams< T > steqr_params=SteqrParams< T >(), size_t cta_wg_size_multiplier=0, Span< int32_t > info=Span< int32_t >())
CTA-optimized symmetric eigen-solver (SYEV-like), n <= 32; real symmetric and complex Hermitian.
size_t syevx_direct_buffer_size(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, size_t neigs, JobType jobz, const MatrixView< T, MatrixFormat::Dense > &V, const SyevxParams< T > &params)
Required workspace, in bytes, for syevx_direct().
Event lanczos(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, Span< std::byte > workspace, JobType jobz=JobType::NoEigenVectors, const MatrixView< T, MatrixFormat::Dense > &V=MatrixView< T, MatrixFormat::Dense >(), const LanczosParams< T > &params=LanczosParams< T >())
All eigenvalues, and optionally eigenvectors, of a batch of symmetric/Hermitian matrices by a full-le...
Event ritz_values(Queue &ctx, const MatrixView< T, MFormat > &A, const MatrixView< T, MatrixFormat::Dense > &V, const VectorView< typename base_type< T >::type > &ritz_vals, Span< std::byte > workspace)
Computes the Ritz values (Rayleigh quotients) of a set of trial vectors.
Event syevx(Queue &ctx, const MatrixView< T, MFormat > &A, Span< typename base_type< T >::type > W, size_t neigs, Span< std::byte > workspace, JobType jobz=JobType::NoEigenVectors, const MatrixView< T, MatrixFormat::Dense > &V=MatrixView< T, MatrixFormat::Dense >(), const SyevxParams< T > &params=SyevxParams< T >(), Span< int32_t > info=Span< int32_t >())
Selected eigenpairs of a batch of symmetric/Hermitian matrices, dense or CSR.
SyevxPreconditioner
Preconditioner of the LOBPCG path of syevx; the other algorithms ignore it.
Definition enums.hh:449
@ Auto
ILU(k) if configured, else the environment default, else None.
Definition enums.hh:450
JobType
Whether an eigensolver computes eigenvectors (LAPACK jobz).
Definition enums.hh:189
Transpose
The operator op() applied to an operand: op(A) = A, A^T or A^H.
Definition enums.hh:171
SyevxSelect
Which part of the spectrum syevx returns.
Definition enums.hh:414
SortOrder
Order in which eigenvalues (and their vectors) are returned.
Definition enums.hh:335
Uplo
Which triangle of a symmetric, Hermitian or triangular operand is referenced.
Definition enums.hh:287
SvdVectors
Which singular vectors gesvd computes for one factor (LAPACK jobu / jobvt).
Definition enums.hh:220
OrthoAlgorithm
Algorithm used by ortho to orthonormalise the columns of a batch of blocks.
Definition enums.hh:471
MatrixFormat
Storage format of a Matrix / MatrixView, used as a template argument.
Definition enums.hh:72
SyevxAlgorithm
Algorithm family of the partial symmetric/Hermitian eigensolver syevx.
Definition enums.hh:383
Side
Side on which a special operand multiplies: op(A) * B or B * op(A).
Definition enums.hh:319
@ NoEigenVectors
Eigenvalues only (jobz = 'N').
Definition enums.hh:191
@ EigenVectors
Eigenvalues and eigenvectors (jobz = 'V').
Definition enums.hh:190
@ NoTrans
op(A) = A.
Definition enums.hh:172
@ Value
Every eigenvalue in the half-open interval (vl, vu].
Definition enums.hh:417
@ Index
SyevxParams::il .. iu inclusive, 0-based in the ascending spectrum.
Definition enums.hh:416
@ Extremal
neigs eigenpairs from one end; SyevxParams::find_largest picks the end.
Definition enums.hh:415
@ Ascending
Smallest first.
Definition enums.hh:336
@ All
Full square factor (LAPACK 'A').
Definition enums.hh:222
@ CGS2
Classical Gram-Schmidt with one reorthogonalisation pass.
Definition enums.hh:476
@ Chol2
CholeskyQR applied twice; the default.
Definition enums.hh:472
@ Dense
Column-major dense storage with leading dimension and batch stride.
Definition enums.hh:73
@ Auto
Heuristic selection on format, n, batch size and jobz (the default).
Definition enums.hh:384
size_t inv_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A)
Required workspace size, in bytes, for inv().
Event inv(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &Ainv, Span< std::byte > workspace)
Computes the explicit inverse of each dense square matrix in the batch, , by LU factorisation (getrf(...
Event lascl(Queue &ctx, const MatrixView< T, MType > &mat, T cfrom, T cto)
Scales mat by cto / cfrom without overflow or underflow (LAPACK ?lascl).
size_t ortho_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, Transpose transA, OrthoAlgorithm algo=OrthoAlgorithm::Chol2)
Required workspace, in bytes, for ortho(); arguments as for the call itself.
Event ormqx_cta(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, const VectorView< T > &tau_in, const MatrixView< T, MatrixFormat::Dense > &c_in, Uplo factorization, Side side, Transpose trans, int32_t k, const Span< std::byte > &ws, size_t cta_wg_size_multiplier=1)
CTA-optimized application of Q from a QR/QL factorization (ORMQx/UNMQx semantics) for very small matr...
Event ortho(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, Transpose transA, Span< std::byte > workspace, OrthoAlgorithm algo=OrthoAlgorithm::Chol2)
Orthonormalises the columns (or rows) of each batch item in place.
Event ormbr(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, const VectorView< T > &tau, const MatrixView< T, MatrixFormat::Dense > &c, char vect, Side side, Transpose trans, const Span< std::byte > &ws, int32_t block_size=32)
ORMBR/UNMBR-style application of bidiagonal reduction reflectors.
size_t bdsqr_buffer_size(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, Span< T > singular_values_out)
Required workspace, in bytes, for bdsqr().
size_t gesvd_blocked_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, Span< typename base_type< T >::type > singular_values, const MatrixView< T, MatrixFormat::Dense > &u_out, const MatrixView< T, MatrixFormat::Dense > &vh_out, SvdVectors jobu, SvdVectors jobvh)
Required workspace, in bytes, for gesvd_blocked().
Event gebrd_cta(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, const VectorView< typename base_type< T >::type > &d, const VectorView< typename base_type< T >::type > &e, const VectorView< T > &tauq, const VectorView< T > &taup, size_t cta_wg_size_multiplier=1)
CTA-parallel small-matrix GEBRD for real square matrices, 1 <= n <= 32, where one CTA cooperatively r...
size_t ormbr_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, const VectorView< T > &tau, const MatrixView< T, MatrixFormat::Dense > &c, char vect, Side side, Transpose trans, int32_t block_size=32)
Required workspace, in bytes, for ormbr(); arguments as for the call.
size_t bdsdc_buffer_size(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, Span< T > singular_values_out, bool want_vectors)
Required workspace, in bytes, for bdsdc(); want_vectors selects the overload.
size_t gesvdj_cta_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, Span< typename base_type< T >::type > singular_values, const MatrixView< T, MatrixFormat::Dense > &u_out, const MatrixView< T, MatrixFormat::Dense > &vh_out, SvdVectors jobu, SvdVectors jobvh, GesvdjParams< T > params=GesvdjParams< T >())
Required workspace, in bytes, for gesvdj_cta() (currently 0: all state is in local memory).
size_t gebrd_blocked_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, const VectorView< typename base_type< T >::type > &d, const VectorView< typename base_type< T >::type > &e, const VectorView< T > &tauq, const VectorView< T > &taup, int32_t block_size=16)
Required workspace, in bytes, for gebrd_blocked(); arguments as for the call.
Event gesvd_blocked(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, Span< typename base_type< T >::type > singular_values, const MatrixView< T, MatrixFormat::Dense > &u_out, const MatrixView< T, MatrixFormat::Dense > &vh_out, SvdVectors jobu, SvdVectors jobvh, const Span< std::byte > &ws, Span< int32_t > info=Span< int32_t >())
Blocked native SVD for real dense matrices: GEBRD-style dense -> bidiagonal, then BDSQR (or BDSDC whe...
size_t gesvd_cta_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, Span< typename base_type< T >::type > singular_values, const MatrixView< T, MatrixFormat::Dense > &u_out, const MatrixView< T, MatrixFormat::Dense > &vh_out, SvdVectors jobu, SvdVectors jobvh)
Required workspace, in bytes, for gesvd_cta().
Event bdsdc(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, Span< T > singular_values_out, const Span< std::byte > &ws, bool sort_desc=true, Span< int32_t > info=Span< int32_t >())
Bidiagonal divide-and-conquer SVD for a real upper bidiagonal matrix.
Event gebrd_blocked(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, const VectorView< typename base_type< T >::type > &d, const VectorView< typename base_type< T >::type > &e, const VectorView< T > &tauq, const VectorView< T > &taup, const Span< std::byte > &ws, int32_t block_size=16)
Blocked GEBRD for real square dense matrices: DLABRD-style panel + GEMM.
Event gebrd_unblocked(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, const VectorView< typename base_type< T >::type > &d, const VectorView< typename base_type< T >::type > &e, const VectorView< T > &tauq, const VectorView< T > &taup)
Unblocked GEBRD-like reduction to real bidiagonal form: reduces each square A in place,...
Event gesvdj_cta(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, Span< typename base_type< T >::type > singular_values, const MatrixView< T, MatrixFormat::Dense > &u_out, const MatrixView< T, MatrixFormat::Dense > &vh_out, SvdVectors jobu, SvdVectors jobvh, const Span< std::byte > &ws=Span< std::byte >(), GesvdjParams< T > params=GesvdjParams< T >(), Span< int32_t > info=Span< int32_t >())
One-sided (Hestenes) Jacobi SVD for batches of small matrices.
Event bdsqr(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, Span< T > singular_values_out, const Span< std::byte > &ws, bool sort_desc=true, Span< int32_t > info=Span< int32_t >())
Bidiagonal QR iteration for a real upper bidiagonal matrix.
Event gesvd_cta(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, Span< typename base_type< T >::type > singular_values, const MatrixView< T, MatrixFormat::Dense > &u_out, const MatrixView< T, MatrixFormat::Dense > &vh_out, SvdVectors jobu, SvdVectors jobvh, const Span< std::byte > &ws, Span< int32_t > info=Span< int32_t >())
CTA-oriented native SVD for small matrices, max(m, n) <= 32: bidiagonal reduction,...
Event sytrd_sy2sb(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, const MatrixView< T, MatrixFormat::Dense > &ab_out, const VectorView< T > &tau_out, Uplo uplo, int32_t kd, const Span< std::byte > &ws)
First stage of two-stage reduction: dense -> band (LAPACK xSYTRD_SY2SB).
size_t stedc_buffer_size(Queue &ctx, size_t n, size_t batch_size, JobType jobz, StedcParams< T > params)
Required workspace, in bytes, for stedc().
size_t sytrd_sy2sb_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, const MatrixView< T, MatrixFormat::Dense > &ab_out, const VectorView< T > &tau_out, Uplo uplo, int32_t kd)
Required workspace, in bytes, for sytrd_sy2sb(); arguments as for the call.
Event sytrd_band_reduction_single_step(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &ab_in, const MatrixView< T, MatrixFormat::Dense > &abw_out, Uplo uplo, int32_t kd, const Span< std::byte > &ws, SytrdBandReductionParams params)
Debug/testing hook: execute exactly one BANDR1 “chase step”.
size_t sytrd_band_reduction_single_step_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &ab_in, const MatrixView< T, MatrixFormat::Dense > &abw_out, Uplo uplo, int32_t kd, SytrdBandReductionParams params)
Required workspace, in bytes, for sytrd_band_reduction_single_step().
Event latrd_lower_panel(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, const VectorView< T > &e_out, const VectorView< T > &tau_out, const MatrixView< T, MatrixFormat::Dense > &w_in, int32_t j0, int32_t ib, int32_t wg_hint=0, bool fuse_trailing_update=false)
LATRD-like panel factorization used by blocked SYTRD, Uplo::Lower only.
size_t sytrd_sb2st_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &ab_in, const VectorView< typename base_type< T >::type > &d_out, const VectorView< typename base_type< T >::type > &e_out, const VectorView< T > &tau_out, Uplo uplo, int32_t kd, int32_t block_size)
Required workspace, in bytes, for sytrd_sb2st(); arguments as for the call.
size_t tridiagonal_solver_buffer_size(Queue &ctx, size_t n, size_t batch_size, JobType jobz)
Required workspace, in bytes, for tridiagonal_solver().
size_t stebz_buffer_size(Queue &ctx, size_t n, size_t batch_size, StebzParams< T > params)
Required workspace size, in bytes, for stebz().
size_t sytrd_band_reduction_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &ab_in, const VectorView< typename base_type< T >::type > &d_out, const VectorView< typename base_type< T >::type > &e_out, const VectorView< T > &tau_out, Uplo uplo, int32_t kd, int32_t block_size)
Required workspace, in bytes, for sytrd_band_reduction() with one block size.
size_t sytrd_blocked_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a, const VectorView< T > &d, const VectorView< T > &e, const VectorView< T > &tau, Uplo uplo, int32_t block_size=tuning::SYTRD_BLOCK_SIZE_MEDIUM)
Required workspace, in bytes, for sytrd_blocked(); arguments as for the call.
Event sytrd_sb2st(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &ab_in, const VectorView< typename base_type< T >::type > &d_out, const VectorView< typename base_type< T >::type > &e_out, const VectorView< T > &tau_out, Uplo uplo, int32_t kd, const Span< std::byte > &ws, int32_t block_size)
Second stage of two-stage reduction: band -> tridiagonal (LAPACK xSB2ST/xHB2ST bulge chasing,...
Event sytrd_cta(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, const VectorView< T > &d_out, const VectorView< T > &e_out, const VectorView< T > &tau_out, Uplo uplo, const Span< std::byte > &ws, size_t cta_wg_size_multiplier=1)
CTA-optimized symmetric tridiagonal reduction (SYTD2-style), n <= 32.
size_t stein_buffer_size(Queue &ctx, size_t n, size_t k, size_t batch_size, SteinParams< T > params)
Required workspace size, in bytes, for stein().
size_t steqr_cta_buffer_size(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, const VectorView< T > &eigenvalues, JobType jobz=JobType::NoEigenVectors, SteqrParams< T > params=SteqrParams< T >())
Required workspace, in bytes, for steqr_cta(); arguments as for the call.
Event steqr_cta(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, const VectorView< T > &eigenvalues, const Span< std::byte > &ws, JobType jobz=JobType::NoEigenVectors, SteqrParams< T > params=SteqrParams< T >(), const MatrixView< T, MatrixFormat::Dense > &eigvects=MatrixView< T, MatrixFormat::Dense >(), Span< int32_t > info=Span< int32_t >())
steqr() pinned to the CTA kernel: one sub-group partition per matrix, for small n (runtime-dispatched...
size_t steqr_buffer_size(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, const VectorView< T > &eigenvalues, JobType jobz=JobType::NoEigenVectors, SteqrParams< T > params=SteqrParams< T >())
Required workspace, in bytes, for steqr(); arguments as for the call.
Event sytrd_band_reduction(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &ab_in, const VectorView< typename base_type< T >::type > &d_out, const VectorView< typename base_type< T >::type > &e_out, const VectorView< T > &tau_out, Uplo uplo, int32_t kd, const Span< std::byte > &ws, int32_t block_size)
Symmetric/Hermitian band -> tridiagonal reduction (BANDR1-style).
EigenRangeType
How stebz() selects a subset of the spectrum.
Event stedc(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, const VectorView< T > &eigenvalues, const Span< std::byte > &ws, JobType jobz, StedcParams< T > params, const MatrixView< T, MatrixFormat::Dense > &eigvects, Span< int32_t > info=Span< int32_t >())
Eigenvalues and eigenvectors of a batch of symmetric tridiagonal matrices by divide and conquer (LAPA...
Event stebz(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, const VectorView< T > &w, Span< int32_t > m, const Span< std::byte > &ws, StebzParams< T > params=StebzParams< T >())
Computes selected eigenvalues of a batch of symmetric tridiagonal matrices by bisection on Sturm sequ...
Event steqr(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, const VectorView< T > &eigenvalues, const Span< std::byte > &ws, JobType jobz=JobType::NoEigenVectors, SteqrParams< T > params=SteqrParams< T >(), const MatrixView< T, MatrixFormat::Dense > &eigvects=MatrixView< T, MatrixFormat::Dense >(), Span< int32_t > info=Span< int32_t >())
Eigenvalues, and optionally eigenvectors, of a batch of symmetric tridiagonal matrices by implicit QR...
SteqrShiftStrategy
Shift used by the CTA STEQR's implicit QR/QL steps.
StedcMergeVariant
How the merge step (secular solve + eigenvector formation) is dispatched.
Event stein(Queue &ctx, const VectorView< T > &d, const VectorView< T > &e, const VectorView< T > &w, size_t k, const MatrixView< T, MatrixFormat::Dense > &Z, const Span< std::byte > &ws, SteinParams< T > params=SteinParams< T >())
Computes eigenvectors of a batch of symmetric tridiagonal matrices by inverse iteration,...
Event sytrd_blocked(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &a_in, const VectorView< T > &d_out, const VectorView< T > &e_out, const VectorView< T > &tau_out, Uplo uplo, const Span< std::byte > &ws, int32_t block_size=tuning::SYTRD_BLOCK_SIZE_MEDIUM)
Blocked symmetric/Hermitian tridiagonal reduction, n > 32: LATRD-style panel plus a BLAS-3 trailing u...
Event tridiagonal_solver(Queue &ctx, Span< T > alphas, Span< T > betas, Span< typename base_type< T >::type > W, Span< std::byte > workspace, JobType jobz, const MatrixView< T, MatrixFormat::Dense > &Q, size_t n, size_t batch_size)
Eigenvalues, and optionally eigenvectors, of a batch of small symmetric tridiagonal matrices by shift...
StedcSecularSolver
Secular-equation root finder used by stedc()'s merge step.
SteqrUpdateScheme
Update scheme of the CTA STEQR's implicit QR/QL steps.
StedcAlgorithm
Which divide-and-conquer driver runs the merge tree.
constexpr Span< const int32_t > stein_all_counts
Sentinel meaning: every batch item has all k eigenvalues valid.
@ All
Every eigenvalue.
@ Wilkinson
Wilkinson shift computed from the trailing 2x2 block.
@ Lapack
LAPACK-style implicit shift (the stable dsteqr formulation).
@ Auto
Use BatchLAS tuning tables for the current problem size.
@ FusedCta
CTA-partitioned root solve using tunable threads per root.
@ Baseline
Serial-per-root path (3 separate kernels).
@ Fused
One kernel: warp-parallel root solve + build/normalize Qprime columns.
@ Legacy
Original BatchLAS solver.
@ Rocm
rocSOLVER-style solver (default).
@ PG
Parlett-Gray style scalar recurrence.
@ EXP
Explicit similarity update mirroring steqr.cc's bulge-chasing math.
@ Auto
Currently: Levels.
@ Recursive
Depth-first: one node per launch (kept for A/B comparison).
@ Levels
Level-synchronous: every node at a tree level merges in one launch.
The batched containers: Matrix (owning), MatrixView (non-owning), KernelMatrixView (device-side),...
void require_arg_accessible(const Queue &ctx, const A &arg, const std::string &what)
constexpr int32_t SYTRD_BLOCK_SIZE_MEDIUM
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 for the one-sided Jacobi SVD (gesvdj_cta).
size_t max_sweeps
Cap on cyclic sweeps; convergence normally takes well under 10.
Span< int32_t > sweep_counts
Optional per-problem diagnostic, the analogue of cusolverDnXgesvdjGetSweeps: when non-empty (size >= ...
typename base_type< T >::type Real
size_t cta_wg_size_multiplier
Multiplies the baseline problems-per-work-group.
Real tol_multiplier
A rotation is applied to pivot pair (p,q) only when over the 2x2 Gram entries of the current columns...
Real zero_sigma_multiplier
means is not determined by A and is filled from the orthogonal complement.
Parameters of an incomplete LU factorization with level-of-fill k, ILU(k).
Definition iluk.hh:28
Owning ILU(k) factor of a batch of CSR matrices sharing one pattern.
Definition iluk.hh:77
Parameters for syev_jacobi_cta() (cyclic two-sided Jacobi).
typename base_type< T >::type Real
Real tol_multiplier
A rotation is applied to pivot pair (p,q) only when .
bool sort
Sort eigenpairs on output.
size_t cta_wg_size_multiplier
Multiplies the baseline work-group size, LCM(P, sub_group_size) for the compile-time partition width ...
SortOrder sort_order
Order used when sort is set.
size_t max_sweeps
Cap on cyclic sweeps; a safety net for pathological inputs.
Parameters for lanczos().
bool sort_enabled
Sort eigenvalues (and eigenvectors) on output.
typename base_type< T >::type float_type
size_t ortho_iterations
Orthogonalisation passes per vector.
OrthoAlgorithm ortho_algorithm
Orthogonalisation of each new Lanczos vector.
SortOrder sort_order
Order used when sort_enabled.
size_t reorthogonalization_iterations
Iterations between full reorthogonalisations.
Options for the option-struct ortho() against an external basis M.
Transpose transM
Vectors of M are its columns (NoTrans) or rows.
OrthoAlgorithm algorithm
Orthogonalisation algorithm.
size_t iterations
Project-then-orthonormalise passes.
Transpose transA
Vectors of A are its columns (NoTrans) or rows.
Options for the option-struct ortho(); fields as the positional arguments.
OrthoAlgorithm algorithm
Orthogonalisation algorithm.
Transpose transA
Columns (NoTrans) or rows.
Where BatchLAS work runs: a device, a backend, an ordered SYCL queue and a workspace arena.
void wait() const
Block until all work enqueued on this Queue has completed.
batchlas::WorkspaceLease workspace(size_t bytes)
Borrow bytes of device scratch from this Queue's arena.
Non-owning view of size() contiguous elements of type T.
Definition sycl-span.hh:39
Parameters for stebz() (bisection on a symmetric tridiagonal matrix).
int64_t iu
Last wanted index (0-based, inclusive), range == Index; < 0 means n-1.
T vl
Lower bound (exclusive), range == Value.
int32_t max_iterations
Safety cap on bisection steps per eigenvalue; the loop also exits on interval convergence.
int64_t il
First wanted index (0-based, inclusive), range == Index.
T abstol
Absolute tolerance on each eigenvalue.
T vu
Upper bound (inclusive), range == Value.
EigenRangeType range
Which eigenvalues to compute.
SortOrder order
Output order of w.
Parameters for stedc().
int secular_threads_per_root
<= 0 uses BatchLAS tuning; otherwise this exact partition width
int64_t recursion_threshold
Leaf size; <= 0 uses BatchLAS tuning, otherwise this exact threshold.
StedcMergeVariant merge_variant
Merge-step dispatch.
int merge_threads
work-group size for fused kernel
StedcSecularSolver secular_solver
Secular root finder.
StedcAlgorithm algorithm
Merge-tree driver.
int max_sec_iter
iteration cap for secular root solver
bool enable_rescale
keep ROCm-style v rescale (disable for perf experiments)
int secular_cta_wg_size_multiplier
<= 0 uses BatchLAS tuning; otherwise this exact multiplier
SteqrParams< T > leaf_steqr_params
Parameters of the leaf steqr() solves.
Parameters for stein() (inverse iteration on a symmetric tridiagonal).
uint32_t seed
Seed of the random starting vectors.
int32_t max_iterations
Fixed number of inverse-iteration steps; two or three suffice for eigenvalues accurate to working pre...
T ortho_threshold
Eigenvalues closer than ortho_threshold * ||T|| form one cluster and have their vectors explicitly re...
Parameters for steqr() and steqr_cta(), and for stedc()'s leaf solves.
size_t max_sweeps
Cap on Francis QR sweeps; 2-3 typically suffice per eigenvalue.
bool back_transform
If false, the eigenvector matrix is set to the identity and the rotations are applied to it; if true,...
bool transpose_working_vectors
Interleave the working copies of d and e across the batch.
SteqrShiftStrategy cta_shift_strategy
CTA STEQR only: shift strategy.
bool block_rotations
Apply the rotations in block_size blocks.
size_t block_size
Rotations are applied in blocks of this size; larger means excess FLOPs but better memory reuse.
size_t cta_wg_size_multiplier
CTA STEQR only: multiplies the baseline work-group size, LCM(N, sub_group_size).
SteqrUpdateScheme cta_update_scheme
CTA STEQR only: update scheme.
bool sort
Sort eigenpairs on output.
SortOrder sort_order
Order used when sort is set.
Optional per-iteration convergence histories recorded by the iterative syevx paths.
Definition extensions.hh:94
bool store_ritz_values
Record ritz_value_history.
bool store_convergence_rate
Record convergence_rate_history.
size_t max_iterations
Rows the history spans can hold.
size_t iteration_stride
Stride between iterations; 0 means batch_size * neigs.
int32_t * iterations_done
Optional per-batch-item iteration count.
Span< float_type > ritz_value_history
Ritz values per iteration (needs store_ritz_values).
Span< float_type > current_residual_history
Residual norm of the current iterate (needs store_current_residual).
Definition extensions.hh:98
size_t store_every
Record every k-th iteration.
Span< float_type > best_residual_history
Best residual norm seen so far, per eigenpair.
Definition extensions.hh:97
bool store_current_residual
Record current_residual_history.
typename base_type< T >::type float_type
Definition extensions.hh:95
size_t batch_stride
Stride between batch items; 0 means neigs.
Span< float_type > convergence_rate_history
Per-iteration residual reduction (needs store_convergence_rate).
Definition extensions.hh:99
Options for syevx(): algorithm family, convergence controls, preconditioning and which part of the sp...
Definition extensions.hh:41
const SyevxInstrumentation< T > * instrumentation
Optional convergence-history sink; nullptr records nothing.
Definition extensions.hh:66
const ILUKPreconditioner< T > * preconditioner
Caller-owned ILU(k) factor used as preconditioner, or nullptr.
Definition extensions.hh:54
size_t filter_degree
Chebyshev degree for SyevxAlgorithm::Filtered; 0 selects a default.
Definition extensions.hh:62
size_t extra_directions
Extra search directions carried beyond neigs.
Definition extensions.hh:47
OrthoAlgorithm algorithm
Orthogonalisation used inside the iterative solvers.
Definition extensions.hh:44
int64_t iu
select == Index: last wanted index, inclusive; iu < 0 means n-1.
Definition extensions.hh:73
size_t iterations
Cap on outer iterations of LOBPCG / Filtered.
Definition extensions.hh:46
SyevxSelect select
Range kind: Extremal (with find_largest), Index or Value.
Definition extensions.hh:68
float_type abstol
Absolute tolerance per eigenvalue; non-positive means .
Definition extensions.hh:81
bool find_largest
Extremal selection: largest (true) or smallest (false) eigenpairs.
Definition extensions.hh:48
SortOrder order
Output order for Index and Value; for Extremal it follows find_largest.
Definition extensions.hh:83
float_type vl
select == Value: lower end of the half-open interval (vl, vu], as in LAPACK.
Definition extensions.hh:76
typename base_type< T >::type float_type
Definition extensions.hh:42
float_type vu
select == Value: upper end (inclusive).
Definition extensions.hh:77
size_t ortho_iterations
Orthogonalisation passes per use.
Definition extensions.hh:45
ILUKParams< T > iluk_params
Parameters of the factor built under build_preconditioner.
Definition extensions.hh:61
int64_t il
select == Index: first wanted index, 0-based and inclusive, into the ASCENDING spectrum.
Definition extensions.hh:72
SyevxAlgorithm method
Algorithm family; Auto defers to syevx_select_algorithm().
Definition extensions.hh:43
bool build_preconditioner
Build the ILU(k) factor inside syevx from the caller's workspace.
Definition extensions.hh:60
T relative_tolerance
Relative residual tolerance of the iterative paths.
Definition extensions.hh:50
T absolute_tolerance
Absolute residual tolerance of the iterative paths.
Definition extensions.hh:49
int init_power_iterations
LOBPCG only: power-iteration steps on the random start block; -1 selects the default,...
Definition extensions.hh:65
SyevxPreconditioner preconditioner_type
Preconditioner family.
Definition extensions.hh:57
Resolved, algorithm-independent description of the caller's request.
bool value_range
true: (vl, vu]; false: the index block [il, iu]
bool reverse
write the selected block in descending order
int64_t iu
valid iff !value_range; 0-based, inclusive
int64_t max_count
Upper bound on eigenpairs per item, clamped to [0, n] so a consumer may index [il,...
int64_t il
valid iff !value_range; 0-based, inclusive
Sweep schedule for sytrd_band_reduction() (successive band reduction).
int32_t max_sweeps
Maximum number of sweeps; < 0 means the implementation default.
std::vector< int32_t > block_size_seq
Block size per sweep (Algorithm 2: nb^(i)).
int32_t kd_work
Working band semibandwidth; <= 0 means the implementation default.
std::vector< int32_t > d_seq
Diagonals eliminated per sweep (Algorithm 2: d^(i)); 0 means the default schedule.
int32_t max_steps
Debug/testing: chase steps for sytrd_band_reduction_single_step(); <= 0 means exactly one step.
Owning, growable array of T in USM shared memory (sycl::malloc_shared).
Owning batch of batch_size() vectors of length size() in USM shared memory.
Definition matrix.hh:1963
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.
Tuning constants, bucketed by problem order n, with runtime BATCHLAS_TUNE_* overrides.