L2 — Matrix-Vector Operations#

General matrix-vector products (gemv), rank-1 updates (ger), the strided / segmented gemv variants, and triangular solve / matvec (trsv / trmv).

gemv#

Functions

template<typename T, bool TRANSPOSE = false, bool ROW_MAJOR = false, bool TRAILING_SYNC = true>
void gemv(uint32_t m, uint32_t n, T alpha, const T *A, const T *x, T beta, T *y)#

Matrix-vector product: y = alpha * A * x + beta * y (GEMV).

Threads are distributed over the output rows of the m×n matrix A. Set TRANSPOSE=true to compute Aᵀ * x and ROW_MAJOR=true for row-major A (A is column-major by default). NumPy equivalent: y = alpha*A@x + beta*y (or alpha*A.T@x + beta*y when transposed).

Unlike gemm — where a row-major operand is just a transpose, so the only layout flag is ROW_MAJOR_C — GEMV keeps a per-matrix ROW_MAJOR flag: TRANSPOSE already selects the mathematical operation (A·x vs Aᵀ·x), so it cannot also stand in for the storage order. TRANSPOSE and ROW_MAJOR are therefore independent. (This flag fully subsumes the former gemv_ex, which was just gemv with the defaults removed and has been deleted.)

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • TRANSPOSE – When true, multiply by Aᵀ instead of A (default false).

  • ROW_MAJOR – When true, A is stored row-major (default false = column-major).

Parameters:

  • m – Number of rows of A.

  • n – Number of columns of A.

  • alpha – Scalar multiplier on the product.

  • A – Input matrix of m*n elements.

  • x – Input vector (length n, or m when transposed).

  • beta – Scalar multiplier on the prior y.

  • y – In/out vector (length m, or n when transposed).

template<typename T, bool TRANSPOSE = false, bool ROW_MAJOR = false, bool TRAILING_SYNC = true>
void gemv(uint32_t m, uint32_t n, T alpha, const T *A, const T *x, T *y)#

Matrix-vector product: y = alpha * A * x (GEMV), no-beta overload.

Same as the full GEMV but overwrites y (no beta * y term). Set TRANSPOSE=true for Aᵀ * x and ROW_MAJOR=true for row-major A. NumPy equivalent: y = alpha*A@x (or alpha*A.T@x when transposed).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • TRANSPOSE – When true, multiply by Aᵀ instead of A (default false).

  • ROW_MAJOR – When true, A is stored row-major (default false = column-major).

Parameters:

  • m – Number of rows of A.

  • n – Number of columns of A.

  • alpha – Scalar multiplier on the product.

  • A – Input matrix of m*n elements.

  • x – Input vector (length n, or m when transposed).

  • y – Output vector (length m, or n when transposed).

template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE, bool ROW_MAJOR_A>
void gemv_impl_ct(uint32_t rank, uint32_t size, T alpha, const T *A, const T *x, T beta, T *y)#
template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE, bool ROW_MAJOR_A>
void gemv_impl_ct(uint32_t rank, uint32_t size, T alpha, const T *A, const T *x, T *y)#
template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE = false, bool ROW_MAJOR = false, bool TRAILING_SYNC = true>
void gemv(T alpha, const T *A, const T *x, T beta, T *y)#

Matrix-vector product: y = alpha * A * x + beta * y (GEMV), compile-time size.

Compile-time-M,N overload; the inner column loop is fully unrolled. Set TRANSPOSE=true for Aᵀ * x and ROW_MAJOR=true for row-major A. NumPy equivalent: y = alpha*A@x + beta*y (or alpha*A.T@x + beta*y).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • M – Number of rows of A (compile-time constant).

  • N – Number of columns of A (compile-time constant).

  • TRANSPOSE – When true, multiply by Aᵀ instead of A (default false).

  • ROW_MAJOR – When true, A is stored row-major (default false = column-major).

Parameters:

  • alpha – Scalar multiplier on the product.

  • A – Input matrix of M*N elements.

  • x – Input vector (length N, or M when transposed).

  • beta – Scalar multiplier on the prior y.

  • y – In/out vector (length M, or N when transposed).

template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE = false, bool ROW_MAJOR = false, bool TRAILING_SYNC = true>
void gemv(T alpha, const T *A, const T *x, T *y)#

Matrix-vector product: y = alpha * A * x (GEMV), compile-time size, no-beta overload.

Compile-time-M,N overload that overwrites y (no beta * y term). NumPy equivalent: y = alpha*A@x (or alpha*A.T@x when transposed).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • M – Number of rows of A (compile-time constant).

  • N – Number of columns of A (compile-time constant).

  • TRANSPOSE – When true, multiply by Aᵀ instead of A (default false).

  • ROW_MAJOR – When true, A is stored row-major (default false = column-major).

Parameters:

  • alpha – Scalar multiplier on the product.

  • A – Input matrix of M*N elements.

  • x – Input vector (length N, or M when transposed).

  • y – Output vector (length M, or N when transposed).

template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE = false, bool ROW_MAJOR = false>
void gemv(T alpha, const T *A, const T *x, T beta, T *y)#

Matrix-vector product within one warp: y = alpha * A * x + beta * y (GEMV), single-warp, compile-time size.

Matrix-vector product on one thread: y = alpha * A * x + beta * y (GEMV), compile-time size.

One 32-lane warp computes the matvec with lanes striding over the output rows of the M×N matrix A (each row an independent inner product). Set TRANSPOSE=true for Aᵀ * x and ROW_MAJOR=true for row-major A. No shared scratch, no __syncthreads; independent warps may run distinct problems concurrently. Full 32 lanes required. y is read only when beta != 0 (BLAS semantics: beta == 0 treats y as write-only). NumPy equivalent: y = alpha*A@x + beta*y (or alpha*A.T@x + beta*y when transposed).

One thread computes the matvec, walking the output rows of the M×N matrix A serially (each row an independent inner product). Set TRANSPOSE=true for Aᵀ * x and ROW_MAJOR=true for row-major A. No shared scratch, no barriers, no threadIdx read; operands may be thread-local register arrays. y is read only when beta != 0 (BLAS semantics: beta == 0 treats y as write-only). NumPy equivalent: y = alpha*A@x + beta*y.

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • M – Number of rows of A (compile-time constant).

  • N – Number of columns of A (compile-time constant).

  • TRANSPOSE – When true, multiply by Aᵀ instead of A (default false).

  • ROW_MAJOR – When true, A is stored row-major (default false = column-major).

  • T – Scalar type (e.g. float, double).

  • M – Number of rows of A (compile-time constant).

  • N – Number of columns of A (compile-time constant).

  • TRANSPOSE – When true, multiply by Aᵀ instead of A (default false).

  • ROW_MAJOR – When true, A is stored row-major (default false = column-major).

Parameters:

  • alpha – Scalar multiplier on the product.

  • A – Input matrix of M*N elements.

  • x – Input vector (length N, or M when transposed).

  • beta – Scalar multiplier on the prior y.

  • y – In/out vector (length M, or N when transposed).

  • alpha – Scalar multiplier on the product.

  • A – Input matrix of M*N elements.

  • x – Input vector (length N, or M when transposed).

  • beta – Scalar multiplier on the prior y.

  • y – In/out vector (length M, or N when transposed).

template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE = false, bool ROW_MAJOR = false>
void gemv(T alpha, const T *A, const T *x, T *y)#

Matrix-vector product within one warp: y = alpha * A * x (GEMV), single-warp, compile-time size, implicit beta = 0.

Matrix-vector product on one thread: y = alpha * A * x (GEMV), compile-time size, no-beta overload.

Overwrites y (no beta * y term — y is never read, so it is safe to write into cold/uninitialized scratch). Otherwise identical to the beta overload above. No shared scratch, no __syncthreads. Full 32 lanes required. NumPy equivalent: y = alpha*A@x (or alpha*A.T@x when transposed).

Overwrites y (no beta * y term). NumPy equivalent: y = alpha*A@x (or alpha*A.T@x when transposed).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • M – Number of rows of A (compile-time constant).

  • N – Number of columns of A (compile-time constant).

  • TRANSPOSE – When true, multiply by Aᵀ instead of A (default false).

  • ROW_MAJOR – When true, A is stored row-major (default false = column-major).

  • T – Scalar type (e.g. float, double).

  • M – Number of rows of A (compile-time constant).

  • N – Number of columns of A (compile-time constant).

  • TRANSPOSE – When true, multiply by Aᵀ instead of A (default false).

  • ROW_MAJOR – When true, A is stored row-major (default false = column-major).

Parameters:

  • alpha – Scalar multiplier on the product.

  • A – Input matrix of M*N elements.

  • x – Input vector (length N, or M when transposed).

  • y – Output vector (length M, or N when transposed; overwritten).

  • alpha – Scalar multiplier on the product.

  • A – Input matrix of M*N elements.

  • x – Input vector (length N, or M when transposed).

  • y – Output vector (length M, or N when transposed).

namespace warp
namespace thread

Contraction-parallel gemv (gemv_reduced)#

One warp owns each output element and its lanes split the contraction — the L2 sibling of gemm_reduced. See Contraction-parallel ops (the *_reduced family).

Functions

template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE = false, bool TRAILING_SYNC = true>
void gemv_reduced(T alpha, const T *A, const T *x, T beta, T *y)#

Contraction-parallel GEMV: y = alpha * op(A) * x + beta * y.

Single-warp contraction-parallel GEMV: y = alpha * op(A) * x + beta * y.

Compile-time-size matrix-vector product that parallelizes the contraction across a warp’s lanes (one warp per output element) rather than one thread summing serially — the L2 analogue of glass::gemm_reduced. Column-major A. Thread-count invariant at any block size.

Warp-per-problem analogue of glass::gemv_reduced; one full 32-lane warp.

Template Parameters:

  • T – Scalar type.

  • M, N – A is M x N (column-major).

  • TRANSPOSE – If true, computes Aᵀ x (output length N, contract over M); else A x (length M, contract over N).

  • TRAILING_SYNC – Emit a trailing __syncthreads() (default true).

  • T, M, N, TRANSPOSE – See glass::gemv_reduced.

  • TRAILING_SYNC – Emit a trailing __syncwarp() (default true).

Parameters:

  • alpha – Scalar on the product.

  • A – Input matrix (M x N, column-major).

  • x – Input vector (length N if !TRANSPOSE else M).

  • beta – Scalar on the existing y (read only when beta != 0).

  • y – In/out result (length M if !TRANSPOSE else N).

  • alpha, A, x, beta, y – See glass::gemv_reduced.

template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE = false, bool TRAILING_SYNC = true>
void gemv_reduced(T alpha, const T *A, const T *x, T *y)#

Contraction-parallel GEMV with implicit beta = 0: y = alpha * op(A) * x.

Single-warp contraction-parallel GEMV, implicit beta = 0: y = alpha * op(A) * x.

Overwrites y (not read). Otherwise identical to the beta overload.

Template Parameters:

  • T, M, N, TRANSPOSE, TRAILING_SYNC – See the beta overload.

  • T, M, N, TRANSPOSE, TRAILING_SYNC – See the beta overload.

Parameters:

  • alpha, A, x – See the beta overload.

  • y – Output (overwritten).

  • alpha, A, x, y – See the beta overload.

namespace gemv_reduced_detail#
namespace warp

Strided / segmented gemv#

Functions

template<typename T, uint32_t M, uint32_t N, uint32_t ROW_STRIDE = M>
void gemv_strided(T alpha, const T *A, const T *x, T beta, T *y)#

Column-major GEMV with an explicit leading dimension: y = alpha * A * x + beta * y.

Compile-time-M,N matrix-vector product where A[i][j] = A[i + j*ROW_STRIDE], letting an M×N matrix be addressed inside a larger array (e.g. a spatial 6×6 embedded in a wider buffer, as in GRiD). When ROW_STRIDE == M this is identical to glass::gemv<T,M,N>. NumPy equivalent: y = alpha*A@x + beta*y.

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • M – Number of rows of A (compile-time constant).

  • N – Number of columns of A (compile-time constant).

  • ROW_STRIDE – Column-major leading dimension of A (default M).

Parameters:

  • alpha – Scalar multiplier on the product.

  • A – Input matrix, addressed at A[row + col*ROW_STRIDE].

  • x – Input vector of length N.

  • beta – Scalar multiplier on the prior y.

  • y – In/out vector of length M.

template<typename T, uint32_t M, uint32_t N, uint32_t ROW_STRIDE = M>
void gemv_strided(T alpha, const T *A, const T *x, T *y)#

Column-major GEMV with an explicit leading dimension: y = alpha * A * x, no-beta overload.

No-beta variant of the strided GEMV that overwrites y, with A[i][j] = A[i + j*ROW_STRIDE]. NumPy equivalent: y = alpha*A@x.

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • M – Number of rows of A (compile-time constant).

  • N – Number of columns of A (compile-time constant).

  • ROW_STRIDE – Column-major leading dimension of A (default M).

Parameters:

  • alpha – Scalar multiplier on the product.

  • A – Input matrix, addressed at A[row + col*ROW_STRIDE].

  • x – Input vector of length N.

  • y – Output vector of length M.

Functions

template<typename T, uint32_t M, uint32_t N, uint32_t ROW_STRIDE = M, bool FUSE_SCALED_ADD = false, bool TRANSPOSE = false, bool ATOMIC_Y = false, typename IDX_T = int>
void gemv_segmented(uint32_t segments, const IDX_T *seg_a_off, const IDX_T *seg_x_off, const IDX_T *seg_y_off, const T *A, const T *x, T *y, T alpha, T beta, const IDX_T *seg_s_off = nullptr, const T *S = nullptr, const T *scalar = nullptr)#

Segmented (batched) column-major GEMV: y_seg = alpha*A_seg*x_seg + beta*y_seg for each segment.

Computes segments independent small GEMVs concurrently in one block. Each A_seg is M×N, column-major with leading dimension ROW_STRIDE, and base offsets into the flat A/x/y arrays come from the descriptor arrays. A single block-stride loop walks the flattened segments * OUT_ROWS outputs (OUT_ROWS = TRANSPOSE ? N : M). Optional flags: TRANSPOSE computes Aᵀ_seg · x_seg (the leaf→root direction); ATOMIC_Y accumulates into y via atomicAdd so segments may overlap (beta is then ignored — pre-scale y yourself); FUSE_SCALED_ADD folds an extra per-segment S * scalar term into the single y store (non-atomic path only).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • M – Rows per segment matrix (compile-time constant).

  • N – Columns per segment matrix (compile-time constant).

  • ROW_STRIDE – Column-major leading dimension of each A_seg (default M).

  • FUSE_SCALED_ADD – When true, add S[seg]*scalar[seg] into the output (non-atomic only).

  • TRANSPOSE – When true, each segment computes Aᵀ_seg · x_seg.

  • ATOMIC_Y – When true, accumulate into y with atomicAdd (overlap-safe).

  • IDX_T – Index type of the offset descriptor arrays (default int).

Parameters:

  • segments – Number of independent GEMVs.

  • seg_a_off – Per-segment base element offsets of A_seg within A.

  • seg_x_off – Per-segment base element offsets of x_seg within x.

  • seg_y_off – Per-segment base element offsets of y_seg within y.

  • A – Flat backing array of all segment matrices.

  • x – Flat backing array of all segment input vectors.

  • y – Flat backing array of all segment output vectors (in/out).

  • alpha – Scalar multiplier on each segment product.

  • beta – Scalar multiplier on the prior y (ignored under ATOMIC_Y).

  • seg_s_off – Per-segment offsets of S_seg within S (FUSE only).

  • S – Flat backing array for the fused scaled-add term (FUSE only).

  • scalar – Per-segment multiplier for the fused scaled-add (FUSE only).

template<typename T, uint32_t M, uint32_t N, uint32_t ROW_STRIDE = M, bool FUSE_SCALED_ADD = false, bool TRANSPOSE = false, bool ATOMIC_Y = false, typename IDX_T = int>
void gemv_segmented(uint32_t segments, const IDX_T *seg_a_off, const IDX_T *seg_x_off, const IDX_T *seg_y_off, const T *A, const T *x, T *y, T alpha, const IDX_T *seg_s_off = nullptr, const T *S = nullptr, const T *scalar = nullptr)#

Segmented (batched) column-major GEMV, no-beta overload: y_seg = alpha*A_seg*x_seg per segment.

No-beta variant of gemv_segmented: each segment overwrites its y_seg (optionally plus the fused scaled-add) or computes the transpose. This also serves as the ATOMIC_Y entry point, since atomic accumulate takes no beta (see the full overload’s note on beta-under-atomic).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • M – Rows per segment matrix (compile-time constant).

  • N – Columns per segment matrix (compile-time constant).

  • ROW_STRIDE – Column-major leading dimension of each A_seg (default M).

  • FUSE_SCALED_ADD – When true, add S[seg]*scalar[seg] into the output (non-atomic only).

  • TRANSPOSE – When true, each segment computes Aᵀ_seg · x_seg.

  • ATOMIC_Y – When true, accumulate into y with atomicAdd (overlap-safe).

  • IDX_T – Index type of the offset descriptor arrays (default int).

Parameters:

  • segments – Number of independent GEMVs.

  • seg_a_off – Per-segment base element offsets of A_seg within A.

  • seg_x_off – Per-segment base element offsets of x_seg within x.

  • seg_y_off – Per-segment base element offsets of y_seg within y.

  • A – Flat backing array of all segment matrices.

  • x – Flat backing array of all segment input vectors.

  • y – Flat backing array of all segment output vectors (out, or accumulated under ATOMIC_Y).

  • alpha – Scalar multiplier on each segment product.

  • seg_s_off – Per-segment offsets of S_seg within S (FUSE only).

  • S – Flat backing array for the fused scaled-add term (FUSE only).

  • scalar – Per-segment multiplier for the fused scaled-add (FUSE only).

ger#

Functions

template<typename T, bool TRAILING_SYNC = true>
void ger(uint32_t m, uint32_t n, T alpha, const T *x, const T *y, T *A)#

Rank-1 update: A += alpha * x * yᵀ (GER).

Adds the scaled outer product of x and y to the m×n column-major matrix A. NumPy equivalent: A += alpha * np.outer(x, y).

Template Parameters:

T – Scalar type (e.g. float, double).

Parameters:

  • m – Number of rows of A (length of x).

  • n – Number of columns of A (length of y).

  • alpha – Scalar multiplier on the outer product.

  • x – Input vector of length m.

  • y – Input vector of length n.

  • A – In/out matrix of m*n elements (column-major).

template<typename T, uint32_t M, uint32_t N, bool TRAILING_SYNC = true>
void ger(T alpha, const T *x, const T *y, T *A)#

Rank-1 update: A += alpha * x * yᵀ (GER), compile-time size.

Compile-time-M,N overload of the rank-1 update. NumPy equivalent: A += alpha * np.outer(x, y).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • M – Number of rows of A / length of x (compile-time constant).

  • N – Number of columns of A / length of y (compile-time constant).

Parameters:

  • alpha – Scalar multiplier on the outer product.

  • x – Input vector of length M.

  • y – Input vector of length N.

  • A – In/out matrix of M*N elements (column-major).

Triangular solve / matvec (trsv / trmv)#

Functions

template<typename T, FillMode FILL = FillMode::Lower, Diag DIAG = Diag::NonUnit, bool TRANSPOSE = false, bool TRAILING_SYNC = true>
void trsv(uint32_t n, const T *A, T *x)#

Triangular solve op(A) x = b in place (TRSV).

Solves the triangular system for x, overwriting the right-hand side x (x holds b on entry, the solution on return). A is an n×n triangular matrix stored column-major; only the triangle selected by FILL is read. Set TRANSPOSE=true to solve Aᵀx = b against that same stored triangle, and DIAG=Diag::Unit for an implicit unit diagonal (the diagonal of A is not read). Column-oriented elimination (forward when op(A) is lower-triangular, i.e. (FILL==Lower) != TRANSPOSE, backward otherwise); ends with a trailing __syncthreads() so it composes without a defensive barrier. SciPy equivalent: x = scipy.linalg.solve_triangular(A, b, lower=(FILL==Lower), unit_diagonal=(DIAG==Unit), trans=(1 if TRANSPOSE else 0)).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • FILL – Which triangle of A holds the data (default FillMode::Lower).

  • DIAG – Diag::Unit for an implicit unit diagonal (default Diag::NonUnit).

  • TRANSPOSE – When true solve Aᵀx = b (default false).

Parameters:

  • n – Dimension (A is n×n, x has length n).

  • A – Triangular matrix (column-major, n*n elements; read-only).

  • x – In/out right-hand side; on return holds the solution.

template<typename T, uint32_t N, FillMode FILL = FillMode::Lower, Diag DIAG = Diag::NonUnit, bool TRANSPOSE = false, bool TRAILING_SYNC = true>
void trsv(const T *A, T *x)#

Triangular solve op(A) x = b in place (TRSV), compile-time size.

Same as the runtime trsv but with the dimension as a template parameter. SciPy equivalent: x = scipy.linalg.solve_triangular(A, b, lower=(FILL==Lower), unit_diagonal=(DIAG==Unit), trans=(1 if TRANSPOSE else 0)).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • N – Dimension (A is N×N, x has length N).

  • FILL – Which triangle of A holds the data (default FillMode::Lower).

  • DIAG – Diag::Unit for an implicit unit diagonal (default Diag::NonUnit).

  • TRANSPOSE – When true solve Aᵀx = b (default false).

Parameters:

  • A – Triangular matrix (column-major, N*N elements; read-only).

  • x – In/out right-hand side; on return holds the solution.

template<typename T, FillMode FILL = FillMode::Lower, Diag DIAG = Diag::NonUnit, bool TRANSPOSE = false, bool TRAILING_SYNC = true>
void trmv(uint32_t n, const T *A, const T *x, T *y)#

Triangular matrix-vector product y = op(A) x, out of place (TRMV).

Computes the triangular matvec into a separate output y (distinct from the input x). A is an n×n triangular matrix stored column-major; only the triangle selected by FILL is read. Set TRANSPOSE=true to compute Aᵀx against that same stored triangle, and DIAG=Diag::Unit for an implicit unit diagonal. No interior barrier: each thread owns disjoint outputs and reads the intact x. NumPy equivalent (lower, non-unit): y = np.tril(A) @ x (upper: np.triu(A) @ x; transposed: op(A).T @ x; unit: diagonal forced to 1).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • FILL – Which triangle of A holds the data (default FillMode::Lower).

  • DIAG – Diag::Unit for an implicit unit diagonal (default Diag::NonUnit).

  • TRANSPOSE – When true compute Aᵀx (default false).

Parameters:

  • n – Dimension (A is n×n, x and y have length n).

  • A – Triangular matrix (column-major, n*n elements; read-only).

  • x – Input vector (length n; read-only).

  • y – Output vector (length n); must not alias x.

template<typename T, uint32_t N, FillMode FILL = FillMode::Lower, Diag DIAG = Diag::NonUnit, bool TRANSPOSE = false, bool TRAILING_SYNC = true>
void trmv(const T *A, const T *x, T *y)#

Triangular matrix-vector product y = op(A) x, out of place (TRMV), compile-time size.

Compile-time-N overload of the out-of-place TRMV. NumPy equivalent (lower, non-unit): y = np.tril(A) @ x.

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • N – Dimension (A is N×N, x and y have length N).

  • FILL – Which triangle of A holds the data (default FillMode::Lower).

  • DIAG – Diag::Unit for an implicit unit diagonal (default Diag::NonUnit).

  • TRANSPOSE – When true compute Aᵀx (default false).

Parameters:

  • A – Triangular matrix (column-major, N*N elements; read-only).

  • x – Input vector (length N; read-only).

  • y – Output vector (length N); must not alias x.

template<typename T>
inline constexpr std::size_t trmv_scratch_bytes(uint32_t n)#

Scratch length (in elements of T) required by the in-place trmv.

The in-place TRMV wrapper computes op(A) x into a temporary then copies it back over x; that temporary is n elements long.

Parameters:

n – Dimension passed to the in-place trmv.

Returns:

Bytes the scratch buffer must hold.

template<typename T, FillMode FILL = FillMode::Lower, Diag DIAG = Diag::NonUnit, bool TRANSPOSE = false, bool TRAILING_SYNC = true>
void trmv(uint32_t n, const T *A, T *x, T *scratch)#

Triangular matrix-vector product x = op(A) x, in place (TRMV).

In-place form: overwrites x with op(A) x. Because trmv reads the whole x while writing each output, the wrapper computes into a caller-provided scratch (length n, see trmv_scratch_bytes) and copies the result back with a single barrier in between. Ends with a trailing __syncthreads(). NumPy equivalent (lower, non-unit): x = np.tril(A) @ x.

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • FILL – Which triangle of A holds the data (default FillMode::Lower).

  • DIAG – Diag::Unit for an implicit unit diagonal (default Diag::NonUnit).

  • TRANSPOSE – When true compute Aᵀx (default false).

Parameters:

  • n – Dimension (A is n×n, x and scratch have length n).

  • A – Triangular matrix (column-major, n*n elements; read-only).

  • x – In/out vector (length n); on return holds op(A) x.

  • scratch – Workspace of length n (see trmv_scratch_bytes).

template<typename T, uint32_t N, FillMode FILL = FillMode::Lower, Diag DIAG = Diag::NonUnit, bool TRANSPOSE = false, bool TRAILING_SYNC = true>
void trmv(const T *A, T *x, T *scratch)#

Triangular matrix-vector product x = op(A) x, in place (TRMV), compile-time size.

Compile-time-N overload of the in-place TRMV. NumPy equivalent (lower, non-unit): x = np.tril(A) @ x.

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • N – Dimension (A is N×N, x and scratch have length N).

  • FILL – Which triangle of A holds the data (default FillMode::Lower).

  • DIAG – Diag::Unit for an implicit unit diagonal (default Diag::NonUnit).

  • TRANSPOSE – When true compute Aᵀx (default false).

Parameters:

  • A – Triangular matrix (column-major, N*N elements; read-only).

  • x – In/out vector (length N); on return holds op(A) x.

  • scratch – Workspace of length N (see trmv_scratch_bytes).

template<typename T, uint32_t N, FillMode FILL = FillMode::Lower, Diag DIAG = Diag::NonUnit, bool TRANSPOSE = false>
void trsv(const T *A, T *x)#

Triangular solve on one thread: A x = b in place, compile-time size.

One thread solves the N×N triangular system by forward or back substitution (direction set by FILL and TRANSPOSE). A is column-major and read-only; x is overwritten with the solution. No shared scratch, no barriers, no threadIdx read; operands may be thread-local register arrays. SciPy equivalent: x = scipy.linalg.solve_triangular(A, b, lower=...).

Template Parameters:

  • T – Scalar type (e.g. float, double).

  • N – Dimension (A is N×N, x has length N).

  • FILL – Which triangle of A holds the data (default FillMode::Lower).

  • DIAG – Diag::Unit for an implicit unit diagonal (default Diag::NonUnit).

  • TRANSPOSE – When true solve Aᵀx = b (default false).

Parameters:

  • A – Triangular matrix (column-major, N*N elements; read-only).

  • x – In/out right-hand side; on return holds the solution.

namespace thread

Cooperative-groups variants (glass::cgrps::)#

Functions

template<typename T, bool TRANSPOSE = false, bool ROW_MAJOR = false, bool TRAILING_SYNC = true>
void gemv(uint32_t m, uint32_t n, T alpha, const T *A, const T *x, T beta, T *y, cgrps::thread_group g = cgrps::this_thread_block())#

Matrix-vector multiply: y = alpha * op(A) * x + beta * y (GEMV, cooperative-groups variant).

Runtime-size, single-block; thread rank/size come from the cooperative group. NumPy equivalent: y = alpha * A @ x + beta * y (or A.T @ x when TRANSPOSE).

Template Parameters:

  • T – Scalar type.

  • TRANSPOSE – If true, computes A^T * x.

  • ROW_MAJOR – Storage order of A (false = column-major).

Parameters:

  • m, n – A is m x n.

  • alpha – Scalar multiplier on the product.

  • A – Input matrix.

  • x – Input vector.

  • beta – Scalar multiplier on the existing y (read only when beta != 0).

  • y – In/out result vector.

  • g – Cooperative thread group (defaults to the whole block).

template<typename T, bool TRANSPOSE = false, bool ROW_MAJOR = false, bool TRAILING_SYNC = true>
void gemv(uint32_t m, uint32_t n, T alpha, const T *A, const T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#

GEMV with implicit beta = 0: y = alpha * op(A) * x (cooperative-groups variant).

Runtime-size overload that overwrites y (the existing y is not read). NumPy equivalent: y = alpha * A @ x (or A.T @ x when TRANSPOSE).

Template Parameters:

  • T – Scalar type.

  • TRANSPOSE – If true, computes A^T * x.

  • ROW_MAJOR – Storage order of A (false = column-major).

Parameters:

  • m, n – A is m x n.

  • alpha – Scalar multiplier on the product.

  • A – Input matrix.

  • x – Input vector.

  • y – Output result vector (overwritten).

  • g – Cooperative thread group (defaults to the whole block).

template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE = false, bool ROW_MAJOR = false, bool TRAILING_SYNC = true>
void gemv(T alpha, const T *A, const T *x, T beta, T *y, cgrps::thread_group g = cgrps::this_thread_block())#

Compile-time-size GEMV: y = alpha * op(A) * x + beta * y (cooperative-groups variant).

Dimensions baked in as template parameters. NumPy equivalent: y = alpha * A @ x + beta * y (or A.T @ x when TRANSPOSE).

Template Parameters:

  • T – Scalar type.

  • M, N – Compile-time dimensions (A is M x N).

  • TRANSPOSE – If true, computes A^T * x.

  • ROW_MAJOR – Storage order of A (false = column-major).

Parameters:

  • alpha – Scalar multiplier on the product.

  • A – Input matrix.

  • x – Input vector.

  • beta – Scalar multiplier on the existing y (read only when beta != 0).

  • y – In/out result vector.

  • g – Cooperative thread group (defaults to the whole block).

template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE = false, bool ROW_MAJOR = false, bool TRAILING_SYNC = true>
void gemv(T alpha, const T *A, const T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#

Compile-time-size GEMV with implicit beta = 0: y = alpha * op(A) * x.

Overwrites y (the existing y is not read). NumPy equivalent: y = alpha * A @ x (or A.T @ x when TRANSPOSE).

Template Parameters:

  • T – Scalar type.

  • M, N – Compile-time dimensions (A is M x N).

  • TRANSPOSE – If true, computes A^T * x.

  • ROW_MAJOR – Storage order of A (false = column-major).

Parameters:

  • alpha – Scalar multiplier on the product.

  • A – Input matrix.

  • x – Input vector.

  • y – Output result vector (overwritten).

  • g – Cooperative thread group (defaults to the whole block).

template<typename T, bool TRAILING_SYNC = true>
void ger(uint32_t m, uint32_t n, T alpha, const T *x, const T *y, T *A, cgrps::thread_group g = cgrps::this_thread_block())#

Rank-1 update: A += alpha * x * y^T (GER, cooperative-groups variant).

Adds the scaled outer product of x and y into the column-major matrix A. NumPy equivalent: A += alpha * np.outer(x, y).

Template Parameters:

T – Scalar type.

Parameters:

  • m, n – A is m x n (x length m, y length n).

  • alpha – Scalar multiplier on the outer product.

  • x, y – Input vectors.

  • A – In/out matrix (column-major).

  • g – Cooperative thread group (defaults to the whole block).

template<typename T, uint32_t M, uint32_t N, bool TRAILING_SYNC = true>
void ger(T alpha, const T *x, const T *y, T *A, cgrps::thread_group g = cgrps::this_thread_block())#

Compile-time-size rank-1 update: A += alpha * x * y^T (GER, cooperative-groups variant).

NumPy equivalent: A += alpha * np.outer(x, y).

Template Parameters:

  • T – Scalar type.

  • M, N – Compile-time dimensions (A is M x N; x length M, y length N).

Parameters:

  • alpha – Scalar multiplier on the outer product.

  • x, y – Input vectors.

  • A – In/out matrix (column-major).

  • g – Cooperative thread group (defaults to the whole block).

template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE = false, bool TRAILING_SYNC = true>
void gemv_reduced(T alpha, const T *A, const T *x, T beta, T *y, cgrps::thread_group g = cgrps::this_thread_block())#

Contraction-parallel GEMV: y = alpha * op(A) * x + beta * y (cooperative-groups variant).

Cooperative-groups form of glass::gemv_reduced. See it for semantics.

Template Parameters:

  • T, M, N, TRANSPOSE – See glass::gemv_reduced.

  • TRAILING_SYNC – Emit a trailing g.sync() (default true).

Parameters:

  • alpha, A, x, beta, y – See glass::gemv_reduced.

  • g – Cooperative thread group (defaults to the whole block; pass a warp-multiple group).

template<typename T, uint32_t M, uint32_t N, bool TRANSPOSE = false, bool TRAILING_SYNC = true>
void gemv_reduced(T alpha, const T *A, const T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#

Contraction-parallel GEMV with implicit beta = 0: y = alpha * op(A) * x (cooperative-groups variant).

Template Parameters:

T, M, N, TRANSPOSE, TRAILING_SYNC – See the beta overload.

Parameters:

alpha, A, x, y, g – See the beta overload.