Status: current · RTX 4090 box and threadripper02 unless a line says otherwise · moved here from the root AGENTS.md on 2026-10-07; .claude/CLAUDE.md imports it.

Terse working rules for agents and contributors: toolchain, build performance, testing policy, routing architecture, measurement rules, GPU kernel facts and repository conventions.

TL;DR. Build dependencies: CMake >=3.17 (>=3.21 for cmake --preset), a C++20 compiler with SYCL 2020 support and a backend for your GPU vendor; optionally netlib LAPACKE + CBLAS (the NETLIB host backend and the LAPACKE test oracles). Runtime needs the same plus a GPU the SYCL runtime exposes. The exercised configuration is a self-built CUDA-enabled DPC++ (/opt/dpcpp-cuda) against CUDA 13.2 and an RTX 4090 (sm_89) on Ubuntu 22.04; see "Tested platforms" in docs/guide/build.md.

Apple GPU (macOS): a separate AdaptiveCpp Metal build (presets acpp-metal / acpp-metal-tests, toolchain from scripts/macos/build_adaptivecpp_metal.sh, M5 Max, macOS 27). GPU = Backend::SYCL, float and complex<float> only; double runs on a CPU Queue via Homebrew OpenBLAS; vendor-free. Recipe: "macOS: Apple GPU through AdaptiveCpp Metal" in docs/guide/build.md; design: docs/design/sycl-implementations.md#sycl-impl-metal; traps below.

Warning: NVIDIA targets need a CUDA-capable DPC++. The stock intel-oneapi-compiler-dpcpp-cpp package has no CUDA adapter: on an NVIDIA machine it configures cleanly and builds a CPU-only library. Use Intel oneAPI plus the Codeplay oneAPI for NVIDIA GPUs plugin, or a self-built intel/llvm configured with --cuda.

sycl-ls must list a [cuda:gpu] entry.

-- Using SYCL targets: spir64_x86_64 in the configure output means a CPU-only build, whatever GPUs are present. A CUDA configure names the architecture, e.g. nvidia_gpu_sm_89.

-DBATCHLAS_ENABLE_CUDA=ON means "require it": without a [cuda:gpu] the configure aborts with a FATAL_ERROR naming the missing entry (see "CMake options" in docs/guide/build.md). The default AUTO silently builds CPU-only.

1. Prerequisite Packages

Component	Debian/Ubuntu (apt)	Fedora/RHEL-like (dnf)	Arch Linux (pacman)	Source build (fallback)
BLAS/LAPACK (Fortran APIs)	libblas-dev liblapack-dev	blas-devel lapack-devel	blas lapack	see §4
C interface (CBLAS & LAPACKE headers)	liblapacke-dev	lapack-devel	lapacke	see §4
Build tools	build-essential cmake git	@development-tools cmake git	base-devel cmake git	—
Git LFS (raw benchmark results; docs/perf/README.md)	git-lfs	git-lfs	git-lfs	—

libopenblas-dev alone is not enough: Ubuntu's OpenBLAS package omits lapacke.h, so install liblapacke-dev too (a packaging decision, not a BatchLAS bug).

2. Installing a SYCL Compiler

2a. Intel oneAPI DPC++/C++ (icpx): Intel GPUs and CPU-only builds

The easy path, and the WRONG path for NVIDIA (no CUDA adapter; see the TL;DR and §2b).

# 1. Add Intel's APT repo and key (root)
wget -O- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB |
  sudo gpg --dearmor -o /usr/share/keyrings/oneapi-archive-keyring.gpg
echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" \
  | sudo tee /etc/apt/sources.list.d/oneAPI.list
 
sudo apt update
 
# 2. Install the minimal SYCL compiler package
sudo apt install intel-oneapi-compiler-dpcpp-cpp      # SYCL 2025.x
 
# (optional) Classic compilers for C/C++ & Fortran
sudo apt install intel-oneapi-compiler-dpcpp-cpp-and-cpp-classic \
                 intel-oneapi-compiler-fortran

Per shell session:

source /opt/intel/oneapi/setvars.sh # sets PATH, LD_LIBRARY_PATH, MKLROOT, etc.

The single compiler package is enough for CPU or Intel GPU; the full intel-basekit is not needed.

2b. CUDA-capable DPC++: NVIDIA GPUs

Both options give a clang-family clang++/icpx that emits NVPTX:

Codeplay oneAPI for NVIDIA GPUs: install oneAPI as in §2a, then the matching Codeplay plugin. It adds the CUDA UR adapter, after which sycl-ls reports [cuda:gpu].
Self-built intel/llvm: clone https://github.com/intel/llvm, configure with --cuda (plus --cmake-opt=-DCMAKE_INSTALL_PREFIX=<prefix>). This machine uses it at /opt/dpcpp-cuda; every number in the repository was measured with it.

Also install a CUDA Toolkit (13.2 here) and pass the compiler explicitly:

cmake -S . -B build -DCMAKE_CXX_COMPILER=/opt/dpcpp-cuda/bin/clang++

A consumer of an installed BatchLAS must configure the whole consuming project with this same compiler ("Building against BatchLAS" in docs/cpp-api.md).

2c. AdaptiveCpp: the second SYCL implementation

BatchLAS also builds with the AdaptiveCpp 25.10 release (BATCHLAS_SYCL_IMPL=ACPP; design Two SYCL implementations: DPC++ and AdaptiveCpp). Presets acpp and acpp-tests set CMAKE_CXX_COMPILER=/opt/adaptivecpp-25.10.0-cuda13.2/bin/acpp (trees build-acpp, build-acpp-tests). The hosted CI job acpp-compile compiles every target with --acpp-targets=generic and CUDA off; it runs nothing. An acpp test tree is gated against its own ledger, tests/known-failures-acpp.txt (compare_failures.py --sycl-impl ACPP, or run_local_checks.sh on an acpp tree).

3. Verifying the Toolchain

icpx --version          # or: /opt/dpcpp-cuda/bin/clang++ --version
cmake --version         # ≥3.17 (≥3.21 for cmake --preset)
sycl-ls                 # must list your GPU, e.g. a [cuda:gpu] entry
pkg-config --exists lapacke cblas && echo "LAPACKE & CBLAS found"

4. Building netlib LAPACKE/CBLAS from Source (if distro packages are unavailable)

git clone https://github.com/Reference-LAPACK/lapack.git
cd lapack
cmake -B build -S . -DCMAKE_BUILD_TYPE=Release -DLAPACKE=ON -DBUILD_SHARED_LIBS=ON
cmake --build build -j $(nproc)
sudo cmake --install build   # installs liblapacke.so, libcblas.so, headers

Add the install prefix (e.g. /usr/local/lib) to LD_LIBRARY_PATH, or run sudo ldconfig.

5. CMake Configuration Hints

There is no find_package for SYCL: the compiler is the SYCL implementation. LAPACKE and CBLAS are found with find_library (find_netlib_libs in cmake/BatchLASDependencies.cmake): lapacke in the Debian multiarch directories first, then the default search path; CBLAS (cblas, else blas) from LAPACKE's own directory first. Both are optional: without them configure warns and disables the NETLIB host backend. A library outside those paths is named explicitly:

cmake -S . -B build -DLAPACKE_LIBRARY=/opt/lib/liblapacke.so -DCBLAS_LIBRARY=/opt/lib/libcblas.so

6. Quick Smoke Test

cmake -B build .
# GPU-only iteration tree (nvptx64 only, 13-25% faster; never the gate, §7):
#   cmake -B build-gpu . -DBATCHLAS_CPU_TARGET=none
 
# Iterating on one algorithm: this builds the library and only this one test
# binary. Do not build the default target while iterating — it also builds every
# other test executable (`ctest -N` lists them), which you are not about to run.
cmake --build build --target stedc_tests -j8
ctest --test-dir build -R '^stedc_tests$' --output-on-failure
 
# Before pushing, build and run everything. The final gate is a full-target
# tree (not BATCHLAS_CPU_TARGET=none), plus the vendor-free tree
# (-DBATCHLAS_ENABLE_VENDOR_BLAS=OFF) built only now. One test per GPU slot (§8)
# where ctest_gpus.sh has been validated (threadripper02); on the 2x4090 box, and
# wherever configure printed "GPU test resources: none", keep it serial.
cmake --build build -j8
scripts/ctest_gpus.sh --test-dir build      # threadripper02
ctest --test-dir build                       # 2x4090 box (OOM history, docs/ci.md)
 
# Also before pushing, if you touched cmake/, include/ or the install rules:
# the packaging gate. These are the same checks CI runs (.github/workflows/ci.yml).
sh .github/ci/run_local_checks.sh
 
# The --package mode is the one CI cannot run, because it needs a real install
# from a real SYCL build. Give run_local_checks.sh a prefix and it adds it:
cmake --install build --prefix /tmp/batchlas-prefix
sh .github/ci/run_local_checks.sh /tmp/batchlas-prefix
 
# The end-to-end version of the same thing — install, then configure, build and
# run examples/consumer/ as a standalone outside project:
ctest --test-dir build -R '^consumer_package_tests$' --output-on-failure

Known Pitfalls

Silent CPU-only build on an NVIDIA box: the compiler has no CUDA adapter; -- Using SYCL targets: spir64_x86_64 is the tell (TL;DR, §2b).
icpx (oneAPI 2026) defaults to -fp-model=fast, which stops LLVM inlining libspirv builtins: sycl::fma, barriers and id queries become out-of-line CALLs in ~98% of kernels (cfloat gemm 6-15x, potrf cfloat 6x slower) with every test green. The build passes -ffp-model=precise for IntelLLVM; ctest -R '^device_calls_tests$' (scripts/check_device_calls.py) scans the built PTX and fails on any __spirv_* call. Never add a later -ffp-model=fast/-ffast-math (the last flag wins). benchviz warns about such a build.
oneAPI 2026 icpx on an NVIDIA box without the Codeplay plugin: oneAPI's libsycl ships no CUDA UR adapter. Put /opt/dpcpp-cuda/lib FIRST on LD_LIBRARY_PATH, ahead of setvars.sh, for configure, build and every run. Otherwise sycl-ls shows no [cuda:gpu] and configure silently picks spir64_x86_64.
Missing lapacke.h: install liblapacke-dev even with libopenblas-dev.
Multiple BLAS providers: sudo update-alternatives --config libblas.so.3.
Device not found at runtime: the GPU driver and its SYCL adapter must match the compiler version (Level-Zero runtime for Intel, CUDA adapter and driver for NVIDIA).
libsycl.so.9: cannot open shared object file: the DPC++ runtime is not on the loader path. Export LD_LIBRARY_PATH=<dpcpp-prefix>/lib, or add it to /etc/ld.so.conf.d and run ldconfig.
Wrong double results from the host (NETLIB) backend: some OpenBLAS builds pick a broken dgemm kernel (0.3.20 on Cooperlake Xeons; errors ~66 at n=128). Configure runs a dgemm health check (cmake/BatchLASBlasHealthCheck.cmake) and ctest sets OPENBLAS_CORETYPE; a bare ./build/tests/foo does not. For a double-only CPU failure, re-run with OPENBLAS_CORETYPE=SKYLAKEX before blaming BatchLAS.
-fopenmp binds to the NVIDIA HPC SDK's libgomp when it is on LIBRARY_PATH (symbol version literally VERSION), then fails at load time. Link with env -u LIBRARY_PATH; readelf -d <lib> | grep gomp should show libgomp.so.1.
nvidia-smi says "Driver/library version mismatch": unattended-upgrades replaced the driver under the running kernel. Compare cat /proc/driver/nvidia/version with modinfo nvidia | grep ^version; only a reboot (or module reload) fixes it. Not a BatchLAS problem.
A CUDA binary that "hangs at startup" at 100% CPU+GPU is usually a stuck kernel, not JIT. Run SYCL_UR_TRACE=2 ./bin > trace.log 2>&1 (no pipe), map urKernelCreate handles to names, read the last urEnqueueKernelLaunch. compute-sanitizer cannot attach (SYCL calls cuInit first); for GPU out-of-bounds, rebuild with -DCMAKE_CXX_FLAGS_RELWITHDEBINFO="-O1 -g -UNDEBUG" to re-enable the device asserts in VectorView::at / KernelMatrixView.

Metal (macOS, AdaptiveCpp; M5 Max). Design and limits: docs/design/sycl-implementations.md#sycl-impl-metal.

ACPP_ADAPTIVITY_LEVEL=0 breaks every float kernel: level 0 JITs a TU's kernels together, double ones included, and Metal cannot compile double. A GPU Queue throws unsupported at level 0 (backend_dispatch_tests_adaptivity0). Leave it unset (1) or higher.
A double call on the Apple GPU throws unsupported by design (no fp64). Use Queue(Device("cpu")). set_backend(NETLIB) on the GPU throws too.
Monolithic only on macOS: split component dylibs have cyclic undefined symbols and ld64 rejects them; configure fails without -DBATCHLAS_MONOLITHIC_LIBRARY=ON. Test trees also need BATCHLAS_HIDDEN_VISIBILITY=OFF (the presets set both).
Homebrew OpenBLAS has no liblapacke: LAPACKE and CBLAS are inside libopenblas.dylib. BATCHLAS_OPENBLAS_ROOT (probed at /opt/homebrew/opt/openblas) finds it; frameworks are off for that search so vecLib's cblas.h never pairs with it.
Keg-only paths: llvm@21, lld@21, libomp and openblas are not on default search paths. Use the /opt/homebrew/opt/<pkg> paths, never Cellar/<version>.
brew upgrade can break acpp: ~/opt/adaptivecpp/etc/AdaptiveCpp/acpp-core.json holds the LLVM and libomp paths; a Cellar/ path there dangles after an upgrade (the build script warns). Rebuild with the script, or brew pin llvm@21 libomp.
Clear ~/.acpp/apps after reinstalling AdaptiveCpp: the JIT cache outlives the toolchain, so a patched translator otherwise runs only on kernels it has not cached. ACPP_APPDB_DIR=<dir> gives a run its own cache.
The patches are local (scripts/macos/adaptivecpp-patches/, branch batchlas-metal). The script reuses an existing branch and stops on a patch-count mismatch: ACPP_REAPPLY=1 re-applies.
make install resets the acpp shebang to #!/usr/bin/env python3. When /usr/bin precedes Homebrew on PATH that is the Command Line Tools shim, which exports CPATH=/usr/local/include and LIBRARY_PATH=/usr/local/lib into every compile: a /usr/local GoogleTest then shadows the fetched one and tests fail to link. The script pins acpp/syclcc/syclcc-clang to Homebrew's python3; re-run it (or its sed) after any reinstall. env -u CPATH -u LIBRARY_PATH is the belt-and-braces for configure, build and ctest.
DYLD_LIBRARY_PATH, not LD_LIBRARY_PATH, selects dylibs (e.g. a test runtime prefix: DYLD_LIBRARY_PATH=<acpp-prefix>/lib). SIP strips DYLD_* from /usr/bin interpreters and shell scripts run through them.
ACPP_VISIBILITY_MASK=omp;metal (set by the presets): double and batchlas::verify need the OpenMP device, and configure fails without it.
Debug: ACPP_DEBUG_LEVEL=3 logs every Metal dispatch ("Dispatching kernel"); __ACPP_PRINT_METAL_CODE=1 prints the generated MSL (with IR comments); __ACPP_PRINT_IR_BEFORE_EMIT=1 prints the LLVM IR the emitter receives (look for <4 x, double, llvm.nvvm). A GPU fault under coarse events covers many ops: BATCHLAS_METAL_FLUSH_EACH_OP=1 commits one command buffer per op to attribute it.
A data-dependent loop around barriers never ends on the OpenMP device (the host NETLIB instantiations run their native kernels there), even at batch 1, unless the loop condition is a group collective: tridiagonal_solver broadcasts it (sycl::group_broadcast) for NETLIB under acpp only. A host test that hangs at 100% CPU in __acpp_sscp_kernel is this.
Float subnormals flush to zero on the Apple GPU, though the device reports denorm support. A test that needs one must probe the device and skip.
Routing on Apple is borrowed: no tuned/*.apple_m5_max.txt, so every op borrows the nearest table (sm_120 here) and warns once. Every window and constant is an sm_89/sm_120 hypothesis there.
Run ctest serially on the Mac (ctest --preset acpp-metal-tests sets -j1); its failure ledger is tests/known-failures-acpp-metal.txt.

7. Build Performance

The build is SYCL device-link-bound, not compile-bound. The unit of device linking is the shared library: changing one object re-runs sycl-post-link + ptxas + CPU-target AOT for every object in that .so, single-threaded, so -j cannot shorten it. Measured on threadripper02 (64 cores, 4x RTX PRO 6000, default target = library + 82 test binaries, -j32, other builds sometimes running): fresh tree, no cache, 284-334 s.

Cap build parallelism at -j8 unless you have measured headroom. Heavy SYCL TUs exceed 14 GB each, and wide builds (-j20, -j28, or two -j builds side by side) have been OOM-killed several times. Run one build at a time; add a second only after checking free -g.
ccache is on by default (BATCHLAS_USE_CCACHE, cmake/BatchLASCcache.cmake) when ccache is on PATH or in ~/.local/bin; configure prints ccache enabled (...). No ccache: drop the static binary from https://github.com/ccache/ccache/releases into ~/.local/bin (no root).
- The generated <build>/batchlas-ccache wrapper sets depend mode, CCACHE_BASEDIR (deepest common parent of source and build dir, BATCHLAS_CCACHE_BASEDIR), CCACHE_NOHASHDIR and sloppiness, so a second tree hits the first tree's cache: 187/196 hits, 203 s vs 284 s cold (another pair: 196/196, 176 s vs 329 s). Misses are TUs with an absolute path in a -D.
- Plain CMAKE_CXX_COMPILER_LAUNCHER=ccache without those settings got 3/196.
- Links are never cached: even an all-hits fresh tree pays every device and test link.
Never configure a fresh tree per stage when an existing tree can be reused: reconfigure and rebuild it.
Iterate GPU-only: -DBATCHLAS_CPU_TARGET=none (or the dev-gpu / dev-gpu-tests presets) builds nvptx64 only: 248 s cold against full-target colds of 284-334 s on the same box (13-25%; not a back-to-back pair), and every relink is cheaper. The NETLIB/CPU instantiations need the CPU target, so a GPU-only tree is not a gate: the final pre-push run uses full targets (dev-tests/cuda or a default configure), and the vendor-free tree (-DBATCHLAS_ENABLE_VENDOR_BLAS=OFF) is built only at that final gate.
Build only the target you work on (--target stedc_tests). Every test target is in the default target, so a plain build links them all.
An incremental relink of libbatchlas_backends after a one-line select.cc edit is ~17 s, whatever BATCHLAS_SYCL_LINK_JOBS is (4: 17 s, 16: 18 s). Do not tune it.
cmake --build --target A B C with many targets can exit 0 and still skip some links. Check test -x build/tests/<name>, not the exit code.
Measured and ruled out (do not retry): Ninja (the Makefile null build is already 0.07 s), lld, PCH (the driver cannot emit one), unity builds (kernel-name collisions), -fno-sycl-rdc, explicit -fsycl-device-code-split, cutting one edge of a header cycle (the umbrella header cuts both edges; docs/design/build-performance.md).
-fsyntax-only together with -fsycl-targets is rejected by this driver. Use real -c compiles as the oracle.
To time a real rebuild, append a line rather than touch: with ccache a touched file is a cache hit and the timing measures only the link.
ROCm TUs can be syntax-checked without an AMD GPU: scripts/rocm_syntax_check.sh (headers are in /opt/rocm/include/roc*/). Run it after touching src/backends/roc*.cc or any signature or instantiation the ROCm TUs also spell out.
Tiny-kernel TUs build with -mllvm -pragma-unroll-threshold=262144. Without it LLVM silently declines #pragma unroll, register arrays go to the stack, and tests stay green.

8. Testing Policy

Do not run the full ctest by default. Scope to what you changed: one case ./build/tests/X --gtest_filter=... -> one binary ctest -R '^X$' (-R is a substring regex, so anchor it) -> one component -L util|blas|ortho|tridiag|eig|sparse -> -LE slow -> full only before pushing or after touching shared code (Queue, Matrix/MatrixView, the mempool, sg_partition, include/batchlas/util).
On threadripper02, run multi-test scopes through scripts/ctest_gpus.sh (same arguments as ctest, e.g. scripts/ctest_gpus.sh -LE slow; --test-dir <build> first, default build/).
- Configure writes <build>/ctest_resources.json from nvidia-smi --list-gpus (BATCHLAS_TEST_GPUS=<n>, =0 disables; BATCHLAS_TEST_GPU_SLOTS, default 2). Every GPU test has RESOURCE_GROUPS gpus:1, and tests/ctest_gpu_env.sh sets CUDA_VISIBLE_DEVICES to its slot's GPU. A pre-set CUDA_VISIBLE_DEVICES=1 trims the run to that one GPU.
- Only when configure's sycl-ls lists a [cuda:gpu]; otherwise no spec is written and you run plain ctest, which ignores all of this.
- syev_cta_tests and consumer_package_tests are RUN_SERIAL (flake; three parallel consumer builds).
- Measured -LE slow (88 tests, threadripper02, 4x RTX PRO 6000): serial with CUDA_VISIBLE_DEVICES=1 469 s, ctest_gpus.sh 1 slot/GPU 126 s, 2 slots/GPU 78 s (critical path is gemm_candidates_tests alone; 115 s for the 96 tests after #147); failing names identical in all modes.
- Validated only on threadripper02 (64 cores, ample RAM). The 2x4090 box (20 cores, 125 GB, zero swap) was OOM-killed twice by a build or second ctest beside ctest (docs/ci.md#oom-kill-during-ctest); keep its full gate serial until ctest_gpus.sh is measured there.
Observed on threadripper02 only, cause unknown: a process that sees every GPU is slow. trsm_candidates_tests took 132 s with 4 GPUs visible and 18 s with one (one case: 0.37 s with 1, 1.2 s with 2, 5.1 s with 4). ONEAPI_DEVICE_SELECTOR=cuda:N did not help, so the cost is below SYCL device selection (CUDA or UR adapter/platform init); not checked on the 2x4090 box. Running a binary by hand there: set CUDA_VISIBLE_DEVICES=<n>.
BATCHLAS_TEST_BACKEND=CUDA skips the NETLIB/CPU instantiations and BATCHLAS_TEST_FLOAT_TYPE=float skips the other types. Both are GTEST_SKIPs (case list unchanged), but the coverage is silently gone.
main is not green. Accepted failures are in tests/known-failures.txt; CI diffs against that ledger (docs/ci.md). Compare failing test names against a baseline, never pass/fail counts. To attribute a failure: save your diff, git checkout <base> -- <files>, rebuild only that test target, re-run, compare names.
syev_cta_tests has a concurrency-dependent wrong-answer flake (a few percent of runs under ctest -j2 on the 4090 box), so it is RUN_SERIAL (BATCHLAS_RUN_SERIAL_TESTS in tests/CMakeLists.txt). 0 of 15 concurrent runs on threadripper02 does not rule out a few-percent rate. sytrd_blocked_tests also flakes on untouched main. Check the baseline before blaming your change.
A CUDA-off (CPU-only) tree shows many fake failures: unbuilt tests appear as "Not Run", and there is no sub-group size 32, so CTA kernels throw.
Before pushing, if you touched cmake/, include/ or the install rules, run sh .github/ci/run_local_checks.sh. It includes the per-file comment-density gate (§12).

Writing tests: guards that cannot fail. More than a dozen recorded cases of a green suite passing a wrong answer. Every new or rewritten test needs:

A deliberate break of the code under test, once per axis, giving a narrow, named red set. A break that turns everything red is too coarse to prove anything.
The break must run where the code is reachable. If the table sends the shape to the vendor, the native kernel is never exercised. Pin it (BATCHLAS_<OP>_ROUTE=<spelling> or a select::ScopedPin) or use the vendor-free build.
Non-natural values for every accessor the op reads (ld, batch stride, inc, rows!=cols), plus complex data with a nonzero imaginary part, ConjTrans, and beta!=0. The bug to catch is a kernel that derives a value instead of reading it; derived values are correct in the natural case.
Poison the code under test will accept: an in-range index with a large finite value, not NaN or an out-of-range index that the kernel's own guard discards. It must also survive the upstream code (cuSOLVER Upper potrf overwrites the lower triangle).
A saturating-batch case for any shared-local-memory kernel: batch >=1024, every item the same matrix, each result bit-identical to item 0. Small batches cannot race, and work-group-size ladders put every small test on wg=32, the one width where cross-sub-group SLM races are hidden.
Thresholds straddled in both directions. A routing test covering one shape proves nothing about the router.
No self-referential reference paths. A reference that goes through the same tuning tables or routes as the code under test agrees with it by construction.
Graded input (e.g. condition numbers ~1e±6) for eigen/tridiagonal accuracy. GOE-style random matrices cannot see relative-accuracy bugs.
For a capacity guard, a launch at the advertised ceiling, not an arithmetic re-derivation.
To verify a revert of an untracked file, use md5sum against a pristine copy (git diff prints nothing).
Verify results with batchlas::verify; in a test the macro is EXPECT_VERIFY(T, Check::kind, n, value) (tests/test_utils.hh; Verification library). No local residual, reference or tolerance helpers.

9. Architecture: Routing and Dispatch

Flat selection (every op that reads BATCHLAS_<OP>_ROUTE; docs/design/flat-kernel-selection.md). There is no RouteTable and no include/batchlas/blas/dispatch/.

src/ops/<op>/choice.hh lists an op's choices (a std::variant of families with int knobs, e.g. lpanel:panel=8). src/ops/<op>/<op>.cc holds key_of, can_run(), launch and workspace. select::run / select::pick (device, choose, the vendor-free NoRouteError, trace, coverage) and the op's select::OpSpec (Op, vendor Lib, last resort) do the rest.
can_run() is correctness only: false means the launch would throw or give a WRONG ANSWER. Never put a speed threshold there, or the shape becomes unservable vendor-free.
Speed lives only in tables. select::choose (src/select/select.hh) takes the first runnable entry of the nearest row of tuned/<op>.<device>.txt: exact keys match by longest prefix (dropped from the right until a row matches); n/batch are log distance, weighted per key (n:log:3: potrf work ~ n^3 x batch). A device with no table borrows the nearest one and warns once; the CPU never borrows; nothing runnable falls to the op's last resort.
Tables, ledgers, tiers, staleness and generation: Tuning data layout (the only normative page).
BATCHLAS_TUNED_DIR=<dir> overrides tables without a rebuild; BATCHLAS_SELECT_TRACE=1 prints every choice and the runner-up.

Per-op choices:

Op	Choices
trsm	cta, sg_left, blocked, vendor (a cta or sg_left pin above order 32 throws)
posv	tiny/cta/blocked, no vendor family; its children select for themselves
gemm	direct, tiled, small, reg:m=..:n=..:k=..:u=.. (float), wide:m=..:n=..:k=.., vendor; the aligned/predicated leg is derived inside the launcher, never a field (docs/perf/gemm.md#choices-flat-selection-p34)
syrk	gram (n <= 128; float and double), triangular (float), vendor; keys form = sq\|tall\|wide and trans
syr2k	triangular (float; a real ConjTrans runs as Trans), vendor
symm	expand (mirrored expansion + public gemm, float and double, needs expansion_fits), vendor; key form = sq\|tall\|wide
trmm	triangular (Side::Left), expand (expansion + public gemm, either side), vendor (the cublas?trmm loop only; the old vendor pin meant expand+gemm, that is now expand)

Level-3 four (docs/design/flat-kernel-selection.md §12 "Level-3 four"): all NoFields, native families CUDA GPU only, homogeneous batch, batch <= 65535. Every level-3 vendor family refuses a heterogeneous batch on every backend (the loops run each item at the top-level extents), so a heterogeneous call has no route. cublasdx and the old level-3 pin-word list are gone: every cublasdx pin throws as an unknown family.

Env override BATCHLAS_<OP>_ROUTE (settings().routing.route("<op>"), 19 ops) takes auto, native, vendor or a spelling (lpanel:panel=8, lpanel:8), case-folded and trimmed.

A bad pin THROWS (unparsable, not a compiled candidate, or cannot run the shape). Only the class words native and vendor with no runnable choice of their class fall back to Auto, with a warning (so vendor in a vendor-free build means Auto).
No aliases: the old native:cta, vendor:auto, batchlas_cta and gemm kernel names (tiled16, 128x128x8) throw; BATCHLAS_<OP>_VARIANT / _PROVIDER / BATCHLAS_GEMM_SYCL_KERNEL are not read.
Coverage chosen_algo is the spelling.
The value comes from the settings snapshot: a raw setenv mid-process is not seen until detail::reload_settings() (ScopedEnvVar does it; docs/design/environment.md#environment-route-pins-are-raw-strings).

hemm, herk and her2k still choose by hand-written rule (cublas.cc, vendor builds only: expand/fold into the public gemm vs the per-item vendor loop; BATCHLAS_EXPAND_ROUTE=expand|loop pins it; herk's gram opt-in is BATCHLAS_SYRK_ROUTE=gram). They have no BATCHLAS_<OP>_ROUTE and throw NoRouteError vendor-free. Migration is a listed follow-up (docs/design/flat-kernel-selection.md §11).

cuBLASDx/MathDx is deleted (code, MathDx probe, BATCHLAS_ENABLE_CUBLASDX, BATCHLAS_HAS_CUBLASDX, BATCHLAS_GEMM_CUBLASDX_KERNEL); it never ran, MathDx was absent on both boxes.

Entry points and vendor-free builds

Public entry points live in src/ops/<op>/<op>.cc (src/ops/level3/level3.cc for hemm/herk/her2k). Vendor-availability constants and throw_no_vendor_route are in src/select/vendor.hh; batchlas::NoRouteError is in <batchlas/no_route.hh>.
-DBATCHLAS_ENABLE_VENDOR_BLAS=OFF builds a vendor-free library; that is the only valid vendor-free measurement. With nothing runnable an op throws NoRouteError and records a coverage miss row.

Verifying routes

Linked != reached. A symbol in nm, a linked coverage row or a green test count does not show a kernel runs. Set BATCHLAS_COVERAGE_OUT and read the reached rows. scripts/route_diff.sh captures and diffs them; coverage_merge.sh merges per-pid shards. Verify a routing change by diffing the chosen route across the input space, not by timing it.
For a migrated op (gemm included) chosen_algo is the full spelling (reg:m=128:n=128:k=8:u=1), so route_diff sees gemm tile changes. It does not see a derived leg (gemm's aligned vs predicated) or a non-migrated op's internal kernel pick: use BATCHLAS_KERNEL_TRACE=1 (never while timing, it inflates ~60%).
tools/tune --gate against a pre-P3.4 parent reads its gemm coverage as native:register_tiled, which now means "best native": gate gemm from a kernel-trace old-choice CSV.
Sub-ops count: a "native" syev, geqrf or getrf can still call cuBLAS gemm underneath. Every library gemm call (blocked drivers' default seams, the hemm/herk/her2k expansions in cublas.cc, symm's and trmm's expand families) goes through the public gemm and its table (gemm_custom and the gemm_use_sycl_custom re-route are deleted).

Routing bug patterns

Highest-yield perf pattern: a finished candidate that no table row ranks first. Before writing a kernel, list the op's candidates and check each can ever be chosen: is it timed in the tables, does can_run ever admit it? Look for hard return false and for terms naming a different family's type list.
Read the predicate, not the comment. Comments describing routing windows drift. docs/perf/<op>.md quotes each can_run term and window with its file and line number and gives bracketing evidence. Several hundred source comments cite evidence: docs/<page>.md#anchor, so renaming a heading breaks them; check_evidence_anchors.py catches this and requires every cited slug to be unique across docs/ (Doxygen numbers repeated headings site-wide). Give a cited heading distinctive text. The generated Evidence index lists what cites each section.
A kernel's fast-*leg* predicate (one the dispatcher re-evaluates to choose aligned vs predicated) used as a routing gate sends the call to a different, slower kernel. Found twice in select_kernel_variant (1.7x and 3-4x; deleted in P3.4). Under flat selection a leg is derived in the launcher and must never become a family field, a key or a can_run term (layout packed|strided is a table key, not a gate).
Guarding on a device family (BATCHLAS_HAS_CUDA_BACKEND) where the code needs a library is a recurring bug; only a vendor-free build exposes it.

API traps

The variadic queue-dispatch overload (BATCHLAS_DISPATCH_ON_QUEUE) must be requires-constrained. Even constrained, it beats option-struct overloads for non-const lvalue and prvalue arguments, so per-call checks that live only in options.hh are unreachable on the ordinary spelling. Test checks through the ordinary spelling.
Write PotrfOptions{}, never a bare {}: {} selects the positional Uplo{} = Upper.
Workspace sizing: BumpAllocator checks the alignment-rounded size but advances by the raw size, so an exactly simulated size is too small. Size with BumpAllocator::measuring() + required_bytes(). A *_buffer_size and its run path must branch identically. Sizing mode hands out unbacked pointers: never test ws.data() != nullptr to mean "the caller passed a workspace"; use separate overloads.
The workspace arena belongs to the Queue, so a per-call Queue defeats pooling. Leak long-lived SYCL objects held in statics: static destruction after the runtime is gone hangs.
linalg:: value-returning wrappers free their scratch while kernels are still enqueued. linalg::qr is deliberately absent because of an unexplained cross-Queue wrong-answer defect (docs/design/known-defects.md#linalgqr-returns-a-wrong-qr-after-an-earlier-call-in-the-process).
Pivots are packed 1-based int32 everywhere. cusolverDnXgetrs wants int64.
Known located-but-unfixed defects: docs/design/known-defects.md.

10. Measurement Rules (benchmarks lie silently here)

Compare algorithms only at saturation (doubling the batch buys <10% per item); below it you measure launch, allocation and clock-ramp overhead. Benchmark and tune at batch >=128; batch=1 does not matter to this project. Still profile across batch sizes: batch-only-parallel kernels (one work-group per matrix) starve at small batch. If a kernel's work-group count equals the batch size, that is the bug.
Run one measuring process on the box at a time, whatever the device pin. Two GPUs on one NUMA node and one UVM driver produced 5.5x errors with stable rel_sd while nvidia-smi showed no foreign process. A display attached to a GPU also slows L2-resident vendor cells: measure on a headless GPU. Use benchmarks/gpu_guard.sh, plus a flock against your own duplicate jobs.
Warm up: idle clocks are ~210 MHz. Use --warmup=3 --min_iters=8. The first process to run a kernel pays JIT (cache ~/.cache/batchlas_sycl, shared): do a throwaway pass over every cell first.
Alternate the A/B order between reps and discard the first pass. Per-rep CV is ~3%, so an effect under 5% needs 14-16 reps; 6 reps is a screen, not a verdict.
The Google Benchmark --name filter is a substring match. Filter on the CSV name column, which is quoted and contains a comma (awk -F, offsets shift).
Admit rows on cross-pass reproduction, not rel_sd alone: a noise filter that deletes a reproducible loss fabricates a window. Bracket every window edge with a measured non-winner.
Before quoting "vs vendor": check which route ran, what "the vendor" is (cuBLAS complex trsm and batched orgqr are per-item loops or our own substitute kernels; cuSOLVER syevjBatched beats the XsyevBatched we call at n<=32), and re-measure the vendor, because versions move.
Rebuild every bench binary after a table or can_run change: harnesses resolve the printed route in their own TU, so relinking only the .so leaves the route column lying.
Check that an argument is used, not just echoed: time should scale with it. Harnesses written before the code (design studies) do not benchmark the shipped code.
Sanity anchors (RTX 4090): real FP32 SGEMM ~47-54 TFLOP/s (~80 means TF32), FP64 ceiling ~1.44 TFLOP/s (1/64 rate), DRAM ~1008 GB/s peak, ~950 achievable. cuBLAS gemv is already at the DRAM roof: do not try to beat it.
A cost table that is non-monotonic in n (a larger n running faster) is a routing bug.
nsys sees cuBLAS kernels; BATCHLAS_KERNEL_TRACE does not (it hid 35% of a cfloat syev). Also useful: BATCHLAS_KERNEL_TRACE_PATH (Chrome trace), BATCHLAS_GESVD_PROFILE=1, BATCHLAS_TUNE_* overrides (A/B any tuning constant without a rebuild), BATCHLAS_BENCH_LD_PAD[_A/_B/_C] and BATCHLAS_BENCH_BETA in the gemm benchmarks.
Always confirm a GEMM at beta=1 and with a strided ld: panel updates always carry their parent's ld. Pin one gemm kernel with BATCHLAS_GEMM_ROUTE=<spelling> (reg:m=128:n=128:k=8:u=1, wide:m=64:n=64:k=16, small, tiled, direct; old kernel names such as 128x128x8 and tiled16 throw); it throws if that tile has no instantiation for the call's transpose form. BATCHLAS_GEMM_SYCL_KERNEL and BATCHLAS_GEMM_EXPERIMENTAL are not read. Compare gemm kernels with tools/tune/batchlas_tune (gemm_spec.cc: beta=1, complex alpha, strided cells padded by --ld-pad, every arm verified).
benchviz (python3 benchmarks/benchviz) measures whatever is in its checkout's build/, not HEAD: run benchviz info first. Its batchlas arm pins the native walk, not Auto. Figures follow the user's style (benchmarks/benchviz/style.py: CM usetex, full box, dotted lines with large markers, 2σ band, viridis maps), not generic defaults.
Raw result grids go in benchmarks/results/ (Git LFS; run git lfs install on every machine, or the commit stores raw files and CI's check_lfs_pointers.py fails). Older evidence is at tag perf-evidence/vendor-independence (git show <tag>:experiments/...). docs/perf is the distilled layer.
Retuning takes ~12 min (evaluation/tuning/). The checked-in include/batchlas/tuning_params.hh is the only copy (nothing is generated); port constants into it by hand. Tuning was float-only: constants measured on float are applied to double and complex, so when another type looks bad, sweep that type first. Check whether a knob has a second consumer (aliasing) or is bypassed on the path (shadowing) before trusting "no change".

11. GPU Kernel Design Facts (sm_89; mostly general)

Register gate: registers are split across four sub-partitions of 16,384. The bound is ceil(warps/4)*32*regs <= 16384, not regs*wg <= 65536. SYCL accepts the launch and the driver rejects it as an uncaught abort. Standalone -Xcuda-ptxas -v register counts differ from the in-library kernel: prefer a measured launch ceiling.
Design against occupancy, not spills: 8x8 double tiles do not spill, but complex<double> x 512 threads exceeds the register file. Shrink the thread tile as the scalar widens. Never pass an accumulator element by T& out-parameter (spills the whole array, 43%); return by value.
std::complex * emits an isnan branch plus __mulsc3 (Annex G, no -ffast-math). Write out the real arithmetic in the hot loop only (1.2-1.3x); converting every site cost occupancy and was slower. Do not add -ffast-math.
The 48 KB SLM launch hole: an attribute sticky per CUfunction means test launch order can hide it. Put the discriminating launch first.
Launch bounds (intel::min_work_groups_per_cu) apply only on a functor's operator(), not a lambda. Worth 1.1-1.9x on tiny kernels, but a cap that is too tight is 2.3x slower, and a change moves routing windows.
SLM staged by a lane-strided loop and read by a different index needs a barrier. NVIDIA hides the race at wg=32.
Column-major coalescing: one work-item per output element. With that mapping NoTrans gemv is coalesced and Trans needs the reduction. Row interchanges (laswp) should be an index gather in SLM, not a per-column walk (7.9x fewer sectors).
GEMM: 128x128x8 tiles with an 8x8 thread tile (64 accumulators), aligned (not +1) shared strides and B staged [k][n] give cuBLAS parity for float NN. The epilogue must vary the m index fastest. Internal GEMM demand is panel updates (large m,n; k = blocking factor 8-136; ~60% transposed), so a min_dim gate never fires on it. The native GEMM is ~2x off on strided ld from exposed B-load latency; double-buffering and packing B were both measured as no gain.
DPC++ chunked_partition masks cost ~4-5 instructions per collective. Full-warp or maskless rewrites lost in real kernels even where microbenchmarks won. Judge partition-primitive changes (src/kernels/common/sg_partition/) by real-kernel A/B (steqr/syev_cta benchmarks, n=5..16, batch 16384), never by microbenchmark.
CTA kernels are capped at n <= 32 (the sub-group width). STEDC_RECURSION_THRESHOLD = 32 is that invariant, not a tuning constant: a leaf of 33+ falls from steqr_cta to steqr_wg (14x cliff).
Traffic models that count bytes miss serial recurrences and lost parallelism. Check the critical path and the work-item count first.

Measured dead ends: do not re-propose without a new idea

Idea	Result
Lifting the CTA eigensolver above n=32	85-211x slower
Block Jacobi as a speed path	2-11x slower (accuracy-only opt-in)
Cooperative TRSM diagonal solve	dead end
Packed sub-group getrf with runtime n in SLM	dead end (register-resident designs are not refuted)
Faster native gemv	dead end
Maskless steqr at P<32	dead end
Givens chase	Householder chase is 5.6x cheaper
GEMM kernel fix for strided ld	the fix was routing

The latrd panel symv is L1-over-fetch-bound (12-16x), not DRAM-bound; a MAGMA-style single-read symv is the open 2.7x.

12. Repository Conventions and Git

Docs are a Doxygen 1.18 + doxygen-awesome site and the project's database of results and decisions: sh scripts/build_docs.sh [out] (or -DBATCHLAS_BUILD_DOCS=ON, target batchlas_docs). The governing page is docs/developer/documentation.md; read it before adding a page. CI's docs job builds with warnings as errors and deploys main to GitHub Pages.
Where things go: measurements, A/B grids, rejected alternatives -> docs/perf/<op>.md; design decisions -> docs/design/; math -> docs/algorithms/; user guides -> docs/guide/; process -> docs/developer/; raw data -> benchmarks/results/ (LFS). A new page opens # Title {#area_topic} plus a one-line status. Math is \f$…\f$ / \f[…\f], not $…$.
No Markdown outside docs/ except a directory README.md (never at the root), .claude/CLAUDE.md (it only imports this page) and .github/, which holds the project README (check_markdown_locations.py). No root Markdown at all, so no root plan files: put conclusions in the area page.
Source layout (table: docs/extending.md): src/ops/<op>/ holds <op>.cc, choice.hh and headers shared by the op's kernels; kernels/ holds its kernel TUs and single-consumer headers. Internal sub-ops nest under the parent (syev/{stedc,steqr,sytrd,tridiag}/, gesvd/{gebrd,bdsqr,bdsdc,ormbr}/). Shared device code: src/kernels/common/. src/sycl/ is only DPC++/AdaptiveCpp portability; src/backends/ only vendor wrappers; src/util/ also holds the batchlas_extra sources. CMake lists (src/ops/CMakeLists.txt) group by device-link cluster, not folder, and keep the old batchlas_extensions_* names. In-tree includes are root-based, #include <src/ops/potrf/potrf_native.hh>, <tests/test_utils.hh>; never quoted or relative (check_no_unprefixed_includes.py fails on #include "...").
Comment density <=18% per file (.github/ci/check_comment_density.py). A code line with a trailing comment counts as code. The repo-wide average hides per-file violations: check per file.
- A comment is for invariants, "looks wrong but is deliberate" notes and traps, plus an evidence: docs/<page>.md#anchor pointer.
- Doxygen doc comments (///, //!, /** */) in include/batchlas/ are exempt, but only for the contract (brief, params, returns, @pre, @throws, @ingroup from docs/pages/api_groups.dox); rationale in /// is still rationale.
- Waivers: .github/ci/comment_density_waivers.txt.
A main ref from session start is a snapshot. git fetch origin before claiming something is absent from a branch.
For stacked PRs, merged can mean merged into the previous stack branch, not into main. Read base.ref.
/code-review ultra rejects diffs over 500 files or 8,000 lines. Split along work-package commit boundaries: a by-op split does not compile, because the dispatch facade includes every op's native header.
nm -DC ... | awk '{print $3}' truncates demangled names at the first space, and nm -C cannot demangle concept-constrained templates. Diff full mangled names (scripts/facade_symbol_check.sh).
In worktree-isolated agent sessions, the Bash guard refuses heredocs, for loops with computed arguments and awk programs. Write scripts to $CLAUDE_JOB_DIR/tmp and run them. Wait on a PID with kill -0, not pgrep -f (it matches itself).

13. Primary-Machine Facts (the RTX 4090 box; re-check on any other machine)

2x RTX 4090 (128 SMs, 72 MB L2, 24 GB, FP64 1/64), Xeon w5-2445, Ubuntu 22.04, DPC++ at /opt/dpcpp-cuda, CUDA 13.2, ROCm 6.2.4 headers (no AMD GPU), MAGMA 2.10 Release/sm_89 at ~/magma/build-rel89. Device 0 drives the display: measure on device 1.
The box is shared with another user who runs GPU sweeps. Check nvidia-smi before timing anything and coordinate before a reboot. perf and gdb -p do not work (perf_event_paranoid, ptrace): use the UR trace.
The system OpenBLAS 0.3.20 picks the broken Cooperlake dgemm (Known Pitfalls).
Unattended-upgrades is not blacklisted for nvidia-*, so driver mismatches recur.

On a new machine, before trusting any result: sycl-ls shows your GPU; configure prints the right SYCL target; nvidia-smi (or equivalent) reports the architecture you think it does; record that architecture's FP32/FP64/DRAM ceilings as sanity anchors. Every routing window and tuning constant in the tree was measured on sm_89 (some tables on sm_120). On a different GPU, an Apple GPU included, especially one with 1:2 FP64 where double GEMM verdicts invert, treat the tuned tables and hemm/herk/her2k's hand-written windows as hypotheses. Re-measure at saturation before quoting any ratio, and record the machine alongside new data in benchmarks/results/.

License

SPDX-License-Identifier: MIT