L1 — Vector Operations#

Single-block vector primitives: scaled sums, copies, scaling, dot products, reductions, norms, and elementwise logic. Pure-SIMT (glass::block::, bare glass::) routines are listed first, followed by the cooperative-groups (glass::cgrps::) variants. Single-warp (glass::warp::) variants, where they exist, render inline beside their block-scoped sibling (e.g. under reduce); see Warp-scoped operations (glass::warp::).

axpy / axpby#

Functions

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

Scaled vector sum: y = alpha * x + y (AXPY).

Each thread in the block strides over the n elements. NumPy equivalent: y += alpha * x.

Template Parameters:

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

Parameters:

n – Number of elements.
alpha – Scalar multiplier.
x – Input vector of length n.
y – In/out vector of length n (overwritten with the result).

template<typename T, bool TRAILING_SYNC = true> void axpy(uint32_t n, T alpha, T *x, T *y, T *z)#

Scaled vector sum into a third buffer: z = alpha * x + y (AXPY).

Out-of-place variant that leaves x and y untouched. NumPy equivalent: z = alpha * x + y.

Template Parameters:

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

Parameters:

n – Number of elements.
alpha – Scalar multiplier.
x – Input vector of length n.
y – Input vector of length n.
z – Output vector of length n (overwritten with the result).

template<typename T, bool TRAILING_SYNC = true> void axpby(uint32_t n, T alpha, T *x, T beta, T *y, T *z)#

Doubly-scaled vector sum: z = alpha * x + beta * y (AXPBY).

Out-of-place generalization of AXPY with an independent scale on y. NumPy equivalent: z = alpha * x + beta * y.

Template Parameters:

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

Parameters:

n – Number of elements.
alpha – Scalar multiplier on x.
x – Input vector of length n.
beta – Scalar multiplier on y.
y – Input vector of length n.
z – Output vector of length n (overwritten with the result).

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

Scaled vector sum: y = alpha * x + y (AXPY), compile-time size.

Compile-time-N overload of AXPY. NumPy equivalent: y += alpha * x.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

alpha – Scalar multiplier.
x – Input vector of length N.
y – In/out vector of length N (overwritten with the result).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void axpy(T alpha, T *x, T *y, T *z)#

Scaled vector sum into a third buffer: z = alpha * x + y (AXPY), compile-time size.

Compile-time-N out-of-place overload. NumPy equivalent: z = alpha * x + y.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

alpha – Scalar multiplier.
x – Input vector of length N.
y – Input vector of length N.
z – Output vector of length N (overwritten with the result).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void axpby(T alpha, T *x, T beta, T *y, T *z)#

Doubly-scaled vector sum: z = alpha * x + beta * y (AXPBY), compile-time size.

Compile-time-N overload of AXPBY. NumPy equivalent: z = alpha * x + beta * y.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

alpha – Scalar multiplier on x.
x – Input vector of length N.
beta – Scalar multiplier on y.
y – Input vector of length N.
z – Output vector of length N (overwritten with the result).

template<typename T> void axpy(uint32_t n, T alpha, T *x, T *y)#

Scaled vector sum within one warp: y = alpha * x + y (AXPY), single-warp.

Scaled vector sum on one thread: y = alpha * x + y (AXPY), single-thread.

One 32-lane warp computes the update with lanes striding over the n elements. Elementwise, no inter-lane comms, no shared scratch, no __syncthreads. Independent warps may run distinct problems concurrently. Full 32 lanes required. NumPy equivalent: y += alpha * x.

One thread walks the n elements serially. No shared scratch, no shuffles, no barriers, no threadIdx read; operands may be thread-local register arrays. NumPy equivalent: y += alpha * x.

Template Parameters:

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

Parameters:

n – Number of elements.
alpha – Scalar multiplier.
x – Input vector of length n.
y – In/out vector of length n (overwritten with the result).
n – Number of elements.
alpha – Scalar multiplier.
x – Input vector of length n.
y – In/out vector of length n (overwritten with the result).

template<typename T, uint32_t N> void axpy(T alpha, T *x, T *y)#

Scaled vector sum within one warp: y = alpha * x + y (AXPY), single-warp, compile-time size.

Scaled vector sum on one thread: y = alpha * x + y (AXPY), single-thread, compile-time size.

Compile-time-N overload of the single-warp AXPY. Elementwise, no shared scratch, no __syncthreads. NumPy equivalent: y += alpha * x.

Compile-time-N overload; the trip count folds and the loop unrolls, so x and y may be thread-local register arrays. NumPy equivalent: y += alpha * x.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).
T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

alpha – Scalar multiplier.
x – Input vector of length N.
y – In/out vector of length N (overwritten with the result).
alpha – Scalar multiplier.
x – Input vector of length N.
y – In/out vector of length N (overwritten with the result).

namespace warp

namespace thread#

Row-strided AXPY: Y[r + c*Y_RS] += alpha * X[r + c*X_RS] over an M×N block.

The AXPY analogue of gemv_strided — adds an M×N (column-major) block of X into a same-shaped block of Y when the two live inside wider buffers with different leading dimensions (X_RS, Y_RS). PDDP packs a 14×14 update into a 21-lead buffer top-left (Y_RS=21, X_RS=14). Each (r,c) element is written by exactly one thread/lane, so there is no race and the op is trivially thread-count invariant. Block + warp::.

Functions

template<typename T, uint32_t M, uint32_t N, uint32_t Y_RS, uint32_t X_RS = M, bool TRAILING_SYNC = true> void axpy_strided(T alpha, const T *X, T *Y)#

Row-strided AXPY Y[r + c*Y_RS] += alpha * X[r + c*X_RS] over an M×N block.

Row-strided AXPY within one warp: Y[r + c*Y_RS] += alpha * X[r + c*X_RS].

Column-major; X is addressed at leading dimension X_RS (default M), Y at Y_RS. When X_RS == M and Y_RS == M this is a plain glass::axpy over the M*N contiguous elements. NumPy: Y[:M,:N] += alpha * X[:M,:N] (col-major lds).

Single-warp form of axpy_strided: one 32-lane warp strides over the M*N block elements. Each element written once; no inter-lane comms, no shared scratch. TRAILING_SYNC gates a closing __syncwarp(). Full 32 lanes required; independent warps may run distinct problems concurrently.

Template Parameters:

T – Scalar type (e.g. float, double).
M – Rows of the block.
N – Columns of the block.
Y_RS – Column-major leading dimension of Y.
X_RS – Column-major leading dimension of X (default M).
TRAILING_SYNC – Emit a trailing __syncthreads() (default true).
T – Scalar type.
M, N – Block shape.
Y_RS – Leading dimension of Y.
X_RS – Leading dimension of X (default M).
TRAILING_SYNC – Emit a trailing __syncwarp() (default true).

Parameters:

alpha – Scalar multiplier on X.
X – Input block, addressed at X[r + c*X_RS] (read-only).
Y – In/out block, addressed at Y[r + c*Y_RS].
alpha – Scalar multiplier on X.
X – Input block, addressed at X[r + c*X_RS] (read-only).
Y – In/out block, addressed at Y[r + c*Y_RS].

template<typename T, uint32_t M, uint32_t N, uint32_t Y_RS, uint32_t X_RS = M> void axpy_strided(T alpha, const T *X, T *Y)#

Row-strided AXPY on one thread: Y[r + c*Y_RS] += alpha * X[r + c*X_RS], single-thread.

One thread adds the column-major M×N block of X (leading dimension X_RS) into the same-shaped block of Y (leading dimension Y_RS), walking the elements serially. No shared scratch, no shuffles, no barriers, no threadIdx read; operands may be thread-local register arrays. NumPy: Y[:M,:N] += alpha * X[:M,:N] (col-major lds).

Template Parameters:

T – Scalar type (e.g. float, double).
M – Rows of the block.
N – Columns of the block.
Y_RS – Column-major leading dimension of Y.
X_RS – Column-major leading dimension of X (default M).

Parameters:

alpha – Scalar multiplier on X.
X – Input block, addressed at X[r + c*X_RS] (read-only).
Y – In/out block, addressed at Y[r + c*Y_RS].

namespace warp

namespace thread

copy#

Functions

template<typename T, bool TRAILING_SYNC = true> void copy(uint32_t n, T *x, T *y)#

Vector copy: y = x (COPY).

Each thread in the block strides over the n elements. NumPy equivalent: y = x.copy().

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n.
y – Output vector of length n (overwritten with a copy of x).

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

Scaled vector copy: y = alpha * x.

Copies x into y while scaling by alpha. NumPy equivalent: y = alpha * x.

Template Parameters:

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

Parameters:

n – Number of elements.
alpha – Scalar multiplier.
x – Input vector of length n.
y – Output vector of length n (overwritten with the result).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void copy(T *x, T *y)#

Vector copy: y = x (COPY), compile-time size.

Compile-time-N overload of copy. NumPy equivalent: y = x.copy().

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Input vector of length N.
y – Output vector of length N (overwritten with a copy of x).

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

Scaled vector copy: y = alpha * x, compile-time size.

Compile-time-N overload of scaled copy. NumPy equivalent: y = alpha * x.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

alpha – Scalar multiplier.
x – Input vector of length N.
y – Output vector of length N (overwritten with the result).

template<typename T> void copy(uint32_t n, T *x, T *y)#

Vector copy within one warp: y = x (COPY), single-warp.

Vector copy on one thread: y = x (COPY), single-thread.

One 32-lane warp copies the vector with lanes striding over the n elements. Elementwise, no inter-lane comms, no shared scratch, no __syncthreads. Independent warps may run distinct problems concurrently. Full 32 lanes required. NumPy equivalent: y = x.copy().

One thread walks the n elements serially. No shared scratch, no shuffles, no barriers, no threadIdx read; operands may be thread-local register arrays. NumPy equivalent: y = x.copy().

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n.
y – Output vector of length n (overwritten with a copy of x).
n – Number of elements.
x – Input vector of length n.
y – Output vector of length n (overwritten with a copy of x).

template<typename T, uint32_t N> void copy(T *x, T *y)#

Vector copy within one warp: y = x (COPY), single-warp, compile-time size.

Vector copy on one thread: y = x (COPY), single-thread, compile-time size.

Compile-time-N overload of the single-warp copy. Elementwise, no shared scratch, no __syncthreads. NumPy equivalent: y = x.copy().

Compile-time-N overload; the trip count folds and the loop unrolls, so x and y may be thread-local register arrays. NumPy equivalent: y = x.copy().

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).
T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Input vector of length N.
y – Output vector of length N (overwritten with a copy of x).
x – Input vector of length N.
y – Output vector of length N (overwritten with a copy of x).

namespace warp

namespace thread

Row-strided COPY: Y[r + c*Y_RS] = alpha * X[r + c*X_RS] over an M×N block.

The COPY (overwrite) sibling of axpy_strided — moves a column-major M×N block of X into a same-shaped block of Y when the two live inside wider buffers with different leading dimensions (X_RS, Y_RS), optionally scaling by alpha. Each (r,c) element is written by exactly one thread/lane (no race, trivially thread-count invariant). Block + warp::.

Functions

template<typename T, uint32_t M, uint32_t N, uint32_t Y_RS, uint32_t X_RS = M, bool TRAILING_SYNC = true> void copy_strided(T alpha, const T *X, T *Y)#

Row-strided COPY Y[r + c*Y_RS] = alpha * X[r + c*X_RS] over an M×N block.

Row-strided COPY within one warp: Y[r + c*Y_RS] = alpha * X[r + c*X_RS].

Column-major; X at leading dimension X_RS (default M), Y at Y_RS. When X_RS == M and Y_RS == M this is a plain scaled glass::copy over M*N contiguous elements. NumPy: Y[:M,:N] = alpha * X[:M,:N] (col-major lds).

Single-warp form of copy_strided: one 32-lane warp strides over the M*N block elements. Each element written once; no inter-lane comms, no shared scratch. TRAILING_SYNC gates a closing __syncwarp(). Full 32 lanes required; independent warps may run distinct problems concurrently.

Template Parameters:

T – Scalar type (e.g. float, double).
M – Rows of the block.
N – Columns of the block.
Y_RS – Column-major leading dimension of Y.
X_RS – Column-major leading dimension of X (default M).
TRAILING_SYNC – Emit a trailing __syncthreads() (default true).
T – Scalar type.
M, N – Block shape.
Y_RS – Leading dimension of Y.
X_RS – Leading dimension of X (default M).
TRAILING_SYNC – Emit a trailing __syncwarp() (default true).

Parameters:

alpha – Scalar multiplier on X.
X – Input block, addressed at X[r + c*X_RS] (read-only).
Y – Output block, addressed at Y[r + c*Y_RS] (overwritten).
alpha – Scalar multiplier on X.
X – Input block, addressed at X[r + c*X_RS] (read-only).
Y – Output block, addressed at Y[r + c*Y_RS] (overwritten).

template<typename T, uint32_t M, uint32_t N, uint32_t Y_RS, uint32_t X_RS = M> void copy_strided(T alpha, const T *X, T *Y)#

Row-strided COPY on one thread: Y[r + c*Y_RS] = alpha * X[r + c*X_RS], single-thread.

One thread moves the column-major M×N block of X (leading dimension X_RS) into the same-shaped block of Y (leading dimension Y_RS), scaling by alpha and walking the elements serially. No shared scratch, no shuffles, no barriers, no threadIdx read; operands may be thread-local register arrays. NumPy: Y[:M,:N] = alpha * X[:M,:N] (col-major lds).

Template Parameters:

T – Scalar type (e.g. float, double).
M – Rows of the block.
N – Columns of the block.
Y_RS – Column-major leading dimension of Y.
X_RS – Column-major leading dimension of X (default M).

Parameters:

alpha – Scalar multiplier on X.
X – Input block, addressed at X[r + c*X_RS] (read-only).
Y – Output block, addressed at Y[r + c*Y_RS] (overwritten).

namespace warp

namespace thread

scal#

Functions

template<typename T, bool TRAILING_SYNC = true> void scal(uint32_t n, T alpha, T *x)#

Scale a vector in place: x = alpha * x (SCAL).

NumPy equivalent: x *= alpha.

Template Parameters:

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

Parameters:

n – Number of elements.
alpha – Scalar multiplier.
x – In/out vector of length n (overwritten with the result).

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

Scale a vector into a second buffer: y = alpha * x (SCAL).

Out-of-place variant that leaves x untouched. NumPy equivalent: y = alpha * x.

Template Parameters:

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

Parameters:

n – Number of elements.
alpha – Scalar multiplier.
x – Input vector of length n.
y – Output vector of length n (overwritten with the result).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void scal(T alpha, T *x)#

Scale a vector in place: x = alpha * x (SCAL), compile-time size.

Compile-time-N overload. NumPy equivalent: x *= alpha.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

alpha – Scalar multiplier.
x – In/out vector of length N (overwritten with the result).

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

Scale a vector into a second buffer: y = alpha * x (SCAL), compile-time size.

Compile-time-N out-of-place overload. NumPy equivalent: y = alpha * x.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

alpha – Scalar multiplier.
x – Input vector of length N.
y – Output vector of length N (overwritten with the result).

template<typename T> void scal(uint32_t n, T alpha, T *x)#

Scale a vector in place within one warp: x = alpha * x (SCAL), single-warp.

Scale a vector in place on one thread: x = alpha * x (SCAL), single-thread.

One 32-lane warp scales the vector with lanes striding over the n elements. Elementwise, no inter-lane comms, no shared scratch, no __syncthreads. Independent warps may run distinct problems concurrently. Full 32 lanes required. NumPy equivalent: x *= alpha.

One thread walks the n elements serially. No shared scratch, no shuffles, no barriers, no threadIdx read; operands may be thread-local register arrays. NumPy equivalent: x *= alpha.

Template Parameters:

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

Parameters:

n – Number of elements.
alpha – Scalar multiplier.
x – In/out vector of length n (overwritten with the result).
n – Number of elements.
alpha – Scalar multiplier.
x – In/out vector of length n (overwritten with the result).

template<typename T, uint32_t N> void scal(T alpha, T *x)#

Scale a vector in place within one warp: x = alpha * x (SCAL), single-warp, compile-time size.

Scale a vector in place on one thread: x = alpha * x (SCAL), single-thread, compile-time size.

Compile-time-N overload of the single-warp scale. Elementwise, no shared scratch, no __syncthreads. NumPy equivalent: x *= alpha.

Compile-time-N overload; the trip count folds and the loop unrolls, so x may be a thread-local register array. NumPy equivalent: x *= alpha.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).
T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

alpha – Scalar multiplier.
x – In/out vector of length N (overwritten with the result).
alpha – Scalar multiplier.
x – In/out vector of length N (overwritten with the result).

namespace warp

namespace thread

swap#

Functions

template<typename T, bool TRAILING_SYNC = true> void swap(uint32_t n, T *x, T *y)#

Swap two vectors element-wise: x ↔ y (SWAP).

Exchanges the contents of x and y.

Template Parameters:

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

Parameters:

n – Number of elements.
x – In/out vector of length n.
y – In/out vector of length n.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void swap(T *x, T *y)#

Swap two vectors element-wise: x ↔ y (SWAP), compile-time size.

Compile-time-N overload of swap.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – In/out vector of length N.
y – In/out vector of length N.

rot / rotg (Givens plane rotation)#

Apply (rot, BLAS SROT/DROT: x = c*x + s*y, y = c*y - s*x; block + glass::warp::) and generate (rotg, BLAS SROTG/DROTG semantics, the overflow-safe scaled form; a __host__ __device__ scalar helper).

Givens plane rotation: apply (rot, BLAS SROT/DROT) and generate (rotg, BLAS SROTG/DROTG).

rot applies the rotation [c s; -s c] to a vector pair in place: x[i] = c*x[i] + s*y[i], y[i] = c*y[i] - s*x_old[i]. Each element index is owned by exactly one thread, which reads BOTH old values into registers before writing — no staging, no cross-thread hazard, trivially thread-count invariant. Block + warp::.

rotg is a __host__ __device__ SCALAR helper (one caller, no cooperation): given (a, b) it computes (c, s, r) with c*a + s*b = r and c*b - s*a = 0 (so applying rot with the returned (c, s) to the pair (a, b) zeroes b), using the overflow-safe scaled form of the reference BLAS DROTG.

Functions

template<typename T, bool TRAILING_SYNC = true> void rot(uint32_t n, T *x, T *y, T c, T s)#

Apply a Givens plane rotation to a vector pair in place (ROT).

x[i] = c*x[i] + s*y[i], y[i] = c*y[i] - s*x_old[i] for all i — the BLAS SROT/DROT update. NumPy equivalent: x, y = c*x + s*y, c*y - s*x.

Template Parameters:

T – Scalar type (e.g. float, double).
TRAILING_SYNC – Emit a trailing __syncthreads() (default true).

Parameters:

n – Number of elements.
x – In/out vector of length n.
y – In/out vector of length n.
c – Rotation cosine.
s – Rotation sine.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void rot(T *x, T *y, T c, T s)#

Apply a Givens plane rotation in place (ROT), compile-time size.

Compile-time-N overload. NumPy equivalent: x, y = c*x + s*y, c*y - s*x.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).
TRAILING_SYNC – Emit a trailing __syncthreads() (default true).

Parameters:

x – In/out vector of length N.
y – In/out vector of length N.
c – Rotation cosine.
s – Rotation sine.

template<typename T> void rotg(T a, T b, T &c, T &s, T &r)#

Generate a Givens plane rotation (ROTG) — __host__ __device__ scalar helper.

Given (a, b), computes (c, s, r) such that [c s; -s c] @ [a; b] = [r; 0] — i.e. c = a/r, s = b/r, r = ±sqrt(a² + b²) with the reference-BLAS DROTG sign convention (the sign of the larger-magnitude input) and its overflow-safe scaling (scale = |a| + |b|; the squares are formed on a/scale, b/scale). (a, b) = (0, 0) returns (c, s, r) = (1, 0, 0). Scalar: call from one thread (or redundantly from all — it is deterministic), or on the host. SciPy equivalent: (c, s), r = scipy.linalg.blas.drotg(a, b), np.hypot(a, b) (up to the sign convention).

Template Parameters:

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

Parameters:

a – First component.
b – Second component (the one the rotation annihilates).
c – Output rotation cosine.
s – Output rotation sine.
r – Output rotated magnitude (c*a + s*b).

template<typename T> void rot(uint32_t n, T *x, T *y, T c, T s)#

Apply a Givens plane rotation within one warp (ROT), single-warp.

Apply a Givens plane rotation on one thread (ROT), single-thread.

One 32-lane warp applies x[i] = c*x[i] + s*y[i], y[i] = c*y[i] - s*x_old[i] (lane-strided; each index owned by one lane, both olds read into registers first). Trailing __syncwarp(). NumPy equivalent: x, y = c*x + s*y, c*y - s*x.

One thread applies x[i] = c*x[i] + s*y[i], y[i] = c*y[i] - s*x_old[i] serially over the n elements (both olds read into registers first). No shared scratch, no shuffles, no barriers, no threadIdx read; operands may be thread-local register arrays. NumPy equivalent: x, y = c*x + s*y, c*y - s*x.

Template Parameters:

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

Parameters:

n – Number of elements.
x – In/out vector of length n.
y – In/out vector of length n.
c – Rotation cosine.
s – Rotation sine.
n – Number of elements.
x – In/out vector of length n.
y – In/out vector of length n.
c – Rotation cosine.
s – Rotation sine.

template<typename T, uint32_t N> void rot(T *x, T *y, T c, T s)#

Apply a Givens plane rotation within one warp (ROT), compile-time size.

Apply a Givens plane rotation on one thread (ROT), single-thread, compile-time size.

Compile-time-N overload of the single-warp rot. NumPy equivalent: x, y = c*x + s*y, c*y - s*x.

Compile-time-N overload; the trip count folds and the loop unrolls, so x and y may be thread-local register arrays. NumPy equivalent: x, y = c*x + s*y, c*y - s*x.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).
T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – In/out vector of length N.
y – In/out vector of length N.
c – Rotation cosine.
s – Rotation sine.
x – In/out vector of length N.
y – In/out vector of length N.
c – Rotation cosine.
s – Rotation sine.

namespace warp

namespace thread

dot#

Functions

template<typename T, bool TRAILING_SYNC = true> void dot(uint32_t n, T *x, T *y)#

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void dot(T *x, T *y)#

Inner product: y[0] = x · y (DOT), in-place, compile-time size.

Compile-time-N overload; the scalar result lands in y[0] (uses y as scratch). NumPy equivalent: np.dot(x, y).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Input vector of length N.
y – In/out vector of length N; the dot product lands in y[0].

template<typename T, bool TRAILING_SYNC = true> void dot_lowmem(uint32_t n, T *x, T *y, T *out)#

Inner product: out[0] = x · y (DOT), low-memory variant.

Writes the element-wise products into out, then thread 0 serially accumulates them into out[0], leaving x and y untouched. NumPy equivalent: np.dot_lowmem(x, y).

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n.
y – Input vector of length n.
out – Length-n scratch/output buffer; the result lands in out[0].

template<typename T> constexpr std::size_t dot_fast_scratch_bytes(uint32_t block_threads)#

Shared-scratch size in bytes for the dot_fast ops.

The warp-shuffle dot combines across warps through one scratch slot per warp: ceil(block_threads / 32) elements of T. Allocate dot_fast_scratch_bytes<T>(block_threads) for the s_scratch argument.

Template Parameters:: T – Scalar type.
Parameters:: block_threads – Number of threads in the launching block.
Returns:: Bytes to allocate for s_scratch.

template<typename T, bool TRAILING_SYNC = true> void dot_fast(uint32_t n, T *x, T *y, T *out, T *s_scratch)#

Inner product: out[0] = x · y (DOT), warp-shuffle variant.

Accumulates the element-wise products with a warp-shuffle reduction plus an inter-warp reduction through shared scratch, leaving x and y untouched. NumPy equivalent: np.dot_fast(x, y).

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n.
y – Input vector of length n.
out – Output buffer; the result lands in out[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void dot_fast(T *x, T *y, T *out, T *s_scratch)#

Inner product: out[0] = x · y (DOT), warp-shuffle, compile-time size.

Compile-time-N overload of the warp-shuffle dot product. NumPy equivalent: np.dot_fast(x, y).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Input vector of length N.
y – Input vector of length N.
out – Output buffer; the result lands in out[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T> T dot(uint32_t n, T *x, T *y)#

Inner product within one warp: returns x · y on every lane, single-warp.

One 32-lane warp forms the element-wise products and reduces them with __shfl_down_sync, then BROADCASTS the scalar total back to all 32 lanes via __shfl_sync (from a lane’s register, never a shared re-read — immune to the __restrict__ stale-cache miscompile). Inputs are left untouched; no shared scratch, no __syncthreads. Full 32 lanes required; independent warps may run distinct problems concurrently. NumPy equivalent: np.dot(x, y).

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n.
y – Input vector of length n.

Returns:

The inner product x · y, identical on every lane.

template<typename T, uint32_t N> T dot(T *x, T *y)#

Inner product within one warp: returns x · y on every lane, single-warp, compile-time size.

Compile-time-N overload of the single-warp dot. Reduces with __shfl_down_sync and broadcasts the total to all 32 lanes from a register. No shared scratch, no __syncthreads. NumPy equivalent: np.dot(x, y).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Input vector of length N.
y – Input vector of length N.

Returns:

The inner product x · y, identical on every lane.

template<typename T> T dot(uint32_t n, const T *x, const T *y)#

Inner product on one thread: returns x · y, single-thread.

Serially accumulates the element-wise products. Inputs untouched; no shared scratch, no shuffles, no barriers. NumPy equivalent: np.dot(x, y). Works well on low DOF problems with small vectors that fit into a register Parallelism can be leveraged by scaling thread count.

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n.
y – Input vector of length n.

Returns:

The inner product x · y.

template<typename T, uint32_t N> T dot(const T *x, const T *y)#

Inner product on one thread: returns x · y, single-thread, compile-time size.

Compile-time-N overload; the trip count folds and the loop unrolls, so x and y may be thread-local register arrays. NumPy equivalent: np.dot(x, y).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Input vector of length N.
y – Input vector of length N.

Returns:

The inner product x · y.

namespace warp

namespace thread

Functions

template<typename T, uint32_t N, uint32_t SX = 1, uint32_t SY = 1> T dot_strided(const T *x, const T *y)#

Per-thread strided inner product: Σ x[i*SX] * y[i*SY] (DOT, strided).

A single thread independently walks the full N-length product with compile-time strides, returning the scalar result. There is NO block-wide reduction — intended for use inside an already thread-parallel outer loop (e.g. GRiD generated kernels). With SX = SY = 1 this is a plain scalar dot; the inner loop is fully unrolled. NumPy equivalent: np.dot(x[::SX], y[::SY]).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of products to accumulate (compile-time constant).
SX – Element stride into x (compile-time, default 1).
SY – Element stride into y (compile-time, default 1).

Parameters:

x – Input vector, accessed at indices 0, SX, 2*SX, ….
y – Input vector, accessed at indices 0, SY, 2*SY, ….

Returns:

The strided inner product Σ x[i*SX] * y[i*SY].

template<typename T, uint32_t N, uint32_t SX = 1, uint32_t SY = 1> void dot_strided(const T *x, const T *y, T *out)#

Per-thread strided inner product, store-to-pointer overload.

Same per-thread strided dot as the value-returning overload, but writes the scalar result to *out. NumPy equivalent: out[0] = np.dot(x[::SX], y[::SY]).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of products to accumulate (compile-time constant).
SX – Element stride into x (compile-time, default 1).
SY – Element stride into y (compile-time, default 1).

Parameters:

x – Input vector, accessed at indices 0, SX, 2*SX, ….
y – Input vector, accessed at indices 0, SY, 2*SY, ….
out – Destination for the scalar result.

Functions

template<typename T, uint32_t N, uint32_t SX = 1, uint32_t SY = 1, bool TRAILING_SYNC = true> void dot_strided_coalesced(const T *x, const T *y, T *out, T *s_scratch)#

Block-cooperative strided inner product with coalesced loads (DOT, strided).

Computes the same value as dot_strided<T,N,SX,SY> (Σ x[i*SX] * y[i*SY]), but the WHOLE block cooperates on one dot product: the i-loop is distributed across threads with a block stride so consecutive ranks touch addresses SX apart, turning a per-warp strided gather into a coalesced burst when SX is the inner contiguous axis. Partial sums are combined via warp-shuffle plus a shared-scratch block reduction; the scalar result is broadcast to *out (visible to all threads after the trailing barrier). Requires a block-wide, __syncthreads-safe launch — use dot_strided instead in a per-thread context. NumPy equivalent: np.dot(x[::SX], y[::SY]).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of products to accumulate (compile-time constant).
SX – Element stride into x (compile-time, default 1).
SY – Element stride into y (compile-time, default 1).

Parameters:

x – Input vector, accessed at indices 0, SX, 2*SX, ….
y – Input vector, accessed at indices 0, SY, 2*SY, ….
out – Destination for the scalar result (broadcast to all threads).
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

reduce#

Functions

template<typename T, bool TRAILING_SYNC = true> void reduce(uint32_t n, T *x)#

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void reduce(T *x)#

Sum reduction: x[0] = Σ x[i] (in-place), compile-time size.

Compile-time-N overload of the halving reduce. NumPy equivalent: np.sum(x).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – In/out vector of length N; the sum lands in x[0].

template<typename T, bool TRAILING_SYNC = true> void reduce_lowmem(uint32_t n, T *x)#

Sum reduction: x[0] = Σ x[i] (in-place), low-memory variant.

Thread 0 serially accumulates all elements into x[0]; uses no scratch. NumPy equivalent: np.sum(x).

Template Parameters:

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

Parameters:

n – Number of elements.
x – In/out vector of length n; the sum lands in x[0].

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void reduce_lowmem(T *x)#

Sum reduction: x[0] = Σ x[i] (in-place), low-memory, compile-time size.

Compile-time-N overload; thread 0 serially accumulates into x[0]. NumPy equivalent: np.sum(x).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – In/out vector of length N; the sum lands in x[0].

template<typename T> constexpr std::size_t reduce_fast_scratch_bytes(uint32_t block_threads)#

Shared-scratch size in bytes for the reduce_fast ops.

The warp-shuffle reduce combines across warps through one scratch slot per warp: ceil(block_threads / 32) elements of T. Allocate reduce_fast_scratch_bytes<T>(block_threads) for the s_scratch argument.

Template Parameters:: T – Scalar type.
Parameters:: block_threads – Number of threads in the launching block.
Returns:: Bytes to allocate for s_scratch.

template<typename T, bool TRAILING_SYNC = true> void reduce_fast(uint32_t n, T *x, T *s_scratch)#

Sum reduction: x[0] = Σ x[i] (in-place), warp-shuffle variant.

Accumulates with a warp-shuffle reduction plus an inter-warp reduction through shared scratch; the total lands in x[0]. NumPy equivalent: np.sum(x).

Template Parameters:

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

Parameters:

n – Number of elements.
x – In/out vector of length n; the sum lands in x[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void reduce_fast(T *x, T *s_scratch)#

Sum reduction: x[0] = Σ x[i] (in-place), warp-shuffle, compile-time size.

Compile-time-N overload of the warp-shuffle reduce. NumPy equivalent: np.sum(x).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – In/out vector of length N; the sum lands in x[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T, bool TRAILING_SYNC = true> T reduce_fast(T partial, T *s_scratch)#

Block-sum of a per-thread register value: returns Σ partial.

Reduces one PER-THREAD contribution partial (one per thread) across the block and returns the total to EVERY thread, with no intermediate x[] buffer. This is the entry point for fused “compute-a-partial-then-sum” patterns (e.g. cost/barrier kernels). The result is also broadcast through s_scratch[0]; the routine ends on a __syncthreads(), so s_scratch is safe to reuse afterwards. Threads with no contribution should pass 0.

Template Parameters:

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

Parameters:

partial – This thread’s contribution to the block sum.
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim); on return s_scratch[0] holds the total.

Returns:

The block-wide total Σ partial, identical on every thread.

template<typename T, bool TRAILING_SYNC = true> T reduce_fast_min(T partial, T *s_scratch)#

Block-min of a per-thread register value: returns min partial.

Min twin of the reduce_fast(partial, s_scratch) sum overload above, with the same shape: one PER-THREAD contribution in, the block-wide minimum returned to EVERY thread, no x[]

buffer. The entry point for fused “compute-a-partial-then

-min” patterns (e.g. a softmax’s max-shift, or a smooth-min collision cost). Threads with no contribution should pass the identity — a large sentinel such as

1e30f, NOT zero.

Kept separate from the sum overload rather than folded into a generic reduce_fast(partial, s_scratch, Op): an operator template would force every existing call site through a functor, and the two identities (0 vs +inf) cannot share a default.

Template Parameters:

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

Parameters:

partial – This thread’s contribution to the block min.
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim); on return s_scratch[0] holds the minimum.

Returns:

The block-wide minimum, identical on every thread.

template<typename T> T reduced_tree32(T p[32])#

Fixed 32-way pairwise tree reduce of a register array (returns the sum in p[0]).

Combines 32 partials in the EXACT pairwise grouping that glass::warp::reduce produces in lane 0 (a __shfl_down_sync tree with offsets 16,8,4,2,1), so a serial caller (e.g. a sub-warp fallback below 32 threads) matches the full-warp caller bit-for-bit. This is the shared primitive that lets the contraction-parallel *_reduced engines stay thread-count invariant across the 32-thread boundary. NumPy equivalent: np.sum(p) (different rounding).

Template Parameters:: T – Scalar type.
Parameters:: p – In/out array of 32 partials; on return p[0] holds the total (p is clobbered).
Returns:: The sum of p[0..31].

template<typename T> void reduce(uint32_t n, T *x)#

Sum reduction within one warp: x[0] = Σ x[i] (in-place), single-warp.

Sum reduction on one thread: x[0] = Σ x[i], single-thread.

One 32-lane warp sums the vector with __shfl_down_sync; the total lands in x[0] (input overwritten). No shared scratch, no inter-warp combine. NumPy equivalent: np.sum(x).

Serially accumulates the elements and writes the total to x[0]. Only x[0] is written — the tail is left untouched (unlike the block halving reduce, which uses x as scratch). No shared scratch, no shuffles, no barriers, no threadIdx read. NumPy equivalent: np.sum(x).

Template Parameters:

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

Parameters:

n – Number of elements.
x – In/out vector of length n; the sum lands in x[0].
n – Number of elements.
x – In/out vector of length n; the sum lands in x[0].

template<typename T, uint32_t N> void reduce(T *x)#

Sum reduction within one warp: x[0] = Σ x[i] (in-place), single-warp, compile-time size.

Sum reduction on one thread: x[0] = Σ x[i], single-thread, compile-time size.

Compile-time-N overload. NumPy equivalent: np.sum(x).

Compile-time-N overload; the trip count folds and the loop unrolls, so x may be a thread-local register array. Only x[0] is written. NumPy equivalent: np.sum(x).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).
T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – In/out vector of length N; the sum lands in x[0].
x – In/out vector of length N; the sum lands in x[0].

template<typename T> T reduce(T partial)#

Warp-sum of a per-lane register value: returns Σ partial on every lane.

Sum of a per-thread register value on one thread: returns partial.

Reduces one PER-LANE contribution across a single warp and broadcasts the total back to all 32 lanes — no x[] buffer, no shared scratch. The entry point for fused “compute-a-partial-then-sum” patterns inside a warp (e.g. row-norm / residual accumulation). Inactive lanes should pass 0.

The single-thread mirror of warp::reduce(partial): with one problem per thread there is exactly ONE contribution, so the “reduction” is the identity. Exists so tier-generic call sites (templated on the interface) compile against every tier with the same shape. No barriers, no shuffles, no threadIdx read. NumPy equivalent: np.sum([partial]).

Template Parameters:

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

Parameters:

partial – This lane’s contribution.
partial – This thread’s (sole) contribution.

Returns:

The warp-wide total Σ partial, identical on every lane.

Returns:

partial, unchanged.

template<typename T> T reduce_min(T partial)#

Warp-min of a per-lane register value: returns min(partial) on every lane.

The min twin of the register warp::reduce (same shuffle ladder + broadcast; no buffer, no scratch). Inactive/idle lanes must pass the identity (e.g. +INFINITY — there is no empty-lane sentinel here; use argmin_pair when a lane can be excluded by index). NaN candidates lose every compare, so a NaN lane never wins unless ALL lanes are NaN. Full 32-lane warp required.

Template Parameters:: T – Scalar type (e.g. float, double).
Parameters:: partial – This lane’s contribution.
Returns:: The warp-wide minimum, identical on every lane.

template<typename T> T reduce_max(T partial)#

Warp-max of a per-lane register value: returns max(partial) on every lane.

See reduce_min (pass -INFINITY from idle lanes).

Template Parameters:: T – Scalar type.
Parameters:: partial – This lane’s contribution.
Returns:: The warp-wide maximum, identical on every lane.

namespace shfl_detail#

Functions

template<typename T> T fold_sum(T v)#

template<typename T, typename Op> T fold(T v, Op op)#

template<typename T> T butterfly_sum(T v)#

template<typename T, typename Op> T butterfly(T v, Op op)#

struct MaxOp#

Public Functions

template<typename T> inline T operator()(T a, T b) const#

struct MinOp#

Public Functions

template<typename T> inline T operator()(T a, T b) const#

namespace warp

namespace thread

prefix_sum#

Functions

template<typename T, bool TRAILING_SYNC = true> void prefix_sum_exclusive(T *s_input, T *s_output, int n)#

Exclusive prefix sum (scan): s_output[i] = Σ_{j<i} s_input[j].

Block-wide Hillis-Steele scan; s_output[0] is 0 and each subsequent entry is the running total of all strictly-earlier inputs. NumPy equivalent: s_output = np.concatenate([[0], np.cumsum(s_input)[:-1]]).

Template Parameters:

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

Parameters:

s_input – Input vector of length n (in shared memory).
s_output – Output scan buffer of length n (in shared memory).
n – Number of elements (must not exceed blockDim.x).

template<typename T, bool TRAILING_SYNC = true> void prefix_sum_inclusive(T *s_input, T *s_output, int n)#

Inclusive prefix sum (scan): s_output[i] = Σ_{j<=i} s_input[j].

Block-wide Hillis-Steele scan; each entry is the running total of all inputs up to and including that index. NumPy equivalent: s_output = np.cumsum(s_input).

Template Parameters:

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

Parameters:

s_input – Input vector of length n (in shared memory).
s_output – Output scan buffer of length n (in shared memory).
n – Number of elements (must not exceed blockDim.x).

Norms#

Functions

template<typename T, bool TRAILING_SYNC = true> void vector_norm(uint32_t N, T *a, T *out)#

Euclidean (L2) norm into a separate buffer: out[0] = ‖a‖₂ (default).

Default halving-tree variant, completing the reduction family’s bare surface (mirrors how the plain reduce/dot/nrm2 engines are built): the block writes the per-element squares into out, tree-reduces them with the shared reduce machinery, and thread 0 takes the square root. Non-destructive — a is left untouched; out doubles as the length-N scratch (no shared scratch needed, unlike vector_norm_fast). Thread-count invariant: the pairwise tree gives byte-identical output at any block size. NumPy equivalent: np.linalg.norm(a).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Template Parameters:

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

Parameters:

N – Number of elements.
a – Input vector of length N (read-only).
out – Length-N scratch/output buffer; the result lands in out[0].

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void vector_norm(T *a, T *out)#

Euclidean (L2) norm into a separate buffer: out[0] = ‖a‖₂, compile-time size.

Compile-time-N overload of the default halving-tree vector norm; leaves a untouched. NumPy equivalent: np.linalg.norm(a).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

a – Input vector of length N (read-only).
out – Length-N scratch/output buffer; the result lands in out[0].

template<typename T, bool TRAILING_SYNC = true> void vector_norm_lowmem(uint32_t N, T *a, T *out)#

Euclidean (L2) norm into a separate buffer: out[0] = ‖a‖₂, low-memory variant.

Writes the per-element squares into out, then thread 0 serially sums them and takes the square root, leaving a untouched. NumPy equivalent: np.linalg.norm(a).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Template Parameters:

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

Parameters:

N – Number of elements.
a – Input vector of length N.
out – Length-N scratch/output buffer; the result lands in out[0].

template<typename T, bool TRAILING_SYNC = true> void vector_norm_fast(uint32_t N, T *a, T *out, T *s_scratch)#

Euclidean (L2) norm into a separate buffer: out[0] = ‖a‖₂, warp-shuffle variant.

Accumulates the sum of squares with a warp-shuffle reduction plus an inter-warp reduction through shared scratch, then takes the square root, leaving a untouched. NumPy equivalent: np.linalg.norm(a).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Template Parameters:

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

Parameters:

N – Number of elements.
a – Input vector of length N.
out – Output buffer; the result lands in out[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void vector_norm_fast(T *a, T *out, T *s_scratch)#

Euclidean (L2) norm into a separate buffer: out[0] = ‖a‖₂, warp-shuffle, compile-time size.

Compile-time-N overload of the warp-shuffle vector norm; leaves a untouched. NumPy equivalent: np.linalg.norm(a).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

a – Input vector of length N.
out – Output buffer; the result lands in out[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

Functions

template<typename T, bool TRAILING_SYNC = true> void nrm2_lowmem(uint32_t n, T *x)#

Euclidean (L2) norm: x[0] = ‖x‖₂ (in-place, destructive), low-memory variant.

Squares each element in place, then thread 0 serially sums them and takes the square root, leaving the result in x[0] (the input is overwritten). NumPy equivalent: np.linalg.norm(x).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Template Parameters:

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

Parameters:

n – Number of elements.
x – In/out vector of length n; the result lands in x[0].

template<typename T, bool TRAILING_SYNC = true> void nrm2_fast(uint32_t n, T *x, T *s_scratch)#

Euclidean (L2) norm: x[0] = ‖x‖₂ (in-place), warp-shuffle variant.

Accumulates the sum of squares with a warp-shuffle reduction plus an inter-warp reduction through shared scratch, then takes the square root; the result lands in x[0]. NumPy equivalent: np.linalg.norm(x).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Template Parameters:

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

Parameters:

n – Number of elements.
x – In/out vector of length n; the result lands in x[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void nrm2_fast(T *x, T *s_scratch)#

Euclidean (L2) norm: x[0] = ‖x‖₂, warp-shuffle, compile-time size.

Compile-time-N overload of the warp-shuffle L2 norm. NumPy equivalent: np.linalg.norm(x).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – In/out vector of length N; the result lands in x[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T> T nrm2(uint32_t n, const T *x)#

Euclidean (L2) norm within one warp: returns ‖x‖₂ on every lane.

Euclidean (L2) norm on one thread: returns ‖x‖₂, single-thread.

Single-warp L2 norm, mirroring warp::dot: one 32-lane warp reduces the per-lane sum-of-squares with __shfl_down_sync, BROADCASTS the total to all 32 lanes via __shfl_sync (from a register — immune to the __restrict__ stale-cache miscompile), and each lane takes the square root. Non-destructive (x untouched), no shared scratch, no __syncthreads; uses type-generic sqrt, like every tier (the block forms’ float-only sqrtf was fixed 2026-07-17). Full 32 lanes required. NumPy equivalent: np.linalg.norm(x).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Serially accumulates the sum of squares and returns its square root, mirroring warp::nrm2 (value-returning, non-destructive — unlike the destructive block _fast/_lowmem forms; one thread has no reduction strategy, so the tier ships no such twins). Uses type-generic sqrt (correct for double). Inputs untouched; no shared scratch, no shuffles, no barriers, no threadIdx read; operands may be thread-local register arrays. NumPy equivalent: np.linalg.norm(x).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n (read-only).
n – Number of elements.
x – Input vector of length n (read-only).

Returns:

‖x‖₂, identical on every lane.

Returns:

‖x‖₂.

template<typename T, uint32_t N> T nrm2(const T *x)#

Euclidean (L2) norm within one warp, compile-time size (returns it on every lane).

Euclidean (L2) norm on one thread, compile-time size: returns ‖x‖₂.

Compile-time-N overload of the single-warp L2 norm. Non-destructive, no shared scratch, no __syncthreads; type-generic sqrt. NumPy equivalent: np.linalg.norm(x).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Compile-time-N overload; the trip count folds and the loop unrolls, so x may be a thread-local register array. Non-destructive; type-generic sqrt. NumPy equivalent: np.linalg.norm(x).

Range contract: naive sum of squares (no LAPACK-style snrm2 scaling), so the intermediate sum overflows to inf for ‖x‖ ≳ 1e19 (f32) / ‖x‖ ≳ 1e154 (f64) and loses tiny components to underflow.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).
T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

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

Returns:

‖x‖₂, identical on every lane.

Returns:

‖x‖₂.

namespace warp

namespace thread

1-norm of a difference: ‖x − y‖₁ = Σ|x[i] − y[i]|.

The asum sibling for a difference — the residual/convergence check GATO’s Schur loop wants without materializing x − y first. Inputs are read-only (non-destructive). Block forms mirror asum.cuh (low_memory serial-tail and high_speed warp-shuffle); warp::nrm1_diff mirrors warp::dot (returns the scalar broadcast to every lane). Thread-count invariant.

Functions

template<typename T, bool TRAILING_SYNC = true> void nrm1_diff_lowmem(uint32_t n, const T *x, const T *y, T *out)#

out[0] = Σ|x[i] − y[i]| (‖x−y‖₁), low-memory variant.

Writes the per-element absolute differences into out, then thread 0 serially accumulates them into out[0]. x/y are untouched. NumPy equivalent: np.sum(np.abs(x - y)).

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n (read-only).
y – Input vector of length n (read-only).
out – Length-n scratch/output buffer; the result lands in out[0].

template<typename T, bool TRAILING_SYNC = true> void nrm1_diff_fast(uint32_t n, const T *x, const T *y, T *out, T *s_scratch)#

out[0] = Σ|x[i] − y[i]| (‖x−y‖₁), warp-shuffle variant.

Reduces the absolute differences with a warp-shuffle reduction plus an inter-warp reduction through shared scratch; the result lands in out[0]. x/y are untouched. NumPy equivalent: np.sum(np.abs(x - y)).

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n (read-only).
y – Input vector of length n (read-only).
out – Output; the result lands in out[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void nrm1_diff_fast(const T *x, const T *y, T *out, T *s_scratch)#

out[0] = Σ|x[i] − y[i]| (‖x−y‖₁), warp-shuffle, compile-time size.

Compile-time-N overload of the warp-shuffle nrm1_diff. x/y untouched. NumPy equivalent: np.sum(np.abs(x - y)).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Input vector of length N (read-only).
y – Input vector of length N (read-only).
out – Output; the result lands in out[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T> T nrm1_diff(uint32_t n, const T *x, const T *y)#

‖x − y‖₁ within one warp: returns Σ|x[i] − y[i]| on every lane.

‖x − y‖₁ on one thread: returns Σ|x[i] − y[i]|, single-thread.

One 32-lane warp forms the absolute differences and reduces with __shfl_down_sync, then BROADCASTS the scalar total back to all 32 lanes via __shfl_sync (from a register, never a shared re-read — immune to the __restrict__ stale-cache miscompile). x/y untouched; no shared scratch, no __syncthreads. Full 32 lanes required; independent warps may run distinct problems concurrently. NumPy equivalent: np.sum(np.abs(x - y)).

Serially accumulates the absolute differences, mirroring warp::nrm1_diff (value-returning, non-destructive; one thread has no reduction strategy, so the tier ships no _fast/_lowmem twins). x/y untouched; no shared scratch, no shuffles, no barriers, no threadIdx read; operands may be thread-local register arrays. NumPy equivalent: np.sum(np.abs(x - y)).

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n (read-only).
y – Input vector of length n (read-only).
n – Number of elements.
x – Input vector of length n (read-only).
y – Input vector of length n (read-only).

Returns:

‖x − y‖₁, identical on every lane.

Returns:

‖x − y‖₁.

template<typename T, uint32_t N> T nrm1_diff(const T *x, const T *y)#

‖x − y‖₁ within one warp, compile-time size (returns it on every lane).

‖x − y‖₁ on one thread, compile-time size: returns Σ|x[i] − y[i]|.

Compile-time-N overload of the single-warp nrm1_diff. No shared scratch, no __syncthreads. NumPy equivalent: np.sum(np.abs(x - y)).

Compile-time-N overload; the trip count folds and the loop unrolls, so x and y may be thread-local register arrays. Non-destructive. NumPy equivalent: np.sum(np.abs(x - y)).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).
T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Input vector of length N (read-only).
y – Input vector of length N (read-only).
x – Input vector of length N (read-only).
y – Input vector of length N (read-only).

Returns:

‖x − y‖₁, identical on every lane.

Returns:

‖x − y‖₁.

namespace warp

namespace thread

Functions

template<typename T, bool TRAILING_SYNC = true> void infnorm(uint32_t n, T *x)#

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void infnorm(T *x)#

Infinity norm: x[0] = ‖x‖∞ = max|x[i]|, compile-time size.

Compile-time-N overload (in-place, destructive). NumPy equivalent: np.max(np.abs(x)).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – In/out vector of length N; the result lands in x[0].

Functions

template<typename T, bool TRAILING_SYNC = true> void asum_lowmem(uint32_t n, T *x, T *out)#

Sum of absolute values: out[0] = Σ|x[i]| (ASUM), low-memory variant.

Writes the per-element absolute values into out, then thread 0 serially accumulates them into out[0]. NumPy equivalent: np.sum(np.abs(x)).

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n.
out – Length-n scratch/output buffer; the result lands in out[0].

template<typename T, bool TRAILING_SYNC = true> void asum_fast(uint32_t n, T *x, T *s_scratch)#

Sum of absolute values: x[0] = Σ|x[i]| (ASUM), warp-shuffle variant.

Computes the absolute-value sum with a warp-shuffle reduction plus an inter-warp reduction through shared scratch. The result is written to x[0] (destructive — overwrites the input). NumPy equivalent: np.sum(np.abs(x)).

Template Parameters:

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

Parameters:

n – Number of elements.
x – In/out vector of length n; result lands in x[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void asum_fast(T *x, T *s_scratch)#

Sum of absolute values: x[0] = Σ|x[i]| (ASUM), compile-time size.

Compile-time-N overload of the warp-shuffle ASUM; the result is written to x[0] (destructive). NumPy equivalent: np.sum(np.abs(x)).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – In/out vector of length N; result lands in x[0].
s_scratch – Shared scratch of ceil(blockDim/32) elements (one per warp) — size with reduce_fast_scratch_bytes<T>(blockDim).

template<typename T> T asum(uint32_t n, const T *x)#

Sum of absolute values within one warp: returns Σ|x[i]| on every lane.

Sum of absolute values on one thread: returns Σ|x[i]| (ASUM), single-thread.

Single-warp ASUM, mirroring warp::dot: one 32-lane warp reduces the per-lane absolute-value partials with __shfl_down_sync and BROADCASTS the total to all 32 lanes via __shfl_sync (from a register — immune to the __restrict__ stale-cache miscompile). Non-destructive (x untouched), no shared scratch, no __syncthreads. Full 32 lanes required; independent warps may run distinct problems concurrently. NumPy equivalent: np.sum(np.abs(x)).

Serially accumulates the absolute values, mirroring warp::asum (value-returning, non-destructive — unlike the block _fast/_lowmem forms; one thread has no reduction strategy, so the tier ships no such twins). Inputs untouched; no shared scratch, no shuffles, no barriers, no threadIdx read; operands may be thread-local register arrays. NumPy equivalent: np.sum(np.abs(x)).

Template Parameters:

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

Parameters:

n – Number of elements.
x – Input vector of length n (read-only).
n – Number of elements.
x – Input vector of length n (read-only).

Returns:

Σ|x[i]|, identical on every lane.

Returns:

Σ|x[i]|.

template<typename T, uint32_t N> T asum(const T *x)#

Sum of absolute values within one warp, compile-time size (returns it on every lane).

Sum of absolute values on one thread, compile-time size: returns Σ|x[i]|.

Compile-time-N overload of the single-warp ASUM. Non-destructive, no shared scratch, no __syncthreads. NumPy equivalent: np.sum(np.abs(x)).

Compile-time-N overload; the trip count folds and the loop unrolls, so x may be a thread-local register array. Non-destructive. NumPy equivalent: np.sum(np.abs(x)).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).
T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

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

Returns:

Σ|x[i]|, identical on every lane.

Returns:

Σ|x[i]|.

namespace warp

namespace thread

Index of max-abs (iamax)#

Functions

template<typename T, bool TRAILING_SYNC = true> void iamax(uint32_t n, const T *x, uint32_t *out, T *s_scratch)#

Index of the max-absolute-value element (BLAS i_amax), into out[0].

Non-destructive: x is read-only and never clobbered. Computes the block-wide argmax over |x| (default threadIdx-strided variant) and writes the winning index (uint32_t) to out[0]. NumPy equivalent: int(np.argmax(np.abs(x))).

The routine ends on a trailing __syncthreads(), so out[0] is block-visible on return.

Tie-break: On EQUAL absolute value the LOWER index wins (the BLAS rule). This tie-break is applied at every combine step, which is what makes the result identical for any block size (1 thread, a partial warp, or many warps).
NaN policy: NaN inputs are SKIPPED (IEEE compares are false, so a NaN is never selected). This DIVERGES from np.argmax(np.abs(x)), which propagates NaN; oracle tests must exclude NaN inputs. An all-zero vector returns index 0.
Scratch sizing: This variant uses no shared scratch (a serial pass on thread 0). For the warp-shuffle iamax_fast variant size scratch via iamax_fast_scratch_bytes.

Template Parameters:

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

Parameters:

n – Number of elements.
x – Read-only input vector of length n (not modified).
out – Output: out[0] receives the argmax index.
s_scratch – Shared scratch of iamax_scratch_bytes elements (key + index lanes).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void iamax(const T *x, uint32_t *out, T *s_scratch)#

Index of the max-absolute-value element (BLAS i_amax), compile-time size.

Compile-time-N overload of the default i_amax; non-destructive, writes the winning index to out[0]. NumPy equivalent: int(np.argmax(np.abs(x))). Tie-break (lower index wins on equal |x|), NaN-skip policy, and the trailing __syncthreads() are as in the runtime-n overload.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Read-only input vector of length N (not modified).
out – Output: out[0] receives the argmax index.
s_scratch – Shared scratch of iamax_scratch_bytes elements.

template<typename T, bool TRAILING_SYNC = true> void iamax(uint32_t n, const T *x, uint32_t *out, T *out_val, T *s_scratch)#

i_amax with the max absolute value also returned: out_val[0] = max|x|.

Like iamax but additionally writes the winning absolute value to out_val[0] (max|x|, NumPy np.max(np.abs(x)) over the non-NaN entries). Non-destructive; lower-index tie-break; NaN skipped; trailing __syncthreads().

Template Parameters:

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

Parameters:

n – Number of elements.
x – Read-only input vector of length n (not modified).
out – Output: out[0] receives the argmax index.
out_val – Output: out_val[0] receives max|x|.
s_scratch – Shared scratch of iamax_scratch_bytes elements.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void iamax(const T *x, uint32_t *out, T *out_val, T *s_scratch)#

i_amax with max|x| returned, compile-time size.

Compile-time-N overload of the value-returning i_amax. NumPy equivalents: out[0] = int(np.argmax(np.abs(x))), out_val[0] = np.max(np.abs(x)).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Read-only input vector of length N (not modified).
out – Output: out[0] receives the argmax index.
out_val – Output: out_val[0] receives max|x|.
s_scratch – Shared scratch of iamax_scratch_bytes elements.

template<typename T> constexpr std::size_t iamax_scratch_bytes(uint32_t block_threads)#

Shared-scratch size in bytes for the default/low_memory iamax.

The default variant stores one abs-key (T) and one index (uint32_t) per thread. Allocate iamax_scratch_bytes<T>(block_threads) bytes of T for s_scratch (the index lanes are packed into the same buffer after the keys).

Template Parameters:: T – Scalar type.
Parameters:: block_threads – Number of threads in the launching block.
Returns:: Bytes to allocate for s_scratch.

template<typename T> constexpr std::size_t iamax_fast_scratch_bytes(uint32_t block_threads)#

Shared-scratch size in bytes for iamax_fast.

The warp-shuffle variant reduces within each warp in registers and combines across warps through scratch, needing one (key,index) slot per warp: ceil(block_threads/32) of each. Allocate iamax_fast_scratch_bytes<T>(block_threads) elements of T for its s_scratch.

Template Parameters:: T – Scalar type.
Parameters:: block_threads – Number of threads in the launching block.
Returns:: Bytes to allocate for the iamax_fast scratch.

template<typename T, bool TRAILING_SYNC = true> void iamax_lowmem(uint32_t n, const T *x, uint32_t *out)#

i_amax, low-memory variant (no scratch).

Thread 0 serially scans x for the argmax over |x|, writing the index to out[0]; all other threads idle. Non-destructive. NumPy equivalent: int(np.argmax(np.abs(x))). Lower-index tie-break on equal |x|; NaN skipped (diverges from np.argmax, exclude NaN in tests); all-zero → 0. Ends on a trailing __syncthreads() so out[0] is block-visible.

Template Parameters:

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

Parameters:

n – Number of elements.
x – Read-only input vector of length n (not modified).
out – Output: out[0] receives the argmax index.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void iamax_lowmem(const T *x, uint32_t *out)#

i_amax, low-memory variant, compile-time size.

Compile-time-N overload of the serial i_amax. NumPy equivalent: int(np.argmax(np.abs(x))). Same tie-break / NaN policy as the runtime form.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Read-only input vector of length N (not modified).
out – Output: out[0] receives the argmax index.

template<typename T, bool TRAILING_SYNC = true> void iamax_lowmem(uint32_t n, const T *x, uint32_t *out, T *out_val)#

i_amax + max|x|, low-memory variant (no scratch).

As above but also writes out_val[0] = max|x|. Non-destructive; thread 0 scans serially; lower-index tie-break; NaN skipped; all-zero → 0.

Template Parameters:

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

Parameters:

n – Number of elements.
x – Read-only input vector of length n (not modified).
out – Output: out[0] receives the argmax index.
out_val – Output: out_val[0] receives max|x|.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void iamax_lowmem(const T *x, uint32_t *out, T *out_val)#

i_amax + max|x|, low-memory variant, compile-time size.

Compile-time-N value-returning serial overload.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Read-only input vector of length N (not modified).
out – Output: out[0] receives the argmax index.
out_val – Output: out_val[0] receives max|x|.

template<typename T, bool TRAILING_SYNC = true> void iamax_fast(uint32_t n, const T *x, uint32_t *out, T *s_scratch)#

i_amax, warp-shuffle variant: index of max|x| into out[0].

Each thread forms a strided per-thread argmax in registers, the warp folds it with __shfl_down_sync (carrying the (key,index) pair), and the per-warp winners are combined through s_scratch. Non-destructive. NumPy equivalent: int(np.argmax(np.abs(x))).

Tie-break / NaN: Lower index wins on equal |x| at every shuffle/scratch combine (so the result is block-size invariant). NaN is skipped (diverges from np.argmax; exclude NaN in tests). All-zero vector → index 0.
Scratch: s_scratch must hold iamax_fast_scratch_bytes<T>(blockDim) elements (one (key,index) slot per warp). Ends on a trailing __syncthreads().

Template Parameters:

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

Parameters:

n – Number of elements.
x – Read-only input vector of length n (not modified).
out – Output: out[0] receives the argmax index.
s_scratch – Shared scratch sized by iamax_fast_scratch_bytes<T>(blockDim).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void iamax_fast(const T *x, uint32_t *out, T *s_scratch)#

i_amax, warp-shuffle variant, compile-time size.

Compile-time-N overload of the warp-shuffle i_amax. NumPy equivalent: int(np.argmax(np.abs(x))). Scratch via iamax_fast_scratch_bytes<T>(blockDim).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Read-only input vector of length N (not modified).
out – Output: out[0] receives the argmax index.
s_scratch – Shared scratch sized by iamax_fast_scratch_bytes<T>(blockDim).

template<typename T, bool TRAILING_SYNC = true> void iamax_fast(uint32_t n, const T *x, uint32_t *out, T *out_val, T *s_scratch)#

i_amax + max|x|, warp-shuffle variant.

As the warp-shuffle iamax but also writes out_val[0] = max|x|. Non-destructive; lower-index tie-break; NaN skipped; all-zero → 0.

Template Parameters:

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

Parameters:

n – Number of elements.
x – Read-only input vector of length n (not modified).
out – Output: out[0] receives the argmax index.
out_val – Output: out_val[0] receives max|x|.
s_scratch – Shared scratch sized by iamax_fast_scratch_bytes<T>(blockDim).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void iamax_fast(const T *x, uint32_t *out, T *out_val, T *s_scratch)#

i_amax + max|x|, warp-shuffle variant, compile-time size.

Compile-time-N value-returning warp-shuffle overload. NumPy equivalents: out[0] = int(np.argmax(np.abs(x))), out_val[0] = np.max(np.abs(x)).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Read-only input vector of length N (not modified).
out – Output: out[0] receives the argmax index.
out_val – Output: out_val[0] receives max|x|.
s_scratch – Shared scratch sized by iamax_fast_scratch_bytes<T>(blockDim).

template<typename T> uint32_t iamax(uint32_t n, const T *x)#

Index of max|x| (BLAS i_amax) within ONE warp, returned on every lane.

A single 32-lane warp computes argmax(|x|) and returns it (register return, broadcast to all lanes via __shfl_sync) — there is NO out[] write and NO shared scratch. Each lane forms a strided per-lane (abskey,index) argmax over for (i = lane; i < n; i += 32), then the warp folds the pair with __shfl_down_sync (lower-index tie-break at EVERY step) and broadcasts lane 0’s index. NumPy equivalent: int(np.argmax(np.abs(x))).

Full-warp requirement: This routine assumes a FULL 32-lane warp is active (mask 0xffffffff); it must be called by all 32 lanes of the warp. Inactive lanes (those whose lane >= n, i.e. that never enter the strided loop) seed key = 0, idx = UINT32_MAX, so they never win a tie. An all-zero vector therefore returns index 0.
Multi-warp independence: Shared-free and __syncthreads-free: many warps may run this concurrently in one block, each on its own x, with no cross-warp interference. The result is also identical to the block-scoped glass::iamax (same (key,idx) lower-index combine).
Tie-break / NaN: On EQUAL |x| the LOWER index wins (the BLAS rule), applied at every shuffle combine — so the result is lane/order independent. NaN inputs are SKIPPED (IEEE compares are false, so a NaN is never selected); this DIVERGES from np.argmax(np.abs(x)), which propagates NaN — oracle tests must exclude NaN.

Template Parameters:

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

Parameters:

n – Number of elements.
x – Read-only input vector of length n (not modified).

Returns:

The argmax index (uint32_t), identical on every lane.

template<typename T, uint32_t N> uint32_t iamax(const T *x)#

Index of max|x| (BLAS i_amax) within ONE warp, compile-time size.

Compile-time-N overload of the single-warp i_amax. Returns argmax(|x|) on every lane (register broadcast). NumPy equivalent: int(np.argmax(np.abs(x))). Same full-warp requirement, multi-warp independence, lower-index tie-break, and NaN-skip policy as the runtime-n overload.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – Read-only input vector of length N (not modified).

Returns:

The argmax index (uint32_t), identical on every lane.

namespace iamax_detail#

Functions

template<typename T> void combine(T &key, uint32_t &idx, T ckey, uint32_t cidx)#

namespace warp

clip#

Functions

template<typename T, bool TRAILING_SYNC = true> void clip(uint32_t n, T *x, T *l, T *u)#

Element-wise clamp in place: x = clamp(x, l, u).

Each element is bounded below by the corresponding l[i] and above by u[i]. NumPy equivalent: np.clip(x, l, u).

Template Parameters:

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

Parameters:

n – Number of elements.
x – In/out vector of length n (overwritten with the clamped values).
l – Per-element lower bounds, length n.
u – Per-element upper bounds, length n.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void clip(T *x, T *l, T *u)#

Element-wise clamp in place: x = clamp(x, l, u), compile-time size.

Compile-time-N overload of clip. NumPy equivalent: np.clip(x, l, u).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

x – In/out vector of length N (overwritten with the clamped values).
l – Per-element lower bounds, length N.
u – Per-element upper bounds, length N.

set_const#

Functions

template<typename T, bool TRAILING_SYNC = true> void set_const(uint32_t n, T alpha, T *x)#

Fill a vector with a constant: x[i] = alpha.

NumPy equivalent: x = np.full(n, alpha).

Template Parameters:

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

Parameters:

n – Number of elements.
alpha – Value to broadcast into every element.
x – Output vector of length n.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void set_const(T alpha, T *x)#

Fill a vector with a constant: x[i] = alpha, compile-time size.

Compile-time-N overload. NumPy equivalent: x = np.full(N, alpha).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

alpha – Value to broadcast into every element.
x – Output vector of length N.

identity#

Functions

template<typename T, bool TRAILING_SYNC = true> void set_identity(uint32_t n, T *A)#

Load the identity matrix: A = I_n (column-major).

Writes the n×n identity into A in column-major order. NumPy equivalent: A = np.eye(n).

Template Parameters:

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

Parameters:

n – Matrix dimension (number of rows/columns).
A – Output matrix of n*n elements (column-major).

template<typename T, bool TRAILING_SYNC = true> void add_identity(uint32_t n, T *A, T alpha)#

Add a scaled identity to a matrix in place: A += alpha * I.

Adds alpha to the diagonal of the n×n (column-major) matrix A. NumPy equivalent: A += alpha * np.eye(n).

Template Parameters:

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

Parameters:

n – Matrix dimension (number of rows/columns).
A – In/out matrix of n*n elements (column-major).
alpha – Scalar added to each diagonal entry.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void set_identity(T *A)#

Load the identity matrix: A = I_N (column-major), compile-time size.

Compile-time-N overload. NumPy equivalent: A = np.eye(N).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Matrix dimension (compile-time constant).

Parameters:

A – Output matrix of N*N elements (column-major).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void add_identity(T *A, T alpha)#

Add a scaled identity in place: A += alpha * I, compile-time size.

Compile-time-N overload. NumPy equivalent: A += alpha * np.eye(N).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Matrix dimension (compile-time constant).

Parameters:

A – In/out matrix of N*N elements (column-major).
alpha – Scalar added to each diagonal entry.

template<typename T, bool TRAILING_SYNC = true> void add_identity_partial(uint32_t n, T *A, T alpha, uint32_t diag_count)#

Add a scaled identity to the leading diagonal block: A[:d,:d] += alpha*I.

Adds alpha to only the first diag_count diagonal entries of the n×n (column-major) matrix A, leaving the trailing n-diag_count diagonal entries untouched. Use to regularize a leading sub-block — e.g. the position block of a stacked [q; v] state, where diag_count = n/2. NumPy equivalent: A[:d, :d] += alpha * np.eye(d).

Template Parameters:

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

Parameters:

n – Matrix dimension (number of rows/columns).
A – In/out matrix of n*n elements (column-major).
alpha – Scalar added to each of the leading diagonal entries.
diag_count – Number of leading diagonal entries to bump (<= n).

template<typename T, uint32_t N, uint32_t DIAG_COUNT, bool TRAILING_SYNC = true> void add_identity_partial(T *A, T alpha)#

Add a scaled identity to the leading diagonal block, compile-time size.

Compile-time <N, DIAG_COUNT> overload of add_identity_partial. NumPy equivalent: A[:DIAG_COUNT, :DIAG_COUNT] += alpha * np.eye(DIAG_COUNT).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Matrix dimension (compile-time constant).
DIAG_COUNT – Number of leading diagonal entries to bump (<= N).

Parameters:

A – In/out matrix of N*N elements (column-major).
alpha – Scalar added to each of the leading diagonal entries.

symmetrize#

In-place symmetrization A := 0.5*(A + Aᵀ) for a square matrix.

Numerically enforces symmetry on an n×n column-major matrix — the standard cleanup after a product chain (e.g. a Schur-complement gemm sequence) whose result is symmetric in exact arithmetic but drifts under floating point. MPCGPU hand-rolls exactly this after its Schur assembly (without it, CG stagnates on the slightly-asymmetric operator); this is that loop, once.

Each strictly-lower (i>j) pair owner reads BOTH mirror elements and writes BOTH — every (i,j)/(j,i) pair is touched by exactly one thread, the diagonal is untouched, so there is no cross-thread hazard and the op is trivially thread-count invariant. Block + warp:: + cgrps::.

Functions

template<typename T, bool TRAILING_SYNC = true> void symmetrize(uint32_t n, T *A)#

Symmetrize a square matrix in place: A = 0.5*(A + Aᵀ).

Averages each strictly-off-diagonal mirror pair of the n×n column-major matrix A; the diagonal is untouched. NumPy equivalent: A = 0.5*(A + A.T).

Template Parameters:

T – Scalar type (e.g. float, double).
TRAILING_SYNC – Emit a trailing __syncthreads() (default true).

Parameters:

n – Matrix dimension (number of rows/columns).
A – In/out matrix of n*n elements (column-major).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void symmetrize(T *A)#

Symmetrize in place: A = 0.5*(A + Aᵀ), compile-time size.

Compile-time-N overload. NumPy equivalent: A = 0.5*(A + A.T).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Matrix dimension (compile-time constant).
TRAILING_SYNC – Emit a trailing __syncthreads() (default true).

Parameters:

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

template<typename T> void symmetrize(uint32_t n, T *A)#

Symmetrize within one warp: A = 0.5*(A + Aᵀ), single-warp.

Symmetrize on one thread: A = 0.5*(A + Aᵀ), single-thread.

One 32-lane warp averages the mirror pairs of the n×n column-major matrix in place; the diagonal is untouched. Each pair owned by exactly one lane; trailing __syncwarp(). NumPy equivalent: A = 0.5*(A + A.T).

One thread averages each strictly-off-diagonal mirror pair of the n×n column-major matrix serially; the diagonal is untouched. No shared scratch, no shuffles, no barriers, no threadIdx read; A may be a thread-local register array. NumPy equivalent: A = 0.5*(A + A.T).

Template Parameters:

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

Parameters:

n – Matrix dimension (number of rows/columns).
A – In/out matrix of n*n elements (column-major).
n – Matrix dimension (number of rows/columns).
A – In/out matrix of n*n elements (column-major).

template<typename T, uint32_t N> void symmetrize(T *A)#

Symmetrize within one warp: A = 0.5*(A + Aᵀ), compile-time size.

Symmetrize on one thread: A = 0.5*(A + Aᵀ), single-thread, compile-time size.

Compile-time-N overload of the single-warp symmetrize. NumPy equivalent: A = 0.5*(A + A.T).

Compile-time-N overload; the trip count folds and the loop unrolls, so A may be a thread-local register array. NumPy equivalent: A = 0.5*(A + A.T).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Matrix dimension (compile-time constant).
T – Scalar type (e.g. float, double).
N – Matrix dimension (compile-time constant).

Parameters:

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

namespace warp

namespace thread

transpose#

Functions

template<typename T, bool TRAILING_SYNC = true> void transpose(uint32_t N, uint32_t M, T *a, T *b)#

template<typename T, bool TRAILING_SYNC = true> void transpose(uint32_t N, T *a)#

template<typename T, uint32_t N, uint32_t M, bool TRAILING_SYNC = true> void transpose(T *a, T *b)#

Out-of-place matrix transpose: b = aᵀ, compile-time size.

Compile-time-N,M overload; transposes the N×M column-major matrix a into the M×N column-major matrix b. NumPy equivalent: b = a.T.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of rows of a (columns of b), compile-time constant.
M – Number of columns of a (rows of b), compile-time constant.

Parameters:

a – Input matrix of N*M elements (column-major).
b – Output matrix of M*N elements (column-major).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void transpose(T *a)#

In-place square matrix transpose: a = aᵀ, compile-time size.

Compile-time-N overload; transposes the N×N column-major matrix a in place. NumPy equivalent: a = a.T.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Matrix dimension (compile-time constant).

Parameters:

a – In/out matrix of N*N elements (column-major).

Elementwise logic#

Functions

template<typename T, bool TRAILING_SYNC = true> void elementwise_max(uint32_t N, T *a, T *b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_min(uint32_t N, T *a, T *b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_less_than(uint32_t N, T *a, T *b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_more_than(uint32_t N, T *a, T *b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_less_than_or_eq(uint32_t N, T *a, T *b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_less_than_scalar(uint32_t N, T *a, T b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_and(uint32_t N, T *a, T *b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_not(uint32_t N, T *a, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_abs(uint32_t N, T *a, T *b)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_mult(uint32_t N, T *a, T *b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_sub(uint32_t N, T *a, T *b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_add(uint32_t N, T *a, T *b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_mult_scalar(uint32_t N, T *a, T b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_max_scalar(uint32_t N, T *a, T b, T *c)#

template<typename T, bool TRAILING_SYNC = true> void elementwise_min_scalar(uint32_t N, T *a, T b, T *c)#

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void elementwise_max(T *a, T *b, T *c)#

Element-wise maximum: c = max(a, b), compile-time size.

Compile-time-N overload. NumPy equivalent: np.maximum(a, b).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

a – Input vector of length N.
b – Input vector of length N.
c – Output vector of length N.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void elementwise_min(T *a, T *b, T *c)#

Element-wise minimum: c = min(a, b), compile-time size.

Compile-time-N overload. NumPy equivalent: np.minimum(a, b).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

a – Input vector of length N.
b – Input vector of length N.
c – Output vector of length N.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void elementwise_abs(T *a, T *b)#

Element-wise absolute value: b = |a|, compile-time size.

Compile-time-N overload. NumPy equivalent: b = np.abs(a).

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

a – Input vector of length N.
b – Output vector of length N.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void elementwise_mult(T *a, T *b, T *c)#

Element-wise (Hadamard) product: c = a ⊙ b, compile-time size.

Compile-time-N overload. NumPy equivalent: c = a * b.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

a – Input vector of length N.
b – Input vector of length N.
c – Output vector of length N.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void elementwise_sub(T *a, T *b, T *c)#

Element-wise subtraction: c = a - b, compile-time size.

Compile-time-N overload. NumPy equivalent: c = a - b.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

a – Input vector of length N.
b – Input vector of length N.
c – Output vector of length N.

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void elementwise_add(T *a, T *b, T *c)#

Element-wise addition: c = a + b, compile-time size.

Compile-time-N overload. NumPy equivalent: c = a + b.

Template Parameters:

T – Scalar type (e.g. float, double).
N – Number of elements (compile-time constant).

Parameters:

a – Input vector of length N.
b – Input vector of length N.
c – Output vector of length N.

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

Functions

template<typename T, bool TRAILING_SYNC = true> void axpy(uint32_t n, T alpha, T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#: AXPY y = alpha*x + y (cgrps variant; see glass::axpy).

template<typename T, bool TRAILING_SYNC = true> void axpy(uint32_t n, T alpha, T *x, T *y, T *z, cgrps::thread_group g = cgrps::this_thread_block())#: Out-of-place AXPY z = alpha*x + y (cgrps variant; see glass::axpy).

template<typename T, bool TRAILING_SYNC = true> void axpby(uint32_t n, T alpha, T *x, T beta, T *y, T *z, cgrps::thread_group g = cgrps::this_thread_block())#: AXPBY z = alpha*x + beta*y (cgrps variant; see glass::axpby).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void axpy(T alpha, T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#: Compile-time-size AXPY y = alpha*x + y (cgrps variant).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void axpy(T alpha, T *x, T *y, T *z, cgrps::thread_group g = cgrps::this_thread_block())#: Compile-time-size out-of-place AXPY z = alpha*x + y (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void copy(uint32_t n, T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#: Vector copy y = x (cgrps variant; see glass::copy).

template<typename T, bool TRAILING_SYNC = true> void copy(uint32_t n, T alpha, T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#: Scaled copy y = alpha*x (cgrps variant; see glass::copy).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void copy(T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#: Compile-time-size vector copy y = x (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void scal(uint32_t n, T alpha, T *x, cgrps::thread_group g = cgrps::this_thread_block())#: In-place scale x = alpha*x (cgrps variant; see glass::scal).

template<typename T, bool TRAILING_SYNC = true> void scal(uint32_t n, T alpha, T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#: Out-of-place scale y = alpha*x (cgrps variant; see glass::scal).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void scal(T alpha, T *x, cgrps::thread_group g = cgrps::this_thread_block())#: Compile-time-size in-place scale x = alpha*x (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void swap(uint32_t n, T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#: Swap two vectors x <-> y (cgrps variant; see glass::swap).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void swap(T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#: Compile-time-size swap x <-> y (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void clip(uint32_t n, T *x, T *l, T *u, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise clamp x = clamp(x, l, u) (cgrps variant; see glass::clip).

template<typename T, bool TRAILING_SYNC = true> void set_const(uint32_t n, T alpha, T *x, cgrps::thread_group g = cgrps::this_thread_block())#: Fill with a constant x = alpha (cgrps variant; see glass::set_const).

template<typename T, bool TRAILING_SYNC = true> void set_identity(uint32_t n, T *A, cgrps::thread_group g = cgrps::this_thread_block())#: Load the identity A = I_n (cgrps variant; see glass::set_identity).

template<typename T, bool TRAILING_SYNC = true> void add_identity(uint32_t n, T *A, T alpha, cgrps::thread_group g = cgrps::this_thread_block())#: Add a scaled identity A += alpha*I (cgrps variant; see glass::add_identity).

template<typename T, bool TRAILING_SYNC = true> void reduce(uint32_t n, T *x, cgrps::thread_group g = cgrps::this_thread_block())#: Sum reduction x[0] = sum(x) (cgrps variant; see glass::reduce).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void reduce(T *x, cgrps::thread_group g = cgrps::this_thread_block())#: Compile-time-size sum reduction x[0] = sum(x) (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void dot(uint32_t n, T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#: Dot product y[0] = dot(x, y) (destructive; cgrps variant; see glass::dot).

template<typename T, uint32_t N, bool TRAILING_SYNC = true> void dot(T *x, T *y, cgrps::thread_group g = cgrps::this_thread_block())#: Compile-time-size dot product y[0] = dot(x, y) (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void transpose(uint32_t N, uint32_t M, T *a, T *b, cgrps::thread_group g = cgrps::this_thread_block())#: Out-of-place transpose b = a^T (cgrps variant; see glass::transpose).

template<typename T, bool TRAILING_SYNC = true> void transpose(uint32_t N, T *a, cgrps::thread_group g = cgrps::this_thread_block())#: In-place square transpose a = a^T (cgrps variant; see glass::transpose).

template<typename T, bool TRAILING_SYNC = true> void symmetrize(uint32_t n, T *A, cgrps::thread_group g = cgrps::this_thread_block())#: In-place symmetrization A = 0.5*(A + Aᵀ) (cgrps variant; see glass::symmetrize).

template<typename T, bool TRAILING_SYNC = true> void elementwise_max(uint32_t N, T *a, T *b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise max c = max(a, b) (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_min(uint32_t N, T *a, T *b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise min c = min(a, b) (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_less_than(uint32_t N, T *a, T *b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise less-than c = (a < b) (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_more_than(uint32_t N, T *a, T *b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise greater-than c = (a > b) (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_less_than_or_eq(uint32_t N, T *a, T *b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise less-than-or-equal c = (a <= b) (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_and(uint32_t N, T *a, T *b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise logical AND c = (a && b) (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_not(uint32_t N, T *a, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise logical NOT c = !a (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_abs(uint32_t N, T *a, T *b, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise absolute value b = |a| (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_mult(uint32_t N, T *a, T *b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise (Hadamard) product c = a * b (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_sub(uint32_t N, T *a, T *b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise subtraction c = a - b (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_add(uint32_t N, T *a, T *b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise addition c = a + b (cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_mult_scalar(uint32_t N, T *a, T b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise scalar multiply c = a * b (scalar b; cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_max_scalar(uint32_t N, T *a, T b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise max against a scalar c = max(a, b) (scalar b; cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void elementwise_min_scalar(uint32_t N, T *a, T b, T *c, cgrps::thread_group g = cgrps::this_thread_block())#: Element-wise min against a scalar c = min(a, b) (scalar b; cgrps variant).

template<typename T, bool TRAILING_SYNC = true> void nrm2(uint32_t n, T *x, cgrps::thread_group g = cgrps::this_thread_block())#: Euclidean (L2) norm x[0] = ||x||_2 (destructive; cgrps variant; see glass::nrm2_impl).

template<typename T, bool TRAILING_SYNC = true> void infnorm(uint32_t n, T *x, cgrps::thread_group g = cgrps::this_thread_block())#: Infinity norm x[0] = ||x||_inf (destructive; cgrps variant; see glass::infnorm).

template<typename T, bool TRAILING_SYNC = true> void asum(uint32_t n, T *x, T *out, cgrps::thread_group g = cgrps::this_thread_block())#: Sum of absolute values out[0] = sum(|x|) (cgrps variant; see glass::asum_impl).

struct GroupBarrier#

Public Functions

inline uint32_t rank() const#

inline uint32_t size() const#

inline void sync() const#

Public Members

cgrps::thread_group g#

L1 — Vector Operations#

axpy / axpby#

copy#

scal#

swap#

rot / rotg (Givens plane rotation)#

dot#

reduce#

prefix_sum#

Norms#

Index of max-abs (iamax)#

clip#

set_const#

identity#

symmetrize#

transpose#

Elementwise logic#

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

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