Add Beverin inference image and multi-role cluster example - #15
Draft
LorenzoPaleari wants to merge 125 commits into
Draft
Add Beverin inference image and multi-role cluster example#15LorenzoPaleari wants to merge 125 commits into
LorenzoPaleari wants to merge 125 commits into
Conversation
Skill bodies cost 1169 lines in EVERY prompt and 1081 of them -- 92% -- were four instrument manuals: perf, nsys, rocprof, opt-reports. A box has at most one GPU vendor, so most of that is a manual for hardware the reader does not have, carried on every task including the ones that never profile. So an instrument skill's BODY is now gated on `prompt.profiling_guidance`, and the `profile_first` strategy turns it on by itself -- needing a second knob to get the manual for the tools that strategy exists to use would be a trap nobody finds. Measured on a gemm task: 1373 lines to 284, a saving of 1089. The INDEX line is deliberately not gated. It is the only thing telling an agent the page exists, and gating it too would make the capability undiscoverable rather than merely absent. That asymmetry is what the first new test pins. INSTRUMENT_SKILLS lists both variants of each instrument. A `-judge` page is the same manual with one section swapped, so it costs the same tokens and gates for the same reason; omitting the five would have inlined ~1900 unconditional lines the day they ship. Also: the comment on GENERAL_SKILL claimed the other skills were "read on demand" and that the prompt "indexes the rest instead of inlining everything". The template has always inlined every body. The claim is true now.
Every pointer PARAMETER in every C/C++ reference now carries restrict: 784 of them across 245 .cpp files, plus the translator, which has always emitted it. The three hand-written kernels that were missing it are fixed. The gate is parse-based, and that is the whole point. `grep restrict <file>` is the obvious check and it is wrong twice: it passes a file that qualifies one parameter of six, and it FAILS all 173 *_reference.c files forever, because those have no buffer pointer parameters at all -- TSVC keeps its arrays as file-scope globals and PolyBench passes them through POLYBENCH_2D(...) declarator macros. Their complete pointer inventory is 141 `struct args_t *func_args`, 32 `char **argv` and 23 scalar out-params of an untimed init_array. An agent chasing that grep to green has exactly one move available: hoist the globals into restrict-qualified parameters. That deletes the benchmark. s242 (a[i] = a[i-1] + ...) and s1113 (a[i] = a[len/2-i] + b[i]) exist to test whether a compiler DETECTS the dependence, and a non-aliasing promise answers the question for it. The docstring says so, so the next reader does not rediscover it. Four supporting tests, because a parse-based gate fails silently: a regex that breaks on a multi-line signature reports zero offenders out of zero parameters and looks exactly like a clean tree. One hand-rolled scan here found 24 parameters where there are 784. So the scanner's own yield is asserted (> 500), the multi-line signature that broke it is pinned, an unqualified parameter is proven to fail, and the constructor member-initialiser that reads as a pointer (`nnr_(size_t(n1) * n2)` is multiplication) is proven not to. Separately: DaCe GPU variants are pinned to one stream. Concurrent streams overlap kernels, and every profiling question assumes they do not -- a per-kernel counter bracket needs a synchronised region to bracket, and an nsys timeline attributes a gap to the wrong launch when the next kernel is already running elsewhere.
Level 3 is whole networks -- ResNet, DenseNet, VGG, AlexNet, GoogLeNet, LSTM and GRU in four variants each, Mamba2, minGPT causal attention -- built out of the level 1 primitives already in the corpus. Each lands as the usual pair: a manifest and a buffer-out NumPy kernel. Two agents did the porting and both were killed by process death before reporting, so their LEARNED notes are lost; the ports themselves were complete on disk and are what is committed here. 8 of the 50 remain: four EfficientNet variants, RegNet, and three vision transformers. Unverified, and the reason is worth recording: neither porter had docs/canonical_numpy_form.md, which is the BINDING spec these files have to satisfy and has a CI gate behind it. They copied neighbouring ports instead, so they may satisfy CNF by imitation, but nothing has checked. An audit against the three invariants -- static shape at declaration, explicit indexing, declare-then-fill -- is the next step before these are trusted.
Each instrument gets one page, and each page a -judge twin that is the same text with only `## How it runs` swapped -- byte-identical elsewhere, enforced by a test that is proven to fire three ways (reword a shared line 250 lines from the seam; rename the swap heading; copy the execution section across so the twin never says who runs it). The pages were reviewed against upstream vendor documentation first: 64 errors found across six pages, every one re-verified before the fix, none rejected. That caught things like PAPI's `:stat=` defaulting to `avg` rather than `sum` -- so a bare `cuda:::dram__bytes_read` is bytes PER DRAM PARTITION, low by the instance count with nothing in the output saying so. Then they were FIELD TESTED: a fresh agent given one page, forbidden the siblings, and a task whose answer was known by construction. Three of four failed. Every failure was a page that is factually correct in every claim and still routes the reader wrong: * optimization-hints told the reader tolerance settles reduction reassociation. The harness verifies BITWISE. The tester's fastest variant, 3.13x, built by following the page, graded `correct: true, verified: false` -- score zero. * papi-cpu summed idle threads' barrier spin into the total: 21.9x and 4.5x over truth in two independent tests, both printing its own `armed N threads, counted N` healthy line. Its guard covered armed < used; the failure is the other direction. A guard that checks one direction of a two-directional error is worse than none. * nsys omitted `--force-overwrite=true` from the stats line, so every rerun silently returned the PREVIOUS run's CSVs and exited 0 -- breaking the only loop an agent runs. ncu's four metric tables became a reading -> action table, with every threshold read out of the shipped `sections/*.py` rules rather than remembered, and the same values confirmed in both installed versions. Its ordering discovery matters as much as nsys's total_ns rule: `Memory Throughput` is the MAXIMUM over its constituents, so 85% there with DRAM at 30% means L1 or L2 is saturated and DRAM work buys nothing. Fixtures under fixtures/ carry ground truth by construction. The CPU one documents a trap worth keeping: `a[i] = c[i] > 0 ? x : y` is not branch-bound at -O3 -march=native, because gcc if-converts it to vcmpgtpd plus masked ops -- measured misprediction 88x BELOW the threshold the phase existed to trip. A branch the compiler can flatten is not a branch. Not shipped into hpcagent_bench/skills/ yet: the five -judge pages describe an /instrument route the service does not have, and papi-gpu's numbers have never been observed because the driver profiling gate is on.
… .so one by experiment The design shipped with ten unanswered questions. Nine are now decided in the doc. The tenth was worth measuring rather than deciding, and the measurement changed the answer. CAN a judge count an arbitrary UNINSTRUMENTED .so it dlopens and calls? Yes -- so requiring agents to ship instrumented libraries is unnecessary. But by PAPI_attach, which binds counters to TIDs, not by pre-arming the OpenMP pool, which binds them to thread NUMBERS. Four-way agreement against perf stat as truth: 1.489e9 instructions truth, 1.476e9 attach, 1.486e9 register, 1.472e9 instrumented. Counting perturbs nothing measurable: 0.0562 s uncounted against 0.0560 s attached. The value is in how it BREAKS, because four of the five ways are silent: * raw pthread_create workers report 0.2% of truth -- 3.0M instructions for 1.53e9 executed, every PAPI return PAPI_OK. papi.py's `appeared` guard does not fire: the threads are created AND joined inside the call, so thread_ids() before and after are identical. * nested parallelism reports exactly 24.8% -- two armed outer threads times a quarter of each inner team. Entirely plausible as a magnitude. * a cross-runtime .so (judge libgomp, agent libomp) reports 13.3%. And two OpenMP runtimes in one process fight over affinity: the judge's OMP_PROC_BIND=close confined libomp's workers to one core and made the parallel kernel slower than serial. * OMP_WAIT_POLICY=active inflates cycles 4.01x by counting barrier spin. Outside-in is not wrong here -- it matches perf stat -- it is counting spin as kernel work. This is LLVM libomp's default, so it fires on real submissions. * register mode only: PAPI_stop from a non-owning thread returns PAPI_OK with k*2^47 garbage, and IPC comes out ~1.000 for those slots, so an IPC sanity check misses it. The check that catches all of them is sampling /proc/self/task DURING the call rather than before and after: unarmed_tids was 0 for every correct case and 6-17 for every wrong one, including the pthread case the before/after guard cannot see.
papi-gpu told the reader to arm the event set once and take a PAPI_read delta per
region, with two cudaDeviceSynchronize calls presented as the thing that made the
delta the kernel's. It does not work. The cuda component flushes the counter
ASYNCHRONOUSLY and a device synchronise does not flush it, so the delta between two
reads is whatever happened to be flushed in between -- which has no relationship to
what ran in between.
An EMPTY bracket -- two reads with nothing at all between them -- reported 374 MB of
DRAM reads. That is the whole bug in one line.
Measured against four kernels of known compulsory traffic, 25 regions each:
truth/rep start/stop read-delta
streams b,c into a 128 MiB 134.26 MB 128.4 MB
touches 64 KB x64 64 KB 77.9 KB 93.5 MB <- 1300x
reads a, 64 FMA, writes 64 MiB 67.08 MB 111.1 MB
reads a and c, divergent 128 MiB 134.27 MB 126.0 MB
PAPI_start/PAPI_stop per region lands on the compulsory traffic to within 0.1% on
every row. The read-delta is wrong on every row, and the true 2100x spread across
those four kernels arrives as 1.2x -- it does not add noise, it flattens the ranking
you are profiling to find. Same result on the SM side: a 7168x instruction-count
ratio, exactly matching 512 warps x 11 against 524288 x 77, reported as 1.66x.
The page argued against start/stop on the grounds that re-arming CUPTI per region is
instrumentation cost inside the region. The cost is real -- 2.37 s against 1.22 s over
20 regions -- but the page also says, correctly, that a counted run's wall clock
belongs to no comparison. It traded correctness for a property it tells you to ignore.
The syncs turn out to be redundant too: removing both changed the answer by 0.008%,
because PAPI_stop synchronises to collect. They were a line that looked load-bearing.
gpu_papi_init now arms and disarms around NOTHING at startup and REFUSES to run if
that reads back non-zero. It surfaces the permission gate early, and it is the one
self-test that catches a counter accumulating device-wide instead of attributing --
the failure mode that produces confident, plausible, wrong numbers on every region at
once with no error anywhere.
The GPU fixture had the same class of defect from the other direction: 6 MB of buffers
against 24 MB of L2, so k_stream was entirely cache-resident and never memory-bound.
dram__bytes_read correctly reported ~0 and read as a broken counter. Now 96 MB, 4x L2,
verified landing on each kernel's algorithmic minimum. Costs 0.44 s for 20 reps.
Also adds the AMD side, standalone and judge variants of each:
rocprofv3 the dispatch trace -- which kernel, which copy, which gap
rocprof-compute kernel-level analysis -- SOL, memory chart, which pipe
papi-gpu-amd the rocp_sdk component, same start/stop discipline
There is no AMD GPU on this box and all three say so at the top. What ports is the
METHOD, not the numbers: PAPI's own rocp_sdk README documents the same lagged-flush
behaviour and recommends adding delays between the kernel returning and PAPI_stop,
which is a race you cannot see losing. The empty-bracket self-test is what makes that
checkable on hardware the page could not be tested on.
Two AMD traps that return silent zeros are on the page for the same reason:
AQLPROFILE_READ_API=0 is required on ROCm >= 6.2.0, and PAPI_library_init must run
BEFORE any HIP call -- the opposite of the CUDA rule, where the component needs a live
context first. Same library, opposite order, each silent when wrong.
Seven jobs upload their coverage data. Every one of them names it `.coverage`, and the download step passed `merge-multiple: true`, which flattens all seven into a single directory -- so seven artifacts raced for one path. Six were discarded and the winner became the published project total. Two consecutive GREEN runs show it plainly: both logged `Found 7 artifact(s)` and then `Combined 1 file`, one reporting 59.96% and the next 13.44%. Same repo, same code; the swing is entirely which job happened to land last. The number in the summary has not been a total. Run 30809753679 lost the race harder: two extractions interleaved rather than cleanly overwriting, leaving a torn SQLite file, and combine died with Couldn't use data file '.../HPCAgent-Bench/.coverage': database disk image is malformed against the REPO ROOT path, which is why this reads as a destination problem and is not one. Coverage 7.15.3 combines via `ATTACH DATABASE` and reports the error on the main connection, so the message names the wrong file; the malformed file is `coverage-data/.coverage`. Every individual artifact passes `PRAGMA integrity_check`. No job produced a corrupt file -- the merge did. The coverage config was never the problem: `parallel = true`, `relative_files` and `concurrency = ["multiprocessing", "thread"]` are all already set, which is exactly why the per-job files are clean under xdist. Dropping `merge-multiple` gives each artifact its own subdirectory, so the seven files no longer share a name. `coverage combine` takes explicit paths of any basename and content-hash-dedups, so same-name-different-directory is fine. The second half matters more than the first: a partial combine prints a perfectly plausible percentage and stays green, which is how this survived every green run. The step now fails unless combine consumed every file it was handed, and a repo gate pins all three properties -- proven to fire by breaking each one in turn.
`--cache-control` defaults to `all`, which invalidates L1 and L2 before every replay pass so that pass 3 sees what pass 1 saw. The cost is that the kernel is measured COLD, which is not how it runs, and that is invisible until the working set fits in cache -- at which point it owns the headline number. Same kernel, same binary, 6 MB of buffers against this part's 24 MB of L2: --cache-control all (default) dram__bytes_read 4.20 MB DRAM Throughput 90.04% --cache-control none dram__bytes_read 2.05 MB DRAM Throughput 0.14% A 640x swing in the number that decides whether you are memory-bound, from a flag nobody sets. At a 96 MB working set the two agree (94.53% against 94.33%), because then the data genuinely does not fit and the flush changes nothing. This also reconciles ncu against an in-situ counter, which is how it was found: PAPI's cuda component does not touch the caches, so on that same kernel it reported near-zero DRAM traffic while ncu reported 90% of peak. Neither is broken -- they answer cold-start against steady-state, and which one you want depends on whether the kernel runs once or a thousand times. A timestep loop is the second case and is exactly where the default misleads. Both of the page's ordering claims were also tested with the gate open, since the metric half was written while it was shut. `Waves Per SM` < 1.0 killing the occupancy chapter is CONFIRMED and is the sharpest rule on the page: the launch-bound kernel reads 0.01 waves/SM with 16.71% warps active, so the occupancy chapter argues for tuning occupancy and the ordering rule correctly overrules it with "widen the grid". The Memory-Throughput-is-a-maximum claim could NOT be tested here -- after the fixture resize every kernel is DRAM-limited, so Memory Throughput and DRAM Throughput are equal on all four. The claim stays, unverified, and now says so. Also files the ablation, tagging and speed-up-plot backlog.
The restrict gate landed in 1389108 ahead of the one file that violates it: lavamd_reference.cpp declares 12 bare pointer parameters at :47 and :90, so test_reference_source_form.py has been failing on main ever since. The twelve __restrict__ qualifiers were already written and simply never committed. The rest is provenance, and it is the part worth reading. Four reference kernels carry OpenMP directives that do NOT match their upstream, and until now nothing said so -- which leaves the next reader unable to tell an adaptation from a transcription error. Each file now states which it is and why: - lavamd: upstream kernel_cpu.c:112-117 carries four private(...) clauses. The adapted directive drops ALL of them and the correct list is EMPTY -- upstream declares those eighteen variables at function scope, whereas this extraction declares each at its point of use inside the outer loop body, so C++ scoping already makes them private. A copied private() list would have been wrong here in a way that still compiles. - xsbench: a reimplementation of the unionized-grid lookup, so the directive was re-derived from the loop in this file rather than copied. - cp2k_density_matrix_trs4: upstream carries NO OpenMP directive at all -- every matrix operation is a library call -- so there was nothing to copy and the directives are adapted from scratch. - cp2k_grid_integrate: hand-written Fortran, not a transcription of the C. - velocity_tendencies: adapted, not copied verbatim. Upstream text is quoted verbatim in each header, tabs included, so a reader can check the adaptation rather than trust it.
Phase 2c runs `hpcagent_bench/benchmarks/`. Every file it measures sits inside the `[tool.coverage.run] omit` pattern, so instrumenting it produces no report data whatsoever. It has been pure cost, and the cost is large. `omit` stops LINE tracing, not the per-call dispatch. sys.settrace fires on every call event even for a file coverage will never record, and this phase is call-dominated: profiling the one cloudsc branch test shows 4.4M function calls in 23.7 s, so it pays that dispatch four million times to discard the result. Measured here, 8.28 s bare against >120 s instrumented. In CI the identical 745 tests went 183.57 s -> 736 s when coverage landed, which pushed the heaviest test past --timeout=600 and has made mpi Phase 2c red for three consecutive runs. Not a flake and not new -- it dates to the commit that turned coverage on. COVERAGE_CORE=sysmon looked like the fix and is not one. coverage 7.13.5 refuses sysmon whenever `branch = true` on Python < 3.14 -- "sys.monitoring can't measure branches in this version", since BRANCH_RIGHT/BRANCH_LEFT arrive in 3.14 -- and refuses it again for `concurrency=`. It then warns and silently falls back to the C tracer. Setting it would have changed the log and not the runtime; confirmed by reading core.py's selection logic and by env.PYBEHAVIOR.branch_right_left being False here. So the phase clears PYTEST_ADDOPTS. That is the whole fix for this job, and it loses nothing, because there was nothing to lose. Phase 5's numba/jax sweep is a different case and gets a different answer. It drives real library code, so its coverage IS signal and cannot be switched off. Its budget was set against a 200-port kernelbench subtrack and an uninstrumented run; the subtrack is now 239 and coverage is on. It took 26:01 when last green and reached 93% before the runner killed it at 35:00, so the budget goes to 55. That is a correction for work added on purpose -- the per-test --timeout=600 is still what catches a hang. Two gates pin it, both proven to fire by breaking them: Phase 2c stays uninstrumented, and the omit pattern that makes that safe stays in place. The second matters more than it looks -- narrowing the omit would quietly make this the one phase where a real library path goes unmeasured, and nothing else would notice.
… ceiling Two failures from 0619ec4, which ported 39 KernelBench level3 networks and grew the kernelbench subtrack 200 -> 239 without landing anything that depends on either number. The resolver globbed ("level1", "level2") only, so all 39 new ports fell into res.skips and three tests ERRORED at collection on the `resolved` fixture. level3 IS vendored -- third_party/KernelBench at 423217d9 holds level1 (100), level2 (100), level3 (50), level4 (20) -- and uses the same <index>_<Name>.py convention, so the existing UPSTREAM_INDEX regex and kernelbench_key fold apply unchanged. Checked before trusting that: grouping all three levels by key produces no new collisions, so adding level3 cannot perturb any level1/level2 resolution. The "no upstream model #1 for key" phrasing is the duplicate-name index, not a level-specific numbering. level4 stays out deliberately -- it holds HuggingFace model+batch+seq configs (16_gpt2_bs1_seq1023.py) and nothing was ported from it. That is now a named constant with the reason attached rather than a tuple literal. All 39 resolve 1:1 (alexnet -> level3/5_AlexNet.py, lenet5 -> level3/4_LeNet5.py, ...), 239 copies, 0 skips. The hardcoded 200 in test_kernelbench_references.py becomes a named PORT_COUNT = 239; the assertion still fails loudly if the subtrack grows again without the resolver learning where those sources live, which is the property it was protecting. Second: ml/shallow_wide_mlp resolved to 16.032 GB against a 16 GB XL ceiling. The interesting part is that shrinking the batch CANNOT fix it -- the three weight matrices alone are 32768*16384 + 32768*32768 + 16384*32768 = 2^31 elements = exactly 16.000 GiB, the ceiling to the byte. Even batch_size 1 leaves it 917,504 B over on the biases. A weight dimension had to move. output_size 16384 -> 8192 at XL only. The width is what "shallow wide" means and both hidden layers stay at 32768; the depth stays at upstream's [32768, 32768] so the provenance link the first half of this commit just established is not broken; batch and input keep upstream's get_inputs() shape. XL_BYTE_CEILING is untouched and no skip was added -- the ceiling is a real memory budget with other kernels tuned right up to it. XL footprint 16.032 -> 14.024 GiB. L already had output_size 8192, so L and XL are now equal there: flat rather than growing, which 76 other corpus kernels already do, and no ladder violation. The manifest carries the 2^31 arithmetic as a comment so the next reader does not re-derive it.
Four tests hardcoded the subtrack size as a bare 200. Porting 39 level3 networks moved the real number to 239 and every one of them broke, which is how five CI jobs came to fail with a NUMBER as their first decisive line rather than a defect: tests/test_e2e_numerical.py:67 UNGATED_COUNT = 200 tests/test_kernelbench_references.py expected 200 kernelbench ports, found 239 tests/test_levels.py:99 assert ... == 200 tests/test_kernelbench_translation.py assert len(kernelbench_stems()) == 200 They all guard different consequences of the corpus growing -- provenance resolution, the level selector, the translation ratchet, the numerical sweep's exclusion set -- so none of them is redundant. What was redundant was the literal. A ratchet needing an update in four places is a ratchet that will be wrong in at least one, so the number now lives in tests/corpus_counts.py and the four import it. UNGATED_COUNT deserves its own note, because raising an exclusion count is the one edit here that could be a weakening. It is not: the exclusion is defined by `spec.subtrack in UNGATED_SUBTRACKS`, and that predicate did not change. No kernel that was gated became ungated -- the SUBTRACK grew, and the count is a pin on the subtrack's size. Deriving it from the shared constant is what keeps that honest: the two can no longer disagree, so this stays a size pin and never becomes somewhere to park a failing kernel. The comment says so, since its predecessor explicitly asked for a reason.
The claim was that COVERAGE_CORE=sysmon silently falls back to the C tracer, argued from coverage's core.py selection logic. Measuring it: the same cloudsc test that runs in 8.28 s bare ran 1500 s under sysmon and was KILLED without finishing -- byte for byte the behaviour of the unset run. So >181x, and that is a floor rather than a figure. The earlier '>120 s' was the point at which the first attempt was interrupted, not a completion. Replacing it with the number that was actually observed.
INSTRUMENT_SKILLS gates the instrument bodies out of prompts that did not ask for them -- worth 1373 -> 284 lines when it landed, because four manuals were 1081 of 1169. It is a hand-written list of names that must stay in sync with a directory, which is the same shape as the four test files that each hardcoded the corpus size until the corpus grew. It drifts, silently, in the direction of paying for it. So the check is derived from what the gate actually protects: token cost. Any page long enough to be a manual (>= 100 body lines) must be classified. The threshold separates cleanly -- the strategy skills are all ~22 lines, the manuals 104-480, and nothing sits near the boundary. Drafts are checked too, since a draft graduates by one `mv` and the failure has to arrive BEFORE the page is in every prompt rather than after. Adding the six AMD pages was the reason to write it; the gate then immediately found three more nobody had classified -- optimization-hints (104), pytorch-to-numpy (138) and static-analysis (141), about 380 lines that would have entered every prompt on graduation. Being big is not the same as being an instrument, though, so the fix is not one list but a forced choice between two: - static-analysis is gated. It is a compile-time tool with the same shape as opt-reports: run it, read the report. - optimization-hints is ALWAYS inlined. It is not an instrument at all -- it is the ORDER of operations, and gating it behind a profiling knob would hide the sequencing from precisely the agent least likely to ask for it. Its worst measured failure routed a reader to a zero score and involved no tool. - pytorch-to-numpy is always inlined only because it is PARKED and is a porting skill; listed explicitly so the size gate cannot quietly absorb it into the instrument set while whether it ships at all is still open. A second gate refuses membership in both sets, since that means nobody decided and the behaviour would depend on which check ran first. Both proven to fire by breaking them.
Two kernelbench ports were refused outright:
NotImplementedError: slice step 'stride' must be a compile-time integer;
a symbolic step is read as 1 and the stride is lost
efficientnet_mb_conv (step `stride`) and resnet_basic_block (step `conv_stride`). A
parse sweep of all 246 benchmarks/ml kernels found exactly those two -- the initial
read of the CI log suggested 19 files and 70 slices, but most of those slices are
inside helpers the inliner absorbs, and they were never the problem.
Nor is the cause helper specialization, which is where the investigation started. In
both kernels the helper IS inlined, so the guard fires at the KERNEL BODY site
(frontend.py:534), not the surviving-helper sites. The step is not a call-site literal:
it is a kernel PARAMETER from the manifest's init.scalars (stride: 2, conv_stride: 1)
that is also an ABI argument.
_FoldStructuralUses already exists for exactly this class of name -- a runtime argument
keeps its name everywhere it can be evaluated at runtime, and folds only where nothing
else can be emitted. It covered a structural call's AXIS slot and not a slice's STEP
slot, which is the same kind of slot, refused by the sibling guard for the same reason.
So this adds visit_Slice, folding node.step ONLY. Bounds keep their name and still reach
the ABI: a bound is an ordinary integer expression, and the trip count comes from the
target's extent.
Placement is the delicate part. The fold runs after inlining (so a helper's ::stride has
already become the body's ::conv_stride), after _FoldConstantSymbols, and BEFORE
desugar_tuples and both structural guards -- which respects the ordering the existing
comments require, since the tuple fold must precede the guards and a literal can only
help it. The rebound-name rule was extracted to _rebound_names and is now shared by both
folds; _FoldConstantSymbols' behaviour is byte-identical.
The guard itself is untouched, message intact. Three of the seven new tests are
REFUSALS, pinning that it did not get weaker: a name absent from the manifest, a name
reachable as an EXTENT (which the harness may scale), and a name the body REBINDS.
The different-strides case is real and is proven on a real port: efficientnet_mb_conv
uses two conv helpers with different strides in one kernel and parses to steps {1, 2}
across 6 slices -- the 1x1 convs keep 1, the depthwise conv takes the manifest's 2. A
numerical test pins it too, since a collapse would read 1 3 5 7 instead of 1 4 7 10.
test_abi_corpus_agreement goes 3 failed / 2 passed -> 5 passed. That is the decisive
one: KNOWN_NON_LOWERING is {} and ratcheted in both directions, so all 578 kernels lower
and every emitted signature matches its binding.
The ports themselves are not wrong. canonical_numpy_form.md allows slices over declared
axes and says nothing against a step; an init.scalars knob used as a step is an ordinary
structural constant. The translator was the limitation.
Known and deliberately not fixed: on the surviving-helper path, _build_helper_kirs keeps
only the FIRST call site, so a second call to an array-returning non-inlinable helper is
left un-rewritten. Reproduced -- it is a COMPILE ERROR in all three native backends, not
a wrong answer, and no corpus kernel reaches it. Fixing it is one KernelIR per distinct
constant tuple, a feature-sized change to a function the whole corpus flows through, for
a construct with zero consumers.
…abled
DaCe's `compiler.command_cache` records the first build of a shape with `ninja -t
compdb` and replays those commands for later SDFGs, skipping CMake entirely. It has
been ON by default on spcl/dace@extended the whole time. It has also been doing nothing
in CI, because DaCe picks its generator with `shutil.which('ninja')` and only replays
when it picked Ninja (codegen/compiler.py:560, :644) -- and .github/actions/setup never
installed ninja.
So the config read True, CMake fell back to Make, and every SDFG paid a full configure.
Nothing reported it. A missing build cache is not an error, it is a slow build, which
reads as "CI is sluggish today" rather than as a defect -- the same shape as a guard
that only checks one direction.
ninja-build and ccache now install in the shared setup action rather than in whichever
job noticed first. ccache is there for the same silent-failure reason: DaCe knows
nothing about it, so it helps only via a compiler launcher or a PATH shim, and neither
can point at a package that is not installed.
On the framework side, pin_build_caching() joins pin_cpp_standard() and
pin_single_stream() at the top of optimize(). Same argument as the C++ standard: a
user's ~/.dace.conf must not be able to change what a graded baseline costs to build.
It pins build_mode=cmake (native writes per-object .o.cmd files, so there is no
compile_commands.json and therefore no command cache at all), configure_cache and
command_cache -- and warns when ninja is absent, since that is the one input the config
cannot describe.
ccache is offered through CMAKE_{C,CXX,CUDA}_COMPILER_LAUNCHER rather than by hoping
/usr/lib/ccache sorts first on PATH. CMake reads those from the environment, so it
covers the build DaCe is about to run without touching DaCe -- which is out of scope
here anyway.
Three gates, each proven to fire by breaking it: the pins survive a hostile conf (every
one is set to the WRONG value first, so a no-op function fails), ccache reaches CMake
without depending on PATH order, and CI installs both tools.
A submission's build is write-heavy and entirely disposable, so a memory filesystem is the right medium for it. It is the wrong medium everywhere the build shares memory with the kernel being timed, which is most places -- and that objection is not new here: harness/recording.py already REFUSES a results DB on a memory filesystem, because "on a compute node the DB would compete with the run for RAM". So the sandbox opts IN rather than defaulting on. HPCAGENT_BENCH_SANDBOX_DIR names a directory outright; otherwise /dev/shm is used only under CI. A workstation keeps the ordinary temp dir, which also keeps the standing rule that builds belong on disk rather than in the tmpfs /tmp already is. The second rule matters as much as the first: a tmpfs that runs out does not degrade, it fails the build with ENOSPC, and that failure is then attributed to the SUBMISSION rather than to the host. Below 512 MB free the sandbox declines and falls back. Checked per call rather than once at import -- free space is a property of the moment and several sandboxes can be live at once. Also files npbench PR#47 (0-init) to the backlog. It converts np.empty/np.empty_like to zeros across 256 sites, because a kernel that writes an `empty` buffer only PARTIALLY leaves stale memory in its output and is therefore read-nondeterministic. That lands harder here than upstream: this repo grades BITWISE, so a stale byte is not a tolerance question, it is a failed verification that reads as a flake. Filed to audit rather than port, since CNF's declare-then-fill invariant may already cover most of it.
The column's own docstring said it: `pluto` built `<base>_fp{64,32}.cpp` -- the same
sources as `llvm`, with the same clang++ -- and never invoked polycc. Meanwhile
polycc_report described a transformation that nothing compiled. Every pluto-vs-llvm
number in the results DB is llvm-vs-llvm.
polycc is source-to-source ONLY (the one compiler-adjacent call in the script is
clang-format, to indent its own output), so making this real is a BUILD PATH, not a flag
preset. hpcagent_bench/pluto_transform.py is now the single place the invocation is
spelled, and the report's args are defined as an EXTENSION of the build's
(POLYCC_REPORT_ARGS = POLYCC_ARGS + --debug) so the two are structurally incapable of
describing different transforms again.
Two things had to change that look cosmetic and are not, both measured:
(1) The column takes the C driver, not clang++. polycc's output is C and only C: rank>=2
arrays arrive as VLA parameters (`const double A[restrict NI][NK]`), and neither
variably-modified types nor the `restrict` KEYWORD exist in C++; polycc also prepends
its own `#define min(x,y)`, which detonates inside libstdc++. Verified on a real
kernel: gramschmidt's transformed output compiles clean under clang -std=c17 and does
not compile at all under clang++ -std=c++20.
(2) The OpenMP spelling had to change. clang ACCEPTS -fopenmp=libgomp and generates no
OpenMP for the pragma AT ALL. Measured on clang 21.1.8 with nm -u, one
`#pragma omp parallel for` loop:
-fopenmp=libgomp GOMP=0 kmpc=0 <- pragma dropped, loop is SERIAL
-fopenmp GOMP=0 kmpc=3
-fopenmp=libgomp -fopenmp GOMP=0 kmpc=0 <- the `=<lib>` form wins in EITHER
-fopenmp -fopenmp=libgomp GOMP=0 kmpc=0 order, so appending cannot rescue it
clang implements OpenMP only against its own libomp. Since `polycc --parallel` PUTS
the pragma in the source, the shared clang baseline would have timed a serial binary
under a parallel label -- the same lie this column was rebuilt to stop telling, one
layer down. CPU_BASELINE_CLANG_PLUTO is written as a substitution on
CPU_BASELINE_CLANG so the two cannot drift in any flag except that one, and
flags.pluto_capability gates the column on the object actually referencing a runtime.
The other clang columns keep libgomp deliberately: their sources carry no OpenMP
pragma (0 of 45 emitted *_fp64.cpp), so the spelling cannot change their codegen, and
test_fork_openmp_safety.py pins libgomp as the runtime whose fork() behaviour the
isolation layer is tested against.
Argument order is the third trap. polycc's signature is symbols, then arrays, then
scalars -- a VLA parameter's extents are themselves parameters and C requires them
declared first -- while every other native column uses the canonical ABI order. The
translator already writes that order as <base>_pluto_binding.json, so call_args reads it
rather than re-deriving it. When it is missing the column DECLINES: a positional ctypes
call cannot detect a permuted argument list, it just returns numbers.
Non-affine scops still decline through NotSupportedByFramework rather than falling back
to the untransformed source. polycc may silently MISCOMPILE a non-affine scop rather
than reject it, so "exited 0" is not evidence the transform was sound, and a silent
fallback here would be the original bug with a better hiding place.
End-to-end on gramschmidt: 3 omp pragmas emitted, compiles as C, and the object carries
3 OpenMP runtime symbols. spmm_csr declines (polycc rc=1).
Also files the skill-page review (all five pages DO NOT SHIP, 20 surviving findings) and
two backlog items: npbench 0-init, and unit tests proving the kernelbench NumPy ports
agree with the PyTorch models they were translated from.
…mula
All five failed an 11-agent review (find, then adversarially re-check). The AMD pages
failed hardest, and the root cause is worth recording because it is not "I was careless":
The ROCm counters page gives DEFINITIONS AND UNITS ONLY -- every formula on it reads
"NOT SHOWN". A web-fetch summariser returned a confident "Derived Metrics with
Formulas" markdown table anyway, and the page repeated it. Three of four derived
formulas were therefore invented by a model and shipped as vendor documentation.
The honest fence those pages carried ("no AMD GPU on this box, nothing here was
executed") did not help at all, because a reader cannot tell a fenced-but-correct claim
from a fenced-and-fabricated one. A fence is not a source.
So the formulas now come from ROCm's own counter_defs.yaml, fetched and PARSED rather
than summarised, and the page cites that file rather than the prose page:
VALUBusy 100*reduce(SQ_ACTIVE_INST_VALU,sum)/CU_NUM/reduce(GRBM_GUI_ACTIVE,max)
SALUBusy 100*reduce(SQ_INST_CYCLES_SALU,sum)/CU_NUM/reduce(GRBM_GUI_ACTIVE,max)
MemUnitStalled 100*TCP_TCP_TA_DATA_STALL_CYCLES_max/reduce(GRBM_GUI_ACTIVE,max)/SE_NUM
VALUUtilization 100*reduce(SQ_THREAD_CYCLES_VALU,sum)/(reduce(SQ_ACTIVE_INST_VALU,sum)*MAX_WAVE_SIZE)
LDSBankConflict 100*reduce(SQ_LDS_BANK_CONFLICT,sum)/reduce(GRBM_GUI_ACTIVE,max)/CU_NUM
L2CacheHit 100*reduce(TCC_HIT,sum)/(reduce(TCC_HIT,sum)+reduce(TCC_MISS,sum))
None of those denominators is SQ_BUSY_CU_CYCLES, which is what the page had used for
three of them. L2CacheHit was the one I had right. They are also ARCHITECTURE-SPECIFIC:
on gfx10 LDSBankConflict is SQC_LDS_BANK_CONFLICT/SQC_LDS_IDX_ACTIVE and L2CacheHit
counts GL2C_* instead of TCC_*, so the page now says to ask the tool by name.
Other corrections, each grounded in an upstream quote:
- `--pmc` twice does NOT give two passes. The option is nargs="*" with no append, so the
second occurrence silently DISCARDS the first. Multi-pass is a two-row input file. The
page had been teaching silent data loss.
- VALUUtilization (rocprofv3) and `VALU Utilization` (rocprof-compute) are near-identical
spellings for opposite quantities -- lane occupancy vs how busy the VALU was.
rocprof-compute's divergence metric is `VALU Active Threads`, in work-items.
- Units differ BETWEEN THE TOOLS: rocprof-compute's gfx942 L2 panel declares Read BW as
`unit: (Bytes + $normUnit)`, while FetchSize is documented in KILOBYTES. The page had
carried the KB rule, called it "the unit trap", and then applied it to the byte tool --
creating the 1024x error it warned about.
- Occupancy: 8 wavefront slots per SIMD and 32 per CU on CDNA, so 2048 work-items fill a
CU, not 256. The old text also asserted a factor that works out to 1.
- GPUBusy is a PERCENTAGE (100*GRBM_GUI_ACTIVE/GRBM_COUNT), so dividing by it inverts the
normalisation. The normaliser is GRBM_GUI_ACTIVE.
- Percentages are no longer accumulated across regions; only counts are.
- rocp_sdk enumerates under `rocp_sdk:::`, not `rocm:::` -- every shell example was wrong.
- AQLPROFILE_READ_API=0 is intercept-mode-only, not an unconditional export.
- Copies DO carry a byte volume (`bytes` in the buffer-tracing record); the page had said
you could not get it from the trace.
- v1 rocprof DOES print launch geometry; LDS column is LDS_Block_Size; the pmc CSV is
pid-prefixed; *_domain_stats.csv was missing from the reports table.
- Replay is not the only distortion: dispatches are serialized across HIP streams, and
replay breaks MPI outright (repeated MPI_Init/MPI_Finalize).
papi-gpu: three events its own steps 3 and 5 CONSUME were never armed -- all three
verified to resolve here first. "within 0.1% on every row" was arithmetically false; row
2 is +18.9% (77.9 KB against 64 KB), which is launch overhead at a scale where it stops
being negligible. A leftover claim that the syncs are the measurement is gone.
ncu: the gate section was stale -- this box now reads RmProfilingAdminOnly: 0, so the
page was routing readers off a working profiler onto cuobjdump, which it calls incapable
of costing anything. Deleted, along with the blanket UNVERIFIED fence. --cache-control
none now carries its precondition (valid only on a single-pass collection; --set basic is
8 passes). Step 3's prose contradicted the page's own action table and the table was the
one matching NVIDIA. The occupancy-gap row now sends you to re-read Waves Per SM first,
since inside 1 <= Waves Per SM < 5 NVIDIA attributes the gap to the tail.
Findings and their upstream quotes: docs/BACKLOG_skill_page_review_20260803.md.
Also backlogs porting the remaining 31 KernelBench kernels.
…t wrong
The three AMD pages were written against no hardware. This box turns out to HAVE an
AMD GPU -- a Radeon 780M (gfx1103, RDNA3 integrated) already exposing /dev/kfd, with
ROCm 7.2.4 and rocprofiler-sdk 1.1.0 in Ubuntu's own archive. So the rocprofv3 page has
now been executed rather than researched, against a real HIP fixture.
Three claims confirmed, and three corrected -- including two that were THEMSELVES
corrections made from documentation earlier today. Source-verified and hardware-verified
are different states, and this is the evidence:
CONFIRMED the LDS column is LDS_Block_Size, not Group_Segment_Size
CONFIRMED *_domain_stats.csv exists and the reports table had omitted it
CONFIRMED the output layout is FLAT on this version, not <hostname>/<pid>/
WRONG "copies carry a bytes field, ask for csv json". The CSV emitter has NO size
field at all: Kind, Direction, Stream_Id, Source_Agent_Id,
Destination_Agent_Id, Correlation_Id, Start_Timestamp, End_Timestamp. The
buffer-tracing record does define `bytes`, which is what the doc search
found, but it does not reach --output-format csv. The page now says check
your emitter before believing any page, including itself.
WRONG "no register count". The trace carries VGPR_Count, Accum_VGPR_Count and
SGPR_Count per dispatch -- the numbers that turn "occupancy is low" into a
cause. Both the draft AND the SHIPPED rocprof skill said otherwise.
MISSING kernel_stats also carries StdDev, which nothing surfaced.
The shipped skill's correction hit a gate that turned out to be right to resist: it pins
the skill to columns the repo's READER really parses, and the reader matches
Group_Segment_Size. So on current ROCm the reader finds no LDS column at all, silently,
because a missing optional column reads as null. That is a live product bug, not a doc
error -- the page now names both spellings and says which one the code matches, and the
reader/fixture fix is backlogged rather than smuggled in here.
One operational finding worth as much as any column: rocprofv3 REQUIRES
hsa-amd-aqlprofile and does not depend on it. Without the package the run dies with
./gpu_phases: error while loading shared libraries: libhsa-amd-aqlprofile64.so.1
prefixed with the CHILD's name, not the profiler's, because the library is injected into
the profiled process. The binary links and runs fine standalone, so this reads as a bug
in your own program and is not one.
Ubuntu also ships two ROCm stacks that do not interoperate: /usr/bin/hipcc links against
/usr/lib/rocm/llvm and fails with undefined symbol __hipUnregisterFatBinary, while
/opt/rocm/bin/amdclang++ has no device bitcode. The build that works crosses them --
Ubuntu's clang for device code, /opt/rocm for the runtime.
Not verified here and still marked so: anything CDNA-specific. An iGPU has no HBM and no
Infinity Fabric, and the MI300 counter expressions are a different architecture's.
NVIDIA needs no sample: install the CUDA Toolkit and nsys/ncu/CUPTI are there. AMD is six separate traps, every one of which cost time today and none of which is obvious: Ubuntu already packages ROCm (so amdgpu-install's version-pinned URL is a dead end), the tools land in /opt/rocm/bin off PATH, rocprofv3 needs hsa-amd-aqlprofile and does not depend on it, rocprof-compute pins astunparse==1.6.2 against a PEP-668 system python, two ROCm toolchains coexist without interoperating, and an unsupported gfx target needs HSA_OVERRIDE_GFX_VERSION. Written down so the sample is a script plus a preflight check rather than prose.
The pmc-twice test could not confirm the pass-splitting claim on hardware:
rocprofiler-sdk 1.1.0 reports no counter metrics for the agent at all.
rocprofiler_iterate_agent_supported_counters failed for agent 1 (gfx1103)
:: Agent HW architecture is not supported, no counter metrics found.
terminate called after throwing an instance of 'std::out_of_range'
what(): unordered_map::at
[rocprofv3_error_signal_handler] rocprofv3 caught signal 6
The unsupported-agent line is only a WARNING, so the run continues and then
aborts on the empty counter map, and hangs a further 10s+ in queue sync. The
observable is SIGABRT plus a hang in a program that runs clean unprofiled --
same trap as the missing aqlprofile library, and it reads as the user's kernel
faulting.
Page now says to run --pmc under a timeout, and that trace support on a part
implies nothing about counter support on it. The fence marks every --pmc
SEMANTIC claim (pass splitting, budget, cross-pass ratios) as source-read
rather than run, since they could not be exercised here.
Checked every flag the AMD pages name against rocprofv3 --help on ROCm 7.2.4. All 7 resolve; no page invents an option. Two other things did not survive. The counter budget was described backwards. The page said an over-long --pmc list makes the kernel replay until every counter is collected. That is ncu's behaviour and rocprof v1's; rocprofv3 says the opposite in its own help -- "job will fail if entire set of counters cannot be collected in single pass". So the cost is the run, not the wall-clock, and splitting passes is the user's job rather than the tool's. Page now says so and points at the input file. rocprof-compute is installed by the distro ROCm packages and still cannot run: it pins astunparse==1.6.2 against an installed 1.6.3 and misses 11 packages. Every subcommand including --help prints the dependency errors and exits 0, so a wrapper checking the return code concludes success and finds no output. The pin is exact and the installed version is newer, so upgrading cannot fix it -- the page now says to build a venv from the shipped requirements.txt, and to confirm --help prints usage before assuming the tool is present. The --pmc repeat-overwrite claim was re-checked against the argparse definition and holds: nargs="*", no append action, so the second occurrence replaces the first.
The framework looked for `<base>_pluto_binding.json`. The emitter has only ever written `<base>_fpNN_pluto_binding.json`, one per precision, so the file was never found and every kernel declined with "no binding" -- with the binding sitting in the same directory. CI reported this as a pluto column that ran nothing; locally, where polycc exists, the same kernels declined silently. Only the ORDER now comes from that file. Shape, dtype and which arguments are output pointers come from the manifest-derived binding every other native column allocates against. That is not tidiness: the pluto binding is emitted per precision and call_args has no way to say which precision is running, so reading a dtype out of it would be reading fp64's declaration during an fp32 run half the time. `_ArgView` is gone with it -- it existed to adapt those per-precision dicts, and defaulted shape to `()` and dtype to None, which allocates a 0-d float64 scalar and hands the kernel a pointer to 8 bytes for a buffer it indexes. `run_polycc` now returns the argv it ran alongside the result. The transformation report echoed a command rebuilt from a second `shutil.which`, so it could print something that was never executed. It also deletes its own partial output on failure: polycc writes as it goes, and a truncated translation unit whose mtime is newer than the scop's is exactly the "fresh enough, reuse it" condition transformed_sources tests -- the next build compiled half a kernel and timed it. `_ensure_built` reuses a cached .so only while it is newer than every source that composes it. An existence check alone made the artifact unfalsifiable: the name says which framework built it and nothing about which sources it compiled, so a tree holding a lib<short>_pluto.so from before this column compiled polycc's output would be returned, timed, and recorded as a Pluto number while being a clang one. test_wrap_kernel_matches_numpy drops pluto and gains a replacement at the right layer. That test calls the built symbol positionally in canonical ABI order; pluto's scop declares its size symbols FIRST, because a VLA parameter's extents must be declared before use in C: C SIG : void tsvc_2_s212_fp64(double *a, double *b, const double *c, const double *d, int64_t LEN_1D) PLUTO SIG: void tsvc_2_s212_fp64(int64_t LEN_1D, double *a, double *b, const double *c, const double *d) Measured: passing the first order to the second segfaults immediately. CI never saw it because polycc is absent there. The new test pins that the framework finds the binding and that polycc's order differs from canonical -- the two halves that were broken.
…ways had
The AMD kernel-trace reader matched `Group_Segment_Size`. rocprofiler-sdk 1.1.0 emits
`LDS_Block_Size`. The symptom was not an absent field, which is what the backlog entry
claimed and what would have been tolerable: `column()` returns "" for an unmatched prefix
and `number("")` is 0.0, so a 16 KB workgroup came back as `shared_memory: 0.0,
shared_memory_unit: "B"`. That is a measurement. It says the LDS budget is free, and an
agent sizes a tile against a budget it has already spent.
Both spellings are matched now, and a trace carrying neither reports null -- which is what
the module's own "Absent is not zero" doctrine promised and the code did not deliver.
Fixtures carry both generations plus a no-LDS trace, so neither direction can regress.
`registers_per_thread` was documented as unavailable on this vendor, in the payload note
shipped with every AMD profile, in two docstrings and in the agent-facing skill page. The
trace carries `VGPR_Count`; it is now the row's register count. `SGPR_Count` stays out: the
scalar file is per wavefront and has no NVIDIA counterpart, so it has no field in a schema
whose whole point is being vendor-independent, and averaging it into one that means
something else would be worse than omitting it.
The shipped skill page contradicted itself on this in three places, saying the trace does
carry VGPR/SGPR in one section and "not in the trace at all" two sections down. One prompt
fragment, both answers. Also corrects the kernel-stats schema there (it was missing `Name`
and `StdDev` and in the wrong order) and restores the LDS unit and the allocation-granule
caveat, which a page that teaches LDS occupancy arithmetic cannot do without.
Adds the aqlprofile trap to the gates table. `rocprofv3` requires `hsa-amd-aqlprofile` and
does not depend on it; missing, the traced run dies with `libhsa-amd-aqlprofile64.so.1`
prefixed with the CHILD's name, so the one refusal that misattributes itself to the
submission had no named cause and no fix.
HSA_OVERRIDE_GFX_VERSION is the standard escape hatch for an unsupported ROCm target, and the page presented "gfx1103 has no counters" as a hardware fact without ever applying it -- so the headline measured claim could have been an unset environment variable. It is not. Measured: with HSA_OVERRIDE_GFX_VERSION=11.0.0 exported and the kernel built --offload-arch=gfx1100, the application runs clean and `--pmc` produces the same std::out_of_range abort, with the warning still naming gfx1103. The override is a ROCr/HIP lie about the ISA; rocprofiler reads the real hardware ID when it enumerates counters. Also measured on that run: `timeout`'s SIGTERM at the deadline is caught by rocprofv3's own signal handler, logged as "caught signal 15", and the process keeps running. It needs SIGKILL. The page said to run `--pmc` under a `timeout`; it now says `timeout -k`, because the bare form does not terminate it. The rocprof-compute banner asserted "No command below was executed and no number below was observed" six lines above a paragraph reporting an executed --help and an observed exit code. On a page whose entire value is letting a reader tell measured from documented, a banner the next paragraph refutes trains them to discount both. Scoped to what is actually true: no profile was collected. Its venv remedy also dropped --system-site-packages, which is the form that works -- rocprof-compute imports the distro's ROCm Python modules, so an isolated venv satisfies every pinned pip requirement and then fails on those instead. The copies section deleted the only lead to byte volume while telling the reader to compute achieved bandwidth. The CSV finding is right and stays; the emitters that may still carry `bytes` are named again rather than gestured at. Backlog: the install sample omitted hsa-amd-aqlprofile from the apt line that calls itself "the whole thing", two bullets above declaring it REQUIRED -- scripted as written it installs the broken configuration. The clang bitcode path is derived from --print-resource-dir instead of hardcoding major version 20. The "verbatim" measured header is one physical line again, since a wrapped copy pasted into a fixture makes csv.DictReader read lines 2 and 3 as data. Item 11 records what landed and keeps StdDev as the open half.
The setup action exported DACE_compiler_build_mode=native for speed. A DACE_* environment variable outranks Config.set, so dace_framework.BUILD_CACHE_PINS asked for `cmake` on every job and got `native` -- which skips CMake, writes per-object .o.cmd files, produces no compile_commands.json, and therefore makes compiler.command_cache inert while it still reports True. Exactly the "config that reads enabled and does nothing" failure the pin's own docstring warns about for ninja. Two mutually exclusive optimizations, each claiming to be the fast one, and the env var was winning silently. The pin is the framework's stated design and lives in library code with a test; the export is a CI-only override whose comment claims the opposite. The export goes. tests/test_dace_flavors.py::test_the_build_cache_pins_are_applied_and_survive_a_hostile_conf was the visible symptom (assert 'native' == 'cmake') and now fails if this inverts again.
…does not share The test demanded `libgomp` in PLUTO_PAR on Linux. PLUTO_PAR was changed to a bare -fopenmp on purpose: measured, `clang -fopenmp=libgomp` accepts the flag, parses the pragma and emits no OpenMP call at all, and pluto is the ONE clang column whose sources carry `#pragma omp parallel for` -- so the runtime pin that is inert everywhere else silently serialises exactly this column under a parallel label. The other clang columns keep libgomp and are still asserted. The pluto leg is now pinned to the spelling that emits OpenMP, in both PLUTO_PAR and the baseline it substitutes into.
…code `coverage combine ... | tee combine.log` discarded combine's status, since a run: block gets `bash -e` with no pipefail. A torn or malformed coverage database exited 0, and the only check that then fired was the file-count one -- which reports a partial merge and sends the reader after the wrong cause entirely, on the exact defect whose diagnosis this workflow already spends four paragraphs untangling. set -o pipefail. test_ci_installs_the_tools_that_fail_silently_when_absent searched the whole action.yml for "ccache", and the same file carries a comment block explaining why ccache is there. Drop the package and the test stays green on the comment. It now reads the apt-get install lines (joining the backslash continuation the package list wraps with). test_ccache_is_offered_to_cmake saved two launcher variables and pin_build_caching sets three, leaking CMAKE_CUDA_COMPILER_LAUNCHER into every later test in the worker -- silently routing a later nvcc through ccache. Test-order dependence, and the pollution direction is toward passing.
Three translator items, and the first of them was mis-diagnosed in the backlog.
A kernel names its own dimensions off a parameter -- `batch, channels, h, w =
x.shape`, folded by tuple-desugar to `batch = batch_size` -- while `init.shapes`
keeps the symbol. String `==` on those tokens then declines a contraction whose
extents match, because `batch` and `batch_size` are the same extent spelled two
ways. `dims_agree` walks a cheap ladder instead: literal equality, then
alias-substituted equality, then sympy (imported lazily, kept off the common
path). It fails CLOSED -- an unresolvable comparison answers False, because a
wrong True contracts over two different extents and a wrong False only declines
a matmul.
The backlog claimed one cause. Instrumenting every decline found two: the second
was `relu_self_attention`'s Call operand, which never reached
`_matmul_result_shape` at all, so no amount of alias resolution would have
helped. `_materialise_call_operands` spills it first, skipping rank-1 so the
scalar dot path is untouched. The corpus ratchet `KNOWN_NON_LOWERING` is now
EMPTY (was 5), and all five kernels verify numerically.
A full float sum/mean now accumulates its innermost axis in blocks of 128 --
numpy's own pairwise cutoff -- rather than one serial chain. Measured through
the op oracle, n = 2**22, emitted float32, gcc -O2 (which does not reassociate
on its own), against numpy's pairwise sum:
one accumulator |d| = 1.09e+02 (5.2e-05 relative)
blocked, 128 |d| = 4.00e+00 (1.9e-06 relative)
Deliberately scoped: the FULL reduction only, and only its innermost axis, which
is the one long dependence chain. Outer axes keep plain loops so an emitted nest
still looks like a nest to the parallelism and isopar recognisers. Integer sums
are untouched -- integer addition is exact and associative, so blocking it is
pure code growth.
The test passes `dtypes=` explicitly. Without it `run_op` emits float64, whose
naive error is ~1e-11, and the test goes green whatever the accumulation does;
the first version of this test had no teeth and looked fine.
Finally, `_assert_ok` accepted a run in which every backend reported
"unsupported", so a case could go green having graded nothing. Skips stay
accepted -- a backend that cannot lower an op should not fail the op's test --
but at least one of c/cpp/fortran must now have reported ok, and the failure
names what each backend actually did.
GitHub throttles anonymous CI egress with 403, not 429, so a container build that happened to be rate-limited read as "the repository is gone" while the repository was public and fine. Two changes. The ref is pinned to a commit rather than tracking a branch, which needs `git init` + `fetch <sha>` + `checkout FETCH_HEAD`, since `--branch` takes a branch or a tag and never a SHA. And the fetch retries four times with doubling backoff, naming the dependency and the `git ls-remote` diagnostic when it finally gives up, instead of failing on the first 403. This raises the odds, it does not remove the dependency: only vendoring would. `HPTT_REF`, `HPTT_REPO` and `HPTT_CLONE_TRIES` override each part.
…ing it libstdc++ picks its parallel-algorithm backend per translation unit from `__has_include(<tbb/tbb.h>)`. A runner without the TBB headers keeps compiling, keeps linking, keeps returning right answers, and quietly runs sequentially -- nothing in the flags, the exit code or the output says so. The repo already engineered against exactly this for Polly; isopar had no equivalent. `languages.isopar_capability()` compiles one `std::execution::par_unseq` call at the flags the harness really builds C++ with and reads `nm` for a TBB runtime call, returning the same three-way AutoparVerdict every other column uses. Measured here: ok, runtime_calls=12. It lives in `languages` rather than beside `flags.polly_capability` because only that module can name the cpp block's compiler and its `-std=`, and `stdpar_link_flags` -- which must agree with it -- is directly above. `flags.probe_autopar` grew `runtime_pattern` and `suffix` rather than forking into a second probe: the evidence is the same (compile, then nm) and only what counts as a runtime call differs. Worth stating plainly, because the backlog entry implied otherwise: cpp_isopar is NOT a scored column today. It has no FRAMEWORK_LANG entry, and its only consumers are correctness oracles, where a serial backend is slow rather than wrong. So the exposure today is zero and this keeps it zero -- the probe is registered in preflight.AUTOPAR_PROBES, so the column cannot be added ungated. The negative case is real rather than monkeypatched. `-D_GLIBCXX_USE_TBB_PAR_BACKEND=0` is what a runner without libtbb-dev compiles, since libstdc++ defines that macro AS the `__has_include`. Same source, same exit code, right answers -- 12 TBB references down to 0, object 22088 B down to 1256 B. Both halves are asserted in ONE test so it cannot pass by measuring nothing, and a second test pins that the header question (which decides -ltbb) and the nm evidence never disagree. Also: `test_gate_is_a_no_op_for_ungated_frameworks` still listed pluto after it joined AUTOPAR_GATED, so it asserted the opposite of what the tree does and passed or failed according to whether the host's clang honours an OpenMP pragma.
…n NaN The GPU skill pages tell an agent that a float-atomic reduction will not pass scoring, and a test pins that the pages SAY so. Nothing had ever made it true: the gpu CI job is disabled for want of a runner, so no float-atomic submission had ever met `_determinism_check` on this repo. It needs no device. The gate is host Python over two output dicts, and "two runs of a float-atomic reduction" is fully characterised by what those dicts hold -- values equal to the last ulp, unequal in their bits. Built with `nextafter` rather than a re-summation in another order, because numpy's pairwise sum may reassociate to identical bits and the fixture would then silently test nothing. The tests are paired so none can pass vacuously: the ulp-apart pair is rejected under bitwise=True and ACCEPTED under bitwise=False, which is what shows the rejection is the bitwise leg's work rather than a fixture whose numbers are simply far apart. A reproducing run is accepted; a run that reproduces a WRONG answer is still rejected; and the default is pinned off inspect.signature rather than off a call site's line number. Pinning it surfaced a real defect. `np.array_equal` reports NaN != NaN, so a kernel whose output legitimately holds NaN -- a masked cell, a log of zero -- failed the gate as NONDETERMINISTIC even though the second run was a bit-for-bit copy of the first, with nothing the agent could do about it. Reproducibility and validity are different questions: whether that NaN belongs there is the oracle leg's, and `compare_arrays` was already NaN/+-Inf-aware, so the two legs had disagreed on what NaN means. Now `equal_nan=True`, with a test pinning that this did not become "NaN matches anything".
ubuntu-latest ships LLVM 18, whose Fortran driver is still spelled `flang-new`,
while the harness is developed and measured against 21 -- flags.POLLY_PAR's
Polly findings and the clang OpenMP-spelling measurements behind flags.PLUTO_PAR
are both 21 numbers. A CI-only major nobody develops against is exactly how a
count comes out 0 on a dev box and nonzero on a runner with nothing in the log
to say the toolchains differed.
CI now installs clang-21, flang-21 and libomp-21-dev from apt.llvm.org, with the
key fetch retried. The unversioned clang/clang++/flang are symlinked into
/usr/local/bin, and that is not cosmetic: `resolve_compiler`'s
highest-`<name>-<major>` fallback does not help here, because
`flags.polly_capability` probes bare `shutil.which("clang")` and compilers.yaml
names the drivers unversioned, so a distro clang left on the box would keep
winning and the LLVM columns would silently stay on 18. toolset.yaml's discovery
list gains clang-22/21 so prefer-latest stays true.
`verify_toolchain` now checks clang++ as well as clang. They ship in one apt
package, but a driver is required here because a TEST requires it:
test_warnings_ratchet skips its ENTIRE -Wall -Wextra count when any one of
gcc/g++/clang/clang++/gfortran is missing, so verifying only clang would let the
C++ half of that ratchet go silently unmeasured -- the exact failure this script
exists to make loud.
Six entries close and one is corrected in place. C claimed a single cause and had two -- the second never reached the shape comparison at all. E's headline no longer holds: run 31001878513's integration job passed the ratchet at count 0, and not vacuously, since the test asserts 20+ real builds. D's entry implied a scored cpp_isopar column to protect; there is none, so the work is to keep the exposure at zero rather than to close one. I closes only its determinism half -- the null-workspace protocol is still unexercised, and "the GPU pages are green" is still not evidence the GPU track works.
…n real hardware Two tracks, both measured on this box rather than read off a docstring. The generated *_dace.py are AUTOGENERATED from the numpy reference by numpyto_c.dace_emit, so a frontend gap is not something a kernel author routes around -- it lands in the generator's output. Parsing 336 of them (to_sdfg, never a C++ build) gives 311 clean and 25 refused across 18 classes, and 8 of those 25 are INTERNAL CRASHES rather than clean refusals: KeyError: npwx, KeyError: pyobject, a bodyless NotImplementedError, _numpy_empty() missing a required argument. An unsupported construct should name itself; those are bugs in the refusal, separate from whether the feature is ever supported. The sweep STALLED at 336/576 on vexx_k_dace.py -- stuck, not slow. A kernel that hangs the frontend parse is its own finding, and it means 25 is a floor rather than a total. Root-caused the two known gaps rather than describing them. The bitwise NameError is ONE bare eval with no globals dict (replacements/utils.py:236) silently borrowing that module's namespace, which contains no symbolic head name at all. Shifts survive the identical path only because left_shift.eval FOLDS to a literal when both operands are numbers, while bitwise_and and friends are bodyless and do not -- so the string stays bitwise_and(1, 3) and the eval raises. The __-prefixed variants are half-wired: produced by the parser, absent from symbolic.py:3299's binop map, so codegen emits __left_shift(N, 1) into C++. Patched to the BARE classes, shifts work end to end including inside a map range. And nqueens shows the reversed form (1 << N) needs covering too. For comprehensions there is no desugarer at all -- four sites mention them and none rewrites anything, one of them commented out. They vanish today only as a side effect of constant folding, so only a compile-time-evaluable comprehension survives. The _DISALLOWED_STMTS entry stays: it is the assertion that preprocessing already handled the node. Both GPU lanes pass here (RTX 4050, cc 8.9, nvcc 13.1), including agentbench-gpu.yml, which is workflow_dispatch + self-hosted and had therefore never executed anywhere: 48 passed for agent_bench, 7 for frameworks. Zero skips, which is what makes it evidence -- _has_gpu() gates the tvm/triton cases. This does NOT close the GPU item: the null-workspace protocol is still unexercised, and the leg that would catch an all-zero result is the oracle, not determinism, since all-zeros reproduces perfectly. Also root-caused: cegterg's 17 skips are one missing apt package (libfftw3-dev; blas/lapack are fine), and the numerical oracle still calls polycc directly at numerical_oracle.py:1109 instead of through run_polycc, so it inherits the ambient environment and misses the aarch64 pet shim -- the Pluto merge did not fix that. Nothing is applied: the dace tree is shared, currently dirty with another session's work, and this session files DaCe bugs rather than fixing them.
The merge-landing doc carried ten resolved entries with their full original text and verification prose; that is git history, not backlog. A-J are now one line of root cause each, and the two items that did not close -- the null-workspace protocol and the HPTT clone dependency -- are the only entries with argument behind them. BACKLOG_ci_reds_20260804.md is deleted: all four causes closed. Its one surviving ask (a run summary that says failed vs timed out vs cancelled vs could-not-fetch) moves to the merge-landing doc, and the durable facts -- a 403 from a runner is throttling, preemption reads as failure -- are carried in the method notes rather than in a doc titled after four fixed reds. Ablations doc: levels 1-3 of KernelBench are ported, so item 10 is now the level4 decision alone; the expression-Tuple emit gap it referenced is fixed; item 11 keeps its two open pieces and drops the fixed schema narrative.
Backlog notes are a working session's state, not the project's. They read as project documentation to anyone else, they rot the moment the work lands, and they carry per-session context nobody can verify from the tree. Removes the four docs/BACKLOG_*.md and docs/skills_draft/BACKLOG.md; the open work they held is tracked outside the repo. Real documentation stays: the DESIGN_* docs, the guides, canonical_numpy_form.md. Also drops the "BACKLOG item 5" cross-reference from plot_speedup.py's module docstring -- the rule it states (one reader, never mix two run tags) is the part that mattered and it stays. .gitignore keeps BACKLOG*.md out.
…egterg CI has installed libfftw3-dev since the 4-runner split, but nothing checked it. tests/ports/cegterg builds its C++ reference against <fftw3.h> and skips cleanly when the probe fails, so a runner provisioned without the package reports 25 passed, 17 skipped -- green, and indistinguishable from a suite that ran. Checked as a pkg-config module rather than a -lfftw3 link row because the header is the half that goes missing: libfftw3-double3 arrives as a transitive dependency of half a desktop, which makes the link resolve while fftw3.h is absent -- exactly the state of the box this was found on, where a link-only probe would have reported present.
…machine_learning
hpc -> scientific_computing, foundation -> loop_level_reasoning, ml -> machine_learning,
everywhere: the benchmarks/<track>/ directories, all 631 manifests, the selectors, the HF
config slugs, the Harbor adapter, the sbatch submission scripts and the docs.
Hard rename, no aliases -- an old selector is an error naming the new token rather than a
translation layer nobody would ever delete. The `foundation:` manifest block and
`spec.foundation` follow the track they are named after, so there is no leftover spelling.
Two classes of site a word-boundary substitution CANNOT reach, both found by tests rather
than by grep:
* a token glued to a word character -- `submit_foundation_alps`, `foundation_tsvc_2_s212`
(the HF config slug), `foundation_blk`. `\b` treats `_` as a word char.
* a token inside a glob -- `glob("*.y*ml")` in tests/test_ci_coverage.py became
`glob("*.y*machine_learning")`, which matched nothing, so the loop body never ran and
the test PASSED. A vacuous glob is invisible to every gate except one that counts its
own matches.
Prose is decided by hand: HPC also names the field, ML also names the corpus in translator
comments, and `foundation-model` is an LLM. Only the spellings where the word IS the track
were rewritten; "HPC login node" stayed.
…ch results
The five DO-NOT-SHIP pages and the language pages, each defect checked against the tool
that has to make it true:
* papi-gpu-amd: the metric layer was REGENERATED. Five derived-metric formulas were
fabricated (SQ_BUSY_CU_CYCLES as the denominator for VALUBusy / SALUBusy /
LDSBankConflict); the counter names resolve, so wrong numbers came back silently.
Deleted -- rocprofv3 computes derived metrics by name. Also: the event prefix is
rocp_sdk, not rocm; GPUBusy is a percentage, so dividing by it inverts the
normalisation; percentages are no longer summed across regions; the empty-bracket
self-test now fails on probe == 0, which is the silent zero the page exists to prevent.
* rocprofv3: `--pmc` twice does NOT mean two passes -- the launcher joins one survivor and
discards the rest, silently. Two rows in an input file is the documented form.
* rocprof-compute: VALUUtilization was inverted, and the imported KILOBYTES rule was itself
unsourced -- Read BW is Gbps. The real trap is dimensional: FetchSize is a KiB volume,
Read BW is a rate.
* ncu: the gate section routed a reader off a WORKING profiler onto cuobjdump, which the
page itself calls incapable of costing anything. One line pointing at
/proc/driver/nvidia/params replaces it.
* lang-c told the reader to use [[nodiscard]] under -std=c17, where the page's own line 119
says the syntax is unavailable. lang-fortran gated on `flang-new`, renamed to `flang` at
LLVM 20, so the gate never ran; its -fanalyzer gate is C-only and always scored green.
lang-hip's ASan gate passed having instrumented nothing when the runtime was absent
(-print-file-name echoes the name back). lang-python counted a pre-commit run that
produced no output. opt-reports named the wrong directory for two capture kinds.
static-analysis analysed at c++17/c11 while the harness builds c++23/c17.
The nsys and profiling pages absorb the draft additions; the draft nsys page stays until its
-judge twin can move with it, since the variant-2 test pairs them inside the draft tree.
An honest "nothing here was executed" fence is not a substitute for a source: a reader cannot
tell a fenced-but-correct claim from a fenced-and-fabricated one. Every retained claim now
carries an upstream quote and a URL; anything unsourceable was deleted.
The CI red on 81174cc was tests/test_papi_gpu.py seeing 'counted run failed (TIMEOUT)' where the child died of SIGSEGV: a 5 s rep_timeout lost a scheduling race on a loaded runner, and terminate() relabeled the crash. Two-level fix: forked.py now classifies a child that already died of its own signal as that signal even when the deadline check fired first, and the two forked-child tests get a SCHEDULING_PATIENCE_S = 30 budget that names what it actually pays for. test_dace_flavors only exercises config pins THIS DaCe declares.
POST /profile now dispatches on a 'tool' field -- linuxperf, papi, nsys, rocprofv3, none -- and POST /instrument is gone: it was the same build, the same data and the same run with a different instrument attached, so it was a parameter wearing a route's clothes. tool='none' is that capability (the judge attaches nothing and hands back the agent's own stdout) and tool='papi' finally reaches count_submission, which counts without perf and so is the only measurement left where perf_event_paranoid forbids sampling. A tool the language cannot use is a 400 that names the one that serves it, checked before anything builds. /score and /submit were two names for one handler. They now differ by the thing that matters: /score grades the PUBLIC seed only and is never recorded, /submit adds the held-out second seed and settles the run. An agent iterating against /score cannot overfit inputs it cannot see. /oracle stays as an alias for /submit -- it has the hidden seed on today, so aliasing it to the public-only route would have quietly weakened the anti-overfit gate for every page that calls it. A prebuilt library is now read from the shared mount and only from there. Its path in the agent's container means nothing in the judge's, and the judge dlopen()s what that field names -- so an absolute path outside the mount was an arbitrary object of the agent's choosing. The mount's installed libraries are listed in /task, and a link that failed on a library nobody installed now says which name, instead of leaving it in the linker's output.
The route change lands in the docs the agent actually reads: the contract doc gains the /score vs /submit split and the /profile tool table, the tool prompts say iterate with score and finalize with submit, and the counters page names tool='papi' as what to reach for when perf_event_paranoid forbids sampling. Five GPU judge pages lose a capability rather than keep a promise the service does not make: every host tool is a 400 for a cuda/hip submission, so the device tracer is the judge's only instrument there and the instrumented artifact stays the agent's to run. An 'instrumented_ns' field that never existed is gone with them. Two frontmatter descriptions carried an unquoted colon, which is a YAML mapping and broke the pages' own gate.
Both land in AUTOGENERATED *_dace.py, so neither is something a kernel author can route around -- the generator emits them and the kernel silently has no DaCe column. A chained comparison hits a bodyless NotImplementedError in visit_Compare: 48 conv/pool kernels die on 'if 0 <= oy < oh'. Split into its links, but only when every repeated operand is a Name or a Constant -- the split evaluates the middle operand twice, which for a call is a duplicated side effect and for a subscript a second memlet, and a refusal is recoverable where a miscompile is not. reshape's -1 means 'infer from the size' to numpy and a negative dimension to dace: 47 kernels broadcast a bias with reshape(1, -1, 1, 1). Spell the extent out from the operand's shape, and leave the call alone when it cannot be computed rather than guessing an extent into the SDFG.
195 of 576 generated programs do not parse, and nothing measured it: *_dace.py is emitted from the numpy reference, so a construct the frontend refuses lands in the generator's output and the kernel silently has no DaCe column. The failure is invisible because the DaCe column simply is not there to be red. A ratchet, not a pass/fail: REFUSED carries the known set with each cause, a new refusal fails the gate, and a kernel that starts parsing fails it too, so the list can only shrink. Parse only -- no C++ compiler runs -- so the whole corpus fits a runner with nothing but python. One subprocess per kernel, because the two interesting failures are the ones an in-process loop cannot survive to report: cloudsc wedges the parse, and DaCe's parse state is process-global. Measured 4-wide, 24 kernels came back as crashes that all parse fine alone -- so the sweep runs at the parallelism its numbers survive. The runner checks out submodules it does not need yet: the torch agreement phase belongs in this job and its test is not written.
The confinement landed one layer too deep: Sandbox.build serves the in-process callers too -- the optimizers and framework runners build a .so in a temp dir in THIS process and hand back its path, which is not a claim anyone needs to check. Two of their tests went red, correctly. The path is only untrusted when it arrived over HTTP, so _submission_from_body resolves it and the request faults with a 400 instead of the build failing later.
The numpy reference is the correctness oracle for every backend, so a port that quietly computes a different function grades every submission against the wrong answer and nothing ever goes red -- the DaCe, C, C++ and Fortran columns all agree with each other and all of them are wrong. 250 kernels were ported from KernelBench and nothing compared them to their originals: collect_reference_sources.py resolves provenance and the collected model is never imported. 215 of the 250 run against their upstream model at preset S on CPU. The other 35 cannot be lined up mechanically and are pinned in UNALIGNED with the cause, a ratchet like the frontend gate: a port that stops agreeing fails, and a pinned port that becomes comparable fails too, so the list can only shrink. Nothing is skipped -- a skip and a pass look identical in a summary. Three of the pinned causes are defects the comparison found and cannot grade around: four manifests give GroupNorm 4 channels and 8 or 16 groups, so the upstream model cannot be built at the size the harness runs; conv2d_avg_pool_sigmoid_sum's manifest pools by 2 where the model pools by 4; regnet's port and model disagree on num_classes. They are named in the map rather than fixed here, because each one changes a manifest the corpus already measured. A loss returning (1,) against a torch 0-d scalar is the same number in two spellings, so the comparison reshapes rather than calling it a mismatch; that alone moved five ports from disagree to agree.
A function's frees are emitted after its body, so a data-dependent early return jumped straight over all of them. Only helpers can return at all -- the kernel is void and its returns are dropped -- which is why no kernel-level test ever saw it, and it is the worst shape the bug could take: the caller is a benchmark loop, so the helper leaks its workspace once per element per rep. Measured on the audit's own example, 32 reps over 8 elements leaked 256 allocations; a run long enough to measure is a run long enough to exhaust the box. Every exit now releases what is live at that point, innermost branch first, and a scalar return latches its value into a temporary before any free runs -- the returned expression is usually a read of the buffer being released. A body that ends in a return no longer emits the closing frees at all, since nothing can reach them. Under AddressSanitizer rather than by reading the emitted text: with the frees stripped back out, LeakSanitizer fails the run and names the 256 allocations, so the gate is known to be able to fail.
The agreement phase peaked at 6.8 GB, which no GitHub runner has. It is not a leak -- three ports account for all of it, and the largest is 3.1 GB on its own. The cause is that a manifest names SCALARS, so it can never bind an upstream __init__ parameter that is a LIST of layer sizes. Falling back to upstream's value built shallow_wide_mlp at [32768, 32768] rather than the manifest's hidden1=24, hidden2=20 -- and the port was going to be reported unalignable anyway, so the whole allocation bought nothing. A sequence the manifest cannot name is now that refusal, stated before the model is built rather than after. The rest is ordinary care with peaks: the model is only rebuilt when a submodule scalar actually changes an argument (it was always built twice, with both alive at once), the torch side is dropped as soon as its outputs and parameters are numpy, and torch runs single-threaded. 6815 MB -> 1832 MB, same 252 results.
Freeing on every exit is only right where the value survives the free. A scalar return latches into a temporary first, and an array return copies into the out-param before anything is released -- both safe. The gap is a helper that returns a heap local BY VALUE: the previous commit would have emitted free(t) immediately before returning t. That shape is already ill-typed C (a pointer returned from a double-typed function) and only arises when a helper's array return was misclassified as scalar -- the array path rewrites returns into an out-param and never gets here. So neither answer is emittable: freeing hands back a dangling pointer, not freeing leaks. It now names the buffer and refuses, which is what an unsupported construct owes its caller. Verified the exits that DO fall through still free: a void kernel frees its locals at the end, and a helper whose branches all return frees on each of them.
The reallocating free was emitted only where a SECOND marker made the reallocation visible in the emitted text. A marker whose one occurrence sits inside a loop is emitted once and runs once per iteration, so every iteration but the last overwrote a pointer nothing had freed. dbcsr is the shape: A is allocated inside the a_pos loop and freed once after it, B and __mm1 inside the nested b_pos loop. At preset XL that loop runs 640k times with ~160k buffers live at up to 8 KB each -- order 1.3 GB lost per call, and the kernel is called once per measurement rep. gmres does the same with __spv3 once per restart iteration. Freeing before every deferred allocation covers both shapes and is less code than tracking whether a realloc was visible: the declaration NULL-initialises the name and free(NULL) is a no-op, so the first pass through the loop is safe.
numpy's two indexing rules change rank differently, and a length-1 slice is where they visibly differ: a[0:N, 0] drops the axis and is (N,), while a[0:N, 0:1] keeps it and is (N, 1). A kept length-1 axis BROADCASTS -- every position along it reads the same source element. The scalarizer mapped every RHS slice axis onto the destination's iteration variable, so a[:, 0:1] came out as a[i][j]: a whole row where one column belongs. out[:, :] = a[:, 0:1] + b returned wrong numbers in C, C++ and Fortran, with no diagnostic anywhere -- and it could not be caught by cross-backend agreement, because all three share this lowering and agreed with each other. Only the numpy oracle disagreed. Reading such an axis at its slice start is correct whichever extent the destination has: where the destination is also length 1, the iteration variable only ever takes that one value. Covered against numpy for both rules, since the emitted C was plausible enough to read as correct. The symbolic k:k+1 form is recognised as well as the constant one.
n, c, h, w = x.shape is what the helper inliner emits, and it is the single biggest reason a generated dace program is refused. Each unpacked name reaches the frontend as an ordinary body local, so it mints a fresh opaque symbol per use and the buffer sized from them cannot be written from x: "could not broadcast [batch_size, 3, 224, 224] into [__sym___inl6_n_0, __sym___inl6_c_in_0, ...]". 137 generated programs carry that spelling, against 132 IndexError refusals measured. Split into n = x.shape[0], the passes that already exist resolve every one: _ShapeToSymbol for a declared array, _inline_transient_shape_scalars for a transient. No new inlining logic, and the subscript spelling was always the supported one. A swap must go through temporaries. a, b = b, a lowered in source order assigns a = b and then reads the NEW a -- a wrong answer rather than a refusal, so every source is latched first whenever the right-hand side reads any name the left-hand side binds.
Ranks of one job compile different SDFGs into the same .dacecache and the same precompiled-header cache, and neither is written atomically. Racing there does not merely fail: a rank can load the .so another rank is halfway through writing, so the run validates WRONG. The reported symptoms -- library load errors, FileExistsError, crashes, and job timeouts where one rank waits on a build another rank is rewriting -- are all that one race. Suffixing both roots with the rank removes the sharing instead of trying to lock it: no coordination, no lock file to leak when a rank is killed, and a crashed rank leaves only its own directory behind. Four launcher variables are read because Open MPI, MPICH/PMI, Slurm and MVAPICH each publish the rank under a different name, and missing one silently returns the whole job to a shared folder. The header cache is partitioned through DACE_BUILD_CACHE_DIR, the knob DaCe reads. Its default is already RAM-backed (/dev/shm, falling back to ~/.cache), so this only partitions what is in memory already; the cost is one ~110 MB header per rank rather than per node, still bounded by the existing LRU budget. A run with no launcher variable is left on DaCe's own defaults -- it has nothing to race with, and suffixing anyway would give every plain run a cold cache.
DaCe has no runtime .shape: an array's extents ARE symbols, so a shape read has to be resolved before the frontend sees it. _ShapeToSymbol did that for declared arguments only, and a read on a TRANSIENT survived -- (h.shape[3] + 2 - kw) // 1 + 1. That is not merely unresolved: it makes the enclosing size expression non-symbolic, and _plan_size_promotion was all-or-nothing, so ONE such read denied every size scalar in the kernel its symbol. The whole conv family refused on that. Three parts, each measured on alexnet/vgg16/resnet18/lenet5/mobilenet_v1/ squeezenet: ResolveShapeReads rewrites every remaining .shape[k]. It is flow sensitive, because h is rebound per layer and its extents change with it, and deliberately conservative: an extent guessed wrong is a miscompile, not a refusal, so it infers only through an alias, an allocation, a reshape, a transpose and an elementwise result whose operands agree, and never through @, whose result shape is neither operand's. Promotion drops the names it cannot promote, transitively, instead of abandoning the kernel. The closure follows every name in a candidate's right-hand side including positions that are not sizes: np.full's dtype argument dragged an array-valued name in, and that one name cost every size scalar its symbol. Compound reshape extents are hoisted to a name, because DaCe names the container it builds after the shape EXPRESSION and then wants a symbol of that same name. Collected from reshape only -- an allocation takes a compound extent happily -- but substituted in EVERY shape, or the allocation and the reshape name different symbols and DaCe cannot see they are the same extent.
Co-authored-by: Lorenzo Paleari <lorenzo.paleari@inf.ethz.ch>
LorenzoPaleari
force-pushed
the
agent/inference-cluster-example
branch
from
August 6, 2026 14:00
881851a to
136c6d8
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
mainbranch intoclusterwithout container conflictscontainers/cluster/ce-images/inferencepassskeletons for score/submit/bench/verifyWhy
The cluster branch contained generic AMD/NVIDIA Container Engine images and an archived working vLLM build, but did not expose a maintainable inference build path or an end-to-end multi-role Slurm example. This change makes the known-good inference inputs directly usable and supplies the orchestration skeleton needed for a Beverin deployment.
Impact
Operators can configure inference, agent, and judge node counts from a trusted environment file, submit one Slurm allocation, and inspect per-job service and agent logs. Web search is functional. Benchmark grading and remote task assignment remain intentionally unimplemented and return explicit errors rather than simulating results.
The history includes the requested clean synchronization from current
mainintocluster, so the PR contains the corresponding upstream changes in addition to the container work.Validation
git diff --checkbeverin.sbatchandrun_cluster.shagent_driver.pyandjudge_service.pyA real Beverin allocation, CE image build, distributed vLLM startup, and GPU execution were not available in this environment and remain to be validated on-cluster.