Unify execution scopes and add NVIDIA thread backend

.github/scripts/coverage_obligations.py

@@ -1,5 +1,5 @@
 #!/usr/bin/env python3
-"""Check the 20 declared behavioral obligations (test/coverage-obligations.json)
+"""Check the declared behavioral obligations (test/coverage-obligations.json)
 have passing evidence in collection and the signed receipt. This is an
 obligation-presence check, not line or semantic coverage."""
 

.gitignore

@@ -18,6 +18,7 @@ docs/HANDOFF.md
 
 # Raw per-GPU sweep dumps (regenerate via bench/run_*sweep.sh; distilled results are the tracked *.md)
 bench/mega_sweep_*.txt
+bench/nvt_valid_*.txt
 bench/warp_block_sweep_*.txt
 bench/paper_*.txt
 

CHANGELOG.md

@@ -6,8 +6,28 @@ All notable changes to GLASS will be documented here. The format follows
 
 ## [Unreleased]
 
+### Added
+
+- `glass::nvidia::thread::` provides one-problem-per-thread cuSOLVERDx
+  implementations for the supported LAPACK subset. The measured sm_120 and
+  sm_87 ladders select it only where it clears the native implementation by
+  the tuning margin.
+
 ### Changed
 
+- **Breaking:** NVIDIA operations now require an explicit execution scope
+  (`glass::nvidia::block::op`, `glass::nvidia::warp::op`, or
+  `glass::nvidia::thread::op`); the ambiguous bare `glass::nvidia::op`
+  re-export is removed.
+- **Breaking:** the flat `backend` and `suggested_*` advisor family is replaced
+  by `glass::recommend<op,T,dims...>()`, which returns one `execution_plan`
+  containing implementation family, execution scope, and launch packing.
+  Measured architectures carry paired MathDx and native-only tables.
+- **Breaking:** advisor operation names use `op::potrf` and
+  `op::ldlt_solve`, and rectangular GEMM shape arguments use conventional
+  `(M,N,K)` order.
+- `GLASS_TARGET_SM` is the single architecture selector for native dispatch,
+  execution plans, and MathDx descriptors; legacy `SMS` remains an input alias.
 - `glass::block::riccati_gain` reuses the symmetric `P·B` intermediate
   (`BᵀPA = (PB)ᵀA`), shrinking its shared-scratch requirement from
   `NU²+NX²` to `NU²+NX·NU` scalars and speeding up the fused gain solve.
@@ -18,9 +38,9 @@ All notable changes to GLASS will be documented here. The format follows
   (numerically identical, fewer barriers per iteration).
 - Compile-time `ger` uses a flat one-thread-per-output work mapping
   (bit-identical results by construction).
-- Documented-overload coverage is measured per overload contract — 662
+- Documented-overload coverage is measured per overload contract — 665
   contracts extracted from Doxygen XML with call-shape evidence
-  (`test/api-contracts.json`) — instead of per public name, plus 19 declared
+  (`test/api-contracts.json`) — instead of per public name, plus 21 declared
   behavioral correctness obligations checked against the signed receipt.
 - The signed GPU receipt is split into eight dependency-scoped shards with
   lazy test-binary compilation; development reruns only affected shards, and

CLAUDE.md

@@ -21,9 +21,10 @@ documentation, start with `README.md`.
   block, warp-0, or thread-0 body for measured compile-time cells.
 - `glass::warp::` owns one problem per full 32-lane warp.
 - `glass::thread::` owns one problem per thread. It is compile-time-only,
-  branch-free, and intended for register-resident sizes up to `N <= 7`.
-- `glass::nvidia::block::` and `glass::nvidia::warp::` expose vendor-backed
-  implementations. MathDx is optional; CUB ships with CUDA.
+  branch-free, and is usually strongest while its operands remain register-resident.
+- `glass::nvidia::block::`, `glass::nvidia::warp::`, and
+  `glass::nvidia::thread::` expose vendor-backed implementations. The thread
+  surface requires cuSOLVERDx 0.4+; MathDx is optional and CUB ships with CUDA.
 - `glass::cgrps::` is a cooperative-groups spelling of the block algorithm, not
   an independently tuned backend.
 
@@ -173,7 +174,7 @@ cd docs && PATH="$(cd .. && pwd)/.venv/bin:$PATH" make all SPHINXOPTS="-W --keep
 
 `release.sh` requires a clean `main`, a changelog entry, a complete
 public-overload manifest (every documented overload compile-covered — an
-overload metric, not line coverage), all 20 declared behavioral obligations
+overload metric, not line coverage), all 21 declared behavioral obligations
 passing, a fresh full GPU receipt, and local receipt verification before it
 tags and pushes.
 

README.md

@@ -38,19 +38,22 @@ 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:
+GLASS separates two choices: **execution scope** (thread, warp, or block) and
+**implementation family** (dependency-free GLASS code or optional NVIDIA
+libraries). Pick a scope based on how many independent problems share a block,
+then a family based on availability and measurement:
 
 | 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` |
+| `glass::thread::` (**Thread**) | **thread** | Sequential branch-free subset, one compile-time problem per thread; usually register-resident around `N≤7`, but correct and measured beyond that point | inline in the base headers (via `glass.cuh`) |
+| `glass::nvidia::block::` (**Nvidia block**) | block | CUB + cuBLASDx + cuSOLVERDx; compile-time-size block implementations and queries | `glass-nvidia.cuh` |
+| `glass::nvidia::warp::` (**Nvidia warp**) | warp | CUB `WarpReduce` L1 reductions, one full 32-lane warp per problem | `glass-nvidia.cuh` |
+| `glass::nvidia::thread::` (**Nvidia thread**) | thread | cuSOLVERDx 0.4+ LAPACK, one packed problem per CUDA thread; no shared scratch or block barrier | `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,
+**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
 `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
@@ -62,8 +65,28 @@ 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
+NVIDIA calls always name their scope: `glass::nvidia::block::op`,
+`glass::nvidia::warp::op`, or `glass::nvidia::thread::op`. There is no bare
+`glass::nvidia::op` alias because those scopes require different launches.
+
+When a launcher or code generator wants the measured placement, include
+`glass-defaults.cuh` and make one query:
+
+```cpp
+constexpr auto plan = glass::recommend<glass::op::potrf, float, N>(
+    glass::dependency_set::mathdx);
+// plan.implementation, plan.execution_scope, plan.block_threads,
+// plan.problems_per_block, plan.shared_bytes
+```
+
+`recommend()` is host/code-generation metadata, not an implicit dispatcher:
+the caller still launches and calls the explicit namespace selected by the
+plan. `native_only` is the default dependency set; measured architectures ship
+a paired native-only table from the same capture. Family and scope are measured;
+the packing fields are ready-to-use legal defaults that callers may retune.
+
+> **Note:** `glass::cgrps::` (header `glass-cgrps.cuh`) is a cooperative-groups
+> *adapter* for the Block interface — the same SIMT loop indexed via a
 > `cooperative_groups::thread_group`, numerically identical and **not** a separately-tuned
 > backend.
 
@@ -123,7 +146,9 @@ need NVIDIA MathDx (cuBLASDx / cuSOLVERDx) and extra flags:
 | `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`).
+The NVIDIA wrappers auto-detect availability (`GLASS_HAVE_CUBLASDX`,
+`GLASS_HAVE_CUSOLVERDX`, and `GLASS_HAVE_CUSOLVERDX_THREAD`; the last requires
+cuSOLVERDx 0.4+).
 Full setup, linking, and the MathDx download are in [`bench/INSTALL.md`](bench/INSTALL.md) and
 the [installation guide](docs/source/user_guide/getting_started/installation.rst).
 
@@ -164,10 +189,10 @@ The README is a landing page; the deep reference lives in the
 
 | Topic | Page |
 |-------|------|
-| API reference (L1 / L2 / L3 / nvidia / warp / banded) | [`api_reference/`](docs/source/api_reference/) |
+| API reference (L1 / L2 / L3 / NVIDIA / thread / warp / execution plans / banded) | [`api_reference/`](docs/source/api_reference/) |
 | Namespaces, naming rules, and the two-axis taxonomy | [`concepts/namespaces.rst`](docs/source/user_guide/concepts/namespaces.rst) |
 | Choosing a backend + tuning for your hardware | [`concepts/tuning.rst`](docs/source/user_guide/concepts/tuning.rst) |
-| `glass::nvidia::gemm` cuBLASDx-vs-SIMT dispatch | [`concepts/backend_dispatch.rst`](docs/source/user_guide/concepts/backend_dispatch.rst) |
+| `glass::nvidia::block::gemm` cuBLASDx-vs-SIMT dispatch | [`concepts/backend_dispatch.rst`](docs/source/user_guide/concepts/backend_dispatch.rst) |
 | `TRAILING_SYNC` and barrier conventions | [`concepts/trailing_sync.rst`](docs/source/user_guide/concepts/trailing_sync.rst) |
 | Contraction-parallel (`*_reduced`) family | [`concepts/contraction_parallel.rst`](docs/source/user_guide/concepts/contraction_parallel.rst) |
 | Block-tridiagonal layout (`bdmv` / `pcg`) | [`concepts/block_tridiagonal.rst`](docs/source/user_guide/concepts/block_tridiagonal.rst) |
@@ -176,21 +201,23 @@ The README is a landing page; the deep reference lives in the
 
 ## 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.
+- **Scope determines placement.** Block forms use one block per problem; warp
+  forms pack one problem per full warp. `glass::thread::` and
+  `glass::nvidia::thread::` use one problem per CUDA thread
+  (`<<<ceil(P/TPB), TPB>>>`).
 - **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
+  per-matrix `ROW_MAJOR` flag (its transpose changes the math op); `glass::nvidia::block::` 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
+- `glass::nvidia::block::*` (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`.
+- cuSOLVERDx TRSM has no native non-1.0 `alpha`; both NVIDIA wrappers pre-scale
+  `B` before `execute` (in shared memory at block scope, in the owning thread at
+  thread scope).

bench/INSTALL.md

@@ -11,7 +11,7 @@ benches. All three are auto-detected by `bench/run_bench.py`.
 
 | Library | Required for | Header-only? |
 |---------|-------------|--------------|
-| CUB    | `bench_reduce` (CUB baseline + `glass::nvidia::reduce` variant) | Yes (bundled with CUDA) |
+| CUB    | `bench_reduce` (CUB baseline + `glass::nvidia::block::reduce`) | Yes (bundled with CUDA) |
 | cuBLASDx | `bench_gemv`, `bench_gemm`, `bench_blockdim`, `bench_gemm_batched`, `bench_lapack` | Yes |
 | cuSOLVERDx | `bench_lapack` (Cholesky / TRSM / posv / etc.) | **No** — links a precompiled device fatbin |
 
@@ -41,8 +41,9 @@ portal.
    (Free NVIDIA Developer account required.)
 
 2. Choose: **MathDx for CUDA 12, Linux x86_64** (`.tar.gz` format).
-   Version 25.12.x or later is recommended (this is the version the GLASS
-   wrappers are tested against).
+   Version 25.12.x or later supports the block wrappers. Use **26.03 or later**
+   for `glass::nvidia::thread::` (cuSOLVERDx 0.4+); the thread correctness suite
+   is validated against 26.03.
 
 ### Install
 
@@ -54,7 +55,7 @@ tar -xzf MathDx_*.tar.gz -C /opt
 ls /opt/nvidia/mathdx/
 
 # Set the environment variable (add to ~/.bashrc to persist)
-export MATHDX_ROOT=/opt/nvidia/mathdx/25.12   # adjust version as needed
+export MATHDX_ROOT=/opt/nvidia/mathdx/26.03   # adjust version as needed
 ```
 
 ### Verify

bench/JETSON.md

@@ -4,7 +4,8 @@ One capture per box (Orin AGX / Orin NX / Orin Nano — all `sm_87`). Each run
 produces a single `bench/jetson_<host>_<ts>.tar.gz` containing the timings
 plus a full device/JetPack provenance bundle; send those back for ingestion
 (figure columns + the `ideal_sm87` ladder table, spliced off-box via
-`tune.py --from-ladder --sm 870`).
+`tune.py --from-ladder --from-nvt-valid --sm 870` when MathDx is available;
+native-only captures need no confirmation companion).
 
 ## 0. Pre-flight: get on the latest JetPack the box supports
 
@@ -124,7 +125,8 @@ methodology). The script:
 
 | Capture | Feeds |
 |---|---|
-| `mega_sweep_*.txt` | `ideal_sm87` ladder table (spliced on the desktop via `python bench/tune.py --from-ladder <txt> --sm 870 --allow-no-mathdx`), paper §portability "which crossovers moved" |
+| `mega_sweep_*.txt` | `ideal_sm87` ladder table. Full MathDx replay: `python bench/tune.py --sm 870 --legs ladder --from-ladder <mega> --from-nvt-valid <nvt>`; a native-only capture needs only `--from-ladder <mega> --allow-no-mathdx`. |
+| `nvt_valid_*.txt` | Required valid-input veto companion when the ladder selects NVIDIA thread; pass it with `--from-nvt-valid` during off-box regeneration. |
 | `paper_hostblas_*.txt` | Jetson columns for the hostblas + latency figures |
 | `paper_fusion_*.txt` | Jetson fusion curves |
 | `robotics_sweep_*.txt` | Jetson robotics tier panels |