GLASS

GLASS is a header-only CUDA C++ library of composable __device__ primitives for small, block-local linear algebra and robotics math. It includes BLAS operations, factorizations and solvers, structured-system routines, spatial algebra, Lie-group operations, projections, and small estimation kernels. GLASS is the foundational linear-algebra layer underneath GRiD, MPCGPU, GATO, HJCD-IK, and other A2R Lab GPU solvers.

📖 Full documentation: https://a2r-lab.github.io/GLASS/ (source under docs/source/).

Overview

GLASS functions operate on data already in shared or device memory. Block-scoped functions run within one CUDA block; the caller normally launches one block per independent problem. Warp- and thread-scoped forms let a block pack smaller independent problems more densely.

It began as hand-rolled SIMT subroutines tuned for the very small matrices where vendor launch/dispatch overhead dominates, and has grown into a unified single-block surface that also wraps NVIDIA's device-side libraries — CUB (L1), cuBLASDx (L2/L3), cuSOLVERDx (LAPACK) — under one __device__ calling convention, so one kernel can mix hand-rolled and vendor-backed primitives without leaving the block. Not every operation exists at every scope; the API reference is the authoritative surface inventory.

Interfaces

GLASS exposes four primary interfaces. Pick one based on how many independent problems should share a block and whether optional vendor dependencies are acceptable:

Interface	Scope	What it is / when to choose it	Header
glass::block:: (Block)	block	Explicit hand-rolled SIMT implementation; no dependencies and never re-dispatched	glass.cuh
glass::warp:: (Warp)	warp	Single-warp SIMT via __shfl_*_sync (selected L1/L2/L3 ops, no __syncthreads). Pack many small independent problems into one block	inline in the base headers (via glass.cuh)
glass::thread:: (Thread)	thread	Sequential branch-free subset, one compile-time problem per thread; intended for register-resident sizes up to N≤7	inline in the base headers (via glass.cuh)
glass::nvidia::block:: (Nvidia)	block	CUB + cuBLASDx + cuSOLVERDx, auto-dispatched against SIMT by size at compile time (compile-time sizes). When a vendor tensor-core kernel wins at your size. Plus glass::nvidia::warp:: — CUB WarpReduce L1 reductions, one full 32-lane warp per problem	glass-nvidia.cuh

Bare glass::op (and bare glass::nvidia::op) is the measured-default face: the same block-scope calling contract, with the implementation body chosen per (op, size, dtype) by glass::dispatch_body() (glass-dispatch.cuh, regenerated by bench/tune.py --legs body). This selection is a constexpr decision inside the device function itself — the compiler resolves each call site to a block, warp-0, or thread-0 body at compile time; there is no host-side dispatcher and no runtime branching or repacking. Measured cells route to a warp- or thread-body executed inside the block; a moved cell matches the block body to tolerance, not bit-exactly. Ops with no moved cell remain the same entities as glass::block:: (all old spellings compile unchanged). Pin glass::block:: explicitly wherever determinism is load-bearing: explicit namespace = contract tier, bare namespace = performance tier.

Note: glass::cgrps:: (header glass-cgrps.cuh) is a convenience cooperative-groups alias of the Block interface — the same SIMT loop indexed via a cooperative_groups::thread_group, numerically identical and not a separately-tuned backend.

Many operations offer both runtime (size as an argument) and compile-time (size as a template argument) overloads. Reductions additionally offer _lowmem (no scratch) and _fast (warp-shuffle) suffixed forms (e.g. glass::reduce_lowmem / glass::reduce_fast). The dense surface covers gemm/gemv/ger, iamax, trsv/trmv, syrk/syr2k, symm/trmm/dimm, inv/potrf (single and K-way fused), ldlt/ldlt_solve, and posv/potrs; plus contraction-parallel *_reduced, tensor_*, and congruence_* families. Every block-scope op takes a TRAILING_SYNC template flag (default true = ends on a barrier, always safe to compose; pass false to elide the tail barrier where separable — a documented no-op where the last barrier is fused into the algorithm; see src/base/barrier.cuh for the exact contract). See the namespace and naming guide and synchronization contract.

Higher-level solvers

Built on the primitives above (single-block) for the block-tridiagonal SPD systems of trajectory optimization / MPC:

Function	What	Header
glass::bdmv	block-tridiagonal matvec ([L|D|R] strips, padded vectors)	src/base/banded/bdmv.cuh
glass::bdsv	block-tridiagonal direct SPD solve (block-Cholesky / block-Thomas sweep; the direct sibling of pcg on the identical [L|D|R] layout)	src/base/banded/bdsv.cuh
glass::pcg	single-block preconditioned conjugate gradient (S x = b). ⚠ Documented launch contract: blockDim must be a multiple of 32 (warp-level dot reductions) — the one exception to thread-count invariance	src/base/pcg/solve.cuh

An internal box-constrained QP solver, glass::internal::box_qp, also lives in the tree but is not part of the public surface (QP is optimization, not linear algebra).

Quick start

#include "glass.cuh"

__global__ void my_kernel(float* A, float* B, float* C, int m, int n, int k) {
    glass::gemm(m, n, k, 1.f, A, B, 0.f, C);   // all block threads cooperate on one problem
}

my_kernel<<<num_items, 256>>>(A, B, C, m, n, k);   // one block per data item

Runnable, self-contained programs (one concept each) live in examples/. GEMM follows the standard BLAS convention — C is M×N, contraction K (A is M×K, B is K×N) — with TRANSPOSE_A / TRANSPOSE_B operand flags and a single ROW_MAJOR_C output flag; a row-major operand is just a transpose. See examples/02_gemm_conventions.cu.

Installation

GLASS is header-only — add the repo root to your include path and #include "glass.cuh". The pure-SIMT surface needs only nvcc -std=c++17. The glass::nvidia:: paths additionally need NVIDIA MathDx (cuBLASDx / cuSOLVERDx) and extra flags:

Surface	Build requirements
glass.cuh, glass-cgrps.cuh	C++17 — no extra deps
glass-nvidia.cuh (L1)	C++17 + CUB (bundled with CUDA 11+)
glass-nvidia.cuh (L2/L3 GEMM/GEMV/batched)	C++17 + --expt-relaxed-constexpr + cuBLASDx
glass-nvidia.cuh (LAPACK)	C++17 + --expt-relaxed-constexpr + -rdc=true -dlto -lcusolverdx -lcublas -lcusolver -lcudart + cuSOLVERDx

The nvidia wrappers auto-detect availability (GLASS_HAVE_CUBLASDX / GLASS_HAVE_CUSOLVERDX). Full setup, linking, and the MathDx download are in bench/INSTALL.md and the installation guide.

CMake consumers get an INTERFACE target (glass::glass) via add_subdirectory, find_package, or:

include(FetchContent)
FetchContent_Declare(glass GIT_REPOSITORY https://github.com/A2R-Lab/GLASS.git)
FetchContent_MakeAvailable(glass)
target_link_libraries(your_target PRIVATE glass::glass)

Build & test

pip install -r test/requirements.txt
pytest test/                 # compiles test/cuda/*.cu once, caches by source hash

Selected tests compile only the requested CUDA binaries. Signed correctness is split into eight functional shards, while releases rerun every shard. A receipt is a signed attestation by a keyholder that the suite ran at that exact source tree; CI re-verifies the signature and source fingerprint — it does not independently prove GPU execution (see pytest-gpu-proof's security model). Details in docs/source/user_guide/tutorials/running_tests.rst. What "tested" means per op — which independent oracle (NumPy/SciPy/Pinocchio), finite-difference identity, or pinned contract validates each family, attested under the signed receipt — is documented in testing_oracles.

Documentation map

The README is a landing page; the deep reference lives in the hosted docs (sources in docs/source/):

Topic	Page
API reference (L1 / L2 / L3 / nvidia / warp / banded)	api_reference/
Namespaces, naming rules, and the two-axis taxonomy	concepts/namespaces.rst
Choosing a backend + tuning for your hardware	concepts/tuning.rst
glass::nvidia::gemm cuBLASDx-vs-SIMT dispatch	concepts/backend_dispatch.rst
TRAILING_SYNC and barrier conventions	concepts/trailing_sync.rst
Contraction-parallel (*_reduced) family	concepts/contraction_parallel.rst
Block-tridiagonal layout (bdmv / pcg)	concepts/block_tridiagonal.rst
Worked examples + quickstart	tutorials/ · examples/
Benchmarks + measured sweep results	tutorials/benchmarks.rst · tutorials/sweep_results.rst

Notes / gotchas

• One block per problem. Every function runs inside a single block; launch <<<num_items, threads>>>. Exception: glass::thread:: is one problem per THREAD (<<<ceil(P/TPB), TPB>>>) — for low-DOF packing (N≲7, compile-time size only). See CLAUDE.md for its constraints.
• Column-major by default (Fortran order, matching cuBLAS). GEMM uses TRANSPOSE_A / TRANSPOSE_B + ROW_MAJOR_C (a row-major operand is just a transpose); GEMV keeps a per-matrix ROW_MAJOR flag (its transpose changes the math op); glass::nvidia:: uses the layout enum per matrix (LA/LB/LC).
• Reductions are destructive. dot / nrm2 / reduction variants write the result to x[0] and may consume the input as scratch; nrm2 squares elements before reducing. The glass::warp:: and glass::thread:: forms return the value instead; the glass::nvidia::warp:: (CUB) forms take an explicit per-warp scratch pointer.
• potrf fills only the lower triangle; the upper retains input values.
• glass::nvidia::* (default form) requires exactly gemm_threads<T,M,N,K>() threads; use the BLOCK_THREADS template parameter (with DEFINE_NVIDIA_<NAME>_BLOCKDIM) to launch any count ≥ gemm_min_block_threads<T,M,N,K>(). Compile without -DNDEBUG for a clean assertion instead of a silent deadlock if the launch is too small.
• glass::nvidia::trsm has no native non-1.0 alpha (cuSOLVERDx limitation); the wrapper pre-scales B in shared memory before execute.