GLASS Code Documentation#
For the research overview, paper, figures, and robotics case studies, visit the GLASS project page. These pages document installation, numerical contracts, tuning, and the API.
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.
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.
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.
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.
glass::nvidia::block::CUB / cuBLASDx / cuSOLVERDx at block scope. Choose this when the measured vendor implementation wins at your size (needs NVIDIA MathDx).
glass::nvidia::warp::CUB WarpReduce L1 reductions, one full 32-lane warp per problem and
explicit per-warp scratch.
glass::nvidia::thread::cuSOLVERDx 0.4+ LAPACK, one independent packed problem per CUDA thread; no dynamic shared scratch or block-wide synchronization.
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 (Historical Archive) for dated, configuration-specific
results.
Header-only install, the single-block execution model, and an optional
MathDx setup for the glass::nvidia:: backend.
The L1 / L2 / L3 and NVIDIA device functions, generated from the header doc-comments via Doxygen + Breathe.
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.
Paper results#
The release evaluates GLASS on Jetson AGX Orin, Jetson AGX Xavier, and RTX 5090. See Paper Results for the current measurements, comparison populations, and release figures. The dated Backend Sweep Results (Historical Archive) archive preserves earlier measurements.