NVIDIA Backend (`glass::nvidia::`)#

Vendor-accelerated paths built on CUB (reductions), cuBLASDx (GEMM/GEMV), and cuSOLVERDx (LAPACK). The block-scope ops live in glass::nvidia::block:: (the contract tier; bare glass::nvidia:: re-exports them as the measured-default face). The entry points auto-dispatch between a pure-SIMT implementation and the vendor backend based on a size heuristic / tuning table; see Backend Dispatch. The L2/L3/LAPACK paths require NVIDIA MathDx (MATHDX_ROOT) — see Installation.

Each call has a companion host-side query helper (*_scratch_bytes, *_threads, *_block_threads_valid) used to size the launch.

L1 (CUB-backed reductions)#

reduce / dot / nrm2 at block scope (cub::BlockReduce) and, in glass::nvidia::warp::, at warp scope (cub::WarpReduce — one FULL 32-lane warp per problem, per-warp scratch via warp_reduce_scratch_bytes<T>(), TRAILING_SYNC emits __syncwarp()).

Defines

_GLASS_ASSERT_THREADS_EQ(THREADS)#

Functions

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

Block-level sum reduction backed by CUB BlockReduce.

Sums all N elements of x across the block; thread 0 writes the total to x[0] (in place). Requires blockDim == (THREADS, 1, 1) exactly — CUB does not tolerate a mismatch. Compile-time sizes only. NumPy equivalent: x[0] = np.sum(x).

Template Parameters:

T – Scalar type.
N – Number of elements.
THREADS – Block thread count (must equal blockDim.x).
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

x – Input array of length N; result lands in x[0].
s_scratch – Shared scratch >= reduce_scratch_bytes<T, THREADS>() bytes.

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

Block-level dot product backed by CUB BlockReduce.

Computes the inner product of the N-element vectors x and y; thread 0 writes the scalar result to *out. Requires blockDim == (THREADS, 1, 1) exactly. Compile-time sizes only. NumPy equivalent: *out = np.dot(x, y).

Template Parameters:

T – Scalar type.
N – Number of elements.
THREADS – Block thread count (must equal blockDim.x).
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

x – First input vector (length N).
y – Second input vector (length N).
out – Output pointer for the resulting scalar.
s_scratch – Shared scratch >= reduce_scratch_bytes<T, THREADS>() bytes.

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

Block-level Euclidean (L2) norm backed by CUB BlockReduce.

Sums the squares of the N elements of x and writes the square root to *out from thread 0. Requires blockDim == (THREADS, 1, 1) exactly. Compile-time sizes only. NumPy equivalent: *out = np.linalg.norm(x).

Template Parameters:

T – Scalar type.
N – Number of elements.
THREADS – Block thread count (must equal blockDim.x).
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

x – Input vector (length N).
out – Output pointer for the resulting scalar norm.
s_scratch – Shared scratch >= reduce_scratch_bytes<T, THREADS>() bytes.

template<typename T, uint32_t THREADS = 256> inline constexpr std::size_t reduce_scratch_bytes()#

Shared-memory bytes needed for the L1 reduce/dot/nrm2 scratch (host-callable).

Returns sizeof(cub::BlockReduce<T, THREADS>::TempStorage) — the size of the s_scratch buffer these CUB-backed reductions require. constexpr.

Template Parameters:

T – Scalar type.
THREADS – Block thread count (must match the reduction call).

Returns:

Required scratch size in bytes.

L2 (gemv)#

Defines

SMS#

_GLASS_GEMV_NO_BD(M, N, LA, LB, LC, CT, ARCH)#

_GLASS_GEMV_BD(M, N, TC, LA, LB, LC, CT, ARCH)#

_GLASS_GEMV_NO_BD_E(M, N, LA, LB, LC, CT, ARCH)#

_GLASS_GEMV_BD_E(M, N, TC, LA, LB, LC, CT, ARCH)#

DEFINE_NVIDIA_GEMV(M, N)#

DEFINE_NVIDIA_GEMV_BLOCKDIM(M, N, TC)#

DEFINE_NVIDIA_GEMV_LAYOUT(M, N, LA, LB, LC)#

DEFINE_NVIDIA_GEMV_BLOCKDIM_LAYOUT(M, N, TC, LA, LB, LC)#

DEFINE_NVIDIA_GEMV_SM(M, N, SM)#

DEFINE_NVIDIA_GEMV_BLOCKDIM_SM(M, N, TC, SM)#

DEFINE_NVIDIA_GEMV_PREC(M, N, CT)#

DEFINE_NVIDIA_GEMV_BLOCKDIM_PREC(M, N, TC, CT)#

DEFINE_NVIDIA_GEMV_LAYOUT_SM(M, N, LA, LB, LC, SM)#

DEFINE_NVIDIA_GEMV_BLOCKDIM_LAYOUT_SM(M, N, TC, LA, LB, LC, SM)#

Functions

template<typename T, uint32_t M, uint32_t N, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860, bool TRAILING_SYNC = true> void gemv(T alpha, T *A, T *x, T beta, T *y, char *smem)#

Block-level GEMV that auto-dispatches between SIMT and cuBLASDx.

Computes y = alpha*A*x + beta*y for an M×N matrix A. The primary template consults should_use_cublasdx_gemv<T,M,N,SM_VAL>() at compile time: small shapes fall through to the dependency-free SIMT path glass::gemv (no scratch, LA maps to its ROW_MAJOR flag); shapes the heuristic flags for the vendor backend require a DEFINE_NVIDIA_GEMV* macro in scope (cuBLASDx, needs the MathDx headers / MATHDX_ROOT) or a static_assert fires. Compile-time sizes only. NumPy equivalent: y = alpha*A@x + beta*y.

Template Parameters:

T – Scalar type.
M – Rows of A / length of y.
N – Columns of A / length of x.
BLOCK_THREADS – Pinned cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B (degenerate for a vector; ignored).
LC – Memory layout of C (degenerate for a vector; ignored).
SM_VAL – Target SM architecture (default = SMS).
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

alpha – Scaling factor for A*x.
A – Pointer to the M×N matrix.
x – Pointer to the input vector (length N).
beta – Scaling factor for the incoming y.
y – Pointer to the output vector (length M).
smem – Shared scratch (cuBLASDx route only; unused for SIMT).

template<typename T, uint32_t M, uint32_t N, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860> constexpr std::size_t gemv_scratch_bytes()#

Shared-memory bytes needed by gemv<...> (host-callable).

Returns the cuBLASDx scratch size for this signature, or 0 when the auto-dispatch routes to SIMT (which needs no scratch). Template parameters match gemv<>. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A.
N – Columns of A.
BLOCK_THREADS – Pinned cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.

template<typename T, uint32_t M, uint32_t N, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860> constexpr uint32_t gemv_threads()#

Thread count cuBLASDx wants for gemv<...> (host-callable).

Returns the block thread count to launch with for the cuBLASDx route (the default 256 is overridden by the matching DEFINE_NVIDIA_GEMV* specialization). Template parameters match gemv<>. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A.
N – Columns of A.
BLOCK_THREADS – Pinned cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.

template<typename T, uint32_t M, uint32_t N, uint32_t ROW_STRIDE = M, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860, bool TRAILING_SYNC = true> void gemv_strided(T alpha, T *A, T *x, T beta, T *y, char *smem)#

Strided-A GEMV that auto-dispatches between SIMT and cuBLASDx.

Computes y = alpha*A*x + beta*y where A is stored with an arbitrary leading dimension ROW_STRIDE (A[i + j*ROW_STRIDE]). On the SIMT route the stride is used directly (no packing, no scratch); on the cuBLASDx route the strided A is packed into compact shared scratch and forwarded to gemv<>. When ROW_STRIDE == M this degenerates to standard gemv. Works under any launch geometry (1D/2D/3D) as long as the total thread count matches. cuBLASDx route needs the MathDx headers / MATHDX_ROOT.

Template Parameters:

T – Scalar type.
M – Rows of A / length of y.
N – Columns of A / length of x.
ROW_STRIDE – Leading dimension (row stride) of A.
BLOCK_THREADS – Pinned cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

alpha – Scaling factor for A*x.
A – Pointer to the strided M×N matrix.
x – Pointer to the input vector (length N).
beta – Scaling factor for the incoming y.
y – Pointer to the output vector (length M).
smem – Shared scratch (cuBLASDx route only; unused for SIMT).

template<typename T, uint32_t M, uint32_t N, uint32_t ROW_STRIDE = M, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860> constexpr std::size_t gemv_strided_scratch_bytes()#

Shared-memory bytes needed by gemv_strided<...> (host-callable).

Returns 0 when the auto-dispatch routes to SIMT (no packing), or the A-packing scratch plus the inner cuBLASDx gemv scratch otherwise. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A.
N – Columns of A.
ROW_STRIDE – Leading dimension (row stride) of A.
BLOCK_THREADS – Pinned cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.

L3 (gemm)#

Defines

SMS

_GLASS_GEMM_NO_BD(M, N, K, LA, LB, LC, CT, ARCH)#

_GLASS_GEMM_BD(M, N, K, TC, LA, LB, LC, CT, ARCH)#

_GLASS_GEMM_NO_BD_E(M, N, K, LA, LB, LC, CT, ARCH)#

_GLASS_GEMM_BD_E(M, N, K, TC, LA, LB, LC, CT, ARCH)#

DEFINE_NVIDIA_GEMM(M, N, K)#

DEFINE_NVIDIA_GEMM_BLOCKDIM(M, N, K, TC)#

DEFINE_NVIDIA_GEMM_LAYOUT(M, N, K, LA, LB, LC)#

DEFINE_NVIDIA_GEMM_BLOCKDIM_LAYOUT(M, N, K, TC, LA, LB, LC)#

DEFINE_NVIDIA_GEMM_SM(M, N, K, SM)#

DEFINE_NVIDIA_GEMM_BLOCKDIM_SM(M, N, K, TC, SM)#

DEFINE_NVIDIA_GEMM_PREC(M, N, K, CT)#

DEFINE_NVIDIA_GEMM_BLOCKDIM_PREC(M, N, K, TC, CT)#

DEFINE_NVIDIA_GEMM_LAYOUT_SM(M, N, K, LA, LB, LC, SM)#

DEFINE_NVIDIA_GEMM_BLOCKDIM_LAYOUT_SM(M, N, K, TC, LA, LB, LC, SM)#

DEFINE_NVIDIA_GEMM_TRANSB(M, N, K)#

DEFINE_NVIDIA_GEMM_BLOCKDIM_TRANSB(M, N, K, TC)#

DEFINE_NVIDIA_GEMM_TRANSB_SM(M, N, K, SM)#

DEFINE_NVIDIA_GEMM_BLOCKDIM_TRANSB_SM(M, N, K, TC, SM)#

_GLASS_GEMM_BATCHED_BD(M, N, K, BATCH, TC, LA, LB, LC, ARCH)#

_GLASS_GEMM_BATCHED_BD_E(M, N, K, BATCH, TC, LA, LB, LC, ARCH)#

DEFINE_NVIDIA_GEMM_BATCHED_BLOCKDIM(M, N, K, BATCH, TC)#

DEFINE_NVIDIA_GEMM_BATCHED_BLOCKDIM_LAYOUT(M, N, K, BATCH, TC, LA, LB, LC)#

DEFINE_NVIDIA_GEMM_BATCHED_BLOCKDIM_SM(M, N, K, BATCH, TC, SM)#

DEFINE_NVIDIA_GEMM_BATCHED_BLOCKDIM_LAYOUT_SM(M, N, K, BATCH, TC, LA, LB, LC, SM)#

Functions

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860, bool TRAILING_SYNC = true> void gemm(T alpha, T *A, T *B, T beta, T *C, char *smem)#

Block-level GEMM that auto-dispatches between SIMT and cuBLASDx.

Computes C = alpha*A*B + beta*C for an M×N×K product. The primary template consults should_use_cublasdx<T,M,N,K,SM_VAL>() at compile time: small shapes fall through to the dependency-free SIMT path glass::gemm (no scratch; LA/LB/LC are mapped onto its TRANSPOSE_B / ROW_MAJOR flags), while shapes the heuristic flags for the vendor backend require a DEFINE_NVIDIA_GEMM* macro in scope (cuBLASDx, needs the MathDx headers / MATHDX_ROOT) or a static_assert fires. Compile-time sizes only. NumPy equivalent: C = alpha*A@B + beta*C.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension (cols of A, rows of B).
BLOCK_THREADS – Pinned cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B (row_major expresses A*Bᵀ).
LC – Memory layout of C.
SM_VAL – Target SM architecture (default = SMS).
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

alpha – Scaling factor for A*B.
A – Pointer to the M×K matrix.
B – Pointer to the K×N matrix.
beta – Scaling factor for the incoming C.
C – Pointer to the M×N output matrix.
smem – Shared scratch (cuBLASDx route only; unused for SIMT).

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860> constexpr std::size_t gemm_scratch_bytes()#

Shared-memory bytes needed by gemm<...> (host-callable).

Returns the cuBLASDx scratch size for this signature, or 0 when the auto-dispatch routes to SIMT (which needs no scratch). Template parameters match gemm<>. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BLOCK_THREADS – Pinned cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860> constexpr uint32_t gemm_threads()#

Thread count cuBLASDx wants for gemm<...> (host-callable).

Returns the block thread count to launch with for the cuBLASDx route (the default 256 is overridden by the matching DEFINE_NVIDIA_GEMM* specialization). Template parameters match gemm<>. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BLOCK_THREADS – Pinned cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t A_RS = M, uint32_t B_RS = K, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860, bool TRAILING_SYNC = true> void gemm_strided(T alpha, T *A, T *B, T beta, T *C, char *smem)#

Strided-A/B GEMM that auto-dispatches between SIMT and cuBLASDx.

Computes C = alpha*A*B + beta*C where A and B are stored with arbitrary leading dimensions A_RS and B_RS. On the SIMT route the strides are used directly (no packing, no scratch); on the cuBLASDx route A and B are packed into compact shared scratch and forwarded to gemm<>. C is written standard column-major (LDC = M; no strided output). When A_RS==M and B_RS==N this degenerates to standard gemm. cuBLASDx route needs MathDx / MATHDX_ROOT.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
A_RS – Leading dimension (row stride) of A.
B_RS – Leading dimension (row stride) of B.
BLOCK_THREADS – Pinned cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

alpha – Scaling factor for A*B.
A – Pointer to the strided M×K matrix.
B – Pointer to the strided K×N matrix.
beta – Scaling factor for the incoming C.
C – Pointer to the M×N output matrix.
smem – Shared scratch (cuBLASDx route only; unused for SIMT).

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t A_RS = M, uint32_t B_RS = K, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860> constexpr std::size_t gemm_strided_scratch_bytes()#

Shared-memory bytes needed by gemm_strided<...> (host-callable).

Returns 0 when the auto-dispatch routes to SIMT (no packing), or the A+B packing scratch plus the inner cuBLASDx gemm scratch otherwise. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
A_RS – Leading dimension (row stride) of A.
B_RS – Leading dimension (row stride) of B.
BLOCK_THREADS – Pinned cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860, bool TRAILING_SYNC = true> void gemm_batched(T alpha, T *const *A, T *const *B, T beta, T *const *C, char *smem)#

cuBLASDx-backed batched GEMM: BATCH independent products in one block.

Performs BATCH independent C = alpha*A*B + beta*C products, one per threadIdx.y, in a single CUDA block (requires a 2D launch dim3(TC, BATCH)). A/B/C are length-BATCH arrays of pointers. The primary template is a static_assert stub — a

DEFINE_NVIDIA_GEMM_BATCHED_BLOCKDIM(


M,N,K,BATCH,TC)

specialization must be in scope. Needs cuBLASDx / MathDx (MATHDX_ROOT). For a 1D launch use gemm_batched_1d in l3_simt.cuh.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BATCH – Number of independent GEMMs.
BLOCK_THREADS – Pinned per-batch cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

alpha – Scaling factor for A*B.
A – Array of BATCH pointers to the M×K matrices.
B – Array of BATCH pointers to the K×N matrices.
beta – Scaling factor for the incoming C.
C – Array of BATCH pointers to the M×N output matrices.
smem – Shared scratch (BATCH copies of per-GEMM cuBLASDx smem).

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860> constexpr std::size_t gemm_batched_scratch_bytes()#

Total shared-memory bytes needed by gemm_batched<...> (host-callable).

Returns BATCH copies of the per-GEMM cuBLASDx scratch (set by the matching DEFINE_NVIDIA_GEMM_BATCHED_* specialization; 0 for the unspecialized primary). Template parameters match gemm_batched<>. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BATCH – Number of independent GEMMs.
BLOCK_THREADS – Pinned per-batch cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t BLOCK_THREADS = 0, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, uint32_t SM_VAL = 860> constexpr uint32_t gemm_batched_threads()#

Total thread count for gemm_batched<...> (host-callable).

Returns the full launch thread count (per-batch BlockDim × BATCH) for the cuBLASDx route; launch as dim3(TC, BATCH). The default 256 is overridden by the matching DEFINE_NVIDIA_GEMM_BATCHED_* specialization. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BATCH – Number of independent GEMMs.
BLOCK_THREADS – Pinned per-batch cuBLASDx BlockDim (0 = vendor picks).
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
SM_VAL – Target SM architecture.

L3 SIMT batched (no cuBLASDx)#

Functions

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t TC, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, bool TRAILING_SYNC = true> void gemm_batched_1d(T alpha, T *const *A, T *const *B, T beta, T *const *C)#

Pure-SIMT batched GEMM for a 1D launch (no cuBLASDx, no scratch).

Runs BATCH independent C = alpha*A*B + beta*C products in one block, giving each batch element TC threads carved out of a 1D launch of >= TC*BATCH threads (canonical dim3(TC*BATCH, 1, 1)). Reuses the well-tested compile-time SIMT gemm core; no shared memory and no DEFINE macro needed, so any T works. Best for small shapes (max(M,N,K) ≲ 8) where cuBLASDx tile-load overhead dominates. Use the cuBLASDx gemm_batched (l3.cuh) for larger shapes.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BATCH – Number of independent GEMMs.
TC – Threads assigned to each batch element.
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

alpha – Scaling factor for A*B.
A – Array of BATCH pointers to the M×K matrices.
B – Array of BATCH pointers to the K×N matrices.
beta – Scaling factor for the incoming C.
C – Array of BATCH pointers to the M×N output matrices.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t TC, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, bool TRAILING_SYNC = true> constexpr std::size_t gemm_batched_1d_scratch_bytes()#

Shared-memory bytes needed by gemm_batched_1d<...> (host-callable).

Always 0 (the SIMT path needs no scratch); provided for API symmetry with the cuBLASDx gemm_batched. Template parameters match gemm_batched_1d<>. constexpr.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t TC, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, bool TRAILING_SYNC = true> constexpr uint32_t gemm_batched_1d_threads()#

Total threads required across the 1D block for gemm_batched_1d<...>.

Returns TC * BATCH — launch with at least this many threads. Host-callable constexpr. Template parameters match gemm_batched_1d<>.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t TC, uint32_t B_STRIDE = N * K, uint32_t C_STRIDE = M * N, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, bool TRAILING_SYNC = true> void gemm_strided_batched_1d(T alpha, const T *A_shared, T *B, T beta, T *C)#

Pure-SIMT batched GEMM (1D launch) with one shared A across BATCH (B,C) pairs.

Like gemm_batched_1d, but a single A matrix is broadcast to BATCH strided (B, C) pairs: batch b reads B + b*B_STRIDE, writes C + b*C_STRIDE. Common in GRiD’s EE-pose-gradient (one transform applied to many destinations). Avoids building pointer arrays caller-side. No cuBLASDx, no scratch, no DEFINE macro. Strides default to tightly-packed batches.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BATCH – Number of (B,C) pairs sharing A.
TC – Threads assigned to each batch element.
B_STRIDE – Element stride between consecutive B matrices.
C_STRIDE – Element stride between consecutive C matrices.
LA – Memory layout of A.
LB – Memory layout of B.
LC – Memory layout of C.
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

alpha – Scaling factor for A*B.
A_shared – Pointer to the single shared M×K matrix (read-only).
B – Base pointer for the BATCH K×N matrices.
beta – Scaling factor for the incoming C.
C – Base pointer for the BATCH M×N output matrices.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t TC, uint32_t B_STRIDE = N * K, uint32_t C_STRIDE = M * N, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, bool TRAILING_SYNC = true> constexpr std::size_t gemm_strided_batched_1d_scratch_bytes()#

Shared-memory bytes needed by gemm_strided_batched_1d<...> (host-callable).

Always 0 (SIMT path needs no scratch); provided for API symmetry. Template parameters match gemm_strided_batched_1d<>. constexpr.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t TC, uint32_t B_STRIDE = N * K, uint32_t C_STRIDE = M * N, layout LA = layout::col_major, layout LB = layout::col_major, layout LC = layout::col_major, bool TRAILING_SYNC = true> constexpr uint32_t gemm_strided_batched_1d_threads()#

Total threads required across the 1D block for gemm_strided_batched_1d<...>.

Returns TC * BATCH — launch with at least this many threads. Host-callable constexpr. Template parameters match gemm_strided_batched_1d<>.

LAPACK (cuSOLVERDx)#

cuSOLVERDx-backed single-block LAPACK wrappers for glass::nvidia::.

Block-level, compile-time-size factorizations and solves: Cholesky (potrf), triangular solve (trsm), SPD solve (posv / potrs), unpivoted LU (getrf_no_pivot / getrs_no_pivot / gesv_no_pivot), QR (geqrf), and least-squares (gels). Each is a static_assert stub by default; instantiate a shape with the matching DEFINE_NVIDIA_<NAME> / _BLOCKDIM / _SM / _BLOCKDIM_SM macro in your .cu (inside namespace glass::nvidia). Each also ships *_scratch_bytes / *_threads host-callable constexpr queries.

Requires cuSOLVERDx / NVIDIA MathDx (MATHDX_ROOT) and linking the precompiled device library (-rdc=true -dlto -lcusolverdx -lcublas -lcusolver -lcudart). All matrices are column-major.

Defines

SMS

_GLASS_CHOL_NO_BD(N, CT, ARCH)#

_GLASS_CHOL_BD(N, TC, CT, ARCH)#

_GLASS_CHOL_NO_BD_E(N, CT, ARCH)#

_GLASS_CHOL_BD_E(N, TC, CT, ARCH)#

_GLASS_TRSM_NO_BD(M, N, CT, ARCH)#

_GLASS_TRSM_BD(M, N, TC, CT, ARCH)#

_GLASS_TRSM_NO_BD_E(M, N, CT, ARCH)#

_GLASS_TRSM_BD_E(M, N, TC, CT, ARCH)#

DEFINE_NVIDIA_CHOL(N)#

DEFINE_NVIDIA_CHOL_BLOCKDIM(N, TC)#

DEFINE_NVIDIA_CHOL_SM(N, SM)#

DEFINE_NVIDIA_CHOL_BLOCKDIM_SM(N, TC, SM)#

DEFINE_NVIDIA_CHOL_PREC(N, CT)#

DEFINE_NVIDIA_CHOL_BLOCKDIM_PREC(N, TC, CT)#

DEFINE_NVIDIA_CHOL_PREC_SM(N, CT, SM)#

DEFINE_NVIDIA_CHOL_BLOCKDIM_PREC_SM(N, TC, CT, SM)#

DEFINE_NVIDIA_TRSM(M, N)#

DEFINE_NVIDIA_TRSM_BLOCKDIM(M, N, TC)#

DEFINE_NVIDIA_TRSM_SM(M, N, SM)#

DEFINE_NVIDIA_TRSM_BLOCKDIM_SM(M, N, TC, SM)#

DEFINE_NVIDIA_TRSM_PREC(M, N, CT)#

DEFINE_NVIDIA_TRSM_BLOCKDIM_PREC(M, N, TC, CT)#

DEFINE_NVIDIA_TRSM_PREC_SM(M, N, CT, SM)#

DEFINE_NVIDIA_TRSM_BLOCKDIM_PREC_SM(M, N, TC, CT, SM)#

_GLASS_POSV_NO_BD(N, NRHS, CT, ARCH)#

_GLASS_POSV_BD(N, NRHS, TC, CT, ARCH)#

_GLASS_POSV_NO_BD_E(N, NRHS, CT, ARCH)#

_GLASS_POSV_BD_E(N, NRHS, TC, CT, ARCH)#

DEFINE_NVIDIA_POSV(N, NRHS)#

DEFINE_NVIDIA_POSV_BLOCKDIM(N, NRHS, TC)#

DEFINE_NVIDIA_POSV_SM(N, NRHS, SM)#

DEFINE_NVIDIA_POSV_BLOCKDIM_SM(N, NRHS, TC, SM)#

DEFINE_NVIDIA_POSV_PREC(N, NRHS, CT)#

DEFINE_NVIDIA_POSV_BLOCKDIM_PREC(N, NRHS, TC, CT)#

DEFINE_NVIDIA_POSV_PREC_SM(N, NRHS, CT, SM)#

DEFINE_NVIDIA_POSV_BLOCKDIM_PREC_SM(N, NRHS, TC, CT, SM)#

_GLASS_POTRS_NO_BD(N, NRHS, ARCH)#

_GLASS_POTRS_BD(N, NRHS, TC, ARCH)#

_GLASS_POTRS_NO_BD_E(N, NRHS, ARCH)#

_GLASS_POTRS_BD_E(N, NRHS, TC, ARCH)#

DEFINE_NVIDIA_POTRS(N, NRHS)#

DEFINE_NVIDIA_POTRS_BLOCKDIM(N, NRHS, TC)#

DEFINE_NVIDIA_POTRS_SM(N, NRHS, SM)#

DEFINE_NVIDIA_POTRS_BLOCKDIM_SM(N, NRHS, TC, SM)#

_GLASS_GETRF_NO_BD(N, ARCH)#

_GLASS_GETRF_BD(N, TC, ARCH)#

_GLASS_GETRF_NO_BD_E(N, ARCH)#

_GLASS_GETRF_BD_E(N, TC, ARCH)#

DEFINE_NVIDIA_GETRF(N)#

DEFINE_NVIDIA_GETRF_BLOCKDIM(N, TC)#

DEFINE_NVIDIA_GETRF_SM(N, SM)#

DEFINE_NVIDIA_GETRF_BLOCKDIM_SM(N, TC, SM)#

_GLASS_GETRS_NO_BD(N, NRHS, ARCH)#

_GLASS_GETRS_BD(N, NRHS, TC, ARCH)#

_GLASS_GETRS_NO_BD_E(N, NRHS, ARCH)#

_GLASS_GETRS_BD_E(N, NRHS, TC, ARCH)#

DEFINE_NVIDIA_GETRS(N, NRHS)#

DEFINE_NVIDIA_GETRS_BLOCKDIM(N, NRHS, TC)#

DEFINE_NVIDIA_GETRS_SM(N, NRHS, SM)#

DEFINE_NVIDIA_GETRS_BLOCKDIM_SM(N, NRHS, TC, SM)#

_GLASS_GESV_NO_BD(N, NRHS, ARCH)#

_GLASS_GESV_BD(N, NRHS, TC, ARCH)#

_GLASS_GESV_NO_BD_E(N, NRHS, ARCH)#

_GLASS_GESV_BD_E(N, NRHS, TC, ARCH)#

DEFINE_NVIDIA_GESV(N, NRHS)#

DEFINE_NVIDIA_GESV_BLOCKDIM(N, NRHS, TC)#

DEFINE_NVIDIA_GESV_SM(N, NRHS, SM)#

DEFINE_NVIDIA_GESV_BLOCKDIM_SM(N, NRHS, TC, SM)#

_GLASS_GEQRF_NO_BD(M, N, ARCH)#

_GLASS_GEQRF_BD(M, N, TC, ARCH)#

_GLASS_GEQRF_NO_BD_E(M, N, ARCH)#

_GLASS_GEQRF_BD_E(M, N, TC, ARCH)#

DEFINE_NVIDIA_GEQRF(M, N)#

DEFINE_NVIDIA_GEQRF_BLOCKDIM(M, N, TC)#

DEFINE_NVIDIA_GEQRF_SM(M, N, SM)#

DEFINE_NVIDIA_GEQRF_BLOCKDIM_SM(M, N, TC, SM)#

_GLASS_GELS_BMAX(M, N)#

_GLASS_GELS_NO_BD(M, N, NRHS, ARCH)#

_GLASS_GELS_BD(M, N, NRHS, TC, ARCH)#

_GLASS_GELS_NO_BD_E(M, N, NRHS, ARCH)#

_GLASS_GELS_BD_E(M, N, NRHS, TC, ARCH)#

DEFINE_NVIDIA_GELS(M, N, NRHS)#

DEFINE_NVIDIA_GELS_BLOCKDIM(M, N, NRHS, TC)#

DEFINE_NVIDIA_GELS_SM(M, N, NRHS, SM)#

DEFINE_NVIDIA_GELS_BLOCKDIM_SM(M, N, NRHS, TC, SM)#

Functions

template<typename T, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860, bool TRAILING_SYNC = true> void potrf(T *A, char *smem)#

In-place Cholesky factorization of an N×N SPD matrix (cuSOLVERDx potrf).

Factors A = L*Lᵀ and overwrites A with the lower-triangular L (upper triangle untouched). Primary template is a static_assert stub — add a DEFINE_NVIDIA_CHOL* specialization for the shape. Requires cuSOLVERDx / MathDx (MATHDX_ROOT). NumPy equivalent: A = np.linalg.cholesky(A).

Template Parameters:

T – Scalar type.
N – Matrix dimension.
BLOCK_THREADS – Pinned cuSOLVERDx BlockDim (0 = vendor picks).
SM_VAL – Target SM architecture (default = SMS).
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

A – Pointer to the N×N column-major SPD matrix; overwritten with L.
smem – Shared scratch (>= potrf_scratch_bytes<…>()).

template<typename T, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr std::size_t potrf_scratch_bytes()#

Shared-memory bytes needed by potrf<...> (host-callable, constexpr).

Template Parameters:

T – T scalar;
N – dim;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr uint32_t potrf_threads()#

Thread count cuSOLVERDx wants for potrf<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t M, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860, bool TRAILING_SYNC = true> void trsm(T alpha, T *L, T *B, char *smem)#

Lower-triangular solve L*X = alpha*B in place (cuSOLVERDx trsm).

Solves with side=left, fill=lower, non-transposed, non-unit-diagonal; B is overwritten with X. cuSOLVERDx has no alpha, so the wrapper pre-multiplies B by alpha in shared memory first. Primary template is a static_assert stub — add a DEFINE_NVIDIA_TRSM* specialization. Requires cuSOLVERDx / MathDx (MATHDX_ROOT). SciPy equivalent: solve_triangular(L, alpha*B, lower=True).

Template Parameters:

T – Scalar type.
M – Rows of L (M×M) and of B.
N – Columns (number of right-hand sides) of B.
BLOCK_THREADS – Pinned cuSOLVERDx BlockDim (0 = vendor picks).
SM_VAL – Target SM architecture (default = SMS).
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

alpha – Scaling factor applied to B before the solve.
L – Pointer to the M×M lower-triangular matrix.
B – Pointer to the M×N right-hand sides; overwritten with X.
smem – Shared scratch (>= trsm_scratch_bytes<…>()).

template<typename T, uint32_t M, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr std::size_t trsm_scratch_bytes()#

Shared-memory bytes needed by trsm<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
M – rows;
N – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t M, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr uint32_t trsm_threads()#

Thread count cuSOLVERDx wants for trsm<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
M – rows;
N – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860, bool TRAILING_SYNC = true> void posv(T *A, T *B, char *smem)#

Fused SPD solve A*X = B (cuSOLVERDx posv: Cholesky factor + solve).

Factors SPD A and solves in one call; A is overwritten with L (lower), B with X. Faster than a separate chol+trsm for N >= ~8. Primary template is a static_assert stub — add a DEFINE_NVIDIA_POSV* specialization. Requires cuSOLVERDx / MathDx (MATHDX_ROOT). NumPy equivalent: X = np.linalg.solve(A, B).

Template Parameters:

T – Scalar type.
N – Matrix dimension (A is N×N).
NRHS – Number of right-hand sides (columns of B).
BLOCK_THREADS – Pinned cuSOLVERDx BlockDim (0 = vendor picks).
SM_VAL – Target SM architecture (default = SMS).
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

A – Pointer to the N×N SPD matrix; overwritten with L.
B – Pointer to the N×NRHS right-hand sides; overwritten with X.
smem – Shared scratch (>= posv_scratch_bytes<…>()).

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr std::size_t posv_scratch_bytes()#

Shared-memory bytes needed by posv<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
NRHS – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr uint32_t posv_threads()#

Thread count cuSOLVERDx wants for posv<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
NRHS – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> void potrs(const T *L, T *B, char *smem)#

SPD solve given the Cholesky factor (cuSOLVERDx potrs).

Solves L*Lᵀ*X = B given the lower factor L from potrf; B is overwritten with X. Primary template is a static_assert stub — add a DEFINE_NVIDIA_POTRS* specialization. Requires cuSOLVERDx / MathDx (MATHDX_ROOT). SciPy equivalent: cho_solve((L, True), B).

Template Parameters:

T – Scalar type.
N – Matrix dimension.
NRHS – Number of right-hand sides.
BLOCK_THREADS – Pinned cuSOLVERDx BlockDim (0 = vendor picks).
SM_VAL – Target SM architecture (default = SMS).

Parameters:

L – Pointer to the N×N lower Cholesky factor (read-only).
B – Pointer to the N×NRHS right-hand sides; overwritten with X.
smem – Shared scratch (>= potrs_scratch_bytes<…>()).

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr std::size_t potrs_scratch_bytes()#

Shared-memory bytes needed by potrs<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
NRHS – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr uint32_t potrs_threads()#

Thread count cuSOLVERDx wants for potrs<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
NRHS – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> void getrf_no_pivot(T *A, char *smem)#

In-place LU factorization without pivoting (cuSOLVERDx getrf_no_pivot).

Factors A = L*U (unit-lower L, upper U) in place, no row pivoting. Primary template is a static_assert stub — add a DEFINE_NVIDIA_GETRF* specialization. Requires cuSOLVERDx / MathDx (MATHDX_ROOT). SciPy (pivot-free) equivalent of lu_factor(A) with A := LU.

Template Parameters:

T – Scalar type.
N – Matrix dimension.
BLOCK_THREADS – Pinned cuSOLVERDx BlockDim (0 = vendor picks).
SM_VAL – Target SM architecture (default = SMS).

Parameters:

A – Pointer to the N×N matrix; overwritten with the LU factors.
smem – Shared scratch (>= getrf_no_pivot_scratch_bytes<…>()).

template<typename T, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr std::size_t getrf_no_pivot_scratch_bytes()#

Shared-memory bytes needed by getrf_no_pivot<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr uint32_t getrf_no_pivot_threads()#

Thread count cuSOLVERDx wants for getrf_no_pivot<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> void getrs_no_pivot(const T *LU, T *B, char *smem)#

LU solve given the (unpivoted) LU factor (cuSOLVERDx getrs_no_pivot).

Solves L*U*X = B given the LU factor from getrf_no_pivot; B is overwritten with X. Primary template is a static_assert stub — add a DEFINE_NVIDIA_GETRS* specialization. Requires cuSOLVERDx / MathDx (MATHDX_ROOT). SciPy (pivot-free) equivalent of lu_solve((LU, ...), B).

Template Parameters:

T – Scalar type.
N – Matrix dimension.
NRHS – Number of right-hand sides.
BLOCK_THREADS – Pinned cuSOLVERDx BlockDim (0 = vendor picks).
SM_VAL – Target SM architecture (default = SMS).

Parameters:

LU – Pointer to the N×N LU factor (read-only).
B – Pointer to the N×NRHS right-hand sides; overwritten with X.
smem – Shared scratch (>= getrs_no_pivot_scratch_bytes<…>()).

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr std::size_t getrs_no_pivot_scratch_bytes()#

Shared-memory bytes needed by getrs_no_pivot<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
NRHS – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr uint32_t getrs_no_pivot_threads()#

Thread count cuSOLVERDx wants for getrs_no_pivot<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
NRHS – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860, bool TRAILING_SYNC = true> void gesv_no_pivot(T *A, T *B, char *smem)#

Fused general solve A*X = B without pivoting (cuSOLVERDx gesv_no_pivot).

LU-factors general A (no pivoting) and solves in one call; A is destroyed, B overwritten with X. Primary template is a static_assert stub — add a DEFINE_NVIDIA_GESV* specialization. Requires cuSOLVERDx / MathDx (MATHDX_ROOT). NumPy equivalent: X = np.linalg.solve(A, B).

Template Parameters:

T – Scalar type.
N – Matrix dimension (A is N×N).
NRHS – Number of right-hand sides.
BLOCK_THREADS – Pinned cuSOLVERDx BlockDim (0 = vendor picks).
SM_VAL – Target SM architecture (default = SMS).
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

A – Pointer to the N×N matrix; destroyed (holds LU on exit).
B – Pointer to the N×NRHS right-hand sides; overwritten with X.
smem – Shared scratch (>= gesv_no_pivot_scratch_bytes<…>()).

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr std::size_t gesv_no_pivot_scratch_bytes()#

Shared-memory bytes needed by gesv_no_pivot<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
NRHS – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr uint32_t gesv_no_pivot_threads()#

Thread count cuSOLVERDx wants for gesv_no_pivot<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
N – dim;
NRHS – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t M, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> void geqrf(T *A, T *tau, char *smem)#

In-place QR factorization of an M×N matrix (cuSOLVERDx geqrf).

Factors A into R (upper triangle) and Householder reflectors (below the diagonal), with the scalar factors written to tau. Primary template is a static_assert stub — add a DEFINE_NVIDIA_GEQRF* specialization. Requires cuSOLVERDx / MathDx (MATHDX_ROOT). SciPy equivalent: qr(A, mode=’raw)`.

Template Parameters:

T – Scalar type.
M – Rows of A.
N – Columns of A.
BLOCK_THREADS – Pinned cuSOLVERDx BlockDim (0 = vendor picks).
SM_VAL – Target SM architecture (default = SMS).

Parameters:

A – Pointer to the M×N matrix; overwritten with R + reflectors.
tau – Caller-provided output array of min(M,N) scalars.
smem – Shared scratch (>= geqrf_scratch_bytes<…>()).

template<typename T, uint32_t M, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr std::size_t geqrf_scratch_bytes()#

Shared-memory bytes needed by geqrf<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
M – rows;
N – cols;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t M, uint32_t N, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr uint32_t geqrf_threads()#

Thread count cuSOLVERDx wants for geqrf<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
M – rows;
N – cols;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t M, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860, bool TRAILING_SYNC = true> void gels(T *A, T *tau, T *B, char *smem)#

Least-squares solve of min ||A*X - B|| (cuSOLVERDx gels).

Solves the over- or under-determined system; cuSOLVERDx picks QR (M >= N) or LQ (M < N) internally. A is destroyed, B overwritten with X. Primary template is a static_assert stub — add a DEFINE_NVIDIA_GELS* specialization. Requires cuSOLVERDx / MathDx (MATHDX_ROOT). NumPy equivalent: X = np.linalg.lstsq(A, B).

Template Parameters:

T – Scalar type.
M – Rows of A.
N – Columns of A.
NRHS – Number of right-hand sides.
BLOCK_THREADS – Pinned cuSOLVERDx BlockDim (0 = vendor picks).
SM_VAL – Target SM architecture (default = SMS).
TRAILING_SYNC – Emit a trailing __syncthreads() before return (default true).

Parameters:

A – Pointer to the M×N matrix; destroyed during the solve.
tau – Caller-provided workspace of min(M,N) scalars.
B – Right-hand sides (max(M,N)×NRHS storage); overwritten with X.
smem – Shared scratch (>= gels_scratch_bytes<…>()).

template<typename T, uint32_t M, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr std::size_t gels_scratch_bytes()#

Shared-memory bytes needed by gels<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
M – rows;
N – cols;
NRHS – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

template<typename T, uint32_t M, uint32_t N, uint32_t NRHS, uint32_t BLOCK_THREADS = 0, uint32_t SM_VAL = 860> constexpr uint32_t gels_threads()#

Thread count cuSOLVERDx wants for gels<...> (host-callable, constexpr).

Template Parameters:

T – scalar;
M – rows;
N – cols;
NRHS – rhs;
BLOCK_THREADS – pinned BlockDim;
SM_VAL – SM arch.

Dispatch & query helpers#

Defines

SMS

Functions

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t SM_VAL = 860> constexpr uint32_t gemm_min_block_threads()#

Smallest BlockDim cuBLASDx will accept for a GEMM (host-callable).

Constructs the cuBLASDx GEMM type inline and reads its natural block_dim product — the minimum thread count to pin via DEFINE_NVIDIA_GEMM_BLOCKDIM. No DEFINE macro required. Requires cuBLASDx / MathDx (MATHDX_ROOT). constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
SM_VAL – Target SM architecture (default = SMS).

Returns:

Natural block thread count cuBLASDx picks.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BLOCK_THREADS, uint32_t SM_VAL = 860> constexpr bool gemm_block_threads_valid()#

True iff BLOCK_THREADS is enough for a cuBLASDx GEMM (host-callable).

Returns whether BLOCK_THREADS >= gemm_min_block_threads<T,M,N,K,SM_VAL>(). Use in a static_assert to validate a pinned launch thread count. Requires cuBLASDx / MathDx (MATHDX_ROOT). constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BLOCK_THREADS – Candidate launch thread count to validate.
SM_VAL – Target SM architecture (default = SMS).

Returns:

true if BLOCK_THREADS meets the minimum.

template<typename T, uint32_t M, uint32_t N, uint32_t SM_VAL = 860> constexpr uint32_t gemv_min_block_threads()#

Smallest BlockDim cuBLASDx will accept for a GEMV (host-callable).

GEMV is modeled as a Size<M, 1, N> GEMM; returns the natural block_dim product. No DEFINE macro required. Requires cuBLASDx / MathDx (MATHDX_ROOT). constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / length of y.
N – Columns of A / length of x.
SM_VAL – Target SM architecture (default = SMS).

Returns:

Natural block thread count cuBLASDx picks.

template<typename T, uint32_t M, uint32_t N, uint32_t BLOCK_THREADS, uint32_t SM_VAL = 860> constexpr bool gemv_block_threads_valid()#

True iff BLOCK_THREADS is enough for a cuBLASDx GEMV (host-callable).

Returns whether BLOCK_THREADS >= gemv_min_block_threads<T,M,N,SM_VAL>(). Requires cuBLASDx / MathDx (MATHDX_ROOT). constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A.
N – Columns of A.
BLOCK_THREADS – Candidate launch thread count to validate.
SM_VAL – Target SM architecture (default = SMS).

Returns:

true if BLOCK_THREADS meets the minimum.

Defines

SMS

Functions

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t SM_VAL = 860> constexpr bool should_use_cublasdx()#

Compile-time backend decision for gemm: cuBLASDx vs SIMT (host-callable).

Returns true iff cuBLASDx is expected to beat the SIMT path for this (T,M,N,K) on SM_VAL. Consults the per-SM tuning_table.cuh, falling back to a conservative shape heuristic for unmeasured shapes. Drives the gemm<> primary template’s dispatch. Only T==float is tuned; other types return false. No cuBLASDx dependency. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
SM_VAL – Target SM architecture (default = SMS).

Returns:

true to route to cuBLASDx, false for SIMT.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t SM_VAL = 860> inline void print_dispatch()#

Print which backend gemm<T,M,N,K,...> dispatches to (host/device).

Diagnostic helper: prints “cuBLASDx (needs DEFINE_NVIDIA_GEMM*)” or “SIMT

fallback” for the given shape. Callable from host or device. Useful when debugging an undefined-symbol or unexpectedly-slow GEMM.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
SM_VAL – Target SM architecture (default = SMS).

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t TC, uint32_t BLOCK_THREADS> constexpr bool gemm_batched_1d_block_threads_valid()#

True iff BLOCK_THREADS suffices for gemm_batched_1d (host-callable).

Returns whether BLOCK_THREADS >= TC*BATCH (every batch element gets TC threads). Unlike gemm_block_threads_valid this does not consult cuBLASDx — the SIMT-only requirement is purely “enough threads”. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BATCH – Number of independent GEMMs.
TC – Threads per batch element.
BLOCK_THREADS – Candidate launch thread count to validate.

Returns:

true if BLOCK_THREADS >= TC*BATCH.

template<typename T, uint32_t M, uint32_t N, uint32_t SM_VAL = 860> constexpr bool should_use_cublasdx_gemv()#

Compile-time backend decision for gemv: cuBLASDx vs SIMT (host-callable).

Sibling of should_use_cublasdx<> for the gemv API; consults tuning_table.cuh. Only T==float is tuned; other types return false. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A.
N – Columns of A.
SM_VAL – Target SM architecture (default = SMS).

Returns:

true to route to cuBLASDx, false for SIMT.

template<typename T, uint32_t M, uint32_t N, uint32_t ROW_STRIDE, uint32_t SM_VAL = 860> constexpr bool should_use_cublasdx_gemv_strided()#

Compile-time backend decision for gemv_strided (host-callable).

Sibling of should_use_cublasdx<> for the strided-GEMV API. Only T==float is tuned; other types return false. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A.
N – Columns of A.
ROW_STRIDE – Leading dimension (row stride) of A.
SM_VAL – Target SM architecture (default = SMS).

Returns:

true to route to cuBLASDx, false for SIMT.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t A_RS, uint32_t B_RS, uint32_t SM_VAL = 860> constexpr bool should_use_cublasdx_gemm_strided()#

Compile-time backend decision for gemm_strided (host-callable).

Sibling of should_use_cublasdx<> for the strided-GEMM API. Only T==float is tuned; other types return false. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
A_RS – Leading dimension (row stride) of A.
B_RS – Leading dimension (row stride) of B.
SM_VAL – Target SM architecture (default = SMS).

Returns:

true to route to cuBLASDx, false for SIMT.

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t SM_VAL = 860> constexpr bool should_use_cublasdx_batched()#

Compile-time backend hint for batched GEMM: cuBLASDx vs SIMT (host-callable).

Sibling of should_use_cublasdx<> reporting whether cuBLASDx would win for a batched GEMM shape (used as a tuning signal; gemm_batched_1d itself is SIMT-only and does not auto-dispatch). Only T==float is tuned. constexpr.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BATCH – Number of independent GEMMs.
SM_VAL – Target SM architecture (default = SMS).

Returns:

true if cuBLASDx is expected to win, false for SIMT.

template<typename T, uint32_t M, uint32_t N, uint32_t SM_VAL = 860> inline void print_dispatch_gemv()#

Print which backend gemv<T,M,N,...> dispatches to (host/device).

Diagnostic sibling of print_dispatch<> for the gemv API.

Template Parameters:

T – Scalar type.
M – Rows of A.
N – Columns of A.
SM_VAL – Target SM architecture (default = SMS).

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t A_RS = M, uint32_t B_RS = N, uint32_t SM_VAL = 860> inline void print_dispatch_gemm_strided()#

Print which backend gemm_strided<...> dispatches to (host/device).

Diagnostic sibling of print_dispatch<> for the strided-GEMM API.

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
A_RS – Leading dimension (row stride) of A.
B_RS – Leading dimension (row stride) of B.
SM_VAL – Target SM architecture (default = SMS).

template<typename T, uint32_t M, uint32_t N, uint32_t ROW_STRIDE = M, uint32_t SM_VAL = 860> inline void print_dispatch_gemv_strided()#

Print which backend gemv_strided<...> dispatches to (host/device).

Diagnostic sibling of print_dispatch<> for the strided-GEMV API.

Template Parameters:

T – Scalar type.
M – Rows of A.
N – Columns of A.
ROW_STRIDE – Leading dimension (row stride) of A.
SM_VAL – Target SM architecture (default = SMS).

template<typename T, uint32_t M, uint32_t N, uint32_t K, uint32_t BATCH, uint32_t SM_VAL = 860> inline void print_dispatch_batched()#

Print the batched-GEMM backend hint for a shape (host/device).

Diagnostic sibling of print_dispatch<> reporting the batched-GEMM tuning signal (gemm_batched_1d is SIMT-only; the 2D gemm_batched is the cuBLASDx path).

Template Parameters:

T – Scalar type.
M – Rows of A / C.
N – Columns of B / C.
K – Inner dimension.
BATCH – Number of independent GEMMs.
SM_VAL – Target SM architecture (default = SMS).

Public types and shared helper macros for the glass::nvidia:: wrappers.

Defines the user-facing glass::nvidia::layout enum (per-matrix memory order, used as the LA/LB/LC template arguments of gemm/gemv/…) plus the private helper macros (_GLASS_CUBLAS_LAYOUT, _GLASS_ASSERT_BLOCKDIM_GEQ) shared by l2.cuh and l3.cuh. Included before l1/l2/l3/lapack so layout is in scope inside the glass::nvidia namespace.

Defines

_GLASS_CUBLAS_LAYOUT(L)#

_GLASS_ASSERT_BLOCKDIM_GEQ(GEMM_T)#

namespace glass#

Matrix memory layout for the cuBLASDx-backed glass::nvidia:: wrappers.

Maps directly to cublasdx::Arrangement<>: col_major (Fortran/cuBLAS default) and row_major (C-style). Passed per matrix as the LA/LB/LC template arguments of gemm/gemv/row_strided_*. The numeric values (0/1) are part of the public ABI: the DEFINE_NVIDIA_*_LAYOUT* macros take integer literals and static_cast them back to this enum in their specializations.

namespace nvidia#

Enums

enum class layout : uint8_t#

Values:

enumerator col_major#

enumerator row_major#

Host helpers#

The umbrella headers also expose host-callable helpers for sizing dynamic shared memory at launch time.

Umbrella header for the hand-rolled SIMT surface (no deps).

Pulls in the full pure-SIMT, single-block BLAS/LAPACK surface — L1 vector ops, L2 matrix-vector (gemv/ger), and L3 matrix ops (gemm, inv, Cholesky, trsm) — all as __device__ helpers that cooperate across one CUDA block using threadIdx / blockDim directly (no cooperative-groups dependency). Every op offers runtime-size and compile-time-size (<T, N, ...>) overloads.

NAMESPACE CONTRACT (2026-07-30 restructure):

glass::block:: — the explicit block-scope SIMT tier. CONTRACT tier: bit-exact, thread-count invariant, never re-dispatched.
glass::warp:: — one problem per warp (alias of block::warp — the warp mirrors live inline in the same base headers).
glass::thread:: — one problem per thread (alias of block::thread).
glass::op BARE — the measured-DEFAULT face: block-scope calling contract, implementation chosen per (op, size, dtype) from the shipped dispatch table. Phase 1 pins every cell to the block body (glass::dispatch_body() in glass-defaults.cuh), so today the bare names ARE the block tier via the using-directive below — identical symbols, bit-identical results. A future measured retune may add shadowing wrappers that route specific cells to a warp- or thread-body executed under the same block-scope contract (all threads enter, result valid after return); such a retune is an attested event, and determinism-sensitive consumers should pin glass::block:: explicitly.

Include glass-cgrps.cuh for the cooperative-groups variants, or glass-nvidia.cuh for the CUB / cuBLASDx / cuSOLVERDx-accelerated paths (glass::nvidia::block:: / glass::nvidia::warp::, with the same bare-name re-export inside glass::nvidia::). Also defines the host helper glass_gemm_dispatch_smem below.

Functions

template<typename T = float> inline std::size_t glass_gemm_dispatch_smem(int m, int n, int block_threads = 256, int tile = 8)#

Host helper: shared-memory bytes needed for glass::gemm_dispatch (tiled path).

Compute on the host at launch time and pass as the kernel’s dynamic-smem argument. Returns 0 when tiling is not warranted (m >= 32 or m*n > block_threads), in which case glass::gemm_dispatch runs the plain (non-tiled) path. Host-callable. Standard convention: C is m×n, contraction k; the tiled path stages an m×TILE A-tile and a TILE×n B-tile.

Usage: size_t smem = glass_gemm_dispatch_smem<float>(m, n); kernel<<<grid, 256, smem>>>(m, n, k, A, B, C); // inside the kernel: extern shared T scratch[]; glass::gemm_dispatch(m, n, k, alpha, A, B, beta, C, (smem > 0) ? scratch : nullptr, (smem > 0) ? scratch + m * 8 : nullptr);

Template Parameters:

T – Scalar type (defaults to float).

Parameters:

m – Rows of A / C.
n – Columns of B / C.
block_threads – Launch thread count used for the tiling heuristic.
tile – Tile width (must match the gemm_tiled<T, TILE> used).

Returns:

Bytes of dynamic shared memory to allocate, or 0 for the plain path.

namespace glass

Matrix memory layout for the cuBLASDx-backed glass::nvidia:: wrappers.

Maps directly to cublasdx::Arrangement<>: col_major (Fortran/cuBLAS default) and row_major (C-style). Passed per matrix as the LA/LB/LC template arguments of gemm/gemv/row_strided_*. The numeric values (0/1) are part of the public ABI: the DEFINE_NVIDIA_*_LAYOUT* macros take integer literals and static_cast them back to this enum in their specializations.

namespace block#

Umbrella header for the glass::nvidia:: backend (CUB / cuBLASDx / cuSOLVERDx).

Include this (instead of, or in addition to, glass.cuh) to access the vendor-accelerated single-block linear-algebra paths. It pulls in:

L1 (l1.cuh) CUB-backed reductions: reduce / dot / nrm2 at block scope (cub::BlockReduce) AND warp scope (glass::nvidia::warp::, cub::WarpReduce — one full warp per problem, per-warp scratch).
query_simt.cuh SIMT-only dispatch + diagnostic helpers (should_use_cublasdx<>, print_dispatch<>, …).
L3 SIMT (l3_simt.cuh) 1D-launch batched GEMMs (no cuBLASDx dependency).
L2 (l2.cuh) cuBLASDx-backed gemv + DEFINE_NVIDIA_GEMV* macros.
L3 (l3.cuh) cuBLASDx-backed gemm / gemm_batched / row_strided_*.
query.cuh Host-side constexpr BlockDim query API.
LAPACK (lapack.cuh) cuSOLVERDx chol/trsm/posv/getrf/gesv/geqrf/gels.

The L2/L3/LAPACK wrappers gate themselves on GLASS_HAVE_CUBLASDX / GLASS_HAVE_CUSOLVERDX, auto-detected from include order. Set MATHDX_ROOT and define GLASS_BENCH_CUBLASDX / GLASS_BENCH_CUSOLVERDX to force-enable them. The glass::nvidia::* primary templates auto-dispatch between pure-SIMT and the vendor backend at compile time via the size heuristic / tuning table.

Defines

GLASS_HAVE_CUBLASDX#

GLASS_HAVE_CUSOLVERDX#

namespace glass

Matrix memory layout for the cuBLASDx-backed glass::nvidia:: wrappers.

Maps directly to cublasdx::Arrangement<>: col_major (Fortran/cuBLAS default) and row_major (C-style). Passed per matrix as the LA/LB/LC template arguments of gemm/gemv/row_strided_*. The numeric values (0/1) are part of the public ABI: the DEFINE_NVIDIA_*_LAYOUT* macros take integer literals and static_cast them back to this enum in their specializations.

namespace nvidia

namespace block#

NVIDIA Backend (glass::nvidia::)#

L1 (CUB-backed reductions)#

L2 (gemv)#

L3 (gemm)#

L3 SIMT batched (no cuBLASDx)#

LAPACK (cuSOLVERDx)#

Dispatch & query helpers#

Host helpers#

NVIDIA Backend (`glass::nvidia::`)#