2#include <batchlas/export.hh>
115 template <
typename T>
144 template <Backend B,
typename T>
170 template <Backend B,
typename T>
178 size_t iterations = 2);
181 template <Backend B,
typename T>
188 template <Backend B,
typename T>
195 size_t iterations = 2);
235 template <Backend B,
typename T, MatrixFormat MFormat>
247 template <Backend B,
typename T, MatrixFormat MFormat>
256 return syevx<B,T,MFormat>(ctx,
MatrixView<T, MFormat>(A), W, neigs, workspace, jobz,
MatrixView<T, MatrixFormat::Dense>(), params, info);
275 template <Backend B,
typename T, MatrixFormat MFormat>
288 template <Backend B,
typename T, MatrixFormat MFormat>
298 return syevx<B,T,MFormat>(ctx,
MatrixView<T, MFormat>(A), W, m, neigs, workspace, jobz,
MatrixView<T, MatrixFormat::Dense>(), params, info);
307 template <Backend B,
typename T, MatrixFormat MFormat>
317 template <Backend B,
typename T, MatrixFormat MFormat>
324 return syevx_buffer_size<B,T,MFormat>(ctx,
MatrixView<T, MFormat>(A), W, neigs, jobz,
MatrixView<T, MatrixFormat::Dense>(), params);
332 template <Backend B,
typename T, MatrixFormat MFormat>
346 template <Backend B,
typename T, MatrixFormat MFormat>
355 return syevx_buffer_size<B,T,MFormat>(ctx,
MatrixView<T, MFormat>(A), W, neigs, jobz,
MatrixView<T, MatrixFormat::Dense>(), params);
399 template <
typename T>
432 bool subset_supported,
434 int64_t batch_size = 1,
451 template <Backend B,
typename T, MatrixFormat MFormat>
468 template <Backend B,
typename T, MatrixFormat MFormat>
479 jobz, V, params, info);
487 template <Backend B,
typename T, MatrixFormat MFormat>
513 template <Backend B,
typename T, MatrixFormat MFormat>
526 template <Backend B,
typename T, MatrixFormat MFormat>
537 jobz, V, params, info);
547 template <Backend B,
typename T, MatrixFormat MFormat>
557 template <
typename T, MatrixFormat MFormat>
571 template <Backend B,
typename T, MatrixFormat MFormat>
583 template <Backend B,
typename T, MatrixFormat MFormat>
601 template <Backend B,
typename T, MatrixFormat MFormat>
613 template <Backend B,
typename T, MatrixFormat MFormat>
638 template <Backend B,
typename T, MatrixFormat MFormat>
648 template <Backend B,
typename T, MatrixFormat MFormat>
655 return lanczos<B,T,MFormat>(ctx,
MatrixView<T,MFormat>(A), W, workspace, jobz,
MatrixView<T, MatrixFormat::Dense>(), params);
659 template <Backend B,
typename T, MatrixFormat MFormat>
694 template <Backend B,
typename T>
706 template <Backend B,
typename T>
721 template <
typename T>
753 template <Backend B,
typename T>
769 template <Backend B,
typename T>
776 template <
typename T>
803 template <Backend B,
typename T>
835 template <Backend B,
typename T>
854 template <Backend B,
typename T>
879 template <
typename T>
936 template <Backend B,
typename T>
951 template <Backend B,
typename T>
960 template <
typename T>
965 template <
typename T>
983 template <Backend B,
typename T>
991 size_t cta_wg_size_multiplier = 1);
1005 template <Backend B,
typename T>
1013 int32_t wg_hint = 0,
1014 bool fuse_trailing_update =
false);
1027 template <Backend B,
typename T>
1033 int32_t wg_hint = 0,
1034 bool fuse_trailing_update =
false);
1051 template <Backend B,
typename T>
1062 template <Backend B,
typename T>
1081 template <Backend B,
typename T>
1091 template <Backend B,
typename T>
1107 template <Backend B,
typename T>
1116 int32_t block_size);
1119 template <Backend B,
typename T>
1127 int32_t block_size);
1156 template <Backend B,
typename T>
1165 int32_t block_size);
1168 template <Backend B,
typename T>
1179 template <Backend B,
typename T>
1189 template <Backend B,
typename T>
1197 int32_t block_size);
1200 template <Backend B,
typename T>
1210 template <Backend B,
typename T>
1224 template <Backend B,
typename T>
1234 template <Backend B,
typename T>
1260 template <Backend B,
typename T>
1268 size_t cta_wg_size_multiplier = 0,
1272 template <Backend B,
typename T>
1287 template <Backend B,
typename T>
1295 size_t cta_wg_size_multiplier = 0,
1299 template <Backend B,
typename T>
1306 template <
typename T>
1337 template <Backend B,
typename T>
1348 template <Backend B,
typename T>
1361 template <Backend B,
typename T>
1372 template <Backend B,
typename T>
1388 template <Backend B,
typename T>
1399 template <Backend B,
typename T>
1421 template <Backend B,
typename T>
1434 template <Backend B,
typename T>
1441 size_t cta_wg_size_multiplier = 1);
1448 template <Backend B,
typename T>
1456 int32_t block_size = 16);
1459 template <Backend B,
typename T>
1466 int32_t block_size = 16);
1480 template <Backend B,
typename T>
1486 bool sort_desc =
true,
1494 template <Backend B,
typename T>
1502 bool sort_desc =
true,
1506 template <
typename T>
1513 template <
typename T>
1520 static_cast<void>(u);
1521 static_cast<void>(vh);
1537 template <Backend B,
typename T>
1543 bool sort_desc =
true,
1547 template <Backend B,
typename T>
1555 bool sort_desc =
true,
1559 template <Backend B,
typename T>
1584 template <Backend B,
typename T>
1593 int32_t block_size = 32);
1596 template <Backend B,
typename T>
1604 int32_t block_size = 32);
1615 template <Backend B,
typename T>
1631 template <Backend B,
typename T>
1639 Uplo hermitian_uplo,
1644 template <Backend B,
typename T>
1654 template <Backend B,
typename T>
1662 Uplo hermitian_uplo);
1678 template <Backend B,
typename T>
1690 template <Backend B,
typename T>
1698 Uplo hermitian_uplo,
1703 template <Backend B,
typename T>
1713 template <Backend B,
typename T>
1721 Uplo hermitian_uplo);
1730 template <
typename T>
1771 template <Backend B,
typename T>
1784 template <Backend B,
typename T>
1805 template <Backend B,
typename T>
1815 size_t cta_wg_size_multiplier = 1);
1846 template <
typename T>
1883 template <Backend B,
typename T>
1889 template <Backend B,
typename T>
1911 template <Backend B,
typename T, MatrixFormat MFormat>
1925 template <Backend B,
typename T, MatrixFormat MFormat>
1930 size_t nRitz = V.
cols();
1945 template <Backend B,
typename T, MatrixFormat MFormat>
1969 template <Backend B,
typename T>
1976 template <Backend B,
typename T>
1984 template <Backend B,
typename T>
1995 template <MatrixFormat MType,
typename T>
2101template <Backend B,
typename T>
2115template <Backend B,
typename T>
2125template <
typename T>
2135 [&](
auto Back) {
return ortho<Back.value, T>(ctx, A, opts, workspace); });
2138template <
typename T>
2144 [&](
auto Back) {
return ortho<Back.value, T>(ctx, A, opts); });
2154enum class OrthoEmptyBracesAreAmbiguous {};
2160template <Backend B,
typename T>
2164template <
typename T>
2173template <Backend B,
typename T>
2184template <Backend B,
typename T>
2195template <
typename T>
2207 [&](
auto Back) {
return ortho<Back.value, T>(ctx, A, M, opts, workspace); });
2210template <
typename T>
2220 [&](
auto Back) {
return ortho<Back.value, T>(ctx, A, M, opts); });
Non-owning view of a batch of matrices: the argument type of every entry point.
int cols() const
Columns per item (the capacity on a heterogeneous batch).
int batch_size() const
Number of batch items.
Owning batch of batch_size() matrices of equal (capacity) shape.
Non-owning view of a batch of strided vectors: the vector argument type of the entry points.
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 > ¶ms)
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 > ¶ms=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 > ¶ms, 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 > ¶ms, 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 > ¶ms=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 > ¶ms, 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 > ¶ms)
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 > ¶ms, 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 > ¶ms)
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 > ¶ms)
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 > ¶ms=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 > ¶ms=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.
@ Auto
ILU(k) if configured, else the environment default, else None.
JobType
Whether an eigensolver computes eigenvectors (LAPACK jobz).
Transpose
The operator op() applied to an operand: op(A) = A, A^T or A^H.
SyevxSelect
Which part of the spectrum syevx returns.
SortOrder
Order in which eigenvalues (and their vectors) are returned.
Uplo
Which triangle of a symmetric, Hermitian or triangular operand is referenced.
SvdVectors
Which singular vectors gesvd computes for one factor (LAPACK jobu / jobvt).
OrthoAlgorithm
Algorithm used by ortho to orthonormalise the columns of a batch of blocks.
MatrixFormat
Storage format of a Matrix / MatrixView, used as a template argument.
SyevxAlgorithm
Algorithm family of the partial symmetric/Hermitian eigensolver syevx.
Side
Side on which a special operand multiplies: op(A) * B or B * op(A).
@ NoEigenVectors
Eigenvalues only (jobz = 'N').
@ EigenVectors
Eigenvalues and eigenvectors (jobz = 'V').
@ Value
Every eigenvalue in the half-open interval (vl, vu].
@ Index
SyevxParams::il .. iu inclusive, 0-based in the ascending spectrum.
@ Extremal
neigs eigenpairs from one end; SyevxParams::find_largest picks the end.
@ Ascending
Smallest first.
@ All
Full square factor (LAPACK 'A').
@ CGS2
Classical Gram-Schmidt with one reorthogonalisation pass.
@ Chol2
CholeskyQR applied twice; the default.
@ Dense
Column-major dense storage with leading dimension and batch stride.
@ Auto
Heuristic selection on format, n, batch size and jobz (the default).
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.
@ 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.
@ 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)
bool pointer_checks_enabled()
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).
Owning ILU(k) factor of a batch of CSR matrices sharing one pattern.
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.
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.
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.
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).
size_t store_every
Record every k-th iteration.
Span< float_type > best_residual_history
Best residual norm seen so far, per eigenpair.
bool store_current_residual
Record current_residual_history.
typename base_type< T >::type float_type
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).
Options for syevx(): algorithm family, convergence controls, preconditioning and which part of the sp...
const SyevxInstrumentation< T > * instrumentation
Optional convergence-history sink; nullptr records nothing.
const ILUKPreconditioner< T > * preconditioner
Caller-owned ILU(k) factor used as preconditioner, or nullptr.
size_t filter_degree
Chebyshev degree for SyevxAlgorithm::Filtered; 0 selects a default.
size_t extra_directions
Extra search directions carried beyond neigs.
OrthoAlgorithm algorithm
Orthogonalisation used inside the iterative solvers.
int64_t iu
select == Index: last wanted index, inclusive; iu < 0 means n-1.
size_t iterations
Cap on outer iterations of LOBPCG / Filtered.
SyevxSelect select
Range kind: Extremal (with find_largest), Index or Value.
float_type abstol
Absolute tolerance per eigenvalue; non-positive means .
bool find_largest
Extremal selection: largest (true) or smallest (false) eigenpairs.
SortOrder order
Output order for Index and Value; for Extremal it follows find_largest.
float_type vl
select == Value: lower end of the half-open interval (vl, vu], as in LAPACK.
typename base_type< T >::type float_type
float_type vu
select == Value: upper end (inclusive).
size_t ortho_iterations
Orthogonalisation passes per use.
ILUKParams< T > iluk_params
Parameters of the factor built under build_preconditioner.
int64_t il
select == Index: first wanted index, 0-based and inclusive, into the ASCENDING spectrum.
SyevxAlgorithm method
Algorithm family; Auto defers to syevx_select_algorithm().
bool build_preconditioner
Build the ILU(k) factor inside syevx from the caller's workspace.
T relative_tolerance
Relative residual tolerance of the iterative paths.
T absolute_tolerance
Absolute residual tolerance of the iterative paths.
int init_power_iterations
LOBPCG only: power-iteration steps on the random start block; -1 selects the default,...
SyevxPreconditioner preconditioner_type
Preconditioner family.
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.
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.