Status: current
Matrices are column-major and batched, every call enqueues work and returns immediately, and the backend, the workspace and the device all come from the Queue. This page teaches that model and the traps around it. Signatures, option fields and per-op detail are in the generated API reference (groups in Public API); environment variables are in Environment variables, exceptions in Error model, spelling rules in API conventions.
Everything is in namespace batchlas: Matrix, MatrixView, gemm, potrf and the rest of the numerical surface, plus Queue, Device, Event, Span, UnifiedVector and BumpAllocator.
Note: For one more release, headers re-export Queue, Device, Event, Span, UnifiedVector and BumpAllocator at global scope. Define BATCHLAS_NO_GLOBAL_NAMES to turn that off today.
Matrix, MatrixView and VectorView default to <float, MatrixFormat::Dense> (<float> for VectorView); Vector<T> has no default. Spell the format only when it is not Dense: Matrix<float, MatrixFormat::CSR> S(n, n, NonZeros{nnz}, batch).
Options worth knowing: BATCHLAS_ENABLE_CUDA (AUTO default; ON makes configure fail when no CUDA device is exposed), BATCHLAS_NVIDIA_ARCH (detected; pass it only when cross-building), BATCHLAS_ENABLE_VENDOR_BLAS (ON; OFF is the vendor-free build), BATCHLAS_ENABLE_ROCM (OFF), BATCHLAS_BUILD_TESTS, BATCHLAS_BUILD_BENCHMARKS, BATCHLAS_BUILD_PYTHON, and BATCHLAS_ALLOW_UNSAFE_ENV (OFF; see Configuration). CMakePresets.json has dev and dev-tests presets.
Configure the whole consuming project with the compiler BatchLAS was built with. A mismatch compiles and then fails at link:
The package compares CMAKE_CXX_COMPILER to the recorded compiler (in the install's build-tree CMakeCache.txt) by realpath and warns on a difference; -DBATCHLAS_REQUIRE_MATCHING_COMPILER=ON makes it an error.
examples/consumer/ is this project, standalone and runnable (it prints PASS).
When the building compiler has no nvidia_gpu_sm_<N> alias for the GPU (DPC++ stops at sm_90a), BatchLAS_SYCL_TARGETS is nvptx64-nvidia-cuda and BatchLAS_SYCL_BACKEND_OPTIONS carries the architecture. Pass it too, or your kernels build for sm_75:
At run time the loader must find the DPC++ runtime and its SYCL adapters: export LD_LIBRARY_PATH=<dpcpp-prefix>/lib, or add that directory to /etc/ld.so.conf.d/ and run ldconfig. Headers install under <prefix>/include/batchlas/ plus <prefix>/include/batchlas.hh.
A Queue names the device, carries the backend, owns a workspace arena and orders the work. Device::default_device() is the first GPU, else the first CPU, else the first host device.
The string constructor takes the first device of that type; it throws batchlas::device_error when there is none and batchlas::invalid_argument for any other string.
Every entry point is batched and applies the same operation to every item; all matrix arguments of one call must have the same batch size. α and β are the alpha and beta option fields, and op(A) is A, Aᵀ or Aᴴ according to trans. Shapes and constraints are in the API groups: BLAS level 2, BLAS level 3, Factorizations and linear solvers, QR and orthogonal transforms, Symmetric and Hermitian eigensolvers, Singular value decomposition, Sparse (CSR) operations.
| call | computes | written |
|---|---|---|
| gemm(ctx, A, B, C, opts) | C := α·op(A)·op(B) + β·C | C |
| gemv(ctx, A, x, y, opts) | y := α·op(A)·x + β·y | y |
| symm / hemm(ctx, A, B, C, opts) | C := α·A·B + β·C (Left) or α·B·A + β·C (Right), A symmetric / Hermitian | C |
| syrk / herk(ctx, A, C, opts) | C := α·A·Aᵀ + β·C (NoTrans) or α·Aᵀ·A + β·C; herk uses ᴴ | uplo triangle of C |
| syr2k / her2k(ctx, A, B, C, opts) | rank-2k update of C | uplo triangle of C |
| trmm(ctx, A, B, C, opts) | C := α·op(A)·B or α·B·op(A) | C; B untouched |
| trsm(ctx, A, B, opts) | solves op(A)·X = α·B or X·op(A) = α·B | B, overwritten with X |
| potrf(ctx, A, opts) | Cholesky, A = L·Lᴴ or Uᴴ·U | A in place, uplo triangle |
| getrf(ctx, A, pivots) | LU with partial pivoting | A and pivots |
| getrs(ctx, A, B, pivots, opts) | solves op(A)·X = B from getrf's factors | B |
| getri(ctx, A, C, pivots) | C := A⁻¹ from getrf's factors | C; A read-only |
| geqrf(ctx, A, tau) / orgqr(ctx, A, tau) | QR with Householder reflectors / expand to explicit Q | A, tau |
| syev(ctx, A, W, opts) | symmetric/Hermitian eigendecomposition of the uplo triangle | W ascending; A gets eigenvectors when jobz == JobType::EigenVectors |
pivots is int64_t (n·batch), tau is T (min(m,n)·batch) and W is the real counterpart of T (n·batch).
Warning: trmm and getri write a second operand and leave the input alone (reference ?trmm works in place on B); trsm is in place. Mixing these up gives a wrong answer, not a compile error.
Warning: syrk, herk, syr2k and her2k write only the uplo half of C. The other half comes back as it went in, including uninitialised memory (a fresh Matrix is not zeroed), so the result is triangular, not symmetric, and no call reports it.
Mirror the triangle before such use: symmetrize(ctx, uplo) copies the named triangle across the diagonal, hermitize(ctx, uplo) its conjugate (the right one for herk/her2k). Both are one kernel on a square MatrixView and return an Event.
In the positional spelling alpha is argument 4, right after the matrices, matching trmm: trsm<B>(ctx, A, B, alpha, side, uplo, trans, diag). Two = delete overloads with the old order (Side, Uplo, Transpose, Diag, T) turn a stale call into a "call to deleted function" error. The option-struct spelling is unaffected; backend::trsm_vendor still takes alpha last.
The extension surface (steqr, stebz, stein, stedc, lanczos, ritz_values) follows the same rule: VectorView<T> for one logical vector per batch item (needs inc/stride/batch_size), Span<T> for a flat array with one entry per item. stebz shows both: d, e, w are VectorView<T>, m is Span<int32_t>. They do not interconvert, because a VectorView demoted to a Span would drop the stride and read wrong elements.
| you hold | parameter is Span<T> | parameter is VectorView<T> |
|---|---|---|
| UnifiedVector<T> | pass directly, or .to_span() | VectorView<T>(v, size, batch, Inc{i}, Stride{s}) |
| Vector<T> | .data(): the whole allocation, so only when inc == 1 and packed | pass directly |
| raw pointer | Span<T>(p, n) | VectorView<T>(p, size, batch, Inc{i}, Stride{s}) |
A Matrix is a MatrixView and a Vector a VectorView, so owning arguments go wherever a view is expected, and every view member (fill_random(), batch_item(), symmetrize(), ...) is called on the owning object directly. Pass .view() only to a generic parameter (auto, a template T), which would otherwise copy the matrix. A bare {} in an argument position does not compile: name the type (OrthoOptions{}).
Warning: Vector names inc and stride; VectorView also takes them positionally, in the opposite order. Write Vector<T>(size, batch_size, Stride{s}, Inc{i}); the bare-int forms are deleted. The tags exist because (inc = n, stride = 1) and (inc = 1, stride = n) fit the same buffer, so a swapped pair compiles and reads the wrong elements. Stride, Inc, Ld and BatchSize live in batchlas/blas/matrix.hh beside NonZeros; none decays to int.
Most entry points take an option struct, so you write only what differs:
The structs are in batchlas/blas/options.hh (group Option structs): GemmOptions<T>, GemvOptions<T>, SymmOptions<T>, HemmOptions<T>, SyrkOptions<T>, HerkOptions<T>, Syr2kOptions<T>, Her2kOptions<T>, TrmmOptions<T>, TrsmOptions<T>, PotrfOptions, GetrsOptions, SyevOptions. The BLAS ones are templated on T; the three LAPACK ones are not. Defaults are alpha = 1, beta = 0, side = Left, uplo = Lower, trans = NoTrans, diag = NonUnit, jobz = EigenVectors. The traps:
Warning: Every uplo defaults to Uplo::Lower. Filling the upper triangle and calling potrf or syev with default options factorises whatever is in the lower one, with no report. Fill the lower triangle or pass {.uplo = Uplo::Upper}.
*Options structs belong to the convenience layer: the backend comes from the Queue, T is deduced from the matrices, the struct carries every non-matrix argument, and the workspace of LAPACK-style calls may be omitted.
*Params structs (batchlas/blas/extensions.hh, batchlas/blas/functions/iluk.hh: SyevxParams, LanczosParams, StebzParams, SteinParams, SteqrParams, StedcParams, JacobiParams, GesvdjParams, SytrdBandReductionParams, ILUKParams) are ordinary arguments to entry points with no convenience layer: you name the backend (syevx<Backend::CUDA, float>) and pass a workspace. Their position varies, so read the declaration: usually last, but second to last for SteqrParams and StedcParams (eigvects follows), and SytrdBandReductionParams replaces the int32_t block_size argument.
gesvd, ormqr, ortho and spmm take positional arguments and a workspace span. Lease it from the arena yourself:
With an empty option struct and an explicit workspace, name the type: potrf(ctx, A, PotrfOptions{}, ws). T is deduced from the matrix arguments, never from the option struct: write syev<B>(ctx, ...) or syev(ctx, ...) on the option-struct spelling and syev<B, float>(ctx, ...) on the positional one.
Matrices are column-major, like LAPACK and unlike NumPy. For a dense MatrixView V, element (i, j) of batch item b lives at
ld, stride and the indices count elements, not bytes. Access is M(i, j, b) on an owning Matrix (the batch index is required) and V.at(i, j, b) or V(i, j, b) on a MatrixView, which bounds-check and throw std::out_of_range.
Matrix(rows, cols, batch) allocates exactly rows * cols * batch elements with ld() == rows() and stride() == rows() * cols(); nothing pads. Matrix(rows, cols, batch, ld, stride) and MatrixView(data, rows, cols, ld, stride, batch) resolve the defaults identically. Element access computes int64_t(b) * stride + j * ld + i; the within-item offset stays int (limit 2³¹ elements, 8 GB of float).
MatrixView throws std::invalid_argument on a negative extent and on a resolved ld < rows. The batch count is the third argument on Matrix and the sixth on MatrixView, so MatrixView<float> V(p, n, n, batch) means ld = batch and throws. A view accepts a null data or a zero dimension (the shape-only view used for workspace-size queries).
Warning: stride = 0 is the packed default, not cuBLAS's broadcast. BatchLAS has no broadcast operand; unequal batch sizes throw. MatrixView(dA, n, k, n, /*stride=*/0, batch) over one matrix reads batch consecutive items, past the end of the buffer. To multiply one matrix against many right-hand sides, fold the batch into columns (packed B with ld == k, stride == k*n is one k × (n·batch) matrix) or replicate the matrix.
For gemm, use the operand swap (no copy). Otherwise convert: Matrix::to_column_major() returns a converted packed copy and to_row_major() converts back. Both synchronise before returning; to_column_major(ctx) / to_row_major(ctx) use your queue.
Every pointer given to MatrixView or Span must be device-accessible for the backend the Queue dispatches to. MatrixView takes a bare T* and cannot check at construction; entry points that take the backend from the queue check every pointer argument at the call and throw std::invalid_argument:
Zero-copy on a GPU backend: sycl::malloc_device, malloc_shared, malloc_host (including on your own sycl::context for the same device), cudaMalloc, cudaMallocManaged. Not: malloc, new, std::vector. On a host/CPU device ordinary host memory works and nothing is rejected, so test on the device you ship on.
Matrix owns USM shared memory, so the host can read and write it directly. Bulk-load with the copying constructor:
(ld, stride) describe the source: element (i, j, b) is read from data[b * stride + j * ld + i], ld = 0 meaning rows and stride = 0 meaning ld * cols. Prefer the Span<const T> overload: the span knows the source length, so an over-reading shape throws std::invalid_argument.
Generated data skips the host and synchronises before returning:
Other factories: RandomTriangular, Ones, Diagonal, Triangular, TriDiagToeplitz. To fill a matrix you own, the fill_* family on MatrixView writes in place and returns an Event (fill, fill_zeros, fill_identity, fill_diagonal, fill_random, ...; group Matrices, views and vectors). Pass your own queue: a default-constructed Queue builds a fresh one on Device::default_device(), which targets the wrong device on a multi-GPU box.
Warning: The seed of Random, RandomTriangular, fill_random and fill_triangular_random defaults to 42, and each element is a pure function of (seed, index). Two default-seeded calls return bit-identical matrices; pass distinct seeds for distinct operands. Random keys on the flat index over the whole allocation (different ld or stride gives different values at the same (i, j, b)); fill_triangular_random keys on the index within a matrix, so every batch item of a RandomTriangular result is the same matrix.
View-to-view copy is asynchronous: MatrixView<float>::copy(ctx, dst, src). Host element access is one indexed store into managed memory per element: use it for a handful of entries and tests, not to load a batch. MatrixView never owns.
Matrix memory migrates between host and device on demand; a large batch written from the host and read by many kernels pays for that traffic. For device-resident storage, allocate with sycl::malloc_device and wrap it in a MatrixView, which copies nothing:
A column-major view of a row-major m x k buffer with row length k is its transpose. Since Cᵀ = Bᵀ Aᵀ, passing the transposed views to gemm in the opposite order computes the row-major product with no copy:
The swap is gemm's. For the symmetric routines, flip uplo. For everything else, including potrf, getrf and syev, convert first (Row-major source data).
The owning constructors line up as shape, format-specific extra, batch size; the CSR non-zero count is the NonZeros strong typedef (a bare int does not compile):
NonZeros{} is a capacity. convert_to<MatrixFormat::CSR>() sizes a whole batch by its largest item, so on a heterogeneous batch nnz() over-counts smaller items and a loop to nnz() walks past that item's row range. Three accessors exist on Matrix, MatrixView and KernelMatrixView:
The from-data constructors are in Matrices, views and vectors and Matrix model.
Every exception BatchLAS raises is a batchlas:: type from batchlas/error.hh. Each derives from both the std:: exception that site threw before and an empty tag base, batchlas::exception, so existing std:: handlers still catch it (and Python raises the same type as before).
| class | derives from | means | retry helps? |
|---|---|---|---|
| invalid_argument | std::invalid_argument | The call violates the contract: mismatched shapes or batch sizes, short span, null or non-USM pointer, bad ld. | No. |
| out_of_range | std::out_of_range | Index outside its container (V.at(i, j, b)). | No. |
| error | std::runtime_error | Base of the six below. | n/a |
| unsupported | error | No kernel or backend in this build on this device serves the request. | Not as asked; try another route pin, backend, type or shape. |
| device_error | error | The device or vendor runtime failed (cuBLAS/cuSOLVER status, launch failure, null allocation, no device of that type). | Sometimes. |
| workspace_error | error | Scratch too small, or the arena ran out. | Yes: re-query *_buffer_size() or halve the batch. |
| convergence_error | error | An iterative kernel did not converge or a factorisation broke down (LAPACK info > 0). | With different parameters. Prefer the per-item info spans below. |
| internal_error | error | BatchLAS is inconsistent. Caller-reachable today: gesv and posv (so linalg::solve, linalg::solve_spd) for an empty problem, n, nrhs or batch below 1, and a heterogeneous batch. | No; report it. |
| api_misuse | error | Right call, wrong state: a Queue used from another thread, configure() after a Queue exists. | No. |
| NoRouteError (<batchlas/no_route.hh>) | std::runtime_error | Nothing in this build serves the call (typically -DBATCHLAS_ENABLE_VENDOR_BLAS=OFF). op(), backend(), scalar() say which. | Not in this build. |
Properties of batchlas::exception that tests/error_model_tests.cc asserts, because each fails silently if broken: it does not derive from std::exception (a second std::exception subobject would make catch (const std::exception&) ambiguous, and an ambiguous catch base is not diagnosed); it is a virtual base; and it has no what() (use e.message() or catch std::exception).
Outside the hierarchy, invisible to catch (const batchlas::exception&): std::bad_alloc (where a sycl::malloc_* returned null), sycl::exception (including everything reported asynchronously at ctx.wait_and_throw()), and kernel selection: a BATCHLAS_<OP>_ROUTE value that does not parse, names a family not compiled for the scalar type, or cannot run the shape throws std::invalid_argument; a malformed table in BATCHLAS_TUNED_DIR, or a call no family can run in a build that has the op's vendor library, throws std::runtime_error. Open item: kernel selection throws outside the hierarchy. A boundary that must let nothing escape needs catch (const std::exception&) behind the BatchLAS one. Device-side errors surface at ctx.wait_and_throw(), not at the enqueuing call.
Eight routines report a per-item status, opt-in, as one Span<int32_t> entry per batch item with LAPACK's convention (0 good, positive otherwise): potrf, getrf, getri report factorisation status; syev, syevx, gesvd, steqr, stedc report convergence status. It is a field on PotrfOptions and a trailing parameter elsewhere:
An empty span (the default) reports nothing and costs nothing. The span must be device-accessible USM, is written in place, and changes no workspace size; a non-empty span shorter than batch_size throws std::invalid_argument. geqrf and orgqr have no status: Householder QR has no numerical failure mode.
Each takes a trailing Span<int32_t> info, default empty. Without it, a single non-converged item in a batch of 16384 is invisible: the call and ctx.wait() return normally.
The solve-style calls (getrs, linalg::solve) report nothing. A near-singular item produces plausible numbers; when inputs are not known to be well-conditioned, check the residual ‖A·X − B‖ or scan the factor's diagonal for zeros and NaNs per item.
Every entry point enqueues work and returns an Event. Contents of a Matrix, MatrixView or UnifiedVector are not readable until the work finishes:
Without the wait you read the output as it was before the call, and a fresh Matrix is uninitialised, not zeroed (Matrix::Zeros(...) or view().fill_zeros(ctx) give a known start).
A Queue is single-threaded. Use one per thread. It owns an unsynchronised arena and a cached "last event"; workspace(), trim_workspace(), submissions, enqueue(), get_event() and create_event_after_external_work() compare std::this_thread::get_id() to the constructing thread and throw batchlas::api_misuse on a mismatch. Queues built for the same Device share a SYCL context, so per-thread queues see each other's USM allocations. To move a Queue to another thread, call attach_to_current_thread() from the new owner before its first use.
Backend::AUTO resolves once on first use and is cached; set_backend resets the cache. On a GPU it picks the vendor stack if compiled in (NVIDIA → CUDA, AMD → ROCM, Intel → MKL), else NETLIB, as every non-GPU device does. set_backend throws batchlas::unsupported if the named backend is not compiled in. The nameable backends are CUDA, ROCM, MKL, NETLIB and AUTO; Backend::MAGMA and Backend::SYCL are unavailable on every build.
This covers the whole surface, extensions included. The exceptions are entry points whose template parameters are not deducible, listed under Which spelling each entry point takes. When you need the backend as a compile-time constant, with_backend (<batchlas/blas/queue-dispatch.hh>) turns the runtime choice into one:
The backend picks the library build; it does not pick the kernel. After argument validation, each op chooses one kernel family for the call: a native kernel (tiny, cta, lpanel:panel=8, blocked, ...) or vendor (the cuBLAS/cuSOLVER/rocBLAS/host LAPACK call). The decision for an op is src/ops/<op>/<op>.cc, with its vocabulary in src/ops/<op>/choice.hh.
| variable | effect |
|---|---|
| BATCHLAS_<OP>_ROUTE | Pins the op: auto, native, vendor, or a family spelling from choice.hh. A bad spelling throws std::invalid_argument; native and vendor fall back to auto with a warning when nothing of that kind can run. |
| BATCHLAS_SELECT_TRACE=1 | One stderr line per call: op, shape, chosen family, its time and the runner-up's, and which table answered. |
| BATCHLAS_TUNED_DIR=<dir> | A same-named table file in <dir> replaces the built-in one. |
Nested ops decide for themselves: a blocked factorisation calls the public gemm and trsm, which consult their own tables. hemm, herk and her2k are not table-selected. Design: flat kernel selection.
Everything BatchLAS reads from the process environment lands in one typed struct, parsed once under std::call_once on first call; settings() is thread-safe and the only place in the library that reads the environment. The complete variable list, with types and defaults, is Environment variables.
Settings has five groups: routing (the 19 BATCHLAS_<OP>_ROUTE pins, via route(op)), selection (algorithm choices outside the route vocabulary; three override an explicit API argument), geometry (launch geometry, block widths, iteration counts), diagnostics (tracing, dumping, profiling, opt-in checks; none changes a numeric result) and unsafe. Field names follow variable names (BATCHLAS_TRSM_OUTER_NB is settings().geometry.trsm_outer_nb). A knob with a bespoke parser is an EnvValue (raw string with is_set(), value(), get()); a default that is a sentinel ("", 0, -1, nullopt) means the real default is computed at the call site. Variables read only by the repository's tests and benchmarks (BATCHLAS_TEST_BACKEND, BATCHLAS_BENCH_*, ...) are not in Settings.
batchlas::ScopedEnvVar (<batchlas/util/env.hh>) sets a variable for a scope and restores it on exit; a null value unsets it for the duration.
Its constructor and destructor call batchlas::detail::reload_settings(). Anything else that writes the environment mid-process (a raw setenv) must call reload_settings() too, or the change is invisible. A reload between a *_buffer_size() query and its call desynchronises the workspace from the block width the call uses: do not let a ScopedEnvVar scope straddle a sizing/solve pair. env_truthy accepts exactly {1, true, TRUE, on, ON} and env_falsy exactly {0, false, FALSE, off, OFF}; an unset variable is neither, which is why some fields are std::optional<bool>.
Most knobs pick a kernel, geometry or dump path; a wrong one costs a measurement. A few remove a check, and live in Settings::unsafe behind a CMake option (-DBATCHLAS_ALLOW_UNSAFE_ENV=ON, default OFF). With it OFF (release and install builds) settings() holds the unsafe fields at their safe values whatever the environment says, and prints one warning at first use naming the variable and the option. It is ON in the dev, dev-tests, fast-dev, dev-gpu, dev-gpu-tests and benchmarks presets and OFF in cuda (the pre-push gate). tests/settings_tests.cc asserts both arms.
| variable | what it disables | how it fails |
|---|---|---|
| BATCHLAS_SKIP_POINTER_CHECKS | the one-USM-query-per-argument reachability check (~70 ns) | host memory reaches the device as a wild address: CUDA_ERROR_ILLEGAL_ADDRESS, then SIGABRT at teardown, which no catch block stops |
| BATCHLAS_LATRD_GRID_FORCE_UNSAFE | the co-residency cap on the latrd grid path | the grid barrier relies on that cap, so the kernel hangs (it looks like slow JIT). Run under timeout |
| BATCHLAS_BLAS_HEALTH=off | the host-dgemm correctness probe | with a known-bad OpenBLAS kernel, every double and complex<double> host result is silently wrong by O(1) |
The gate refuses only the unsafe direction (BATCHLAS_BLAS_HEALTH=error is allowed, off is refused). configure() is not gated.
LAPACK-style entry points need scratch. Leaving the workspace argument out leases it from a per-Queue arena, sized by the matching *_buffer_size:
To size and own the buffer yourself:
When to keep managing the workspace yourself (pass a span explicitly, e.g. potrf(ctx, A, {.uplo = Uplo::Lower}, my_span)):
Full lifetime rules: batchlas/util/workspace.hh, Workspace.
Queue::native_handle() returns the backend-native stream as void*: a CUstream (cudaStream_t) when the queue's device runs on the CUDA SYCL backend, a hipStream_t on HIP, nullptr otherwise. It keys off the device, not ctx.backend(). The stream belongs to the Queue: do not destroy it or outlive the Queue. On the default in-order queue, work you push to it runs after everything BatchLAS has submitted. To make BatchLAS wait for your work, call ctx.create_event_after_external_work() after enqueueing it.
For SYCL-typed interop include <batchlas/sycl_interop.hh>, the one BatchLAS header that pulls in <sycl/sycl.hpp> (not reachable from <batchlas.hh>; do not re-export it from your headers):
Both queues must share a SYCL context. Pointers from cudaMalloc, cudaMallocManaged, sycl::malloc_device and sycl::malloc_host wrap into Span/MatrixView zero-copy when reachable from that context.
batchlas::linalg (batchlas/blas/linalg-ops.hh, group Convenience layer (batchlas::linalg)) offers value-returning and elementwise free functions; there are no operator overloads. Each takes its backend from the queue and workspace from the arena.
These enqueue and do not wait, except linalg::norm, linalg::cond and linalg::svd. add, subtract, multiply, divide and scaled allocate their result; the _into forms write into storage you own (use them in inner loops); scale works in place.
Design of the layer and option structs: API conventions.
Adding an entry point rather than calling one? See extending.md.
1.18.0
