Skip to content

feat(BACKEND-ROCM-BF16-MOE): run BF16 grouped experts on ROCm - #3096

Merged
localai-org-maint-bot merged 33 commits into
mudler:mainfrom
VikashLoomba:row/BACKEND-ROCM-BF16-MOE
Sep 12, 2026
Merged

feat(BACKEND-ROCM-BF16-MOE): run BF16 grouped experts on ROCm#3096
localai-org-maint-bot merged 33 commits into
mudler:mainfrom
VikashLoomba:row/BACKEND-ROCM-BF16-MOE

Conversation

@VikashLoomba

@VikashLoomba VikashLoomba commented Sep 9, 2026

Copy link
Copy Markdown
Contributor

Row

BACKEND-ROCM-BF16-MOE. Closes #3094. Closes #1928. Closes #3103. Closes #3116. Closes #3115.

Before starting

The missing providers and shared callers were verified at 6db4bef906859e864c82523c01107473f7dcca29. Committed spec 9202e4c4edc4cf6ef9b3e8da66431effb0fbcee5 precedes implementation 94b8bb0ec67eff82d2860cdda1069e4b0a64bac8. Both remain on this unlanded row branch. One pull request carries the spec, implementation, and test repair.

What changed

ROCm lacked the BF16 grouped expert and fused gate/up providers, preventing the shared resident expert sequence from being selected. This change registers deterministic HIP providers and adds explicit native gate/up, weighted-down, and preweighted-combine operations. The Qwen3 MoE forward path selects native execution only when all five required providers are available. Existing typed calls retain their legacy arithmetic.

The native sequence preserves the pinned vLLM BF16 conversion boundaries and multiplies routing weights before down-output narrowing. Combine sums preweighted BF16 values in FP32. The implementation uses the queue's stream and device and allocates no scratch.

Fresh review found three missing test witnesses. The repair adds all 32 provider subsets, malformed-descriptor cases, and weighted/shared numeric modes. Exact fixtures distinguish F32 retention from BF16 rounding. Product code, original fixtures, and tolerances remain unchanged by the repair. The work is independent of #2782 and changes no CI configuration.

Evidence

  • The clean CPU build passes all six declared regression executables: 72 cases and 2910 assertions. The registry retains one previously disabled case.
  • Both legacy grouped GPU suites pass. The original native gate passes eight cases and 2804 assertions, including streams, graph replay, tails, and two-device execution.
  • All 60 original upstream component cases pass on both runtimes at e126687a9a828d513c01a07cd69f025f27d63280. Their parameters, tolerances, physical-format records, and generated kernels are retained.
  • The repaired native contract passes 219 assertions. The new GPU weighted and shared cases pass 72 and 42 assertions. All 14 review-surviving mutations now fail their intended assertions. An additional F32-store narrowing mutation fails both precision witnesses.
  • The coordinator independently reran the CPU regressions, CPU mutations, and focused GPU controls and mutations. Original source and archive hashes remain unchanged. HIP test relinking uses the original production archive; no new full HIP product build is claimed.
  • Full staged preflight exits 0 with 619/619 host units compiled. The five argument-required skips are reconciled by the CPU/HIP ISA audits, exact-range path classification, and inapplicable ARM/CUDA artifact checks. The pre-edit temporary-repository fixture failure passes all 22 tests with the parent Git lookup isolated. No checker changed.
  • Fresh scoped review passes with no findings at 9083a5fb35822e86b31cf786e0c80150dfae5478. Its independent rebuilds detect all ten CPU and five GPU mutations. The coordinator also reran the full preflight on this exact commit: exit 0, with the generic skips separately resolved or marked inapplicable. Exact-range record, scope, style, and attribution checks pass.
  • The commands, binary hashes, mutation receipts, and gate dispositions are recorded in .agents/specs/rocm-bf16-moe.md. The reviewer report and coordinator receipts are retained with the row evidence.

Speed claims

No performance result is accepted, and this change makes no performance claim. The fidelity work tracked in #3115 and #3116 is judged by the reference's own cross-configuration band rather than by reproducing one configuration's tie-break.

Token gate correction (heads 61e8e8d1b through 691b7af30)

The production token gate previously compared the native sequence with the single captured record at the run's own concurrency. The pinned reference's greedy decode is not unique across its own legitimate configurations: at length 33 the identical prompt yields [66,1,70,57,33,81,63,69] at concurrency 1 and [66,1,70,57,33,81,118,66] at concurrency 2, in all three repeats each, because the two executions resolve an exact BF16 tie differently. Comparing against one capture therefore fitted one noise realization of the reference rather than the reference function.

The developer ratified the corrected definition on 2026-09-09: a workload passes when the native sequence equals, for that request, a whole sequence the reference itself emits under one of its captured configurations; positions at which the captured configurations agree remain exact, because membership implies they match; the gate reports the reference set, the matched configuration, every same-configuration outcome, and every position where the reference disagrees with itself. It never mixes positions from two reference sequences, never drops a workload, and never relaxes a stable position.

Evidence at 691b7af30: the operator GPU gate exits zero with 28,598 of 28,598 assertions passing (the 18 beyond the pre-requirement receipt are the request-count requirement, one per matched record), and the operator's mutation study restores the superseded same-configuration-only comparison and reddens exactly the three recorded failures. The tied workload reports reference set {concurrency 1 [66,1,70,57,33,81,63,69], concurrency 2 [66,1,70,57,33,81,118,66]}; matched configuration 1; same-configuration match false; reference disagreement positions [6,7], and every other workload reports same-configuration match true with no disagreements. Five predicate mutations (per-position mixing, a comparison truncated to the stable prefix, a dropped disagreement report, a single retained configuration, and an ignored stable position) each fail their intended case, and the focused CPU predicate suite passes. The superseded head's red is the recorded three-failure whole-sequence comparison.

LM-head output boundary (head 6d6be4820, child spec rocm-lmhead-bf16.md)

The native head projection stored F32 logits where the compiled primary stores BF16 and widens through its sampler. On the retained primary diagnostic (hidden [2,128], weight [128,128], primary logits [2,128], all BF16, decode step 6 of L33/C2/R0) the mirrored boundary now reproduces the primary word for word: the BF16-store-plus-CastF32 arm equals the primary's logits 0 of 128 words on both rows, where the previous F32 store differed by up to half a BF16 ulp (max 9.2e-04). The arms and both devices agree value-for-value; the focused suite runs with and without its fixture, and the fixture-free production case always executes (a reddened case can no longer be reported as skipped). The 18-workload production gate is unchanged: exit 0, 28,598 of 28,598 assertions, all 216 generated tokens identical, and all 27,648 returned logits are now BF16-representable where none were before. This is a correctness mirror, not a token fix: on the primary's own head input the argmax is 118 either way, and on the native step-6 logits the narrowing turns the 63/118 margin into an exact tie that the lowest-index tie-break resolves to the token native already emits. The tied-embedding arm is changed by the same store and has no artifact replay; that debt is recorded in the child spec.

Attention parity (#3115, landed at this head)

The attention measurement landed first and is operator-verified: at L33/C2/R0 layer 0 step 0 the native output differed from the primary at 2902 of 8448 BF16 words, decomposing into kernel arithmetic 2918 of 8448 on the primary's own Q/K/V (self-consistency guard 0) and the Q/K preamble boundary 1569 Q / 1542 K, with the cos/sin table and qkv slices exact. The obvious mirror was falsified by measurement: narrowing the probability at the accumulate moved the kernel term the wrong way, 2918 to 3011, because a BF16 rounding is relative and the primary's exponent reference is a contiguous tile maximum, not a per-warp running one.

The landed repair mirrors the primary's three executing geometries per row: a 32-key context tile anchored at key zero then 64/32-key chunk tiles anchored at the chunk start for query lengths above one, and min(block_size,128)/32 decode tiles anchored at key zero for single-token steps, with the CTA-wide maximum, the f32 correction, and the narrowed probability in the primary's order. Measured after: native output versus primary 2902 to 0, kernel on the primary's own Q/K/V 2918 to 0, post-RoPE Q/K 1569/1542 to 1/0, and the production tokens unmoved. Both switch settings of the fused-chain lever now run identically — the first review round caught that the new hand-call realization threw on row-major rank-3 tensors because of a copied stride check, which broke VT_FUSED_CHAIN_ADOPT=0; both that and the hardcoded tile geometry were repaired, and the second review round passed. The operator independently re-ran the instrument under both settings: 8 of 8 cases, 30,146 of 30,146 assertions, with every dumped native byte identical between them.

Owed and recorded in the child spec: the decode and cached-context arms are synthetic-device-measured but have no primary capture (the primary attention capture is step-0-only and its recipe is gone), mixed-batch and sliding-window arms are unmeasured, and no throughput axis was measured — the kernel now pays one CTA synchronisation per tile instead of per 64 keys. Decode-step word parity additionally requires the two sides to share a physical block size, which the rule mirrors but the oracle capture cannot.

Honest gaps

The production token gate certifies the reference function at this head: the prior six-position signature was three repeats of one two-token chain at the reference's own exact BF16 tie, and the native sequence matches a reference-emitted sequence exactly at every workload. The gate's pass rule is whole-sequence membership with stable positions exact; the separate fidelity criterion for #3115 (decode attention and Q/K preamble arithmetic) and #3116 (the BF16 LM-head output boundary) is that the native-versus-reference logit difference stays inside the reference's own cross-configuration band, measured from the retained BF16 per-step logit diagnostics as one BF16 ulp (2^-9 = 1.953125e-3) at agreeing-context steps. No performance claim is made here. Issue #3103's residual-normalization portion is repaired at this head: the committed BACKEND-ROCM-RESIDUAL-NORM spec and fixes preserve the compiled primary's FP32-then-narrow residual boundary, match all 128 captured primary values, and pass all 264 pinned upstream export cases plus 528 native/reference executions with 22,576 assertions; the complete component gate passes 33,795 assertions on gfx1100, the stream-fixture repair passes with its wrong-stream mutation red, and the full preflight exits zero. Attribution of the remaining six positions is recorded with the row evidence: they are three repeats of the same two-token chain of request 0, rooted at a decode step where the pinned primary itself flips between its own cc1 (63) and cc2 (118) executions over an exact BF16 tie, while native is bit-identical across batch widths and matches the primary's cc1 answer; native-vs-oracle logit noise at the flip step sits in the same <=2e-3 band as all 210 passing positions, and residual causation is excluded. The attention and Q/K preamble parity gap is filed as #3115 and the BF16 LM-head output boundary as #3116, both assigned to this row; shared resource-device issue #3100 remains separate.

The Qwen3 fixture proves production reachability of the shared expert operations. This change does not establish complete DeepSeek-V2 or dots3-note execution through their currently refused router modes. Shared resource-device issue #3100 remains separate. The row remains ACTIVE; this session publishes for user review and does not merge.

Repairs after the review above, verified at this head. 506832d32 moves the residual-norm repair evidence up into its run directory, because pr-size admits exactly one directory level under docs/bench-evidence/<run>/; all 17 files are byte-identical renames and the only other change is the single reference in .agents/specs/rocm-residual-norm.md. c7b0ebf21 gives the four ROCm-only cases in this branch's test_rocm_moe_bf16 the file's own doctest::skip(...) decoration and REQUIRE_MESSAGE backstop, so the CPU-only lane reports CTest Skipped (exit 77) at the same test number CI uses instead of fataling, while a ROCm host still runs them (4 passed | 4 skipped, 33 assertions). 8af736e67 makes the shared skip note name an absent prerequisite rather than an unset environment, which is the fact both gating kinds share. A fresh reviewer mutated each guarantee in a scratch copy and returned PASS, including a classifier mutation that reproduces CI's unclassified-path error verbatim and a forced-predicate mutation that reddens the four cases; the operator reran the affected gates.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:gpt-6-astra [Codex]

ROCm lacks both shared BF16 expert providers. Record the pinned oracle's
rounding and routing-weight contract before implementation, with a
reproducible production fixture and scoped issue ownership.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:gpt-6-astra [Codex]
ROCm cannot select the existing grouped BF16 expert path because its
providers are absent. Register the legacy operations and add explicit
native gate/up, weighted down, and preweighted combine contracts. Route
the complete provider set through the shared Qwen3 MoE forward path.
The native providers keep intermediate BF16 rounding and use no scratch.

Both legacy suites, eight native boundary cases, all 60 pinned upstream
component cases, and six CPU regressions pass. Twenty HIP mutations and
nine CPU descriptor mutations detect their intended defects. The staged
gate compiles all 619 affected host units; supplementary checks account
for its reported skips.

Exact model parity remains open: the native and pre-implementation legacy
paths emit identical tokens, but both differ from the oracle at six
positions across three repeated workloads. Keep the row ACTIVE and all
performance acceptance open. Record the independent residual-normalization
and resource-device gaps under BACKEND-ROCM without changing their paths.

Refs mudler#3094, mudler#1928, mudler#3100, mudler#3103.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:gpt-6-astra [Codex]
The native MoE tests could miss partial provider sets, malformed descriptors,
and accepted weighted and shared modes in mudler#3094. Check every provider subset
and add exact numeric witnesses that distinguish FP32 retention from BF16
rounding. These tests detect all fourteen review survivors and an additional
FP32-store narrowing mutation.

Product arithmetic, original fixtures, and tolerances remain unchanged.
The established production token failure still prevents performance acceptance.
This repair stays on row/BACKEND-ROCM-BF16-MOE for fresh scoped review.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:gpt-6-astra [Codex]
The native Qwen3 MoE path rounds a residual that the compiled primary
retains as separate BF16 operands across MoE. Define the ordered shared
operation and its materialization points before implementation. Keep the
production token gate mandatory and assign the repair to its child row.

The contract includes captured-input failures, exact association witnesses,
shared fusion routing, and production lifetime mutations. Attention and
head repairs remain explicit dependencies for the enclosing parity work.

Refs mudler#3103.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:gpt-6-astra [Codex]
The compiled primary retains separate BF16 attention and residual operands
across MoE. Normalize their ordered FP32 expression through the shared
ResidualRmsNorm operation and FusedChain, then materialize only the next
layer's BF16 residual. Keep CPU and CUDA production defaults unchanged.

The production red failed 33 of 128 norm words on identical inputs. The
operator's first repaired production witness passes all 128 primary words,
and the native/reference component gate passes 10,700 assertions. Actual
next-residual, next-norm, and final-norm captures come from unchanged
compiled launches with exact paired-control tokens and logprobs. Six CPU
regression suites pass, including the existing Qwen3 MoE forward control.

This immutable implementation starts fresh review while the full staged
gate runs. Expanded hardware, pinned upstream fixtures, both complete
production fusion modes, and fresh mutation review remain pending. The
existing ROCm platform excludes production decode graphs; component graph
coverage does not change or certify that policy. Complete model token
parity and performance acceptance remain open in the enclosing campaign.

Refs mudler#3103.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:gpt-6-astra [Codex]
The stream-order fixture requests a backend queue whose default stream
flags are zero. Its nonblocking assertion fails before kernel dispatch.
Own an explicit nonblocking HIP stream and drain its callback before
releasing callback state or operands. Preserve both assertions and the
real provider mutation that launches on the wrong stream.

The residual descriptor include moves the Backend definition by one line.
Repair its BACKEND-PLATFORM citation and prove every other matrix byte
unchanged. Product source and default queue behavior remain unchanged.

The repaired test compiles for CPU and HIP. Six CPU tests pass 8,985
assertions, including the original 264 upstream cases. The operator's
unchanged native implementation passes all 528 CPU/ROCm upstream cases.
The record and symbol gates pass. Hardware repair and mutation execution,
full preflight, fresh review, and operator verification remain required.
The six original L33/C2 token mismatches remain open under the parent gate.

Refs mudler#3103

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:gpt-6-astra [Codex]
…ead issues

The six-position disposition analysis at cc9d4f5 filed mudler#3115 (decode
attention and Q/K preamble parity) and mudler#3116 (the BF16 LM-head output
boundary) on the BF16 MoE row. Refs mudler#3103.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
… gate

The pinned primary's greedy decode is not unique across its own legitimate
configurations. At length 33 the concurrency-1 and concurrency-2 captures of the
same prompt differ at positions 6 and 7 of request 0, so demanding one
configuration's tokens fits a sample instead of mirroring the function. This
records the developer-ratified definition that a workload passes by
whole-sequence membership in the reference's captured configuration set, and
restates the exact-token row's measured result under that rule. Two statements
the corrected rule falsifies are aligned with it: the paired-performance row's
prerequisite and the repair section's "remains failing" reading of the same
comparison.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
…figuration set

The production gate compared each native sequence with the single oracle record
of its own concurrency, which fails at length 33/concurrency 2 because the pinned
primary's own configurations disagree at positions 6 and 7 of request 0. Build
the reference set from every captured concurrency of the same workload and accept
a request only as a whole-sequence member of it, and report the matched
configuration, the same-configuration outcome, and the reference's own
disagreement positions. The membership predicate is pure host code gated by a new
CPU case, so lengths 1 and 3 keep their exact comparisons through their singleton
reference sets.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The corrected production token gate's pass sentence folded in a fidelity band
that no gate in this change measures, and the acceptance row cited a superseded
head's failure log as the head's own receipt. Keep the pass rule to
whole-sequence membership plus the stable-position consequence, move the band
into a separate fidelity sentence recomputed from the retained BF16 head-logit
captures (1.953125e-3, one BF16 ulp at the observed magnitude) that names itself
the acceptance criterion for mudler#3115 and mudler#3116, add the same-configuration outcome
the gate prints to the report sentence, and cite the head's operator receipt
(exit 0, 28580 of 28580 assertions) while labelling the earlier log as the
superseded-head red history the corrected gate answers. The phantom
production-native-1.log reference is removed because no such file exists. The
Now and Owed records now say the gate passes under the corrected rule and name
the paired decode and prefill measurement as the live performance blocker.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The production gate's request loop was driven by native.size(), so a native run
that returned the wrong number of request sequences would compare fewer requests
than the captured record holds and still pass. Require the native request count
to equal the record's concurrency before the loop, and comment why. This changes
the gate binary, so the operator's existing GPU receipt no longer applies to
these bytes and the gate must be rerun.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The acceptance row cited the pre-REQUIRE receipt; record the operator's receipt
at this head (28,598 of 28,598 assertions, including the 18 request-count
requirements) together with the mutation reddening exactly the three recorded
failures, and name every step that attains the maximum per-logit difference.
Refs mudler#3115. Refs mudler#3116.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
…value

"At most one BF16 unit in the last place" is magnitude-relative, so it is not a
well-defined acceptance number for mudler#3115 and mudler#3116: one ulp is 3.90625e-3 above
a logit of 0.5 and half that below it. State the bound as the absolute
`1.953125e-3` (`2^-9`) per logit that was measured, and keep the ulp reading as
the interpretation of that maximum at the magnitude where it occurs.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
…ifact

The acceptance row cited a file that the campaign cleanup deleted with the
superseded evidence tree. Re-derive the claim at this head: the legacy and
fused native paths both exit zero for all 18 records and emit the same 216
generated tokens, recorded in a fresh receipt. Refs mudler#3094.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
Issue mudler#3116 asks whether the native forward should store its LM-head output in
BF16 the way the pinned primary does, and whether that boundary alone moves the
decode step-6 argmax. Measure it before changing anything: replay the primary's
captured head-6 hidden, weight and logits bytes through the operator the forward
calls (vt::Matmul) in three arms. A BF16 store followed by vt::CastF32
reproduces the primary's logits element for element (0 of 128 differing words on
both rows), the F32 store the tree ships today does not (max 9.23157e-04 and
9.72956e-04, half a BF16 ulp), and neither arm moves the step-6 argmax on the
primary's head input (118 before and after) or on the recorded native production
logits (63 before and after, where narrowing collapses the 5.71579e-04 margin to
an exact tie at 0.3359375). The spec records that measurement, scopes the repair
to the head-output dtype boundary, and states that hidden-state parity (mudler#3115)
is still required and that this change must not be claimed to fix the tie.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The native forward stored its LM-head output in F32 while the pinned primary
stores BF16 and widens it for the sampler, so the two heads disagreed by up to
9.23157e-04 on row 0 and 9.72956e-04 on row 1 of the primary's captured head-6
input. Replay the primary's exact hidden and weight bytes through the projection
the forward calls: a BF16 store followed by the shared vt::CastF32 reproduces
the primary's logits element for element (0 of 128 differing words on both rows,
max_abs 0), where the F32 store differs in 256 of 256. Route the head through
one seam (`lm_head::Project`) so a checkpoint-free test executes the boundary
the forward runs, and keep the returned buffer F32 so the device-logits views,
the host download and the captured graph slot stay valid. No step-6 token moves:
both boundaries choose 118 on the primary's head input and 63 on the recorded
native production logits, where narrowing collapses the 63/118 margin to an
exact tie at 0.3359375 that the lowest-index tie-break resolves the same way.
Hidden-state parity (mudler#3115) is still required. Refs mudler#3116.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The head-boundary spec names `ISSUE-mudlerGH-3116`, and the record gate resolves that
identifier only against a canonical local record, which the issue did not have:
it existed on GitHub alone. Import it so the spec's issue reference resolves to
exactly one record in this branch. Refs mudler#3116.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
…boundary

The row records 2902 differing BF16 words at the first attention output and
cannot say whether they come from the Q/K preamble, from the attention kernel's
arithmetic, or from both, because no native byte of that boundary was captured.
Three provider-seam observers on the live ROCm ops now capture it from the
unchanged production path, and the same process replays the primary's own
captured Q/K/V through vt::PagedAttention at the primary's cache geometry, so
the kernel term is measured on identical inputs instead of inferred. Nothing
under src/ or include/ changes: every observer is the wrapper pattern this file
already uses for the residual row, and all of it is gated on VT_ATTN_DUMP, so
an ordinary run registers nothing and both new cases skip with exit 77.

The verdict on gfx1100 is BOTH, kernel-dominant: 2918 of 8448 words on
identical inputs against 1569 Q and 1542 K words from the preamble, with the
cos/sin table and the qkv projection both byte-exact, which also localizes the
preamble term to the pre-RoPE BF16 norm store. A CPU-only companion case checks
the replay's cache geometry against the recorded capture, so the device run
cannot silently measure the wrong addresses. Refs mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The capture section claimed no existing artifact holds the attention boundary's
inputs and output, which the preserved evidence falsifies: the record and its
ten binary siblings under the attention-operator results directory hold exactly
those bytes. The sentence is replaced with the artifact path and with the fact
that the validated CPU transcription of prefix_prefill._fwd_kernel separated
the preamble term from the kernel term on it, and the native measurement that
followed is recorded beside it so the section no longer reads as blocked on a
capture that already exists. Refs mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The row's replay instrument measured both candidate terms on identical inputs and
this spec designs their repair. The ROCm softmax keeps its probability in f32
where the primary narrows it to the value dtype before the value dot, and the Q/K
preamble rounds the normalized value to bf16 before RoPE where the primary
carries f32 and narrows once at the store. It scopes all eight probability sites
by disposition, names the two rocWMMA kernels and the Tier-0 composite as owed
with their reasons, and fixes the two numbers the repair is measured against: the
primary's own boundary at 1 word of 8448 and the kernel replay at 0. Refs mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
…s focused cases

The row's instrument measures the attention boundary but asserts nothing about
the repaired arithmetic, so a repair could land with no case that reddens when it
is missing. Two focused checks now do: the kernel replay on the primary's own
Q/K/V must reach zero differing words, and the native run's post-RoPE Q and K
must reach the CPU model's one-word residue against the primary. The preamble
boundary is read from whichever realization executed, the Tier-0 composite at
this head and a registered fused qk-norm-rope op once one exists, and a new lane
measures that op directly on the primary's captured qkv. Nothing under src/ or
include/ changes. Red at this head: C 2918 of 8448 words, B 1569 Q and 1542 K,
with every pre-existing measurement line reproduced byte-for-byte. Refs mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
…reference max

The first implementation of design 1 narrowed the softmax probability to the
value dtype exactly as prefix_prefill.py:471 does, and measured row C move the
wrong way on gfx1100: 2918 to 3011 of 8448 words. The preserved CPU transcription
of the native kernel reproduces that 3011, and parameterizing it shows why: a bf16
rounding is relative, so it only mirrors the primary when the exponent's
reference max is the primary's own, and that max is maximum(m_i, tile_max) over a
contiguous 64-key tile, not a per-warp per-key running max. The spec now carries
the falsification table, redefines design 1 as the value-dtype probability taken
at the tiled reference max, states the sync and register cost of that structure,
and moves every unmeasured sibling from changed to owed with its reason. Refs mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
…rimary's reference max

The pinned primary narrows the softmax probability to the value dtype before the
value dot (vllm/v1/attention/ops/prefix_prefill.py:471) and keeps only the running
sum in f32 (:445 against :475), and this kernel kept both in f32. Narrowing at the
accumulate alone was implemented and measured first and moved the replay row the
wrong way, 2918 to 3011 of 8448 words, because a bf16 rounding is relative: it
mirrors the primary only when the exponent's reference max is the primary's own,
which is `maximum(m_i, tile_max)` over a contiguous 64-key tile (:442) and not a
per-warp per-key running max. The key walk is therefore tiled at 64 keys, each
tile's max is reduced across the CTA, and every warp exponentiates, narrows and
accumulates against that shared max, with alpha applied to the f32 accumulator
exactly as :449 applies it. Row C on the primary's own Q/K/V goes 2918 -> 0 of
8448 words on gfx1100 and the recorded A0 gap falls 2902 -> 1798. The other
attention kernels keep their arithmetic and are named owed in
.agents/specs/rocm-attn-parity.md. Refs mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The native preamble stored RmsNorm's bf16 result and rotated that, so the
normalized value was rounded twice; the pinned primary keeps it in f32 through
RotaryEmbedding and narrows once, at the store, reading a bf16 cos|sin cache
(vllm/model_executor/models/qwen3.py:150-167 and
vllm/model_executor/layers/rotary_embedding/base.py:105-131). ROCm registered no
fast realisation of the kAttnQkNormRope recipe, so vt::FusedChain ran the Tier-0
composite and the hand-call ran the same three ops. Both now dispatch to a ROCm
AttnQkNormRope kernel -- per-head standard RMSNorm with the shipped reduction, an
in-place pair-wise rotation, one narrowing at the store -- while a backend that
registers no fast op, and every f32-attention caller, keeps the composite bytes
it has today. Measured on gfx1100 at L33/C2/R0: the native run's post-RoPE Q and
K against the primary go 1569/1542 -> 1/0 of 8448 words, the production preamble
on the primary's own qkv lands at 1/0, and the first attention output closes the
recorded 2902-word gap to 0. Refs mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The two replay cases called `std::exit(77)` when `VT_MOE_HEAD_FIXTURE` was
unset, and doctest runs the cases in file order, so the fixture-free
production call-site case below them never executed: on CI and in a fresh
checkout the binary reported Skipped and the one case in it that needs no
capture never ran. That is the fresh review's first LOW finding.

Skip the two replay cases instead of the process. `doctest::skip(...)` makes
doctest report them skipped in its own summary, so the call-site case runs and
is reported in every configuration, and `--no-skip` over that decorator fails
the case loudly rather than letting an absent variable be read as a path.

A fixture-absent run still exits 77 -- CTest reports Skipped, the convention
the suite's registered SKIP_RETURN_CODE documents -- but only when nothing
failed. A reddened call-site case keeps doctest's EXIT_FAILURE, and the
mutation proves it: with the seam reverted to its pre-change F32 body the
fixture-free run exits 1 with 2 failing assertions, while the same mutation
plus a disabled failure guard reports a masked 77.

Refs mudler#3116.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
Corrections the fresh review of `row/BACKEND-ROCM-BF16-MOE-head-bf16` asked
for, in the records that describe the head boundary.

The child spec listed "the tied-embedding arm's arithmetic" as out of scope and
unchanged. It is not unchanged: `lm_head::Project` sends both arms through one
BF16 store and one `vt::CastF32`, so the tied `vt::MatmulBT` narrows to BF16
where it stored F32 directly before. No artifact replays that arm and the
fixture this spec measures is untied, so the bullet now states changed,
unreplayed, and unreachable from the measurement, which is what `## Owed`
already recorded.

Its `## Now` said the token gate, the mutation review, and the preflight
remained required. The receipts show the opposite: the 18-workload gate passes
at `8e43d18bd` under the corrected whole-sequence-membership rule (28598 of
28598 assertions) and the frozen-head preflight exits 0 there, so the section
records those receipts and the repair this branch adds.

The parent spec named mudler#3116 without the child spec that owns it. It now links
`rocm-lmhead-bf16.md`, and its `## Now` credits the head-output boundary to that
child spec instead of listing it among the differences that still require
resolution.

Refs mudler#3116.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
…ings

The fresh review of the mudler#3115 attention repair found five defects. Three of them
kept a documented lever or a gate from working at all, and none is in the
arithmetic the row measured.

The op-layer AttnQkNormRope validation asserted "row stride must be the inner
dimension" on stride[1] of a rank-3 [T,H,Dh] tensor, where the row stride IS Dh.
It contradicted the IsContiguous() check two lines above it and refused every
Dh > 1 operand, so the hand-call realization the bf16 branch takes under
VT_FUSED_CHAIN_ADOPT=0 threw on every forward and the row's same-binary A/B lever
could not run at all: three of the four cases in the production binary threw
"row stride must be the inner dimension" before this change. The check now names
stride[2], and the instrument runs both settings of the switch to the same
30146 assertions with byte-identical dumps, the adoption flag the only differing
field in its report.

The kernel's reference-max tile was hardcoded at 64 keys. That is the tile of one
arm of the primary. chunked_prefill_paged_decode.py:317 dispatches the call on the
batch's max query length and the decode kernel it launches returns for a row with
query_len > 1 (:93-97, :503), so the arm is per ROW: query_len > 1 runs
prefix_prefill, which walks the cached context in TRITON_BLOCK_SIZE = 32 key tiles
anchored at key 0 (:1007, :379-414) and then the chunk in BLOCK_N-key tiles
anchored at the chunk start (:426-478, BLOCK_N 64 for a power-of-two physical
block size and 32 otherwise, :955-966); query_len == 1 runs the Triton decode
kernel tiled by min(block_size, 128), or 32 for a non-power-of-two physical block
size (:444-445). The kernel now reads the arm off the row and walks that grid, so
a decode step whose physical block size is not 64 narrows against the reference
max the primary uses instead of one it never used. Nothing moves on the measured
workload: row C stays 0 of 8448, and the production gate's 27 token records are
identical with this file reverted and with it in place.

The focused case could not tell a 32-key tile from a 64-key one: for this
capture's rows max(keys 0..31) == max(keys 0..32), so a 64 -> 32 mutation stayed
green. Three device cases now synthesize the data that makes the width
load-bearing -- each key's value row is a single bf16-exact integer in its own
output lane, q = e0, and one high-scoring K row sits just past a tile boundary --
and compare the device against a host transcription of the primary's own key walk
for the arm AND against the same transcription under the neighbouring widths:
0 differing words against the arm's tiles, 1408 / 192 / 64 against the wrong
ones, against a noise floor of 8 and a signal floor of 32. A fourth case runs the
recipe's fast realization and the hand-call over one synthesized q/k and requires
byte identity, which is the ADOPT lever as a gate rather than a claim.

The instrument's cases called std::exit(77) on an absent environment variable,
and doctest runs its cases in file order, so the first absent variable ended the
PROCESS before its summary and every later case silently did not run. Each gated
case is now decorated doctest::skip(...) with an in-case guard, the process exits
77 only when something was skipped and nothing failed, and a fixture-free run
reports the four environment-free cases and its summary. The spec records the
same-binary A/B result and the arms that are implemented but unmeasured.

Two arms the source pins have no capture to measure them against: the primary
attention capture is step-0 only, so every recorded row has query_len > 1 and an
empty context, and the capture recipe that produced it is gone (gap G2). The spec
carries them, the sliding-window arms and the mixed batch under ## Owed, and
re-anchors its sibling-site line numbers on this file. Refs mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The scoped re-review of the attention-parity repair found two of the anchors that
repair added wrong. Each names a range that does not hold what the text beside it
claims, so a reader checking the mirror lands on the wrong loop. This is a
documentation change. No code, test or measured value moves.

`prefix_prefill.py:379-414` was labelled "the context loop from key 0". That range
is the CHUNK loop's KV-load region. At `:379` the comment reads "cache at their
absolute positions [ctx_len, ctx_len + query_len)," and `:381` computes
`cache_token_idx = cur_batch_ctx_len + start_n + offs_n`, the current chunk read
out of the paged cache. The context loop is `:231-343`. It opens at `:231` with
`for start_n in tl.range(`, steps at `:232` with `0, cur_batch_ctx_len,
BLOCK_SIZE`, and last writes `m_i = m_ij` at `:343`. Its arithmetic inside is the
running max at `:312` (`m_ij = tl.maximum(m_i, tl.max(qk, axis=1))`), the narrowed
probability at `:338` (`p = p.to(v.dtype)`) and the value dot at `:340`
(`acc = tl.dot(p, v, acc=acc, input_precision=IN_PRECISION)`). Both files now cite
`:231-343`, two lines past the review's own `:231-341`, because `l_i = l_i * alpha
+ l_ij` at `:342` and `m_i = m_ij` at `:343` are inside the loop. The spec's mask
citation `:291` is unchanged.

The kernel comment's anchor list `(:1007, :1014, :267-295, :379-414)` cited
`:1014` for the context walk's geometry. `:1014` is `SKIP_DECODE=skip_decode,`, an
unrelated launch argument. The tile this kernel mirrors is
TRITON_BLOCK_SIZE = 32. It is defined at `:965` (`TRITON_BLOCK_SIZE = 32`) and
bound as the kernel's BLOCK_SIZE at `:1007` (`BLOCK_SIZE=TRITON_BLOCK_SIZE,`). The
list now reads `(:965, :1007)`, and the spec names that pair where it labelled
`:1007` alone.

The chunk loop is now cited by its header at `:369` (`for start_n in tl.range(`)
instead of the body range `:426-478`. That range also ran one line past the loop
into the epilogue at `:478` (`acc = acc / (l_i[:, None] + 1e-10)`). The arithmetic
anchors `:442`, `:471` and `:473` are unchanged and still hold their quotes. Refs
mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The first re-anchor fixed the spec and the kernel comment. Three comment blocks in
the focused device suite and the mask citation in both files still named the wrong
lines, so the same reader trap survived in the cases that gate the geometry. This
is a documentation change. No code, test or measured value moves.

The mask citation is the one substantive correction. `prefix_prefill.py:291` is
`# qk *= sm_scale`, a comment line, so a reader following it lands one statement
past the mask. The context mask itself is `:288-290`. `:288` is `qk = tl.where(`,
`:289` is `(start_n + offs_bs_n[None, :]) < cur_batch_ctx_len, qk, float("-inf")`
and `:290` closes the call. Both files now cite `:288-290`. The arithmetic anchors
`:312` (`m_ij = tl.maximum(m_i, tl.max(qk, axis=1))`), `:338` (`p = p.to(v.dtype)`)
and `:340` (`acc = tl.dot(p, v, acc=acc, input_precision=IN_PRECISION)`) are kept.

`tests/vllm/models/test_rocm_moe_bf16.cpp` carried the two anchors this row had
already corrected elsewhere. At `:1111` and `:2082` the cached-context tile was
`(:1007, :1014)`. `:1014` is `SKIP_DECODE=skip_decode,`, while the tile is defined
at `:965` (`TRITON_BLOCK_SIZE = 32`) and bound at `:1007`
(`BLOCK_SIZE=TRITON_BLOCK_SIZE,`). At `:1142` and `:2082` the loops were
`(context loop :379-414, chunk loop :426-478)`. `:379-414` is the chunk loop's
KV-load region, and `:426-478` is a body range that also runs one line past the
loop into the epilogue at `:478`. They now read `(:965, :1007, :231-343, :369)`
and `(context loop :231-343, chunk loop :369)`. The context loop opens at `:231`
(`for start_n in tl.range(`) and closes at `:343` (`m_i = m_ij`).

The suite's own transcription range `:442-478` at `:1142` is unchanged. It names
the chunk loop's arithmetic and epilogue, and those anchors hold. Refs mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
Brings the measured attention-parity lane: the step-0 decomposition (kernel
2918 of 8448 on identical inputs; preamble 1569 Q / 1542 K), the per-arm
tile-geometry repair that reaches zero on the measured arm and mirrors the
decode and cached-context arms from the pinned source, the repaired fused-chain
hand-call realization, and the re-anchored primary citations. Refs mudler#3115.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The pr-size check could not classify six raw artifacts of this repair, because
`docs/bench-evidence/<run-id>/<file>.<ext>` is the only per-run evidence shape
`BENCH_EVIDENCE_RUN` admits and those six sat one directory deeper. The check
refused the whole change with `unclassified repository path
'docs/bench-evidence/rocm-residual-norm/stream-fixture-repair/build-fixture.py'`.

Every other per-run evidence directory here is flat: raw run artifacts sit
directly in the run directory, and only `.md` and `.json` documents appear in
nested subdirectories, which the `DOC` rule classifies at any depth. This makes
the directory match that convention. `rocm-residual-norm` is new on this branch,
so no earlier layout is lost.

Move all 17 files up into `docs/bench-evidence/rocm-residual-norm/`, remove the
empty directory, and update the one reference in
`.agents/specs/rocm-residual-norm.md`. Every move is a rename with no content
change, and the receipt JSON files name their siblings by basename, so no
recorded path claim breaks.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The CPU-only lane configures `VLLM_CPP_HIP=OFF`, so no ROCm backend is
registered there. The four cases that measure the ROCm arm on data the file
builds itself opened with a bare `REQUIRE` on that backend and killed the test:
the lane reported four FATAL errors instead of the CTest Skipped that every
other ROCm-only case in this file reports. A lane that cannot host the backend
has nothing to measure, so those cases must say so rather than fail.

Decorate them `doctest::skip(RocmBackendAbsent())` and replace the bare
`REQUIRE` with the same `REQUIRE_MESSAGE` backstop the file's environment-gated
cases use. The predicate asks whether the backend is registered rather than
whether a variable is set, so the cases still run on a ROCm host with no
environment set.

Every assertion, expectation, comment and geometry value in the four case bodies
is unchanged.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
The listener's note said a skipped case's "environment is not set". That was true
for the captured-oracle cases, which skip on an unset environment variable, and
false for the four synthesized-data cases added earlier, which skip because no
ROCm backend is registered on a CPU-only lane. `doctest::TestCaseData` carries no
reason, so this listener cannot tell the two apart and must not claim one cause.

Reword the note, and the comment above it, to name the fact both kinds share: a
prerequisite the case needs is absent. This changes a message and a comment only.
The `m_skip` guard, the `g_env_cases_skipped` flag, the listener registration
name, the exit-77 handler, the predicate and the four decorations are unchanged.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:deepseek-v4-flash [DeepSeek Harness]
@VikashLoomba

Copy link
Copy Markdown
Contributor Author

CI triage after the repair push — head 8af736e67.

Two branch-local failures were repaired on this branch:

  • pr-size rejected six evidence paths that sat one directory too deep (docs/bench-evidence/rocm-residual-norm/stream-fixture-repair/*.log|.py). Commit 506832d32 moves all 17 files of that directory up into docs/bench-evidence/rocm-residual-norm/; every move is a byte-identical rename, the only other change is the single reference in .agents/specs/rocm-residual-norm.md. The checker now reports OK: every explicit path class is within its review budget. (rc=0), and all 69 changed paths classify under origin/main's own copy of the checker as well.
  • build-test-cpu failed on this branch's own new case: test_rocm_moe_bf16 fataled at :2076, :2087, :2098, :2115 with REQUIRE( vt::TryGetBackend(vt::DeviceType::kROCM) != nullptr ) on the CPU-only lane, while the file's four environment-gated cases reported SKIPPED. Commit c7b0ebf21 gives those four cases the file's own doctest::skip(...) decoration and a REQUIRE_MESSAGE backstop, so a CPU build now reports CTest Skipped (exit 77) and a ROCm host still runs them — verified by a local CPU build (Test #111 ... ***Skipped, the same number CI uses) and by a HIP run where the four cases execute, 33/33 assertions, with only the pre-existing env-gated cases skipping. The four case bodies are unchanged. 8af736e67 then makes the shared skip note name an absent prerequisite rather than an unset environment, which is the fact both gating kinds share.

The remaining red checks are red in the same way on origin/main itself, in main's own run 34538060205:

Check Failure Main evidence
sanitize-cpu (address,undefined) test_dots3_note_vision, test_dots3_note_audio, test_openai_api_server_dots3_mm_forward job 103074166410 fails the same three tests
sanitize-cpu (thread) test_gemma4_moe_fp8_native_arm_guard job 103074166517 fails the same test (numbered 162 there, 163 here)
windows-msvc-cpu test_openai_api_server.exe exits -1073740791 (0xC0000409) at localiser cases 58/59/61 job 103074166371 is the same crash and exit code
windows-msvc-vulkan the same crash and exit code job 103074166311

pr-size passes. mergeStateStatus is UNSTABLE because of the four pre-existing checks above.

localai-org-maint-bot added a commit that referenced this pull request Sep 11, 2026
…radicts the canonical record (#3144)

`GATE-PR-BODY-TRAILERS`. Refs ISSUE-LOCAL-01M298Q7PPM145DQA5C30GRYJA.

## What changed

`squash_merge_commit_message = PR_BODY` makes a pull request body the
landed
commit message, so `Closes #N` in a body closes the GitHub mirror the
moment the
squash lands. Local files under `.agents/issues` are the issue
authority, and
nothing checked that the branch also carries `ISSUE-GH-N.md` reading
`State: CLOSED`. The two authorities diverge at the merge, and the
authoritative
half is the one left reading OPEN.

`scripts/agent-pr-body.py` now refuses such a body. It already fetches
the live
body and already runs before a merge, which is exactly when this
matters.

## Why here, and why only closing keywords

`check-agent-record.py` cannot catch it: it is offline and never sees a
body,
and `check_canonical_issue_references` silently `continue`s on an
unresolvable
bare `#N`. Making that strict was the obvious repair and is the wrong
one.
Measured over four open pull requests, bare citations lacking a local
record run
6 of 7, 4 of 4, 8 of 8 and 10 of 11 — so failing them would fire the
gate on
ordinary work, which AGENTS.md names as the defect rather than the
discipline.
The rule therefore keys on the closing keyword alone.

## Evidence

Red before: a body closing an issue with no record exits 1 naming the
missing
path. Green after: a bare `Refs #N` still exits 0. Both are committed as
cases
in `tests/scripts/test_agent_pr_body.py` (suite: 21 passed, 16
subtests), and
the spec's `## Tests` table is updated because that table is compared
with the
loaded suite rather than sampled.

Found live on four pull requests in one review pass: #3101 carried a
closing keyword for #3098, whose
record read OPEN; #3095 did the same for #3092 while that row's own gate
table reads
FAILING; #3096 carries five closing keywords against four OPEN records
and one
issue with no local record at all; #3097 likewise for #3093, against a
spec saying the
row is not ready to land. On its first live run against #3095 it also
caught a closing keyword for #3106, which the human review pass and I
had both missed — that record
reads OPEN and its Resolution says it stays open until the reviewed
production
repair lands.

The spec's `Kind` line is amended in the same change, because it claimed
this
row edits no checker rule and now it does.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5-1m [claude-code]

Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
@localai-org-maint-bot
localai-org-maint-bot merged commit 0b5783e into mudler:main Sep 12, 2026
27 of 31 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment