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, architecture-tuned for edge robotics and beyond. 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

GLASS ships measured placement tables for three architectures — RTX 5090 (sm_120), Jetson AGX Orin (sm_87), and Jetson AGX Xavier (sm_72, native-only) — because placement is worth measuring: on the Orin’s raw candidate ladder the best and worst placements for a cell differ by a median of 4.9× (up to 81×), and the recommended placement changes between the Orin and the RTX 5090 in 145 of 396 measured cells (162 of 396 versus the Xavier). The stakes are highest on embedded GPUs, where per-call host dispatch is most expensive relative to the work — GLASS was tuned edge-first.

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

See Backend Sweep Results for the RTX 5090 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.