GPU-callable inference, full operator coverage, and the retirement of fLibrary by sbryngelson · Pull Request #12 · comp-physics/roseNNa

sbryngelson · 2026-09-14T14:11:06Z

Summary

Makes the generated code callable from a GPU, covers every operator the old runtime library covered, and retires that library.

Three things happened in sequence, each one a consequence of the last:

The GPU path was validated on real hardware (A100, NVIDIA HPC SDK 25.11). rosenna gpu-gate --backend cuda PASSes and records zero cudaMemcpy inside the timed infer_batch call, embedded and file-loaded. The HIP path has since been validated the same way (AMD Instinct MI210, gfx90a): --backend hip PASSes under both ROCm 7.2.0 and the TheRock AFAR 23.2.1 drop, with zero hipMemcpy inside the roctx-scoped timed call
Running it found a 30x cliff on the per-point path, and fixing that required a change to the generated arithmetic.
Operator coverage went from 5 of 21 golden models to 21 of 21, at which point fLibrary/ had nothing left to do.

The contract (unchanged)

init is the plan step and the only routine that allocates or transfers; embedded models have no init at all. Nothing in the loop path (infer, infer_batch) allocates, transfers or synchronizes. A structural test enforces this over every backend in both languages.

A model with several graph inputs — an LSTM's initial hidden and cell state — takes them concatenated in x in declaration order, and a model with several graph outputs — that LSTM's Y, Y_h, Y_c — writes them concatenated in y. rosenna info prints both layouts. That keeps infer(x, y), and with it infer_batch, the native kernel and the whole device contract, unchanged, and lets a solver keep a recurrent model's state on the device: y's state slices feed x next step.

Three more entry points: <name>_sync(stream) waits for infer_batch (a stream synchronize in the cuda/hip archive, a no-op in the omp one), for a host with no stream of its own; <name>_init_dev binds a Fortran host to the C archive's init; <name>_infer_one(x, y, stream) runs one sample as a launch per op with the activations in static device buffers, for a model whose activations do not fit a thread (a conv net over a field), absent when the plan has an LSTM.

Operator coverage

	ops
dense	`Gemm`, `MatMul`, `Relu`, `Tanh`, `Sigmoid`
spatial	`Conv`, `MaxPool`, `AveragePool` (2-D, incl. `auto_pad`)
recurrent	`LSTM` (forward, default activations, optional initial states)
shape	`Reshape`, `Transpose`, `Squeeze`, `Unsqueeze`, `Flatten`, `Identity`, broadcast `Add`, `Concat`

Everything statically knowable is resolved at generation time. A folding pass evaluates every constant-only node away — which is also what removes the int64 shape tensors the emitters could never carry. Relabelling ops, and any Transpose that only moves size-1 axes, become buffer aliases: no code, no copy.

All 21 models in goldenFiles/ generate and verify against onnxruntime on both backends, and the two backends agree with each other.

Measured on an A100

Per point, over a million distinct points with data mapped outside the timed window in every harness:

	ns/point
`infer` from a C `target teams distribute parallel for`	1.73
`infer` from the Fortran equivalent	1.75
`infer_batch`, native CUDA kernel	1.47–1.52

Every route through the library costs about the same, which is what the same arithmetic over the same data should cost.

Measured on an MI210

Same harnesses, amdclang/amdflang -fopenmp --offload-arch=gfx90a as the host compilers, hipcc as the device compiler:

	ROCm 7.2.0	AFAR 23.2.1
`infer` from a C `target teams distribute parallel for`	1.98 (embedded) / 5.67 (file-loaded)	3.08 / 5.78
`infer` from the Fortran equivalent	4.73 / 4.62	4.66 / 4.64
`infer_batch`, native HIP kernel	4.04 / 6.62	4.05 / 6.89

Nothing in the generated arithmetic changed to get there. Three things did:

__HIP__ is not a HIP-compilation signal: clang's OpenMP AMDGPU device pass defines it from openmp_wrappers/math.h, so a header that accepted it next to __HIPCC__ emitted static __device__ const into a plain OpenMP host build. hip-clang defines __HIPCC__ itself, so that is the one macro the guards test (the HIP twin of the nvptx __CUDA_ARCH__-without-__CUDACC__ case). A test forces __HIP__ on under a plain compiler.
hipcc does not include its runtime implicitly the way nvcc does; the gate's device harness includes rosenna_rt.h.
hipcc injects -x hip ahead of a .cu and it applies to every later input, so a bare lib.a after the .cu was compiled as source. The gate links it as -L/-l.

No transfers across a time-step loop

The property a solver needs is that the model goes to the device once -- in init, or baked into the device image when embedded -- and is never moved again, on any step. The gate now measures exactly that instead of a single call: every harness runs a 4-step loop over the same resident points with a profiler range around the whole loop, and all three harnesses a configuration runs (per-point C, per-point Fortran, native infer_batch) are profiled. The count of transfers inside the range must be zero.

HIP: rocprofv3 --hip-trace --marker-trace --memory-copy-trace, counting hipMemcpy* API calls plus MEMORY_COPY operations cut to the roctx range (both are needed: a small hipMemcpy is host-staged and never a copy operation; OpenMP offload's copies go over HSA and are never an API call). Zero in all six harness/configuration pairs on the MI210, under both ROCm 7.2.0 and AFAR 23.2.1, with the drivers' setup copies and init's upload visible in the same traces outside the range.
CUDA: the same loops bracketed with nvtx, counting cudaMemcpy* and cuMemcpy* (the driver API nvc's offload uses) in nsys's cuda_api_sum. Zero in all six harness/configuration pairs on the A100 (HPC SDK 25.11), so both device paths are now measured the same way.

The first run of this check caught something: the Fortran per-point harness made two small copies per step, embedded and file-loaded alike -- so not the model. amdflang re-maps an allocatable's descriptor on every target-region entry. The harness now reaches its resident arrays through explicit-shape dummies (no descriptor, no copy), and python/README.md tells a Fortran solver author to do the same or pass raw c_ptrs as the C path does.

Review fixes

Seven findings from a review of this branch are fixed, each pinned by a regression test in python/tests/test_regressions.py on a model shape the golden set cannot express:

A tied initializer (one weight used by two nodes) was loaded once in C (the second copy stayed zero) and was a duplicate case in Fortran. build_plan registers each initializer once.
AveragePool with an end-only pad divided every window by the full kernel; the divisor is now decided from the last window's reach (Spatial.every_window_is_inside).
A real Transpose in a model with no Add used loop counters the Fortran emitter never declared.
Unsqueeze with several negative axes folded to the wrong shape (resolved against ndim+1, not the output rank).
A (1,) Gemm bias passed validation and was read past its end; validate now requires one value per output.
A weights file whose table of contents omitted a tensor returned 0; both loaders track which tensors were filled and return 9 for a missing or duplicated one.
A Constant LSTM initial state, which fold.py claimed to support, was refused by the plan; it is now lowered to two weights.

Also: gpu-gate ran the golden generator in test/, which this branch deletes, and FATALed on a fresh clone. The generator invocation now lives in rosenna/golden.py, shared with the test fixture, and seeds torch so the golden models are the same on every machine (an unseeded gemm_small draw was sometimes dead, and CI failed on one runner and passed on the other). An LSTM whose Y output is omitted is refused by name instead of a KeyError. The two tests whose evidence is "libgomp refused a target region under OMP_TARGET_OFFLOAD=MANDATORY" skip, naming the toolchain, where libgomp has no offload plugins (a plain gcc 12).

Examples: surrogates inside PDE solvers

examples/surrogates/ has four self-contained solvers, each in C and Fortran, with a network inside the time-step loop, organised by where the network sits and the code structure that forces:

	PDE	Structure
`burgers_closure`	coarse-grid Burgers, per-cell closure	embedded model, header-inline `infer` from the solver's offload loop
`reaction_patch`	2-D FitzHugh-Nagumo, learned time step on 3×3 patches	gather → one `infer_batch` → scatter; file-loaded, `init`, native-backend archive, `_sync`
`bubble_lstm`	acoustics through bubbles, LSTM per cell replacing an 8-bin Rayleigh-Plesset population	state `(h, c)` resident on the device; multi-input/multi-output
`poisson_guess`	periodic Poisson, conv-net initial guess for Jacobi	whole-field input, one call per step

Each runs its reference and its surrogate from the same held-out initial condition, prints an error and timings, and exits 0 only if the surrogate did what its README says; both languages give the same numbers on the MI210. make TOOLCHAIN=amd|nvidia|gnu; tests/test_examples.py runs all four on the host in CI, and on nvidia/amd wherever those compilers and a GPU are present (the arch is detected from nvidia-smi/rocminfo rather than assumed). All four also pass on an A100 in both languages. The trained models are checked in.

Building them changed the generator: infer_batch's OpenMP fallback used target teams loop, which amdclang (like nvc) maps one point per team, 3.5 µs per point; teams distribute parallel for/do is 28× faster. <name>_sync and <name>_init_dev came from reaction_patch, and two kernels:

Register-blocked dense layers. A batched GEMM with weights staged in shared memory was built and measured slower than thread-per-point on reaction_patch: the bound was each thread's loads of its own activation vector (scratch memory), not the weights. Computing 8 output columns per pass over the input amortises those loads, 7.9 → 2.3 ms per big step in both languages. It is in the gemm emitters, so infer, the omp fallback and the native kernel all get it; layers narrower than 96 keep one column per pass (gemm_big was 3× slower blocked). The gate's per-point numbers are unchanged.
infer_one for whole-field models: a launch per op, activations in static device buffers; infer_batch runs it per point when the per-point locals would exceed 64 KB. poisson_guess runs its 660 KB-activation conv net on the device at 3 ms per step with no copy in the loop (it ran on the host at 8 ms plus an upload).

examples/cns_closure/ was a patch.md: a documented diff against microfd, a compact compressible solver this repository does not contain. A patch against a file we do not control cannot be compiled, run or tested, so every claim in it was unverifiable by construction -- it said as much itself ("device path unvalidated"). It is now cns.c, our own 3D compressible Navier-Stokes solver (finite volume, MUSCL+minmod, HLLC, Newtonian viscous stress and Fourier conduction, SSP-RK3, no MPI), with the closure called once per cell inside its own offload region and its turbulent viscosity added to the molecular one at every face. The program checks itself -- mass and energy conserved to round-off, density and pressure positive, and the nut the offloaded solver computed equal to a host evaluation of the same model -- and BATCHED=1 additionally requires closure_infer_batch to agree with the header-inline closure_infer directly. It is in test_examples.py on the host and on the GPU toolchains, which is the whole point of vendoring the solver. NO_CLOSURE=1 builds the plain solver as a baseline: on an idle A100 it saturates at 308 Mcell-updates/s (3.25 ns per cell per SSP-RK3 step at 256^3), and the closure costs 25% of wall-clock -- 0.81 ns/cell-step, which is three closure_infer calls at the 0.28 ns measured separately, so the difference and the attribution agree.

Writing it turned up a second nvc codegen bailout: selecting between two local arrays through a pointer (const double *S = SM >= 0 ? L : R, the obvious way to write HLLC's star state) miscompiles under nvc -O2 -mp=gpu -- the kernel dies with CUDA_ERROR_LAUNCH_FAILED at the first write to a mapped array, -O1 is correct, and compute-sanitizer reports no out-of-bounds access. hllc() copies the chosen side instead.

examples/run_basic.sh is the shortest path through the whole tool -- export gemm_small, generate both backends, call the result from C (cAPI.c) and from Fortran (capiTester.f90), verify against onnxruntime -- and it was run by nothing, with both callers hand-written against the generated API. tests/test_examples.py now drives the script itself rather than a copy of its commands, since a copy keeps passing after the script rots. That found two bugs that were real for anyone on an HPC module system: -J for the .mod directory is gfortran's and flang's spelling while nvfortran and ifx want -module (with an SDK loaded FC is already nvfortran, so the script failed on exactly the machine this branch targets), and step 1 hardcoded python3 rather than the interpreter holding torch.

What the wider CI caught

Adding the jobs above was not free of findings -- four defects, none of which the existing suite saw:

Dead LSTM output buffers. The warning-free check was parametrized over a hardcoded list of the five dense models, which stopped being "the generated code" once Conv, pooling, LSTM and the shape ops landed. gcc had been reporting the defect in every run that never compiled an LSTM model: an LSTM always produces Y, Y_h and Y_c, and the emitters wrote all three whatever the model read -- one or two dead buffers and a dead copy in the loop path, per thread on a device. The plan now drops outputs nothing consumes, Y included.
An AveragePool counter set but never read, found by clang on its first run: the divisor is a literal whenever no window can reach outside the input, so the count was pure waste.
The examples did not link under TOOLCHAIN=nvidia -- that branch of common.mk had never been run. -lcudart needs a -L the HPC SDK does not put on the default search path, and nvcc compiling <model>_kernel.cu as C++ leaves __cxa_guard_acquire undefined when a C or Fortran driver links the archive. Both are now nvc's own flags (-cuda -c++libs), so neither hardcodes a path or a C++ runtime.
Two macOS-only failures: the sanitizer job assumed "compiler exists" meant "can link a sanitized binary" (Homebrew gcc has no ASan runtime there; it probes now), and burgers_closure's 200 steps of target regions timed out, so its step count is overridable as poisson_guess's already was.

flang and ifx found nothing, which is its own result: the generated C11 and Fortran 2008 builds and runs identically under five independent toolchains.

One deliberate change to the generated arithmetic

A dense layer now adds its bias after the dot product rather than seeding the accumulator with it. This reorders the sum by one term and changes generated results in the last ulp on every backend, CPU included. That trade was considered and accepted: ulp-level movement is not meaningful for this library's use, and the alternative is a 30x slower per-point GPU path.

No bug has been filed with NVIDIA, and the reproducer is kept for the next reader rather than for a report; it is worth re-checking against a later HPC SDK, since a fixed compiler would make the reordering unnecessary.

It is there because nvc will not compile a target teams distribute parallel for body whose accumulator is seeded from a declare-target array — it emits a kernel that traps at runtime, with no diagnostic beyond an unrelated-looking warning. Without the reordering the only form nvc accepts is target teams loop, which maps one point to one team (1,000,000 blocks of 32 threads, one active lane each) and costs ~30x more.

Both emitters changed together, so the backends stay bit-comparable, and every tolerance in the suite passes. The bailout is an NVIDIA compiler defect and should be filed; the reordering is a workaround, not a fix. doc/nvhpc_teams_mapping/ has the PTX, the ncu launch geometry, the stall counters, and a 45-line self-contained reproducer.

Also in that directory: embedded weights went to __constant__ up to 48 KB, a bound taken from the 64 KB per-module bank — which is a correctness limit, not a performance one. ncu showed 71% of warp-issue stalls on constant-cache misses. The threshold is now 2 KB, worth 2.5–2.8x on the two largest golden models.

fLibrary retired

fLibrary/, modelParserONNX.py, modelCreator.fpp and the test/ shell suite are removed. python/tests/test_golden_suite.py replaces run.sh: the same 21 models, both backends, compared against onnxruntime rather than against recorded output. A recorded golden file pins whatever the library did the day it was recorded, so a wrong-but-stable implementation records its own error as the expectation.

Root README, python/README.md, doc/methodology.md and doc/opensource.md are rewritten for the single remaining path; doc/opensource.md is now a guide to adding an operator to the generator. The root examples/ build generated code instead of libcorelib.a.

Two layout cleanups came with it. python/examples/ was a second directory named "examples" holding two things that were not peers: a performance writeup, now doc/nvhpc_teams_mapping/ beside the methodology it supports, and an integration example, now examples/cns_closure/. And .gitignore was a whitelist -- *, then !*.py, !*.c, and a fresh exception every time something new needed to ship. Under it git check-ignore matched the root README.md and LICENSE, which survived only by already being in the index, and a new .h, .txt, .json or Dockerfile would have been invisible to both git status and git add .; sixteen of its 42 lines named paths this branch had already deleted. It now ignores what is generated, verified equivalent -- git ls-files byte-identical at 115 files, nothing newly untracked, every previously-ignored artifact still ignored.

OpenACC

Both offload families work, and they are not equal. A per-point OpenACC host against the one-archive cuda build is correct on an A100 (1.11e-16, NVCOMPILER_ACC_NOTIFY confirms a real device kernel), but:

#pragma acc parallel loop over points lets nvc map one point per gang (num_gangs=4096, vector_length=32) and then auto-vectorise infer's own loops across the lanes -- the OpenACC shape of the teams loop cliff in doc/nvhpc_teams_mapping/. gang vector gives one point per thread and was 1.9x faster.
Over a million points, gemm_big per-point is 6.03 ns with OpenACC against 1.74 ns with OpenMP target.
rosenna gpu-gate cannot measure OpenACC: its per-point harnesses carry omp target pragmas, so --flags "-acc=gpu" compiles them away and times the host. Confirmed pre-existing by running the same gate at 922526a.

Known limitations

Whoever's device code is in the archive does the final link. Built with offload flags, the archive holds the host compiler's OpenMP fatbin, which nvcc's device-link step cannot register, so the link goes through the host compiler (nvc -cuda) -- which is what a solver uses anyway. Built without them there is no fatbin and nvcc links it directly. (This replaces the old limitation, where a per-point host needed a second, omp-backend archive.)
A model with large intermediate activations declares them as per-thread locals in infer — mnist's come to ~100 KB at f64, poisson_guess's whole-field conv net 660 KB, which amdclang refuses for a device thread. Above 64 KB, infer_batch runs infer_one per point instead; the per-point infer itself is host-only for such a model, and a Fortran module containing it is built without offload.
Spatial ops are 1-D or 2-D; 3-D (rank-5) would need the global rank cap raised. Softmax normalises the last axis only. Pad takes constant, edge and reflect with constant pads (crops included), and reflect is limited to one reflection, so a pad must be narrower than its axis.

TODO

https://claude.ai/code/session_01QoLws1dnkpAkCvnywpPuGc

…d weights by default

…pies weights to the device The batched entry point <name>_infer_batch(n, x, y, stream) now exists in both forms of the C library. Under ROSENNA_BACKEND=cuda|hip it is a one-thread-per- point kernel in <name>_kernel.cu calling the header-inline infer, built by nvcc or hipcc through the runtime macros in rosenna_rt.h; under the default omp backend it is a target teams loop over device pointers (is_device_ptr) in <name>.c. init is the plan step: it loads the file, copies each weight to the device (status 10 on failure) and publishes the addresses to the kernel's translation unit; nothing in the loop path allocates, transfers or synchronizes (controller ruling R5). ROSENNA_CONST is __constant__ under 48 KB of embedded weights and __device__ const above (ruling R4). The generated Makefile selects the backend; generate writes the kernel and the runtime header; a compile-only nvcc job is added to CI. Device path unvalidated until that job and the GPU gate run.

…no-transfer sweep Review round 1 of the batched kernel. The header now carries <name>_device_bind_here(), the per-translation-unit bind of the __constant__ weight table, which the library's own device_bind forwards to and which any user kernel's translation unit must call after init (ruling R8). The host instantiation of an embedded model's __host__ __device__ infer is an assert stub in a CUDA/HIP build instead of a read of device storage (R9). Launches are checked with a peek, status 11 (R10). infer_batch's pointers are restrict-qualified (R7). The structural no-transfer test sweeps the whole .c outside init and its upload tail, with the OpenMP/OpenACC transfer tokens added (R6), and the stubbed C++ test now compiles the kernel unit with __CUDA_ARCH__ defined so infer reads through the table. Macros are #undef'd before definition; HIP_SYMBOL is explained.

…nit contract ROSENNA_LAUNCH_STATUS() now maps to cudaGetLastError()/hipGetLastError() instead of the Peek variants: Peek leaves a stale error in the thread slot, so every later call after one real failure would report status 11 forever, while Get reports once and clears (both stay non-synchronizing). Also fixes the device_bind_here header comment, which said "once after init" -- a repeated init reallocates the device buffers, so every call to init needs its own device_bind_here to follow it.

…opy in init, infer_batch Gives the Fortran module the same device contract Tasks 1-3 gave the C header: infer carries !$omp declare target / !$acc routine seq; an embedded plan's weights are parameter arrays (reshape'd from a flat, C-order literal list so the memory layout matches the file-loaded `protected` form) with no init at all; a file-loaded plan's init makes its weights device-resident with !$omp target update to / !$acc update device right after the load, in the same host call (Fortran has no runtime-API path without CUDA Fortran). A new infer_batch(n, x, y, status) is the OpenMP-target fallback over device-resident arrays (has_device_addr(x, y), no map/copyin/copyout/update -- ruling R5), verified accepted on an explicit-shape dummy under gfortran 15 with no fallback needed. A declared-only bind(C) interface exposes Task 3's native <name>_infer_batch to Fortran hosts holding device pointers. emit_fortran_recipe/<name>_fortran.mk mirrors the C recipe, building lib<name>.a from <name>_model.o. A weight tensor large enough to need more than 255 continuation lines in one array constructor (gemm_big's 1200-element layer) is split into several <=500-element parameter chunks and reassembled with a second, short reshape statement, since F2008's per-statement continuation cap is independent of the 132-column limit already handled. Also flips ruling R3 (Task 2): now that Fortran embeds by default too, `generate` writes a .rwt only for embed=False plans, in every language; verify.py's Fortran build now compiles the module to an object, archives it, and links the driver against the archive, matching the C backend's library-form path. Existing Fortran tests that are specifically about the file-loaded contract (`_init`, `protected` declarations, the load_tensor case-label select) now build with embed=False explicitly, mirroring the same fix already applied on the C side when embedding was introduced there.

…in one directory Ruling R13. generate --lang both writes both <name>.mk (C) and <name>_fortran.mk (Fortran) into one output directory, and both recipes archived into lib<name>.a: ar rcs appends, so building both there in sequence silently merged gemm_small_model.o into the C archive, and either recipe's `clean` then deleted the shared file. Reproduced and confirmed with a forced rebuild (ar t showed both gemm_small.o and gemm_small_model.o in one lib<name>.a). The Fortran archive is now lib<name>_f.a; a Fortran host that also links the native CUDA/HIP kernel links both archives: -l<name>_f -l<name>. verify.py's Fortran build follows the same name even though its fortran/c backends already build in separate directories, so the two names are never the same anywhere a caller might build both recipes in one place. Adds the two-archive test that would have caught this (build both recipes into one directory via `generate --lang both`, assert two distinct archives each containing only its own object per `ar t`, and that cleaning one leaves the other's archive and object alone), plus test_generated_fortran_is_warning_free_under_openacc mirroring the existing C test -fsyntax-only under -fopenacc for all five dense models, both embed states.

Adds `rosenna gpu-gate`, the script a user runs on a real GPU machine to validate the device library: it generates gemm_big embedded and file-loaded in both languages, builds the C library with the chosen batched backend and the Fortran library with the host compiler, then runs three harnesses (C per-point, Fortran per-point, and infer_batch over device-resident data), each checked against onnxruntime and timed per point, and writes every command and its output to gate-report.md. --host-fallback drops the OMP_TARGET_OFFLOAD=MANDATORY requirement so the omp backend, and the script itself, can be exercised end to end on a machine with no accelerator; that is how tests/test_gate.py runs it here. Moves _live_reference from tests/test_emit_fortran.py into rosenna/verify.py so gate.py can reuse it without importing test code; every existing import of it keeps working via a re-export. Adds the microfd_closure worked example: closure.py exports a 9-16-1 Tanh MLP closure model, and patch.md documents (as an unapplied diff against a reference copy of microfd.c) the per-cell closure() kernel, the muf blend in face(), the g.nut allocation and device mapping, and the batched alternative via closure_infer_batch. README.md states the device path is host-validated only until gpu-gate has actually run on a GPU machine. .gitignore gains an allowlist entry for python/examples/**/*.md.

… check, time only the call Fixes from review round 1 of gate.py (two Critical, three Important; nothing CUDA-related has ever run on this machine, so these are traced-by-reading fixes, not verified-by-running): - Ruling R14 (Critical): _run_c_harness1 now compiles gate_harness1.c with the host compiler and its offload flags as before, but links it with the device compiler (--devcc) for --backend cuda|hip instead of the host compiler and -lm. A file-loaded plan's lib<name>.a was built by nvcc/hipcc for those backends (init's upload/bind calls cudaMalloc/cudaMemcpy/cudaMemcpyToSymbol or the hip equivalents), so the plain host compiler left those undefined on every real run. --backend omp keeps the host compiler as the link driver. - Ruling R15 (Critical): _run_nsys_check no longer counts "cudaMemcpy" over the whole profiled harness (which legitimately memcpys in its untimed setup and would fail an R5-compliant infer_batch). It now brackets only the timed infer_batch call with an nvtx range (-DROSENNA_GATE_NVTX=1 in the .cu driver), profiles with --capture-range=nvtx --nvtx-capture=rosenna_timed, and reads the count from `nsys stats --report cuda_api_sum --format csv`. Skips (without failing the gate) when nsys is absent, when the nvtx header does not compile, or for --backend hip (rocprof scoping is a follow-up). - Ruling R16 (Important): the omp infer_batch harnesses (C and Fortran) now time only the infer_batch call -- target enter/exit data move outside the t0/c0..t1/c1 window, matching what the cuda/hip driver already did. - Important: --help gains an epilog with the three concrete host-compiler pairings (nvc/nvfortran+cuda, amdclang/amdflang+hip, gcc/gfortran+omp+host-fallback); --devcc's help text no longer says "required" (it has a default). README's example swaps the gcc-15/--backend cuda pairing (gcc-15 cannot offload) for the nvc one. - Ruling R17 (Important): closure()'s ghost-layer gap in patch.md is fixed, not just documented -- it now writes nut over the padded block minus one layer on each side (the widest range where its central difference has both neighbours in bounds, which the reference microfd.c's halo() has already filled), not FOR3's interior-only range, so face()'s read one ghost cell into the low boundary is no longer silently zero. The batched alternative's gather loop is fixed the same way; its own remaining edge case (closure_infer_batch is called over all nc cells, one layer wider than the gather loop fills) is called out explicitly instead of reusing the old blanket caveat. Minor: patch.md's no-op -/+ diff pair in section 2 is now a context line; --flags/--devflags carry a comment pointing at _join_dash_valued_options.

…ink, drop the live-reference assert Fix round 2 (two new rulings added after round 1 went out, plus one pre-existing test defect the reviewer found; nothing CUDA-related has run on this machine, so R18/R19 are traced-by-reading fixes): - Ruling R18: _sum_cudamemcpy_calls now returns a structured NsysParseResult(parsed, count) instead of a bare int. `parsed` is True only when the Name/Num Calls columns were recognised AND at least one data row was actually read; an unparseable CSV (no fieldnames, no recognised columns, or a header with zero data rows) comes back as parsed=False, never as a bare 0. _run_nsys_check now treats parsed=False as a gate FAILURE ("nsys check inconclusive: could not parse cuda_api_sum"), with the first lines of what nsys actually produced in the report, rather than the previous silent pass. New unit tests in tests/test_gate.py exercise all three cases from the ruling: two cudaMemcpy* rows summing to 5 plus a launch row (parsed, count=5); zero memcpy rows plus a launch row (parsed, count=0); and empty/unrelated CSV text (not parsed, in both cases). - Ruling R19: _run_c_harness1's cuda/hip link line now forwards the HOST offload flags (e.g. nvc's -mp=gpu -gpu=cc80) into the device-compiler link, not just --devflags -- new _host_flags_for_devcc_link wraps each host flag as -Xcompiler <flag> for cuda (nvcc otherwise treats an unrecognised flag as compiler-only) and passes them through directly for hip (hipcc is clang-based, like the expected amdclang host pairing). Without this, gate_harness1.o's own OpenMP-target runtime (from the host compiler) could be unresolved at a link driven purely by devflags. README.md's example now also documents the alternative -- host compiler as the link driver with the runtime named explicitly (-L$CUDA_HOME/lib64 -lcudart / -L$ROCM_PATH/lib -lamdhip64) -- for a toolchain that rejects the device-compiler-as-driver form. - Ruling R20 (pre-existing defect): test_device_fortran.py:: test_fortran_target_regions_are_real called _live_reference and asserted `inputs is not None` instead of skipping on a dead model, unlike every other _live_reference call site in the test suite -- it failed deterministically once gemm_small's current (unseeded) weights happened to produce a dead model. The MANDATORY check only needs one point and a libgomp refusal, not a live onnxruntime reference, so it now uses a constant input (np.full(..., 0.5)), mirroring test_device_c.py::test_target_regions_are_real, which already did this correctly. The file's one other _live_reference call site already used the pytest.skip dead-model guard; no other fix was needed there. Full suite: 165 passed, 18 skipped, 0 failed (162 + 3 new _sum_cudamemcpy_calls unit tests; the previous round's 161/18/1 is now clean).

…he host or a GPU Every command and code block in python/README.md is checked by tests/test_docs.py: the C and Fortran examples are generated from gemm_small and compiled/run as printed, the status table is rendered from abi.STATUS_CODES, and every `rosenna ...` line parses. Also fixes two test_device_c.py tests that used _live_reference's result without the dead-model pytest.skip guard the sibling tests already have.

…ance Review round 1: documents the silent __constant__ vs __device__ const cutoff at 48 KB (ruling R4) in "Call it from C", states that `verify` compares at the ONNX model's own dtype tolerance rather than --precision's in "Verify", and extends the Limits bullet on unset file-loaded weights to cover the omp infer_batch fallback (no status code there; only cuda/hip returns 10).

…ase device copies on a failed bind - The device-pass guard is now (__CUDACC__ && __CUDA_ARCH__) || ((__HIPCC__ || __HIP__) && __HIP_DEVICE_COMPILE__) (ruling R23): clang's OpenMP nvptx device pass defines __CUDA_ARCH__ without __CUDACC__ and selected the __constant__ table that only the CUDA/HIP guard declares. - Every CUDA/HIP guard in the header, source and rosenna_rt.h accepts __HIP__ next to __HIPCC__, so the header does not depend on the hipcc wrapper's flag order. - infer is emitted as 'static inline ROSENNA_DEVICE_FN', the order CUDA's own headers use. - A failed device_bind (or allocation, or copy) in init releases and nulls every device copy, so infer_batch returns 10 rather than launching over a stale table. - The kernel's launch-check comment names GetLastError (ruling R11); _emit_upload's docstring is reflowed. - Tests: the device-pass guard under a plain compiler with __CUDA_ARCH__ forced on; the three-macro header test also covers the file-loaded header; the release helper is asserted.

…er-point harness, host compiler drives its link - Ruling R21 (C1): under --backend cuda|hip the gate builds two C archives per configuration: the omp-backend lib<name>.a with --cc and --flags in the configuration directory, linked by the per-point C harness, and the cuda|hip lib<name>.a with --devcc in <backend>_lib/ (a copy of the sources), linked by the infer_batch .cu driver. Both build lines are recorded. --help states the underlying library limitation. - Ruling R22 (C2): the host compiler compiles and links harness 1 in every backend, against the omp archive; _host_flags_for_devcc_link and the -Xcompiler forwarding are gone. The device compiler compiles and links only the .cu driver. - I2: the Fortran infer_batch harness (and the test host) wrap the call in target data use_device_addr, so has_device_addr sees device addresses. - I3: no -x cu|hip on the .cu driver lines (a -x before the archive makes clang-based hipcc compile the archive as source). - I4: nsys stats runs with -q and the parser starts at the first line naming Num Calls and Name, skipping the stdout preamble; test added. - I6 (ruling R24): -Wall -Wextra -std=c11|f2008 only when the compiler's basename starts with gcc, gfortran, cc or clang; CFLAGS=/FFLAGS= are passed to both recipes explicitly; test added. - I9: nsys profile gets --capture-range-end=stop. - --devcc is shlex-split wherever it becomes argv; a missing compiler in the version probe is recorded, not raised; report command lines are shell-quoted; the per-point timing loops state the benign output race.

…nly fails on compiler-absence skips - I5: patch.md section 5 adds nut=g.nut to face()'s locals line and reads nut[c], nut[c+s] inside the FOR3 region, matching microfd's rule that kernels copy g.* to locals before offloading (closure() in section 3 already did). - I7 (ruling R25): the device-path CI step sets pipefail so pytest's exit status survives the tee, and greps only for the compiler-absence skip reasons; a dead-model skip is not a CI failure.

… hosts, the runtime link flags, and what CI has and has not run - python/README.md: the opening says generate writes sources and recipes and make builds the archives, and points at Verify's residency caveat; ROSENNA_BACKEND / --backend are a recipe variable and a gate flag, not a generate flag; the Call it from C section states that a cuda/hip archive does not serve the per-point path of a file-loaded model and gives the omp-backend build line for the host compiler; the bind(C) route names the runtime link flags (-lcudart / nvfortran -cuda, -lamdhip64); the Verify section no longer claims nvcc/hipcc compile in CI; Limits carries the limitation and its planned resolution. - examples/microfd_closure/README.md: the two-archive gate build replaces the retired device-compiler link and -Xcompiler forwarding text.

…rtran -fopenacc can compile the module gfortran 15 -fopenacc materializes a parameter array read inside a routine seq as a static, demands an OpenACC declare for it, and refuses a declare on a named constant, so an embedded Fortran module did not compile at all under -fopenacc (-c; the diagnostic comes after the front end, and the suite's -fsyntax-only check let it through). The embedded weights are now initialized protected module arrays with declare copyin, the same declare-target line as before; the chunk arrays stay parameter since only the initializer names them. The OpenACC test compiles for real, the module test asserts the new declarations, README wording follows.

…of aborting the run _sh caught only FileNotFoundError; a produced-but-not-executable output (PermissionError) escaped to the FATAL handler and skipped every later configuration. Any OSError from the exec is now a 127 step in the report.

…the C tests already do

… libgomp ignores MANDATORY On the macOS CI runner (Homebrew GCC 13.4), test_fortran_target_regions_are_real ran a real !$omp target region to completion and returned 0 instead of being refused under OMP_TARGET_OFFLOAD=MANDATORY, while the identical C test passed on the same runner and compiler. MANDATORY enforcement is therefore not portable evidence on its own. Both tests now compile their host source to a standalone object (-c) and assert via nm that it references GOMP_target_ext -- a real target region cannot be compiled without a call to it -- as the primary, platform-independent evidence. The MANDATORY run is kept as corroborating evidence: it still passes the test where libgomp enforces it, and is downgraded to a pytest.skip (naming the platform and compiler) rather than a failure where it does not.

…retire fLibrary Three pieces of work, all driven by running the code on a real A100 (NVIDIA HPC SDK 25.11, nvc/nvfortran/nvcc 13.0) rather than reasoning about it. GPU gate, first run on real hardware ------------------------------------ `rosenna gpu-gate --backend cuda` now PASSes, and the R5 device-residency claim is measured rather than asserted: zero cudaMemcpy inside the timed infer_batch call, embedded and file-loaded. Four defects it found: * nsys never opened its capture range (a plain nvtxRangePushA needs NSYS_NVTX_PROFILER_REGISTER_ONLY=0), resolved its -o path twice, and refused its own --stats export. All three exited 0 while producing nothing, so each is now checked explicitly. * nvfortran does not implement `has_device_addr` at all. The module is emitted as .F90 and the clause is chosen by the preprocessor, so gfortran keeps the standard OpenMP 5.1 spelling. Per-point offload was 30x slower than the batched kernel -------------------------------------------------------- Not IPO/LTO and not register pressure -- nvc's per-thread code is the better of the two. nvc maps one point to one *team*: 1,000,000 blocks of 32 threads with a single active lane each. `distribute parallel for`, which would use the whole block, made nvc emit a kernel that traps. The trigger is an accumulator seeded from a declare-target array, which is exactly how a dense layer is written, so both emitters now add the bias *after* the dot product. That reorders the sum by one term and changes results in the last ulp. Embedded weights also went to __constant__ up to 48 KB, chosen against the 64 KB bank -- a correctness bound, not a performance one. ncu showed 71% of warp-issue stalls on constant-cache misses. The cut is now 2 KB. Per point, every route through the library now costs the same: 1.5-1.8 ns. Operator coverage: 5 of 21 golden models to 21 of 21 ---------------------------------------------------- * rank 1-4 throughout; buffers stay flat and row-major, and a Spatial spec carries literal extents so no ONNX attribute reaches the emitters. * Conv, MaxPool, AveragePool, including auto_pad resolved at generation. * A folding pass evaluates every constant-only node away, which is also what removes the int64 shape tensors the emitters cannot carry. * Reshape/Squeeze/Unsqueeze/Flatten/Identity, and any Transpose that only moves size-1 axes, become buffer aliases: no code, no copy. Other Transposes gather. Add broadcasts a constant. * LSTM: forward, ONNX i/o/f/c gate order, optional initial states. Several graph inputs arrive concatenated in x, which keeps infer(x, y) and with it the batched entry point and the whole device contract. * Gemm gained a row count: an LSTM sequence feeding a Gemm applies it per timestep, which the single-row lowering got wrong by 3e-2. * verify allows a cancellation term in its tolerance (n*eps*scale). onnxruntime blocks and vectorises, so two correct implementations differ by more than rtol*|expected| when a sum cancels. mnist forced this: its f32 build is 4.6e-6 off ORT, and its f64 build matches an independent float64 reference to 2.4e-15. fLibrary retired ---------------- Every golden model now generates and verifies against onnxruntime on both backends, so the runtime library, its parser and the shell suite that drove it are removed. test_golden_suite.py replaces run.sh: 21 models, both backends, compared against onnxruntime rather than recorded output. Both READMEs and both doc/ files rewritten; the root examples/ now build generated code instead of libcorelib.a. 235 tests pass. CUDA gate and gcc/gfortran host-fallback gate both PASS.

Deleting test/ took nnLSTM.py with it, and lstm_cell, lstm_gemm, lstm_gemm_hid and lstm_output all `import nnLSTM`. Local runs did not catch it because their .onnx files already existed, so the generators never ran; CI regenerates from a clean checkout and failed on both platforms with ModuleNotFoundError. The helper now lives beside the models it serves, in goldenFiles/, and conftest puts that directory on PYTHONPATH for the scratch cwd the generators run in. Also: goldenFiles/mnist/mnist.py reads its .onnx rather than writing one -- that model is checked in -- so the fixture now says that plainly instead of returning a path that does not exist. Verified by deleting every generated goldenFiles/*.onnx and running the suite from scratch: 235 passed.

No bug is being filed. The reproducer stays as the evidence for why a dense layer adds its bias where it does, and as the thing to re-run against a later HPC SDK -- a fixed compiler would make the reordering unnecessary.

…ped it rosenna gpu-gate --backend hip PASSes on an AMD Instinct MI210 (gfx90a) under ROCm 7.2.0 (amdclang, amdflang -fopenmp --offload-arch=gfx90a, hipcc) and under the TheRock AFAR 23.2.1 drop; both reports are in gate-reports/. Nothing in the generated arithmetic changed. What did: - __HIP__ is dropped from every CUDA/HIP guard. clang's OpenMP AMDGPU device pass defines it from openmp_wrappers/math.h, so the header took the __device__ branch inside a plain OpenMP host build. hip-clang defines __HIPCC__ itself for any HIP compilation, so that is the one macro tested (the HIP twin of ruling R23). A test forces __HIP__/__AMDGCN__ on under a plain compiler and asserts the host branch compiles. - The gate's device harness includes rosenna_rt.h: nvcc includes its runtime implicitly, hipcc does not. - The gate links the device archive as -L/-l: hipcc injects -x hip ahead of a .cu input and it applies to every later input, so a bare lib.a was compiled as source. - The gate ran the golden generator in test/, which this branch deletes. How a generator is run now lives in rosenna/golden.py, shared with the test fixture that already did it right. The rocprof-scoped transfer count is still a follow-up; the report says so. Claude-Session: https://claude.ai/code/session_01QoLws1dnkpAkCvnywpPuGc

Every model in tests/test_review_regressions.py is a shape the golden set cannot express; each test failed before its fix and passes after. - A tied initializer (one weight used by two nodes) was one WeightSpec per use: C's loader filled the first and returned, leaving the second all zeros; Fortran's select-case had two labels for the same name and did not compile. build_plan now registers each initializer once and hands a second user the first symbol. - AveragePool with an end-only pad divided every window by the full kernel: the emitters tested only the begin pads for "no pad cell can fall in a window". Spatial.every_window_is_inside reads the last window's reach off the output extent, and both emitters ask it. - A real (non-alias) Transpose loops over c0, c1, ... which the Fortran emitter declared only when the model also had an Add. - Unsqueeze with several negative axes resolved them against ndim+1 rather than the output rank, folding c[3] with axes [-1,-2] to (1,1,3) instead of (3,1,1): a silently wrong constant wherever both shapes broadcast. - A (1,) Gemm bias, a legal ONNX broadcast, passed validation and the emitted b[i] read past it. validate now requires one value per output. - A weights file whose table of contents omitted a tensor returned 0 and infer ran on zeroed arrays. Both loaders keep a per-weight seen flag: a tensor named twice or not at all is status 9, as the ABI documents. - fold.py said a Constant LSTM initial state was supported, and the plan then refused it as "must be a rank-3 value". It is now lowered to two weights (init_syms), read by the emitters the way a caller-supplied state in x is; validate accepts either form, not a mix. The HIP gate still PASSes 6/6 on the MI210 after the loader change. Claude-Session: https://claude.ai/code/session_01QoLws1dnkpAkCvnywpPuGc

…nores it, LSTM without Y is refused - The gate's ruling-R15 check now has a HIP twin. rocprofv3 has no nsys-style capture range, so the infer_batch harness is rebuilt with a roctx range around the timed call, run under `rocprofv3 --hip-trace --marker-trace -f csv`, and the HIP API trace is cut to the range's timestamps (_scope_hipmemcpy_calls, with the same parsed-vs-inconclusive contract as the nsys parser and unit tests for both halves). On the MI210 under ROCm 7.2.0 and AFAR 23.2.1: zero hipMemcpy inside the timed call, embedded and file-loaded, with the driver's setup copies and init's weight upload visible in the same trace outside it. The harness template carries one generic marker bracket (nvtx or roctx) instead of an nvtx-only one; _run_transfer_check picks the profiler by backend. - test_gate's MANDATORY test and test_kernel's target-loop test assumed libgomp refuses a target region under OMP_TARGET_OFFLOAD=MANDATORY with no device; a libgomp built without offload plugins (gcc 12 here) ignores it. Both now go through one conftest probe that skips, naming the toolchain, as test_device_* already did; the kernel test also gains the GOMP_target_ext symbol check as its platform-independent evidence. - An LSTM whose Y output is omitted (outputs ["", "yh"]) raised KeyError('') in build_plan; validate now refuses it by name. Claude-Session: https://claude.ai/code/session_01QoLws1dnkpAkCvnywpPuGc

The file opened with `*` and `!*/`, then un-ignored one extension at a time. That inverts the failure mode: a file whose kind nobody thought to whitelist is not just untracked, it is invisible -- absent from `git status`, skipped by `git add .`, and gone at the next clean checkout. `git check-ignore` matched the root README.md and LICENSE; both survive only because they are already in the index, and a new .h, .txt, .json, .cpp or Dockerfile would not be seen at all. Sixteen of the 42 lines named paths that no longer exist (fLibrary/, test/, the .fpp files of the retired runtime library, three golden models). The replacement ignores what is generated. Verified equivalent: `git ls-files` is byte-identical at 115 files, `git status` reports nothing newly untracked, and every previously-ignored artifact still checks as ignored -- the entries that moved are .venv/, __pycache__/ and .egg-info/, which `!*/` had been un-ignoring as directories and which now collapse under one rule.

…un_basic.sh `python/examples/` was a second directory named "examples" holding two things that were not peers of each other. nvhpc_teams_mapping is the writeup of the per-point performance investigation -- PTX, ncu counters, a reproducer -- so it goes to doc/ next to methodology.md, which already linked it. microfd_closure is an integration example, so it joins examples/. Every cross-reference is rewritten: doc/methodology.md, doc/opensource.md, python/README.md (two), the emit_c docstring that points a reader at the bias-ordering evidence, and the nvhpc README's own mention of patch.md. microfd_closure's README dropped its `cd python` and now uses the installed `rosenna` CLI, as patch.md already did. examples/run_basic.sh -- generate gemm_small, call it from C (cAPI.c) and from Fortran (capiTester.f90), verify against onnxruntime -- was run by nothing. It is the shortest path through the whole tool and the first thing a reader tries, and both callers are hand-written against the generated API: the header's signature, the module name, whether an `_init` is needed. A change to what `generate` emits would have left them stale with only a reader to notice. The test drives the script itself rather than a copy of its commands, which is the point -- a copy keeps passing after the script rots. To make that possible the script takes PYTHON/ROSENNA/CC/FC overrides, defaulting to what a reader would type. Two things surfaced while wiring it up, both real for a user on an HPC module system rather than artifacts of the test: - `-J` for the .mod directory is gfortran's and flang's spelling; nvfortran and ifx want `-module`. With an HPC SDK loaded FC is already nvfortran, so the script failed on exactly the machine this repository targets. It now picks the flag from the compiler's version output -- matched against the whole output, since nvfortran prints a blank first line. - step 1 hardcoded `python3`, which is not the interpreter holding torch when the project is installed in a virtualenv. The test asserts both callers print the same three numbers and that `verify` reports `ok` for each language. The second assertion carries the weight: gemm_small is re-exported unseeded every run and sometimes comes out all zeros, which makes "the two agree" a comparison of 0 with 0. 430 passed, 26 skipped, coverage unchanged at 97.99%.

patch.md was a documented diff against microfd, a compact compressible solver this repository does not contain. A patch against a file we do not control cannot be compiled, run or tested, so every claim in it was unverifiable by construction -- it said so itself ("device path unvalidated"), and it would have drifted silently as either side changed. cns.c is our own implementation: 3D compressible Navier-Stokes, finite volume on a periodic box, MUSCL with a minmod limiter, HLLC, full Newtonian viscous stress plus Fourier conduction, SSP-RK3, no MPI. The closure is a rosenna MLP mapping the nine velocity-gradient components to a turbulent viscosity, called once per cell inside the solver's own offload region and added to the molecular viscosity at every face, so its output feeds real numerics. microfd still gave the example its shape (padded blocks, one array per quantity, the g/LOCALS/IDX conventions); no microfd source is used. The program checks itself and `make` succeeding is the assertion: mass and total energy conserved to round-off, density and pressure positive and finite, and -- the roseNNa claim -- the nut field the offloaded solver computed equal to a host-side evaluation of the same model on the same primitives. BATCHED=1 builds the gather-and-one-call path, links lib<model>.a, and additionally requires closure_infer_batch to agree with the header-inline closure_infer directly rather than each against the host. Measured, NX=16 NSTEPS=5: gnu 4.2e-15 mass drift, nvidia 1.1e-16, both to the same min density and pressure in every printed digit; nut vs host 0 and 1.4e-19; batched vs per-point exactly 0 on the A100. Two things found while writing it, both recorded where they bit: - Selecting between two local arrays through a pointer -- the obvious way to write HLLC's star state, `const double *S = SM >= 0 ? L : R` -- miscompiles under nvc -O2 -mp=gpu. The kernel dies with CUDA_ERROR_LAUNCH_FAILED at the first write to a mapped array, -O1 is correct, and compute-sanitizer reports no out-of-bounds access. hllc() copies the chosen side instead. - A pointer pre-offset on the host into the middle of a mapped region (`g.F + d*NV*nc`) is not translated on the device: the runtime maps a base, so face() captures the base and puts the offset in the index. common.mk gains LANGS/GENLANG so a C-only example can borrow its toolchain block rather than restate the offload and NVIDIA link flags. Its `run` target now chains the drivers with && instead of `;` -- make checks one exit status per recipe line, so the semicolon form reported success whenever the last language passed, whatever the earlier ones did. tests/test_examples.py covers cns_closure on the host (both call paths) and on the GPU toolchains, which is the point of vendoring the solver: patch.md could not be tested at all. 433 passed, coverage unchanged at 97.99%.

CI caught what I did not: `closure.onnx` was in .gitignore, and unlike the four surrogates -- whose models are checked in -- nothing regenerated it, so `make` in a fresh tree died with "No rule to make target 'closure.onnx'". All three jobs that run the full suite failed. My local runs passed only because I had run closure.py by hand earlier in the session, which is exactly the trap the deleted test/nnLSTM.py sprang on this branch before. The model is now tracked, as the surrogates' are. It is deterministic (torch.manual_seed, 1.2 KB), and `make train` reproduces the tracked file byte for byte -- verified by md5. The surrogates' regenerator is train.py and this example's is closure.py, so common.mk takes TRAINER rather than hardcoding the name; the default is unchanged. Verified the way the first attempt should have been: built the example from a tree exported from the git index alone, so nothing untracked could satisfy it.

I said the example worked without ever measuring it: cns.c printed no timings, unlike the four surrogates. It now reports loop time, the closure's share, and ns per cell per call, so the numbers in the README come from the program. A100 80GB, NVHPC 25.11, 20 steps. Per cell per closure call: per-point 64^3 0.30 ns (14.4% of the step) per-point 128^3 0.28 ns (21.7%) batched 64^3 0.56 ns (23.9%) batched 128^3 0.43 ns (29.4%) host, gcc, 128 threads 65 ns host, gcc, 1 thread 305 ns The per-point path beats the batched one here. Splitting the batched path at 128^3: gather 42.7 ms, infer_batch and its sync 13.4 ms, rescale 1.6 ms. The native kernel is the cheapest part -- 0.097 ns per cell, about a third of A100 fp64 peak for a net with 16 tanh -- and the gather feeding it costs three times the inference, writing nine doubles per cell and reading them straight back. The fused per-point path never materialises the features. Batching pays once a network is big enough that per-call overhead outweighs that traffic; 177 parameters is not it. Measuring found a real bug in the example. closure_batched() never called closure_sync(), and closure.h is explicit that the launch "is asynchronous on it" -- the rescale loop right after it reads what the kernel wrote. It printed the correct answer every time anyway, because nvc's OpenMP target regions happen to serialize against the CUDA default stream: an implementation accident, not a guarantee, and exactly the latent race _sync exists for. It also made the timing lie, which is how I noticed. Unsynced, infer_batch measured 0.65 ms for 60 launches over 2.2M cells -- 68 TFLOP/s of fp64, seven times the card's peak -- because the cost landed in the next synchronizing region and showed up as an impossibly slow rescale. One untimed rhs_eval runs before the timed loop so module load and JIT do not land in step 1. It writes only w, nut and F, never q, so the answer is unchanged.

Running the case at full size showed the example printed no fluid quantity at all -- conservation, positivity, closure agreement and timings, but nothing about the flow. It now reports kinetic energy start to finish and the closure's mu_t against the molecular viscosity, and asserts that kinetic energy decreases: a periodic box with no forcing can only lose it, so this is what would catch a closure wired in with the wrong sign, where negative turbulent viscosity adds energy. 128^3, 2000 steps, t = 1.19 (Re 1000, M 0.1): kinetic energy 31.006 -> 30.375, -2.04%, mass drift 1.3e-13, min rho 0.9976, nothing non-finite, closure within 2.2e-19 of a host evaluation. mu_t is 7.4e-5 mean and 3.1e-4 peak against a molecular 1e-3, so the closure moves the dissipation by something like 7-31% rather than rounding off under it. The conservation bound now grows with the step count (1e-12 + 1e-15*NSTEPS). The fixed 1e-12 was fine for the runs in CI but left only 2.7x margin over the host's 2.2e-14 at 5 steps once I had widened nothing, and would trip a long run for no reason; a non-conservative update is wrong by many orders more. Also a caveat on the speed table that I should have had before quoting it. The machine is shared. The same 64^3 case measured 2.0 ms/step on a quiet card and 63 ms/step with four other jobs saturating all four A100s -- and the giveaway was 64^3 and 128^3 reporting the same total time, which is per-launch stall, not compute. The published figures are from an idle GPU; the README now says so and says to check nvidia-smi and pin CUDA_VISIBLE_DEVICES.

… cost The example could only report the closure's share of a run that always included it. NO_CLOSURE=1 drops the network entirely -- nut stays zero, mu_eff = mu -- so the plain compressible Navier-Stokes solver is measurable on its own and the closure's cost is a difference rather than an attribution. A100 80GB, 100 steps, fp64, idle card, ms/step (Mcell-updates/s): plain with closure slowdown closure cost 64^3 1.565 (168) 1.833 (143) 1.17x 1.02 ns/cell-step 128^3 6.941 (302) 8.680 (242) 1.25x 0.83 ns/cell-step 256^3 54.50 (308) 68.00 (247) 1.25x 0.81 ns/cell-step The solver saturates at ~308 Mcell-updates/s, 3.25 ns per cell per timestep, where a step is SSP-RK3: three full RHS evaluations, each MUSCL + HLLC + the full Newtonian stress + conduction. The closure costs 25% of wall-clock, and its 0.81 ns/cell-step is three closure_infer calls at the 0.28 ns already measured -- the attribution and the difference agree. The same plain solver on this machine's host toolchain at 64^3 is 281 ms/step on one thread and 411 ms/step on all 128: slower with more threads. These loops are written for offload and gcc's host fallback oversubscribes them, the same pathology test_examples.py already documents for the surrogates. The README now says to treat the host path as a correctness fallback rather than a CPU baseline. The host test parametrizes over per_point / batched / plain, so the new #ifdefs cannot rot: nothing else compiles that variant. It also asserts the plain build reports no host comparison, since with nut fixed at zero the old "nut vs host: 0.000e+00" line read like a check that had passed. Both variants build warning-free under -Wall -Wextra.

…a pointer faulted on the MI210 use_device_addr takes the device address of the list item itself, and for a pointer variable that is not the pointee: closure_infer_batch got a garbage address and the kernel faulted (exit 139). The host build passed because the clause is inert there. use_device_ptr is what the other batched examples use. Claude-Session: https://claude.ai/code/session_01QoLws1dnkpAkCvnywpPuGc

The comments said every embedded weight is a `parameter` array. They are deliberately not: gfortran -fopenacc refuses a `declare` on a named constant ("not a variable"), so _emit_embedded_weights uses initialized `protected` module arrays instead, and only the chunk arrays a long literal list is split into stay `parameter`. Last of the PR's "small ones" that was a real defect rather than a missing feature.

Four places, per doc/opensource.md: SUPPORTED plus a _validate_softmax, a Softmax spec in plan.py resolving the axis to a flat [outer, axis_len] view, and a loop nest in each emitter written line-for-line parallel. The validator accepts only a last-axis Softmax, and that is about opsets rather than laziness. ONNX changed this operator at 13: before, `axis` coerced the input to 2-D and normalised every trailing axis together with a default of 1; from 13 it normalises along that one axis with a default of -1. The two readings coincide exactly when the axis is the last one, so requiring that makes the emitted loop right under either opset instead of silently picking a reading. An absent `axis` is accepted only at rank 2, where both defaults land on the last axis anyway. Two things in the loop are deliberate and now have tests rather than assertions in a comment: - The row maximum is subtracted before exponentiating. A logit of 800 overflows to inf otherwise, and inf/inf is nan; the shift cancels exactly in the ratio. test_softmax_does_not_overflow_on_a_large_logit pins it. - The maximum is written `v > mx`, which doc/opensource.md's NaN rule says not to do -- and this is the one op where it is correct. A NaN must LOSE the maximum: `mx` then stays a real number, exp(nan - mx) is nan, the row sum is nan, and the whole row comes out nan. A NaN-sticky maximum would give nan - nan everywhere and lose why. The test checks that the NaN's row is all nan AND that the other row still sums to 1. goldenFiles/softmax_head is a classifier head (4 -> 8 ReLU -> 5 Softmax), so the new loop nest also goes through the device tests, the ASan/UBSan job and the flang/ifx runs, not only the regression test. It pins opset 13 where the other golden models use 10, to exercise the semantics validate.py checks rather than relying on the two readings coinciding. Three rejection tests had used Softmax as their example of an unsupported op and failed for this coverage win -- the same trap test_cli.py's own docstring warns about one level down, where it tells you not to pin such a test to a golden model. They now use `Erf`, with a note to move them again rather than delete the check if Erf is ever implemented. 447 passed, coverage 97.98%.

…ODO item 3) At inference BatchNormalization is a per-channel affine map, y = scale * (x - mean) / sqrt(var + eps) + B and a Conv or Gemm already applies an affine map, so the two compose into one: s = scale/sqrt(var+eps) multiplies into the weight along the output-channel axis, and the shift pushes through the bias as (b - mean)*s + B. The operator disappears in fold.py before validate.py ever sees it, which is why neither emitter gains a loop nest -- a runtime BatchNormalization would read five extra arrays per channel to compute what is by then a constant. A producer with no bias gains one; the shift is not optional. transB decides which axis the scale broadcasts along: a Conv's [OC,IC,KH,KW] and a transB=1 Gemm's [OUT,IN] lead with the output channel, a transB=0 Gemm's [IN,OUT] trails with it. Both layouts are tested, because getting this wrong produces a model that runs and is quietly wrong. The pass declines rather than guesses, and each refusal is a case where folding would change the model's meaning: training_mode set, a graph that reads the running-stat outputs, non-constant scale/B/mean/var, wrong lengths, the intermediate value read by anything else or exported as a graph output, a Gemm with alpha/beta != 1 or transA. A declined BatchNormalization stays in the graph, where validate.py now refuses it with a message that says it is supported only by folding and what folding needs -- the generic "not supported, this generator handles [...]" would have been actively misleading, since the op IS supported in the shape that actually turns up. Verified on real PyTorch exports (nn.Conv2d+BatchNorm2d+ReLU and nn.Linear+BatchNorm1d, exported with do_constant_folding=False so the op survives): both match onnxruntime on both backends. test_fold.py checks the rewritten weights against the formula directly; test_regressions.py compiles a Conv->BN->Relu graph and compares to onnxruntime, since a sign error in the broadcast axis would pass the first and fail the second. float32 there: onnxruntime has no float64 Conv kernel, so a f64 model has no reference. 454 passed, coverage 97.62%.

Output channel oc belongs to group oc/c_out_per_group and reads only that group's c_in_per_group input channels, so the accumulation loop bound is the per-group count and the input channel gains a group offset. Spatial carries the two derived extents rather than `group` itself, which keeps the emitters interpolating numbers as doc/opensource.md asks. At group=1 c_in_per_group == c_in and every expression collapses to what it was, so an ordinary convolution emits byte-identical code. The group offset is hoisted into a named local computed once per output channel rather than spelled into the innermost index. It is invariant there, and gfortran's -Winteger-division had a fair complaint about `oc / 3` sitting inside the subscript -- a warning about the shape of the code, not a false positive. It is clearer and one division cheaper this way. validate.py now refuses the malformed group instead of every group: a group that divides neither channel count evenly (some channel would be in no group, and the loop would read across a boundary), a weight whose channel axis is not c_in/group, and group < 1. Three existing rejection assertions moved to those, since "grouped Conv is not supported" was the thing that changed. Verified against onnxruntime on both backends for depthwise (groups == channels), uneven groups (4->6 in 2 groups, so c_in_per_group 2 and c_out_per_group 3 differ) and square groups, plus real nn.Conv2d exports including a bias-free one. The uneven case is the one that matters: with the two per-group counts equal, an off-by-one in the offset is indistinguishable from a correct one, and with them different it lands inside the buffer and returns plausible numbers. goldenFiles/conv_grouped is that shape, so it also reaches the ASan/UBSan job, the device tests and the flang/ifx runs. 462 passed, coverage 97.64%.

A loop over the output that reads the input where the shifted index is in range and writes the pad value everywhere else. Only axes that are actually padded get a bounds test; on an unpadded axis the output index is the input index, so a test there would always pass and only obscure the nest. The interesting part is not the loop but the operands. ONNX moved `pads` and `constant_value` from attributes (opset 2) to inputs (opset 11), and the input form is a problem for a reason unrelated to Pad: `pads` is int64, and validate.py refuses any initializer that is not floating-point because everything surviving to the emitters gets laid out as a weight. So fold.absorb_pad_inputs normalises the operand form back into attributes before validate runs, exactly as auto_pad is resolved early -- plan.py then reads literal integers and the int64 array never reaches the layout pass. It also normalises the two opset spellings into one, which the emitters never see. An `axes` operand (opset 18) is deliberately NOT absorbed: it changes which axes `pads` counts, so leaving it in place makes validate refuse the node rather than mis-read it. Refused, each because the loop would otherwise compute something else: non-constant mode ('reflect'/'edge'), negative pads (a crop -- the nest only writes the output and reads inside the input, so a crop would silently be a no-op on that axis), a pads length that does not match twice the rank, and pads that disagree with the inferred output shape. A runtime `pads` is refused even earlier, by the frontend's dtype check on the int64 graph input; the test pins the message that actually fires rather than the one _validate_pad would have given. Verified against onnxruntime on both backends, exactly (a pad is a copy): asymmetric pads on two axes, so a begins/ends mix-up cannot pass, and a Pad->Conv chain, which is what an export emits for asymmetric padding a Conv cannot express itself. 467 passed, coverage 97.27%.

The op list and the limits both claimed less than the generator now does: grouped Conv was listed as unsupported, and Softmax, Pad and folded BatchNormalization were absent.

Refused at first out of caution, on the theory that the nest only writes the output and reads inside the input so a crop would be a silent no-op. That was wrong: the output index maps to ck - b, which for b < 0 reads FURTHER into the input and is in range by construction. The nest already computed it. The only thing actually broken was the spelling. `(c2 - -1)` is legal C and a syntax error in Fortran, so a crop emitted a module that would not compile -- which is why this surfaced as a build failure rather than as wrong numbers. A negative begin is now spelled as an addition, and the redundant `>= 0` test is dropped on cropped axes, where it cannot fail. The mixed case is what the test pins: cropping the front of an axis while padding its back, where the read can still run off the end, so the upper bounds test has to survive and the pad value has to appear there. Exact against onnxruntime on both backends.

…erator guide doc/opensource.md is now doc/adding-an-operator.md. The content was already current -- it is what the four new operators were written against, and every file and symbol it names checks out -- but the name said "open source" while the file says "Adding an operator". Two references updated. `axes` (opset 18) names which axes `pads` counts. Expanding it to a full-rank pads in fold.absorb_pad_inputs is the whole of its meaning, so nothing downstream learns the operand existed; a negative axis counts from the end. Absorbing it declines rather than guesses when the operand is not constant, when the lengths disagree, or when an axis repeats. edge and reflect are index maps, not bounds tests: every output element lands on a real input element, so neither writes the pad value and neither emits a condition. reflect is (IN-1) - |(IN-1) - |e||, spelled without abs() so the two emitters stay parallel -- a ternary in C, merge() in Fortran, where both arms are evaluated and that is free because it is integer index arithmetic. reflect is limited to one reflection, which validate enforces: the map covers e in [-(IN-1), 2*(IN-1)] and no further, so a pad at least as wide as its axis would need repeated reflection and would otherwise fold quietly to the wrong element rather than failing. The transformed index goes into a named local per padded axis. Inlined, the expression appears three times per axis, and four axes of it in one subscript ran past Fortran's 132-column limit -- a compile error, which is the good failure mode, but the local is also the clearer and cheaper code. Verified exactly against onnxruntime on both backends for edge, reflect and axes, with asymmetric pads on both spatial axes so an edge/reflect mix-up or an off-by-one in the mirror cannot cancel. The mode rejection test moved from 'reflect' to 'wrap', which is the mode that is still refused. 471 passed, coverage 97.23%.

A file-loaded model built with the cuda/hip backend could not serve a per-point OpenMP/OpenACC host; such a host had to link a second, omp-backend archive built by its own compiler. That is the first entry under Known limitations, and it is now gone. The cause was sharper than "the OpenMP pragmas are off". <name>.c held two regions needing two different compilers: the declare-target weights and the `target update` that puts them on the device, which only a host compiler with offload flags can act on; and the device memory for the native kernel (ROSENNA_MALLOC/MEMCPY_H2D and the <symbol>_dev pointers), which nothing but nvcc or hipcc can compile -- rosenna_rt.h is #error for anything else. The recipe had to pick one compiler for the file and picked nvcc, which silently disabled the declare-target on the weights that same file defines. So the split is now by role rather than by compiler. <name>.c keeps the host half and is ALWAYS built by the host compiler with its offload flags; <name>_kernel.cu takes every CUDA/HIP symbol, including the device copies and the upload that fills them, beside the kernel that reads them. Which half a translation unit provides is a recipe-defined -DROSENNA_NATIVE_KERNEL instead of a sniffed __CUDACC__. init still drives the device copies, through <name>_upload_device(). Verified on an A100 with the thing that was impossible: one libgemm_big.a, built with nvc -mp=gpu for the .c and nvcc for the kernel, linked by BOTH a per-point OpenMP-target host (under OMP_TARGET_OFFLOAD=MANDATORY, so no silent host fallback) and a native infer_batch driver. Identical to all 12 printed digits. One new rule, which the docs will carry: whoever's device code is in the archive does the link. Built with offload flags the archive holds nvc's OpenMP fatbin, which nvcc's device-link step cannot register, so the final link goes through the host compiler (nvc -cuda) -- which is what a solver uses anyway. Built without them there is no fatbin and nvcc links it directly; both were checked. Five structural tests asserted the old shape and now assert the new one, with their invariants intact: the transfer check follows <name>_upload to the kernel TU, the recipe check requires that nothing is handed to nvcc as CUDA source, and the C++-compile check builds the .c as host C with the role macro -- which is also what stopped it defining <name>_sync twice. 471 passed.

The suite has no hot test to fix. The slowest 15 come to ~100s of 264, and the other ~456 share the rest at about a third of a second each: it is hundreds of small compile-and-run tests, which is exactly the shape that parallelises. serial 264s -n 2 124s -n 4 90s (about what a 4-vCPU GitHub runner will see) -n 16 47s Coverage under -n is byte-identical at 97.22%, so this is a straight win rather than a trade. CI takes -n auto everywhere except the two compile-only jobs, which pass -s to show the compiler's own output and so cannot be captured by xdist; their test steps come to ~758s across the matrix today. Two things had to be fixed first, and both were real bugs rather than concessions to parallelism: - examples/run_basic.sh regenerates gemm_small.onnx IN THE REPOSITORY'S goldenFiles with an unseeded generator, so the test I added for it was rewriting a model the golden-suite tests read. It had already produced a dead (all-zero) gemm_small, which fails verify by design -- serially, not just in parallel. The script now takes GOLDEN_DIR and the test points it at a private copy. - Golden-model generation now takes an O_EXCL lock. Each xdist worker is its own process with its own session fixtures, so the session-scoped fixture was no protection on a cold tree: two workers would run the same generator onto the same path. Verified by deleting every generated .onnx and running -n 16 from cold.

…strict as CI CI caught a real defect in the one-archive change. The prototype for <name>_upload_device sat inside `#if defined(__CUDACC__) || defined(__HIPCC__)` with the device pointers, but its only caller is a HOST compilation of <name>.c -- which is the whole point of the change and does not take that branch. So the call was implicit: a warning under gcc, an error under the clang CI runs, which is why three jobs failed while the suite passed here. The declaration is now unguarded, with a comment saying why. A declaration nothing calls costs nothing; the guard was protecting against nothing. The local test compiled the .c without -Werror=implicit-function-declaration, so gcc's warning went unread. It passes the flag now, which makes this class of mistake fail locally instead of three minutes into CI. 471 passed, coverage 97.22%, 50s on 16 workers.

Replaces the limitation it removed: a per-point OpenMP/OpenACC host no longer needs a second omp-backend archive.

infer_one launched one kernel per op. Measured on an A100, batchnet's eleven ops -- output lengths 8 to 68, about 200 elements of work in total -- cost 60.07 us a call, 5.46 us per launch. That is almost entirely launch overhead. A run of consecutive ops that each fit in one block now share a kernel, with __syncthreads() between them. That is what makes it sound: the barrier is a full barrier over a single BLOCK, so the launch is one block and every element the next op reads has been written. An op larger than a block keeps its own kernel and its own grid, where it gets the parallelism it needs; squeezing it into one block to fuse it would trade 5 us of launch for far more compute. batchnet 11 launches -> 1 60.07 us -> 13.39 us per call mnist 10 launches -> 8 (the 256/10/10 tail fuses; the 6272s do not) Each op's loop is strided over the block rather than one element per thread. On the device with lengths at or under FUSE_THREADS that is one element each, exactly as before -- but it makes the kernel correct for ANY block size, including a block of one, which is how the stubbed host build in test_kernel.py emulates a launch. Without the stride that test computed only element 0 of each op and failed with a wrong answer, and the fix is better code rather than a concession: FUSE_THREADS is now a performance choice, not a correctness precondition. Verified numerically against the header-inline infer computed in a host-compiled translation unit (under nvcc, infer is the R9 device-only stub, so the reference cannot come from the same file): batchnet 5.55e-17, mnist 3.61e-16. `rosenna gpu-gate --backend cuda` PASSes on the A100. The gate also caught a real regression from the one-archive change, which is why it is fixed here: nvcc links PIE by default and nvc does not emit PIC by default, so once the host compiler owned <name>.c the archive failed to link with "relocation R_X86_64_32 against `.rodata' can not be used when making a PIE object". The recipe passes -fPIC, held in its own ROSENNA_PIC variable so that overriding CFLAGS cannot drop it. Objects in a static library get linked into executables the generator knows nothing about; before this, nvcc compiled the .c and the question never came up.

…aunch status Multi-device was not an optimization, it was broken. One set of <sym>_dev pointers meant a second init replaced the first device's addresses, and the first device's per-translation-unit __constant__ table then pointed at memory belonging to another device -- a kernel reading a valid-looking address on the wrong one. The pointers, the bind, the null checks and infer_one's event and stream state are all indexed by device now (the activation buffers are __device__ and so were already per device); the current device comes from a new ROSENNA_GET_DEVICE, and an index at or past ROSENNA_MAX_DEVICES is status 13 rather than a silent overrun. Verified on two A100s: init on device 0, upload_device on device 1, then infer_batch on both against a host reference -- 1.11e-16 on each. Before, the second device's upload would have left device 0 pointing into device 1's memory. The other two items are answered rather than implemented, both by measurement: - ASYNCHRONOUS UPLOAD is not worth doing. gemm_big's init costs 188.6 ms, of which the whole alloc-copy-bind is 0.286 ms -- 0.15%. The rest is CUDA context creation, which happens once whatever we do. And from pageable host memory cudaMemcpyAsync stages through a pinned buffer anyway, so "async" would be largely a fiction without registering the weight arrays. - A LAUNCH-STATUS CHECK FOR THE OMP FALLBACK cannot be written soundly, and the generated source now says why instead of leaving the asymmetry with status 11 looking like an oversight. OpenMP has no launch-status API; omp_target_is_present takes a HOST pointer while ruling R5 says x and y are device pointers, so it would answer about the wrong thing; and a deviceless run cannot be an error because --backend omp --host-fallback is a documented configuration of this same code. OpenACC, which I had verified nothing about: a per-point OpenACC host against the one-archive cuda build works on the A100 (1.11e-16, and ACC_NOTIFY confirms a real device kernel rather than a host fallback). Two findings, both now in python/README.md: - `#pragma acc parallel loop` over points lets nvc map ONE POINT PER GANG (num_gangs=4096, vector_length=32) and then auto-vectorises infer's own loops across the lanes -- the OpenACC shape of the teams-loop cliff in doc/nvhpc_teams_mapping. `gang vector` gives one point per thread and was 1.9x faster. Over a million points: 6.03 ns per point with OpenACC against 1.74 with OpenMP target, so the families are not equal in speed even when both are right. - `rosenna gpu-gate` cannot measure OpenACC at all: its per-point harnesses carry omp target pragmas, so `--flags "-acc=gpu"` compiles them away and times the host (743 ns per point, and the nsys check then has nothing to parse). I confirmed this is pre-existing by running the same gate at 922526a, before item 1 -- identical failures, 802 ns per point. 475 passed, coverage 97.26%; `gpu-gate --backend cuda` PASSes on the A100.

The badge states the FLOOR (>=97%), not a snapshot, because the floor is what CI actually enforces: pytest fails below it, so the badge cannot be wrong while the build is green. It needs no third-party service, no token and no write access to the repository -- the three things a live-percentage badge costs, since shields.io has nothing to read unless something publishes it. The exact number is not lost: the coverage job now writes it to $GITHUB_STEP_SUMMARY, so it is on the run's own page, one click from the badge, on every push. A test ties the badge to pyproject.toml's fail_under. A badge that drifts from the value it advertises is worse than no badge, and the floor does move -- it has gone 96 -> 97 during this branch -- so something has to fail when the two disagree.

CI caught a flake that is a genuine defect in the example. The script regenerates gemm_small, and it ran the generator RAW -- bypassing rosenna.golden, which exists precisely because these draws are unseeded and a tiny model is sometimes dead. gemm_small is 2 -> 2 -> ReLU -> 3 -> ReLU; when every pre-activation into the last ReLU comes out negative the model is all zeros, and step 5's verify refuses to compare against it, so the script exits non-zero. That is exactly the failure golden.py's docstring describes, on the one path that did not use golden.py. It now seeds through rosenna.golden's own GOLDEN_SEED rather than repeating the number, so there is still one source of truth. Three consecutive runs give a byte-identical model (same md5), live outputs, and both backends verifying. This is the second thing this script's regeneration has cost -- the first was writing into the repository's shared goldenFiles, fixed with GOLDEN_DIR. Both came from the same root: it regenerates rather than reusing, which is the point of the demo, so it has to regenerate the same way the suite does.

No new loop nest. On a flat row-major buffer a rank-3 (N,C,W) value and a rank-4 (N,C,1,W) one are THE SAME BYTES -- ((n*C+c)*1+0)*W+w is (n*C+c)*W+w -- and so are a weight's (OC,IC,KW) and (OC,IC,1,KW). So a 1-D op is the existing 2-D nest with h_in = h_out = kh = 1, sh = 1, ph = 0, dh = 1, which the compiler folds away. Only the attributes are lifted: kernel_shape, strides and dilations gain a leading 1, and pads (begin, end) becomes (0, begin, 0, end). The one thing that is not free is the weight's REGISTERED rank. emit_fortran declares a weight with the ONNX shape reversed and the nest writes four subscripts, w(kw, kh, ic, oc), so a rank-3 declaration would not match. The weight is registered at the lifted rank; the bytes are identical either way. validate now takes the arity from the input's rank rather than hardcoding 2, so a 1-D op is refused for a 2-entry strides or a 4-entry pads by the same checks that refuse the mirror-image mistakes in 2-D. Verified against onnxruntime on both backends: Conv1d with padding, stride AND dilation together, grouped and depthwise Conv1d (the group offset composes with the lift for nothing), MaxPool1d and AvgPool1d. goldenFiles/conv1d_stack is a conv/pool/conv chain, so the equivalence is also checked by the sanitizer, device and flang/ifx jobs rather than only by a unit test -- it is exactly the kind of claim that is either right or silently off by a stride.

test_spatial_shape_rejections still expected "only 2-D is supported" for a 3-axis kernel_shape. The arity now comes from the input's rank, so that case is refused for disagreeing with the input rather than for not being 2-D. I pushed the previous commit with this test failing: the suite and the commit were in one command, so the commit ran regardless of the result.

Three gates (z update, r reset, h new) and no cell state, so it carries only H. B is 6*hidden: the three W biases then the three R biases. The LSTM was the template and most of the plumbing came free -- multi-input/multi-output, positional `outs` with "" holes, an initial state that is either a graph value or a folded constant. The one thing that is not a smaller LSTM is the h gate. It cannot share the z/r loop, because it needs r, which is only known once that loop has finished; and `linear_before_reset` changes its ARITHMETIC, not its spelling: 0 (the ONNX default): h~ = g(Xt.Wh + (r . Ht-1).Rh + Rbh + Wbh) 1 (what PyTorch exports): h~ = g(Xt.Wh + r . (Ht-1.Rh + Rbh) + Wbh) The reset gate multiplies the state before the recurrent matmul in the first and the matmul's result in the second, and they agree only where r is 1. Both are emitted rather than one assumed, because choosing either would have been wrong for half the models people actually have -- PyTorch's exporter emits 1 and the ONNX default is 0. A test asserts the two readings really do disagree on the model used to test them, so the matrix is not checking one thing twice. The state update is a separate pass from the gates: the gates read the PREVIOUS state for every j, so updating it in place inside that loop would feed j's new value to the gates of every j after it. Verified against onnxruntime on both backends across all eight combinations of linear_before_reset x bias x initial_h. goldenFiles/gru_cell is a PyTorch export, so it also reaches the sanitizer, device and flang/ifx jobs; its state is an explicit forward() argument because, left to itself, PyTorch builds h0 from the input's batch with an Expand and exports a symbolic dimension, which this generator refuses by name. 506 passed, coverage 97.19%.

python/README.md still said "No batch norm, no Softmax, no Pad node, no GRU" -- all four of which this branch added. Flatly wrong rather than merely stale, which is the worst kind of documentation to leave in a merge. The four LSTM generators put "../test" on sys.path, a directory this branch DELETES. They kept working only because golden.py also puts goldenFiles/ there, so the line was dead and pointed a reader at somewhere that no longer exists. They now take the path from __file__, which is where nnLSTM.py actually is and does not depend on the working directory -- these run from a scratch cwd. Checked and clean: no other doc claim contradicts validate.SUPPORTED, no untracked files, nothing generated is tracked, and the golden list matches the 25 directories on disk exactly. The branch is a superset of master's content (the two commits it is "behind" are the merges of comp-physics#10 and comp-physics#11, whose content is already here), and the PR is MERGEABLE/CLEAN. 506 passed, coverage 97.19%.

sbryngelson and others added 25 commits September 13, 2026 20:45

feat: c library form with the per-point infer inline in the header

2e60a5d

fix: prefix every external weight symbol with the model name

2ca5720

feat: device decoration for openmp, openacc, cuda and hip hosts; embe…

5ae8d5f

…d weights by default

fix: explain the device macros and _dev pointers in the emitted header

c26df4f

test: library-form test asserts the storage-class-first infer signature

55d7cf4

test: filter driver-level toolchain notices in every device test, as …

933909d

…the C tests already do

sbryngelson changed the title ~~GPU library: device-callable per-point inference, native CUDA/HIP batched kernel, gpu-gate~~ GPU-callable inference, full operator coverage, and the retirement of fLibrary Sep 14, 2026

sbryngelson and others added 4 commits September 14, 2026 18:41

sbryngelson added 28 commits September 15, 2026 12:55

Document the four new ops in the READMEs

055e407

The op list and the limits both claimed less than the generator now does: grouped Conv was listed as unsupported, and Softmax, Pad and folded BatchNormalization were absent.

Document the one-archive linking rule in the README

657ac0a

Replaces the limitation it removed: a per-point OpenMP/OpenACC host no longer needs a second omp-backend archive.

sbryngelson merged commit fc4a195 into comp-physics:master Sep 16, 2026
9 checks passed

sbryngelson deleted the device-library branch September 16, 2026 02:27

Navigation Menu

GPU-callable inference, full operator coverage, and the retirement of fLibrary - #12

GPU-callable inference, full operator coverage, and the retirement of fLibrary#12
sbryngelson merged 84 commits into
comp-physics:masterfrom
sbryngelson:device-library

sbryngelson commented Sep 14, 2026 •

edited

Loading

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

sbryngelson commented Sep 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

The contract (unchanged)

Operator coverage

Measured on an A100

Measured on an MI210

No transfers across a time-step loop

Review fixes

Examples: surrogates inside PDE solvers

What the wider CI caught

One deliberate change to the generated arithmetic

fLibrary retired

OpenACC

Known limitations

TODO

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

sbryngelson commented Sep 14, 2026 •

edited

Loading