GLASS: GPU Linear Algebra Simple Subroutines

GLASS is a header-only CUDA C++ library of composable __device__ primitives for small, block-local linear algebra and robotics math. It is the foundational linear-algebra layer underneath GRiD, MPCGPU, GATO, HJCD-IK, and other A2R Lab GPU solvers.

The library covers BLAS, factorizations and solvers, structured systems, spatial algebra, Lie groups, projections, and small estimation kernels. Block-scoped operations compose inside a caller-owned kernel; selected warp- and thread-scoped forms pack smaller independent problems more densely. See the API reference for the exact surface and robotics conventions for the load-bearing layout choices.

Interfaces

GLASS separates execution scope (thread, warp, or block) from implementation family (dependency-free GLASS code or optional NVIDIA libraries). The spellings intentionally overlap, but do not all expose every operation.

Block — glass::block::

One block per problem; the block’s threads cooperate over shared/global data. Pure SIMT, no dependencies (#include "glass.cuh"). This explicit implementation is never re-dispatched.

Library Overview

Warp — glass::warp::

One warp per problem (__shfl_*_sync, no __syncthreads), so warps run independently. Choose this to pack many small independent problems into one block for intra-block parallelism. Requires a full 32-lane warp.

Warp-scoped operations (glass::warp::)

Thread — glass::thread::

One problem per thread, 32 packed per warp. This is a compile-time, branch-free subset. It is usually register-resident around N≤7; larger measured sizes remain correct and can still win despite spills.

Thread-scoped operations (glass::thread::)

Nvidia — glass::nvidia::block::

CUB / cuBLASDx / cuSOLVERDx at block scope. Choose this when the measured vendor implementation wins at your size (needs NVIDIA MathDx).

Backend Dispatch

Nvidia warp — glass::nvidia::warp::

CUB WarpReduce L1 reductions, one full 32-lane warp per problem and explicit per-warp scratch.

NVIDIA Backends (glass::nvidia::*)

Nvidia thread — glass::nvidia::thread::

cuSOLVERDx 0.4+ LAPACK, one independent packed problem per CUDA thread; no dynamic shared scratch or block-wide synchronization.

NVIDIA Backends (glass::nvidia::*)

Note

glass::cgrps:: is a cooperative-groups adapter for the Block interface — identical numerics (the same SIMT loop, indexed via a thread_group handle), for callers already in a cooperative-groups context or tiling arbitrary sub-block groups. It is not a separately-tuned backend. #include "glass-cgrps.cuh".

Note

Bare glass::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 the measured tune.py --legs body sweep). The selection is a constexpr decision inside the device function — resolved per call site at compile time, with no host-side dispatcher and no runtime branch. Cells with a robust measured win route to a warp- or thread-body inside the block; every other name is the same entity as glass::block::. Determinism-sensitive callers pin glass::block:: explicitly; see Namespaces, suffixes, and flags.

NVIDIA calls always name block, warp, or thread; there is no bare glass::nvidia::op re-export.

Measured defaults

glass::recommend<op, T, dims...>() exposes architecture-specific measured execution plans for operations included in the tuning ladder. Measurements are not a promise that interfaces have identical coverage or reduction order. See Tuning for Your Hardware for the selection policy and Backend Sweep Results for dated, configuration-specific results.

Get started

Header-only install, the single-block execution model, and an optional MathDx setup for the glass::nvidia:: backend.

Installation

API reference

The L1 / L2 / L3 and NVIDIA device functions, generated from the header doc-comments via Doxygen + Breathe.

API Reference

Quick start

#include "glass.cuh"

// One block solves one problem; threads stride over the data.
__global__ void saxpy_kernel(uint32_t n, float a, float *x, float *y) {
    glass::axpy(n, a, x, y);          // y = a*x + y
}

saxpy_kernel<<<1, 256>>>(n, 2.0f, d_x, d_y);

See Quickstart for a complete, compilable example, and Worked Examples for a worked program per concept.

Measured performance

The measured native and NVIDIA thread / warp / block ladder on an RTX 5090 (sm_120) — each op’s fastest interface across problem size, in ns/problem (the data behind glass::recommend<>), shown here in the NPROB=8192 throughput regime:

See Backend Sweep Results for the same ladder across the NPROB=64 / 1024 / 8192 batch regimes (the winner shifts with batch size), a narrowly configured host-batched cuBLAS/cuSOLVER comparison, the fused riccati_gain case study, and the per-(op, N) winner table; see Tuning for Your Hardware to regenerate everything for your own GPU with bench/tune.py.