linalg-ops.hh
Go to the documentation of this file.
1#pragma once
2
3#include <batchlas/export.hh>
4#include <algorithm>
5#include <stdexcept>
6#include <utility>
7
9// Needed here: the qualified ids in the forwards below bind at definition context.
12// Called directly by svd(), solve() and solve_spd(); never rely on them transitively.
21
22/// @file
23/// @brief The batchlas::linalg convenience layer: elementwise operations and value-returning wrappers.
24/// @ingroup api_linalg
25
26/**
27 * @addtogroup api_linalg
28 * @details
29 * Free functions only, no operator overloads. Membership rule: value-returning,
30 * backend from the Queue, workspace from the Queue's arena => `linalg::`;
31 * out-parameter, workspace yours => `batchlas::`.
32 *
33 * Every function here **enqueues and returns without waiting**, including the
34 * value-returning ones: a returned matrix is readable only after `ctx.wait()`.
35 * The exceptions are norm(), cond(), svd(), solve(), solve_spd() and eigvalsh(),
36 * which wait because they free local scratch on return. Inputs are never
37 * modified, except by the `_into` forms (which write their last matrix argument)
38 * and scale(). A value-returning call allocates its result, so prefer the `_into`
39 * forms in an inner loop. User guide: @ref md_docs_2cpp-api (the `linalg` section);
40 * design notes: @ref design_api_conventions.
41 */
42
44
45/// @addtogroup api_linalg
46/// @{
47
48// ---- elementwise -----------------------------------------------------------
49
50/// @brief Elementwise operation selected by elementwise_into().
51// BATCHLAS_API is load-bearing: an enum used as a template argument must be
52// exported or every instantiation over it is hidden (see blas/enums.hh).
53// evidence: docs/design/symbol-visibility.md#symbol-visibility-binaryop-is-a-template-argument-enum-too
54enum class BATCHLAS_API BinaryOp { Add, Subtract, Multiply, Divide };
55
56// Signature aliases for explicit instantiation; see BATCHLAS_INSTANTIATE in
57// src/util/template-instantiations.hh. Keep in sync with the declarations below.
58namespace sig {
59template <typename T, BinaryOp Op>
60using elementwise_into = Event(Queue&,
64
65template <typename T>
66using axpby_into = Event(Queue&,
67 T,
69 T,
72
73template <typename T>
75
76template <typename T>
77using triangular_mask_into = Event(Queue&,
80 Uplo,
81 int64_t);
82} // namespace sig
83
84/// @brief Elementwise \f$C := A \circ B\f$ for \f$\circ\f$ = Op, per batch item.
85/// @tparam T scalar type
86/// @tparam Op Add, Subtract, Multiply (Hadamard) or Divide
87/// @param ctx queue the kernel is enqueued on
88/// @param A first operand
89/// @param B second operand, same shape and batch size as `A`
90/// @param C output of the same shape; may alias `A` or `B`
91/// @return event of the enqueued kernel
92template <typename T, BinaryOp Op>
93BATCHLAS_API Event elementwise_into(Queue& ctx,
97
98/// @brief \f$C := \alpha A + \beta B\f$, elementwise, per batch item.
99/// @param ctx queue the kernel is enqueued on
100/// @param alpha scale of `A`
101/// @param A first operand
102/// @param beta scale of `B`
103/// @param B second operand, same shape and batch size as `A`
104/// @param C output of the same shape; may alias `A` or `B`
105/// @return event of the enqueued kernel
106template <typename T>
107BATCHLAS_API Event axpby_into(Queue& ctx,
108 T alpha,
110 T beta,
113
114/// @brief In-place scaling \f$A := \alpha A\f$.
115/// @return event of the enqueued kernel
116template <typename T>
117BATCHLAS_API Event scale(Queue& ctx, const MatrixView<T, MatrixFormat::Dense>& A, T alpha);
118
119/// @brief \f$C :=\f$ A with everything outside the requested triangle zeroed.
120/// @param ctx queue the kernel is enqueued on
121/// @param A input, m x n per item
122/// @param C output of the same shape; may alias `A`
123/// @param uplo Upper keeps the part on and above diagonal `k`, Lower on and below
124/// @param k diagonal offset, as NumPy `triu`/`tril`: 0 is the main diagonal,
125/// > 0 toward the upper right, < 0 toward the lower left
126/// @return event of the enqueued kernel
127template <typename T>
131 Uplo uplo,
132 int64_t k = 0);
133
134/// @brief \f$C := A + B\f$ elementwise; see elementwise_into().
135template <typename T>
142
143/// @brief \f$C := A - B\f$ elementwise; see elementwise_into().
144template <typename T>
151
152/// @brief \f$C := A \odot B\f$ elementwise (Hadamard), not the matrix product; see matmul().
153template <typename T>
160
161/// @brief \f$C := A \oslash B\f$ elementwise; see elementwise_into().
162template <typename T>
169
170/// @}
171
172namespace detail {
173template <typename T>
177} // namespace detail
178
179/// @addtogroup api_linalg
180/// @{
181
182/// @brief Returns a new matrix \f$A + B\f$ (elementwise).
183template <typename T>
187 auto C = detail::like(A);
188 // (void) on an Event is deliberate: the in-order Queue already orders the next call.
189 (void)add_into<T>(ctx, A, B, C);
190 return C;
191}
192
193/// @brief Returns a new matrix \f$A - B\f$ (elementwise).
194template <typename T>
198 auto C = detail::like(A);
199 (void)subtract_into<T>(ctx, A, B, C);
200 return C;
201}
202
203/// @brief Returns a new matrix \f$A \odot B\f$ (Hadamard, not matmul()).
204template <typename T>
208 auto C = detail::like(A);
209 (void)multiply_into<T>(ctx, A, B, C);
210 return C;
211}
212
213/// @brief Returns a new matrix \f$A \oslash B\f$ (elementwise).
214template <typename T>
218 auto C = detail::like(A);
219 (void)divide_into<T>(ctx, A, B, C);
220 return C;
221}
222
223/// @brief Returns a new matrix \f$\alpha A\f$; `A` is not modified (scale() works in place).
224template <typename T>
227 T alpha) {
228 auto C = detail::like(A);
229 (void)axpby_into<T>(ctx, alpha, A, T(0), A, C);
230 return C;
231}
232
233// ---- value-returning wrappers ---------------------------------------------
234
235/// @brief Options for matmul(). Deliberately has no `beta`: C is freshly allocated.
236template <typename T>
238 T alpha = T(1); ///< scale of the product
242};
243
244/// @brief Returns a new matrix \f$C = \alpha\,\mathrm{op}(A)\,\mathrm{op}(B)\f$.
245/// @param ctx queue; supplies the backend
246/// @param A `op(A)` is m x k per item
247/// @param B `op(B)` is k x n per item, same batch size
248/// @param opts alpha, transA, transB, precision
249/// @return m x n x batch matrix, fully written once the queue is waited on
250template <typename T>
254 const MatmulOptions<T>& opts = {}) {
255 const bool ta = opts.transA != Transpose::NoTrans;
256 const bool tb = opts.transB != Transpose::NoTrans;
257 const auto m = ta ? A.cols() : A.rows();
258 const auto n = tb ? B.rows() : B.cols();
260 (void)gemm(ctx, A, B, C,
261 GemmOptions<T>{.alpha = opts.alpha,
262 .beta = T(0), // C is fresh
263 .transA = opts.transA,
264 .transB = opts.transB,
265 .precision = opts.precision});
266 return C;
267}
268
269/// @brief Default triangle of the linalg wrappers that read one.
270inline constexpr Uplo kDefaultUplo = Uplo::Lower;
271
272/// @brief Returns the Cholesky factor of a Hermitian positive-definite A; A is not modified.
273/// @param ctx queue; supplies the backend and the arena
274/// @param A n x n per item; the `uplo` triangle is read
275/// @param uplo Lower gives \f$L\f$ with \f$A = LL^H\f$, Upper gives \f$U\f$ with \f$A = U^HU\f$
276/// @return new matrix holding the factor in the `uplo` triangle; the other
277/// triangle holds a copy of A's
278/// @note No status is reported; a non-positive-definite item yields an undefined factor.
279template <typename T>
282 Uplo uplo = kDefaultUplo) {
283 auto L = detail::like(A);
285 (void)potrf(ctx, L, {.uplo = uplo});
286 return L;
287}
288
289/// @brief Returns the eigenvalues of a symmetric/Hermitian A, ascending per item; A is not modified.
290/// @param ctx queue; supplies the backend and the arena
291/// @param A n x n per item; the `uplo` triangle is read
292/// @param uplo triangle of A that is read
293/// @return real vector of `n * batch_size` values, item b at offset `b * n`
294template <typename T>
297 Uplo uplo = kDefaultUplo) {
298 UnifiedVector<typename base_type<T>::type> W(static_cast<size_t>(A.rows()) *
299 static_cast<size_t>(A.batch_size()));
300 auto work = detail::like(A);
302 (void)syev(ctx, work, W.to_span(), {.jobz = JobType::NoEigenVectors, .uplo = uplo});
303 ctx.wait(); // ~Matrix frees `work` without waiting, as in svd()
304 return W;
305}
306
307/// @brief Result of eigh().
308template <typename T>
309struct Eigh {
310 UnifiedVector<typename base_type<T>::type> values; ///< ascending, `n` per batch item
311 Matrix<T, MatrixFormat::Dense> vectors; ///< eigenvectors as columns, n x n per item
312 /// Per-item convergence status: 0 converged, > 0 LAPACK-like. Always filled;
313 /// check it, since a non-converged item otherwise looks like a converged one.
315};
316
317/// @brief Eigenvalues and eigenvectors of a symmetric/Hermitian A; A is not modified.
318/// @param ctx queue; supplies the backend and the arena
319/// @param A n x n per item; the `uplo` triangle is read
320/// @param uplo triangle of A that is read
321/// @return values, vectors and per-item info, valid after the queue is waited on
322/// @throws batchlas::invalid_argument if `A` is not square
323template <typename T>
324inline Eigh<T> eigh(Queue& ctx,
326 Uplo uplo = kDefaultUplo) {
327 UnifiedVector<typename base_type<T>::type> W(static_cast<size_t>(A.rows()) *
328 static_cast<size_t>(A.batch_size()));
329 UnifiedVector<int32_t> info(static_cast<size_t>(A.batch_size()));
330 // The positional syev below does not check squareness. Qualified: `detail::`
331 // here would find batchlas::linalg::detail.
333 auto V = detail::like(A);
335 // Positional because SyevOptions has no `info` field.
336 auto lease = ctx.workspace(
338 (void)syev(ctx, V, W.to_span(), JobType::EigenVectors, uplo, lease.span(),
339 info.to_span());
340 // No wait needed: `info` is moved into the result (unlike svd()'s local scratch).
341 return Eigh<T>{std::move(W), std::move(V), std::move(info)};
342}
343
344/// @brief Solves \f$\mathrm{op}(A_i)X_i = B_i\f$ by LU with partial pivoting; returns X.
345///
346/// Neither A nor B is modified; pivots and workspace come from the arena. No
347/// status is reported: a singular or near-singular item yields a plausible-looking
348/// but wrong X, so check the residual where conditioning is not known.
349/// @param ctx queue; supplies the backend and the arena
350/// @param A n x n per item
351/// @param B n x nrhs per item, same batch size
352/// @param trans op(A)
353/// @return new n x nrhs x batch matrix
354/// @throws batchlas::internal_error for NoTrans when an extent (n, nrhs or batch)
355/// is below 1: `gesv` has no kernel for an empty problem
356template <typename T>
361 auto LU = detail::like(A);
363
366
367 const size_t n_pivots = static_cast<size_t>(A.rows()) * static_cast<size_t>(A.batch_size());
368 auto pivot_bytes = ctx.workspace(n_pivots * sizeof(int64_t));
369 Span<int64_t> pivots(reinterpret_cast<int64_t*>(pivot_bytes.data()), n_pivots);
370
371 // Not stylistic, do not flatten: gesv has no Transpose parameter. No shape gate
372 // here either; gesv selects its own kernel. Degenerate extents throw (deliberate).
373 // evidence: docs/design/api-conventions.md#api-conventions-linalgsolve-keeps-its-transpose-branch
374 if (trans == Transpose::NoTrans) {
375 // Released before the pivot lease (reverse order): the arena reuses it at once.
376 auto lease = ctx.workspace(gesv_buffer_size(ctx, LU, X));
377 (void)gesv(ctx, LU, X, pivots, lease.span(), Span<int32_t>{});
378 ctx.wait(); // ~Matrix frees `LU` without waiting, as in svd()
379 return X;
380 }
381
382 (void)getrf(ctx, LU, pivots);
383 (void)getrs(ctx, LU, X, pivots, {.trans = trans});
384 ctx.wait();
385 return X;
386}
387
388/// @brief Solves \f$A_iX_i = B_i\f$ for Hermitian positive-definite A by Cholesky; returns X.
389///
390/// About half the arithmetic of solve(), which stays the safe default: the SPD
391/// claim is the caller's. A matrix that is not positive definite gives an
392/// undefined X for that item, as LAPACK `?POSV`, and nothing is reported.
393/// Neither A nor B is modified; the workspace comes from the arena.
394/// @param ctx queue; supplies the backend and the arena
395/// @param A n x n per item; the `uplo` triangle is read
396/// @param B n x nrhs per item
397/// @param uplo triangle of A that is read
398/// @return new n x nrhs x batch matrix
399/// @throws batchlas::internal_error when an extent (n, nrhs or batch) is below 1:
400/// `posv` has no kernel for an empty problem
401template <typename T>
405 Uplo uplo = Uplo::Lower) {
406 auto F = detail::like(A);
408
411
412 auto lease = ctx.workspace(posv_buffer_size(ctx, F, X, uplo));
413 (void)posv(ctx, F, X, uplo, lease.span(), Span<int32_t>{});
414 ctx.wait(); // ~Matrix frees `F` without waiting, as in svd()
415 return X;
416}
417
418/// @brief Returns a copy of A with everything below diagonal `k` zeroed (NumPy `triu`).
419template <typename T>
422 int64_t k = 0) {
423 auto C = detail::like(A);
424 (void)triangular_mask_into<T>(ctx, A, C, Uplo::Upper, k);
425 return C;
426}
427
428/// @brief Returns a copy of A with everything above diagonal `k` zeroed (NumPy `tril`).
429template <typename T>
432 int64_t k = 0) {
433 auto C = detail::like(A);
434 (void)triangular_mask_into<T>(ctx, A, C, Uplo::Lower, k);
435 return C;
436}
437
438// ---- forwarding aliases ----------------------------------------------------
439// Every call below must stay qualified with ::batchlas::: linalg is nested inside
440// batchlas, so `return inv(ctx, A);` here is infinite recursion.
441
442/// @brief Returns \f$A^{-1}\f$ as a new matrix; A is not modified.
443/// @pre `A` is square
444template <typename T>
447 return ::batchlas::inv(ctx, A);
448}
449
450/// @brief Returns \f$A^T\f$ as a new matrix. Real `T` only.
451///
452/// Plain transpose; a complex call is a compile error because nothing conjugates.
453template <typename T>
454 requires RealScalar<T>
457 return ::batchlas::transpose<T, MatrixFormat::Dense>(ctx, A);
458}
459
460/// @brief Returns one norm per batch item. **Blocks** until the result is ready.
461/// @see batchlas::norm for the asynchronous out-parameter form
462template <typename T>
464 Queue& ctx,
466 NormType norm_type = NormType::Frobenius) {
467 return ::batchlas::norm<T, MatrixFormat::Dense>(ctx, A, norm_type);
468}
469
470/// @brief Returns \f$\|A_i\|\,\|A_i^{-1}\|\f$ per item (eigenvalue-magnitude ratio for Spectral).
471///
472/// **Blocks** until the result is ready. Real `T` only.
473template <typename T>
474 requires RealScalar<T>
477 NormType norm_type = NormType::Frobenius) {
478 return ::batchlas::cond(ctx, A, norm_type);
479}
480
481// ---- composites ------------------------------------------------------------
482
483/// @brief Result of svd().
484template <typename T>
485struct Svd {
486 Matrix<T, MatrixFormat::Dense> U; ///< m x m (All) or m x k (Thin)
487 UnifiedVector<typename base_type<T>::type> values; ///< k = min(m, n) per batch item
488 Matrix<T, MatrixFormat::Dense> Vh; ///< n x n (All) or k x n (Thin)
489 /// Per-item convergence status, as Eigh::info. Always filled.
491};
492
493/// @brief Singular value decomposition \f$A_i = U_i\Sigma_iV_i^H\f$; A is not modified.
494///
495/// **Blocks** before returning (its internal copy of A must outlive the kernels).
496/// @param ctx queue; supplies the backend and the arena
497/// @param A m x n per item
498/// @param vectors All or Thin
499/// @return U, singular values, Vh and per-item info
500/// @throws batchlas::invalid_argument for SvdVectors::None; call batchlas::gesvd
501/// directly for values only
502/// @note Complex input the one-sided Jacobi kernel cannot take (beyond its size
503/// limit, or on a device without sub-group 32) needs the vendor library: no
504/// other native gesvd kernel takes complex general input,
505/// so a build without it throws batchlas::NoRouteError.
506template <typename T>
507inline Svd<T> svd(Queue& ctx,
509 SvdVectors vectors = SvdVectors::All) {
510 if (vectors == SvdVectors::None) {
512 "linalg::svd: SvdVectors::None would leave U and Vh empty; use "
513 "::batchlas::gesvd directly for a values-only decomposition");
514 }
515 const int64_t m = A.rows();
516 const int64_t n = A.cols();
517 const int64_t k = std::min(m, n);
518 const int batch = A.batch_size();
519
520 auto work = detail::like(A);
522
523 Matrix<T, MatrixFormat::Dense> U(static_cast<int>(m),
524 static_cast<int>(svd_u_cols(vectors, m, k)), batch);
525 Matrix<T, MatrixFormat::Dense> Vh(static_cast<int>(svd_vh_rows(vectors, n, k)),
526 static_cast<int>(n), batch);
527 UnifiedVector<typename base_type<T>::type> S(static_cast<size_t>(k) *
528 static_cast<size_t>(batch));
529
530 UnifiedVector<int32_t> info(static_cast<size_t>(batch));
531 // Positional because GesvdOptions has no `info` field.
533 ctx, work, S.to_span(), U, Vh, vectors, vectors));
534 (void)::batchlas::gesvd(ctx, work, S.to_span(), U, Vh, vectors, vectors,
535 lease.span(), info.to_span());
536
537 // ~Matrix frees `work` without waiting; returning early frees it under the kernels.
538 ctx.wait();
539 return Svd<T>{std::move(U), std::move(S), std::move(Vh), std::move(info)};
540}
541
542/// @brief Result of lu().
543template <typename T>
544struct Lu {
545 Matrix<T, MatrixFormat::Dense> factors; ///< L and U packed; L's unit diagonal is implicit
546 UnifiedVector<int64_t> pivots; ///< rows * batch_size, in getrf()'s pivot format
547};
548
549/// @brief LU factorisation with partial pivoting, \f$A_i = P_iL_iU_i\f$; A is not modified.
550/// @pre `A` is square
551/// @return packed factors and pivots, valid after the queue is waited on
552// Pivots are a UnifiedVector, not an arena lease as in solve(): they outlive the call.
553template <typename T>
555 auto LU = detail::like(A);
557 UnifiedVector<int64_t> pivots(static_cast<size_t>(A.rows()) *
558 static_cast<size_t>(A.batch_size()));
559 (void)::batchlas::getrf(ctx, LU, pivots.to_span());
560 return Lu<T>{std::move(LU), std::move(pivots)};
561}
562
563// linalg::qr is deliberately absent: an unexplained cross-Queue wrong answer.
564// evidence: docs/design/known-defects.md#linalgqr-returns-a-wrong-qr-after-an-earlier-call-in-the-process
565
566
567/// @}
568
569} // namespace batchlas::linalg
Non-owning view of a batch of matrices: the argument type of every entry point.
Definition matrix.hh:485
int rows() const
Rows per item (the capacity on a heterogeneous batch).
Definition matrix.hh:605
int cols() const
Columns per item (the capacity on a heterogeneous batch).
Definition matrix.hh:611
static Event copy(Queue &ctx, const MatrixView< T, MType > &dest, const MatrixView< T, MType > &src)
Enqueues dest := src on ctx, honouring each side's ld and stride.
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
The call violates the API contract.
Definition error.hh:77
A real floating-point scalar (float, double).
Definition enums.hh:57
Enumerations, scalar traits and scalar concepts shared by every public header.
Matrix norms, condition numbers, conditioned test-matrix generators and transpose.
Batched general linear solve (gesv) by LU with partial pivoting.
Event gemm(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, const MatrixView< T, MatrixFormat::Dense > &C, T alpha, T beta, Transpose transA, Transpose transB, ComputePrecision precision=ComputePrecision::Default)
Batched general matrix-matrix multiply.
size_t syev_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, Span< typename base_type< T >::type > eigenvalues, JobType jobtype, Uplo uplo)
Workspace, in bytes, that syev() needs for the same arguments.
Event syev(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &descrA, Span< typename base_type< T >::type > eigenvalues, JobType jobtype, Uplo uplo, Span< std::byte > workspace, Span< int32_t > info)
Eigenvalues, and optionally eigenvectors, of a batch of symmetric/Hermitian matrices (LAPACK ?...
constexpr int64_t svd_u_cols(SvdVectors job, int64_t m, int64_t k)
Number of columns of U that job writes.
Definition enums.hh:243
constexpr int64_t svd_vh_rows(SvdVectors job, int64_t n, int64_t k)
Number of rows of V^H that job writes.
Definition enums.hh:253
Transpose
The operator op() applied to an operand: op(A) = A, A^T or A^H.
Definition enums.hh:171
NormType
Matrix norm computed by norm and used by the condition estimators.
Definition enums.hh:545
ComputePrecision
Internal compute precision of a GEMM (the cuBLAS cublasComputeType_t).
Definition enums.hh:505
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
@ EigenVectors
Eigenvalues and eigenvectors (jobz = 'V').
Definition enums.hh:190
@ NoTrans
op(A) = A.
Definition enums.hh:172
@ Frobenius
sqrt(sum |a_ij|^2).
Definition enums.hh:546
@ Default
The precision of the scalar type.
Definition enums.hh:506
@ Upper
The upper triangle (i <= j).
Definition enums.hh:288
@ Lower
The lower triangle (i >= j).
Definition enums.hh:289
@ None
Do not compute this factor.
Definition enums.hh:221
@ All
Full square factor (LAPACK 'A').
Definition enums.hh:222
Event scale(Queue &ctx, const T &alpha, const MatrixView< T, MType > &mat_view)
In-place A := alpha * A on every element of every item (dense).
size_t posv_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, Uplo uplo)
Workspace, in bytes, that posv() needs for these operands on this queue.
size_t gesv_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B)
Workspace, in bytes, that gesv() needs for these operands on this queue.
Event potrf(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &descrA, Uplo uplo, Span< std::byte > workspace, Span< int32_t > info)
Batched Cholesky factorization of Hermitian positive-definite matrices.
Event getrs(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, Transpose transA, Span< int64_t > pivots, Span< std::byte > work_space)
Batched solve of using the LU factors from getrf().
Event getrf(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, Span< int64_t > pivots, Span< std::byte > work_space, Span< int32_t > info)
Batched LU factorization with partial pivoting.
Event posv(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, Uplo uplo, Span< std::byte > work_space, Span< int32_t > info)
Batched solve of for Hermitian positive-definite A.
Event gesv(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, Span< int64_t > pivots, Span< std::byte > work_space, Span< int32_t > info)
Batched solve of by LU with partial pivoting.
Eigh< T > eigh(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, Uplo uplo=kDefaultUplo)
Eigenvalues and eigenvectors of a symmetric/Hermitian A; A is not modified.
Matrix< T, MatrixFormat::Dense > transpose(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A)
Returns as a new matrix.
Matrix< T, MatrixFormat::Dense > subtract(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B)
Returns a new matrix (elementwise).
Matrix< T, MatrixFormat::Dense > divide(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B)
Returns a new matrix (elementwise).
Event triangular_mask_into(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &C, Uplo uplo, int64_t k=0)
A with everything outside the requested triangle zeroed.
Event divide_into(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, const MatrixView< T, MatrixFormat::Dense > &C)
elementwise; see elementwise_into().
Matrix< T, MatrixFormat::Dense > multiply(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B)
Returns a new matrix (Hadamard, not matmul()).
Matrix< T, MatrixFormat::Dense > solve_spd(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, Uplo uplo=Uplo::Lower)
Solves for Hermitian positive-definite A by Cholesky; returns X.
Event elementwise_into(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, const MatrixView< T, MatrixFormat::Dense > &C)
Elementwise for = Op, per batch item.
Event scale(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, T alpha)
In-place scaling .
Lu< T > lu(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A)
LU factorisation with partial pivoting, ; A is not modified.
Matrix< T, MatrixFormat::Dense > matmul(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, const MatmulOptions< T > &opts={})
Returns a new matrix .
Matrix< T, MatrixFormat::Dense > cholesky(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, Uplo uplo=kDefaultUplo)
Returns the Cholesky factor of a Hermitian positive-definite A; A is not modified.
Event add_into(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, const MatrixView< T, MatrixFormat::Dense > &C)
elementwise; see elementwise_into().
Event axpby_into(Queue &ctx, T alpha, const MatrixView< T, MatrixFormat::Dense > &A, T beta, const MatrixView< T, MatrixFormat::Dense > &B, const MatrixView< T, MatrixFormat::Dense > &C)
, elementwise, per batch item.
Svd< T > svd(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, SvdVectors vectors=SvdVectors::All)
Singular value decomposition ; A is not modified.
UnifiedVector< typename base_type< T >::type > eigvalsh(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, Uplo uplo=kDefaultUplo)
Returns the eigenvalues of a symmetric/Hermitian A, ascending per item; A is not modified.
UnifiedVector< typename base_type< T >::type > norm(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, NormType norm_type=NormType::Frobenius)
Returns one norm per batch item.
Matrix< T, MatrixFormat::Dense > solve(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, Transpose trans=Transpose::NoTrans)
Solves by LU with partial pivoting; returns X.
Matrix< T, MatrixFormat::Dense > triu(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, int64_t k=0)
Returns a copy of A with everything below diagonal k zeroed (NumPy triu).
Matrix< T, MatrixFormat::Dense > inv(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A)
Returns as a new matrix; A is not modified.
Matrix< T, MatrixFormat::Dense > scaled(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, T alpha)
Returns a new matrix ; A is not modified (scale() works in place).
BinaryOp
Elementwise operation selected by elementwise_into().
Definition linalg-ops.hh:54
Event multiply_into(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, const MatrixView< T, MatrixFormat::Dense > &C)
elementwise (Hadamard), not the matrix product; see matmul().
Matrix< T, MatrixFormat::Dense > add(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B)
Returns a new matrix (elementwise).
Matrix< T, MatrixFormat::Dense > tril(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, int64_t k=0)
Returns a copy of A with everything above diagonal k zeroed (NumPy tril).
constexpr Uplo kDefaultUplo
Default triangle of the linalg wrappers that read one.
Event subtract_into(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, const MatrixView< T, MatrixFormat::Dense > &B, const MatrixView< T, MatrixFormat::Dense > &C)
elementwise; see elementwise_into().
UnifiedVector< T > cond(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, NormType norm_type=NormType::Frobenius)
Returns per item (eigenvalue-magnitude ratio for Spectral).
size_t gesvd_buffer_size(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, Span< typename base_type< T >::type > singular_values, const MatrixView< T, MatrixFormat::Dense > &U, const MatrixView< T, MatrixFormat::Dense > &Vh, SvdVectors jobu, SvdVectors jobvh)
Workspace, in bytes, that gesvd() needs for the same arguments.
Event gesvd(Queue &ctx, const MatrixView< T, MatrixFormat::Dense > &A, Span< typename base_type< T >::type > singular_values, const MatrixView< T, MatrixFormat::Dense > &U, const MatrixView< T, MatrixFormat::Dense > &Vh, SvdVectors jobu, SvdVectors jobvh, Span< std::byte > workspace, Span< int32_t > info)
Singular value decomposition of a batch of matrices (LAPACK ?gesvd).
The batched containers: Matrix (owning), MatrixView (non-owning), KernelMatrixView (device-side),...
void require_square(const char *fn, const char *name, const MV &A)
Definition options.hh:98
Matrix< T, MatrixFormat::Dense > like(const MatrixView< T, MatrixFormat::Dense > &A)
Option-struct spellings of the dense BLAS and LAPACK-style entry points.
Batched Hermitian positive-definite solve (posv) by Cholesky.
Completion handle of enqueued work, returned by every entry point.
Event()
An empty Event, already complete.
Options for gemm(): .
Definition options.hh:144
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
Owning, growable array of T in USM shared memory (sycl::malloc_shared).
constexpr Span< T > to_span() const
A Span over the first size() elements.
Result of eigh().
UnifiedVector< typename base_type< T >::type > values
ascending, n per batch item
UnifiedVector< int32_t > info
Per-item convergence status: 0 converged, > 0 LAPACK-like.
Matrix< T, MatrixFormat::Dense > vectors
eigenvectors as columns, n x n per item
Result of lu().
Matrix< T, MatrixFormat::Dense > factors
L and U packed; L's unit diagonal is implicit.
UnifiedVector< int64_t > pivots
rows * batch_size, in getrf()'s pivot format
Options for matmul(). Deliberately has no beta: C is freshly allocated.
ComputePrecision precision
internal compute type
T alpha
scale of the product
Result of svd().
UnifiedVector< typename base_type< T >::type > values
k = min(m, n) per batch item
UnifiedVector< int32_t > info
Per-item convergence status, as Eigh::info. Always filled.
Matrix< T, MatrixFormat::Dense > Vh
n x n (All) or k x n (Thin)
Matrix< T, MatrixFormat::Dense > U
m x m (All) or m x k (Thin)
batchlas::Device, batchlas::Event and batchlas::Queue: where work runs and how it is ordered.
batchlas::Span, the non-owning pointer + length every entry point takes for 1-D data.
batchlas::UnifiedVector, an owning growable array in USM shared memory.