fix(kb): address one kernel's records by one address, however its source and its environment spell it

[chennyiiis]

The defect

A forge-loop record is addressed by the operator as the campaign declared it, and the campaign declares the entry point as the source it read spells it. One tree spells a kernel both ways:

where	spelling	normalized to	page
the header that declares the kernel	KdaPackedDecodeKernel	kdapackeddecodekernel	A
the module that binds it	kda_packed_decode_kernel	kda_packed_decode	B

normalize_operator_name lowercased, replaced ., collapsed separators and stripped a trailing _kernel — but never split camel case. So the two spellings addressed two pages, and the second addresses one nothing ever wrote to.

The read is exact, so the miss surfaced as no_prior_record: indistinguishable from an operator nobody has ever ported. A validated 1.44x port stayed on disk, still applying cleanly, while the campaign that should have started from it re-derived the kernel from scratch.

Why this is the common case, not a rare one

The task contract asks for the camel-cased, namespace-qualified spelling:

Give operator_name the entry point as the source writes it, keeping camel case and any namespace prefix — aiter::fusedAddRmsNorm, not fused_add_rms_norm. Upstream pull-request search splits that spelling into terms, and a name normalized before it arrives has no boundaries left to split on.

That instruction is right about search and was silently wrong about addressing. Every kernel whose Python binding is snake-cased has both spellings within reach of the agent that reads it.

The change

normalize_operator_name splits camel-case boundaries before lowercasing.

An acronym run stays one word. Splitting on every case change would cut MoE into mo_e and QKV into q_k_v — inventing a difference between spellings of one kernel instead of removing one. The tail of a capitalized run only starts a new word when a lowercase letter follows it, and a one-character piece is given back to the word it was cut from only when that word is itself a fragment of at most two letters — which is what an acronym written with a lowercase letter inside looks like once the boundaries have cut it. After a whole word a capital is a word of its own.

KdaPackedDecodeKernel    -> kda_packed_decode
kda_packed_decode_kernel -> kda_packed_decode
aiter::fusedAddRmsNorm   -> fused_add_rms_norm
FusedMoE                 -> fused_moe          (Mo|E rejoined: Mo is a 2-letter fragment)
ChunkFwdKernelO          -> chunk_fwd_kernel_o (Kernel|O kept: Kernel is a word)
paged_attention_ll4mi_QKV_mfma16_kernel -> paged_attention_ll4mi_qkv_mfma16
unified_attention_with_output           -> unified_attention_with_output   (unchanged)

No other identity dimension moves, and the write and read sides share this function, so they cannot disagree.

Validated against every real kernel name reachable from here

The hand-picked examples above are not evidence that the rule is right, so it was property-tested over three corpora rather than over examples: the 107 distinct page names in the production store, the 201 implementation_symbols recorded across its 144 pages, and the 632 kernel symbols declared in this tree (@triton.jit defs, __global__ entry points, *_kernel defs, C++ camel entry points) — 926 distinct names. For every snake-cased name in that union, both its Pascal and its camel spelling were generated and required to normalize to the same page.

This found a real defect in the first version of the rule. Giving a trailing capital back unconditionally kept MoE whole but also welded the dimension letter these kernels end in onto the word before it: ChunkFwdKernelO became chunk_fwd_kernelo while the source declaring it spells it chunk_fwd_kernel_o. 82 of 769 kernel names, 230 of 1538 generated cases. The two-letter fragment bound above is the fix; the two classes are separable because an acronym cut in half leaves a fragment and a word does not.

property	result
idempotence — read side and write side normalize independently, so a drifting rule splits a page	pass, 926/926
a name already written as words is a fixed point	pass
all 107 live page names are their own normal form	yes
camel agreement — Pascal and camel spellings reach the snake name's page	330 / 359 = 91.9%
injectivity over live page names	no merges beyond the 2 pre-existing _kernel pairs below

The 29 disagreements are all one ambiguity and it is not decidable: a digit beside a letter. mxfp4_moe_2stage says 2stage is one word, Mxfp4Moe2Stage says it is two, and gemm_a8w8_blockscale says a8w8 is one — a rule that split digits from letters would re-address that page to gemm_a_8_w_8, which is worse than the gap. 28 of the 29 are of this shape; the last, recompute_w_u_fwd_kernel, spells two adjacent single letters that Pascal writes WU, indistinguishable from an acronym. 9 live page names fall in this set, and each costs a page only if a later campaign supplies the camel spelling of one. test_a_digit_run_keeps_whatever_boundary_the_spelling_gave_it pins the gap so a change in it cannot pass unnoticed.

Measured against the production store

Read-only scan of producer=forge-loop: 144 pages, 106 distinct names. Every page's champion record was fetched and compared; nothing was written.

No stored page is rewritten, and there is no migration. What changes is the page a later campaign addresses. Two of the 106 names land elsewhere:

name	goes to	cost
kdapackeddecodekernel	kda_packed_decode — already exists	none, the two are one kernel
topkgatingsoftmax	topk_gating_softmax — nothing there yet	that page's history stops being read
the other 104	unchanged	—

The merge is the fork this PR reports, and the store proves the two pages are the same kernel rather than two that happen to share a name:

field	kda_packed_decode	kdapackeddecodekernel
implementation_signature	527b02951b281dac…	527b02951b281dac…
implementation_symbols	["kda_packed_decode_kernel"]	["kda_packed_decode_kernel"]
source_paths	sglang/kernels/…/kda_packed_decode.{cuh,py}	identical
champion speedup	1.4439x	1.4439x

Two campaigns, one kernel, the same result derived twice.

The orphan is conditional rather than certain. topkgatingsoftmax's record carries the source symbol topkGatingSoftmax, which the old rule squashes to the current page name and the new rule splits — so the page moves only if a later campaign again names the operator from that symbol. Its history is one record.

Eight further names read like squashed camel (hipbsolgemm, hipgraphlaunch_*, gemm_a8w8_*_bpreshuffle, rocm_unquantized_gemm) but no symbol in their records shows a camel spelling, and the new rule is the identity on a lowercase input. They move only if a campaign supplies the camel spelling, and no page exists at the name they would move to, so nothing merges and nothing is overwritten.

No two stored pages can collide. Every stored name is already lowercase and therefore its own normal form, so the new rule maps the 106 names injectively.

What the same scan ruled out

The scan found 18 groups where one implementation_signature spans several names. Only the kda pair is a rename. The other 17 are distinct operators that share a source file and therefore share a file-level signature — six different operators in vllm/_aiter_ops.py alone, flydsl_moe_stage1 beside stage2, ck_moe_stage1 beside stage2, mxfp4_moe_2stage_t16/t32/t64/t128. That ruled out the two designs this PR previously carried and one considered after it:

• Probing other names on a signature match (the earlier revision of this PR, reverted here): the signature is file-level, so a new operator in a shared file would warm-start from a sibling's patch at match_mode='exact' and then file its result on the sibling's page. @xiaofei-zheng raised exactly this; the store says it is the majority structure, not a corner case.
• Deriving the name from the file's kernel symbols: the symbols are GPU kernels, while the operator a campaign targets is usually the host-side API. Re-deriving the 144 pages that way renames unified_attention_with_output to apply_softcap, collapses mxfp4_moe_2stage_t16/t32/t64/t128 into one name and custom_all_reduce_tp2/tp4 into another.

operator_name names the right thing. Only its normalization was lossy.

The same defect in the version dimension

The same scan found the address splitting a second way, for the same reason: a dimension was whatever string the campaign happened to hold rather than the fact it names.

Three code paths answer "which framework version?" and they answer in three different words.

who answers	what it says	what it means
framework_version() via importlib.metadata	0.24.0+rocm723, 0.5.15.post1.dev20260724+g3d91a569ce	the build the wheel was compiled as
a campaign reading the package tag or its image	v0.24.0, v0.5.15.post1-rocm720-mi35x-20260724	the same release, written the way a tag writes it
a path that observed nothing	none, unspecified, unknown, ""	one absence, under four names

A port is written against a release. It does not depend on the machine that compiled the framework, nor on the tag convention of whoever wrote the version down, nor on which of three code paths failed to find a version. Each of those spellings opened its own page.

What it cost, measured

Read-only scan of the same 146 producer=forge-loop pages, replaying the canonicalizer this PR ships and fetching each page's champion record:

146 pages → 135. Eleven merge, nine of them provably one kernel by identical implementation_signature. Every one of those nine held a different champion on each page:

kernel	the two spellings	champion on each
sparse_attn_prefill_ragged	v0.24.0 / 0.24.0+rocm723	13.7293x / 8.2320x
mxfp8_grouped_gemm	0.5.15.post1.dev20260724+g3d91a569ce / v0.5.15.post1-rocm720-mi35x-20260724	4.1771x / 3.3885x
mhc_fused_post_pre	0.24.0+rocm723 / v0.24.0	3.3633x / 3.2760x
unified_attention_with_output	unspecified / unknown	2.3303x / 2.0238x
paged_attention_rocm	unspecified / unknown	1.5456x / 1.4935x
rmsnorm	unknown / unspecified	1.4946x / 1.1582x
fused_moe	0.24.0+rocm723 / v0.24.0	1.4650x / 1.4435x
aiter_mxfp4_moe_2stage	unknown / unspecified	1.3203x / 1.3097x
moe_flydsl_stage2	unspecified / unknown	1.0345x / 1.0283x

Each row is one kernel ported twice, by two campaigns that could not see each other. The remaining two groups share an operator but not a signature, and both elect the faster of the two: dynamic_per_tensor_quant 10.4489x over 2.8170x, unified_attention_with_output@0.24.0 11.5837x over 2.5712x.

The change

canonical_framework_version resolves a version string to the release it names, and every word for not having observed one to a single word.

0.24.0+rocm723                              -> 0.24.0
v0.24.0                                     -> 0.24.0
0.5.15.post1.dev20260724+g3d91a569ce        -> 0.5.15.post1
v0.5.15.post1-rocm720-mi35x-20260724        -> 0.5.15.post1
0.1.dev19253+g5f76ae224.d20260727.rocm723   -> 0.1
none | unspecified | unknown | "" 	        -> unknown
not-a-version                               -> not-a-version   (unchanged)
0.5.15 / 0.5.15.post1 / 1.0.0rc1 / 1.0.0    -> distinct, as they are distinct releases

Parsing is PEP 440 (packaging.version.Version, already a hard dependency of warmstart_identity). An image tag joins the build to the release with the same character a pre-release uses, so it is only distinguishable by failing to parse whole — hence the one retry on the first -. A string naming no release is returned exactly as written, so the read side and the write side, which canonicalize independently, cannot turn it into some other page's name.

It is applied at the three points that write the dimension and once on the read side:

	where	what it fixes
write	rewrite_by_flydsl/identity.py: framework_version()	the build suffix, and the two words this function itself used for absence — NO_FRAMEWORK_VERSION = "none" and UNKNOWN_VERSION = "unspecified", both deleted
write	kernel_identity.py: KernelRecipeIdentity.from_mapping	the tag spelling an agent declares in task.json
write	task_publisher.py: _normalize_agent_task_payload	makes the published task.json state the address the controller went on to use, beside the kernel_name it already re-derives
read	warmstart_identity.py: rank_fallback_identities	reaches pages written before this, as the release they name

No stored page is rewritten and there is no migration. The fuzzy tier reads a legacy spelling as the release it names rather than as a neighbouring release, so the histories above stay reachable from the canonical address.

This also un-strands the pages that never had a version. 55 of 146 name no release — 42% — and the fuzzy tier used to end before it started for all of them, because unknown fails to parse. Under one word for absence, 12 of them now reach a donor that nothing could reach before.

What the ranking still refuses is unchanged: how far a known release sits from an unknown one is not a question the strings can answer, so that comparison stays rejected rather than being given an invented distance. test_fuzzy_ranking_rejects_unknown_or_unparseable_dimensions passes untouched.

Where the six dimensions stand after this

#	dimension	distinct values in store	what pins it
1	producer	1	whitelist, plus a hard == forge-loop check in task.py	pinned
2	kernel_name	106	normalize_operator_name, re-derived host-side from operator_name	this PR
3	framework	5	folded to the override the run resolves, then checked against the task's own paths	this PR
4	framework_version	up to 4 spellings of one release	canonical_framework_version	this PR
5	backend	5, all registered	membership in KERNEL_BACKENDS, the set campaign setup resolves against	this PR
6	gpu	1 (mi355x)	uniformly lowercased	see below

The same defect in the backend dimension

identity.backend is one of the five dimensions the agent types into task.json, and it becomes --kernel-backend on the forge-loop command line. campaign_setup.resolve_kernel_backend_override then resolves it against KERNEL_BACKENDS and, for a name that is not there, substitutes _FALLBACK_KERNEL_BACKEND = "flydsl", warns on stderr, and continues.

So a draft saying tilelang produces a run that builds with flydsl and files its result on the flydsl page, while the task directory, the --experience-id and the published pointer all say tilelang. One run, two addresses — the shape this PR exists to remove.

The comment standing at that spot in task.py justified not checking, and stated the reason as fact:

forge-loop answers an unregistered one with no layer rather than an error

It does not. It answers with a different layer, silently. The rule that comment defended is what let the substitution through.

This commit checks membership in KERNEL_BACKENDS — the same frozen set campaign_setup resolves against, and the same one opportunity_agent.py:418 already injects into the agent's prompt. The name the agent has to spell is the one it was handed.

registered	assembly, ck, flydsl, triton, gluon, aiter, hip, hipblaslt, fusion
present in the store (146 forge-loop pages)	triton 81, hip 32, flydsl 19, ck 11, aiter 3
stored values this refusal would reject	none — every one is registered

The refusal costs nothing that campaigns write today. It only closes the path by which a name the agent invents becomes a page nothing writes to.

The same defect in the framework dimension

detect_framework(kernel_path, framework_override) infers the owning package from the path — but a non-empty override short-circuits the inference entirely. Since the agent's identity.framework travels to forge-loop as exactly that override, the dimension reaches the knowledge-base page having never been examined by anything. Two consequences:

1. A kernel that belongs to no package is addressed two ways. _NO_FRAMEWORK_SENTINELS = {"standalone", "none", "unknown"} all fold to unknown inside the run, but the draft keeps its own spelling in its directory name and its published pointer. The store is the same gap seen from the far end:

spelling	pages in store	controller fixtures declaring it
standalone	0	4
unknown	1	0

Every campaign that declared standalone filed its result under unknown and pointed at standalone.

2. A framework naming a package the source does not live in is never contradicted. The override suppressed the one check that could have noticed.

This commit folds the override at publication the way the run will, and refuses a framework that none of the task's own paths sit under. The paths were just verified to exist at the base commit, which makes them the only witness available at that point.

Why it refuses rather than re-derives

I first wrote this to derive framework from the paths and overwrite whatever the agent said. Scanning the store killed that design:

verdict over 118 pages that recorded source paths	count
path inference agrees with the stored framework	102
no single owner — declared value kept	13
inference "disagrees"	3

All 3 disagreements are the same path: sglang/aiter/ops/flydsl/kernels/mixed_moe_gemm_2stage.py — under both packages at once. detect_framework answers aiter only because _FRAMEWORKS = ("aiter", "sglang", "vllm") starts there. Re-deriving would have moved those three pages on the strength of a tuple's ordering.

So the check refuses a contradiction and never re-derives. A declared framework appearing anywhere in the paths stands. A path under no known package witnesses nothing and is left alone — which is what keeps the 4 pages stored under torch (a value _FRAMEWORKS does not know) addressable.

Not covered

• gpu is the one dimension of the six with an alias the write side does not fold. warmstart_identity._GPU_ISA already knows mi355x and amd_instinct_mi355x are one card, and mi300x/amd_instinct_mi300x likewise — but only on the read side, for ranking. Nothing folds them when the address is written. All 146 stored pages say mi355x, so there is no split to repair today; the fix shape is a canonical_gpu() fed by that same alias table, and it is a separate concern from the three here. It is the only one of the six this branch leaves unpinned.
• A second split-history defect, pre-existing and not from camel case. Seven stored names end in _kernel — mxfp8_quant_kernel, mxfp8_grouped_gemm_kernel, fwd_kernel, llgemm1_kernel, dynamic_per_group_scaled_quant_kernel, recompute_w_u_fwd_kernel, tp2_fused_ar_rms_kernel — which normalize_operator_name has stripped since before this PR, so no campaign can address them today either. mxfp8_quant_kernel sits beside mxfp8_quant and mxfp8_grouped_gemm_kernel beside mxfp8_grouped_gemm: the same two-half-histories shape, but it implies a write path that bypasses the normalizer, which is a different defect and not traced here.
• match_mode means "same files", not "same kernel". Given the file-level signature, implementation_match is true for siblings in one file. It labels the reference handed to the model and gates nothing — warm start is decided by whether the patch applies cleanly and benchmarks faster — so this is a naming problem, not a behavioural one, and fixing it needs an operator selector in the signature payload.
• Pages written before implementation_identity was recorded (74 of 144) carry no symbols; nothing here reaches them, and nothing needs to.
• Two campaigns that target the same operator at different levels — one naming the host API, one naming the Triton kernel it calls — still address two pages. That is a contract question (operator_name is not required to appear in target_functions), separate from spelling.

Test

• test_the_same_kernel_spelled_either_way_is_one_operator — the reported pair.
• test_namespaced_camel_case_matches_its_snake_case_spelling — the spelling the contract asks for costs no page.
• test_an_acronym_run_is_one_word — MoE, KVCache, QKV, HGEMV_WFPerRow.
• test_a_capital_after_a_whole_word_is_its_own_word — the counterpart, and the regression for the defect the validation above found.
• test_a_digit_run_keeps_whatever_boundary_the_spelling_gave_it — the known gap, pinned.
• test_names_already_written_as_words_are_left_alone — real page names from the store normalize to themselves.
• test_distinct_operators_sharing_a_prefix_stay_distinct — stage1 ≠ stage2, rmsNorm ≠ addRmsNorm.
• test_either_spelling_of_one_operator_reaches_the_same_record — end to end through resolve_loop_identity: written under one spelling, found under all three.

Version dimension:

• test_the_build_a_release_was_compiled_as_is_not_part_of_the_release — 0.24.0+rocm723, and a dev build with three suffixes.
• test_a_release_reads_the_same_however_it_was_written_down — the tag spelling and the image spelling reach the release.
• test_every_word_for_not_knowing_the_version_is_the_same_word — all four absences.
• test_distinct_releases_stay_distinct — .post1, rc1 and a patch bump are not folded away.
• test_a_version_naming_no_release_is_left_exactly_as_written — idempotence, which the independently-canonicalizing read and write sides depend on.
• test_two_spellings_of_one_release_rank_as_that_release — the fuzzy tier reads a legacy page as the release it names.
• test_runs_that_both_observed_no_version_can_reach_each_other — the 42%, with the known-vs-unknown comparison still refused.

Backend dimension:

• test_an_unregistered_backend_is_refused_rather_than_silently_substituted — replaces test_an_unregistered_backend_is_accepted, which pinned the old behaviour and repeated the false premise in its docstring.
• test_every_backend_the_agent_is_offered_is_one_it_may_declare — every name in KERNEL_BACKENDS survives the check, so the refusal cannot cost an operator a name its prompt offered.

Framework dimension:

• test_a_kernel_belonging_to_no_package_is_published_under_one_word_for_that — standalone is published as unknown, which is where the run files its result.
• test_a_framework_alias_is_published_as_the_package_it_names — aiter_meta → aiter.
• test_a_framework_none_of_the_paths_sit_under_is_refused — vllm declared over a sglang/ path leaves a rejection the agent can read.
• test_a_path_under_two_packages_keeps_whichever_the_task_declared — the real sglang/aiter/ops/flydsl/... case: declared sglang stands.
• test_a_path_under_no_known_package_witnesses_nothing — torch, which no path names, is left alone.

The first commit was verified on crsuse2-m2m-069 (8x MI355X): 4351 passed; the 8 failures in that run are *_gpu.py tests that fail identically on an unpatched tree (No module named 'torch' in the interpreter). That run predates the two-letter-fragment amendment and all three later commits, and the box has been unreachable since — it owes a re-run before merge.

Locally, the two commits added here land entirely inside the controller package, which Windows could not collect at all before: fcntl is absent, so its conftest fails to import and every test in it is unreachable. For this round I stubbed fcntl.flock and the directory fsync out of process — neither stub is in the repository, and neither proves the lock contention or the durability the real calls exist for — which made the package runnable for the first time. Run on this branch and on a detached worktree at the parent commit:

controller package	parent bd77313df	this branch
passed	244	250 (+6, exactly the new tests)
failed	27	27
set of failing test ids	—	identical

Those 27 are the platform's: os.open on a directory and os.fsync on a read-only descriptor are POSIX abilities Windows lacks, symlink creation needs a privilege this account does not hold, and the four test_controller_full_path cases fail on both trees before reaching any assertion of mine.

That last point is the one gap worth naming: test_controller_full_path is where the agent's draft and the published address are both visible, and it hardcoded standalone in the operator id. This branch updates that constant to unknown, and the local run cannot confirm the update — the test fails for its own reasons on both trees. What does confirm it is test_a_kernel_belonging_to_no_package_is_published_under_one_word_for_that, which passes and drives publication over the identical six-tuple.

Outside the controller package, the modules these commits import from were diffed the same way — test_assembly_backend (the only test file outside the package that imports the controller), test_implementation_identity, test_warmstart_identity, test_experience_sink and test_experience_sink_helpers: 106 passed, 2 failed, byte-identical on both trees. The earlier version commit was diffed over the whole suite (1886 → 1893 passed, 223 failed / 116 errors unchanged, identical failing ids).

The whole suite was then diffed the same way. The stubs make it a much larger run than the figures above: the 69 files fcntl used to block are collected, so the total goes from 1893 to 4016.

whole suite, stubbed	parent bd77313df	this branch
passed	4010	4016 (+6, exactly the new tests)
failed	292	292
errors	2	2
skipped	40	40
set of failing test ids	—	identical

The 292 are overwhelmingly the platform's — directory os.open, os.fsync on a read-only descriptor, symlink privilege — and being identical on both trees is what says so. None of this substitutes for the devbox re-run, which is what actually decides these commits.

[xiaofei-zheng]

Reviewed the alias tier end to end against the merge base (9c1d4d757), with a local repro on the in-memory KB store. Three blocking issues.

1. Two operators declared in the same source file collide on the alias tier (correctness, regression vs main)

implementation_signature hashes {source_paths, implementation_symbols} (implementation_identity.py:186-208) — every symbol in the declared files, with no selector for the operator actually under optimization. target_functions is not part of the payload. So two forge-loop tasks scoped to different kernels in one file produce an identical signature, which is exactly the equality aliased_identity accepts (identity_alias.py:96), and alias_kernel_names feeds the sibling symbol in as a probe (identity_alias.py:58).

Repro — one file declaring fwd_grouped_kernel and decode_kernel, two tasks with distinct --operator-name / --target-functions (the controller passes both per task, forge_runner.py:135-154):

A written:      kernel:forge-loop:fwd_grouped:unknown:none:triton:mi300x   speedup=2.0
B read status:  read_reason='alias_hit'  match_tier='alias'
                requested=kernel:forge-loop:decode:...
                selected =kernel:forge-loop:fwd_grouped:...
B match_mode:   ['exact']        <-- auto-apply grade
B written:      kernel:forge-loop:fwd_grouped:...  speedup=5.0, champion=True

Operator decode warm-starts from fwd_grouped's patch at match_mode='exact', then recorded_identity files decode's own result onto fwd_grouped's page and takes the champion slot there. decode's page stays empty forever, so every later decode run repeats the redirect, and fwd_grouped's history is now a mix of two kernels that future fwd_grouped runs read back as exact matches.

Same script on the merge base: read_reason='no_prior_record', and B writes to kernel:forge-loop:decode:.... This is introduced here.

test_a_shared_name_does_not_hand_over_another_operator_s_record does not cover it — it uses a different signature, which is the easy case. The hard case is a shared signature with different operators.

2. A transport error on the run's own page is read as "page is empty", and the write is redirected (silent failure)

KernelRecipeKB.list_candidates catches every exception into self.reason and returns [] (agent_kb.py:148-153). recorded_identity (identity_alias.py:116-118) checks only the truthiness of that list and never inspects kb.reason, so a 5xx/timeout on the run's own page is indistinguishable from a cold page — and the fall-through is alias probing, i.e. the record gets filed under a different operator's name even though its own page has history.

Verified with a store that raises on the own-page candidates() call: the exception is dropped and recorded_identity proceeds past the list_candidates guard with no log, no status field, nothing in the returned write result.

_read_top_solutions_impl is careful to turn kb.reason into read_error on every other read (experience_reader.py:303-308, 330-333, 350-357); the two list_candidates call sites in identity_alias.py (95, 117) are the only ones that don't.

3. The candidates handed back from an alias page are not signature-filtered, contrary to the PR body and CHANGELOG

aliased_identity only proves some record in the alias page's top 16 carries the signature. The reader then re-reads that page with alias_kb.read_top_n(destination, limit=top_k) (experience_reader.py:329) and returns whatever ranks highest by speedup, unfiltered. On a page shared with a genuinely different implementation that happens to normalize to the same name, the top candidates can all be foreign; they are labelled match_mode='reference' correctly, but they are returned under read_reason='alias_hit' and injected as warm-start reference material.

That contradicts:

• PR body: "accepts a record only when the candidate's implementation_signature equals the consumer's";
• CHANGELOG.md: "takes it only when a record there carries the implementation signature in hand" and "match_mode is unaffected: the implementation matched, so an aliased candidate is as applicable as any other".

Either filter the returned candidates by signature, or correct both texts.

---

Checked and found nothing blocking: the alias_hit verdict surviving kb_warmstart's normalization (experience_integration.py:1117-1119), the alias-before-fuzzy ordering, match_tier / requested_canonical_id / selected_canonical_id reporting, the MAX_ALIAS_PROBES bound, the docs update in autonomous-loop.md, and the CHANGELOG entry's presence and placement. No tests were run beyond the two targeted repro scripts above; CI covers the rest.

[xiaofei-zheng]

Blocking

identity_alias.py swallows KB read failures, and both call sites act on the empty result as if it were an answer.

KernelRecipeKB.list_candidates reports transport failures by setting self.reason and returning [] (agent_kb.py:144-165). aliased_identity and recorded_identity never look at kb.reason, so a 5xx/timeout on a single metadata read is indistinguishable from an empty page. The exact-address read two lines away in experience_reader.py:302-307 deliberately does the opposite ("surface it as a read error so a cold start is not mistaken for an empty store"), and so does the alias read_top_n at experience_reader.py:331-334.

Two consequences, both verified against this branch with KernelRecipeKB stubbed to set reason and return []:

1. identity_alias.py:97 (read side) — a failing probe skips that alias. The reader falls through to fuzzy/cold start and reports read_reason='no_prior_record' with read_error='': the exact "a validated solution stops being found without anything reporting a loss" failure this module's docstring is written against.

2. identity_alias.py:124 (write side) — recorded_identity treats a failed probe of the run's own page as "this page is empty" and goes on to resolve an alias. Repro: own page newname holds a record and its probe raises; alias page oldname holds a record with the same implementation_signature → recorded_identity returns oldname. The write lands on the alias page while later reads hit the populated newname page on the exact tier and never run the alias probe, so the improvement this run just made is stranded — the split this PR exists to close, opened by a single transient error.

Suggest propagating kb.reason out of aliased_identity / recorded_identity so the reader can set read_error as it does for the other tiers, and so the write side keeps its own address rather than relocating on an unanswered probe.

Checked, no issue found

• Read/write alias derivation agrees: the sink is handed the pristine implementation_signature_override / implementation_identity_override (experience_integration.py:1442-1464), so implementation_symbols, target_functions and the kernel stem are the same inputs the reader used pre-optimization, and alias_kernel_names sorts each tier, so both sides probe in the same order.
• alias_kernel_names ordering and the MAX_ALIAS_PROBES bound exercised directly (['fwd_grouped','attention','zz'], 16 for 32 symbols).
• read_status['read_reason'] is always populated before experience_integration.py:1119 reads it (_set_read_status sets both keys on every path), and cli.py:1529 uses setdefault, so alias_hit survives to the persisted result.
• Grepped every consumer of the warm-start read_reason: nothing outside experience_reader/experience_integration/cli branches on "hit". rewrite_by_flydsl/kb.py:157,405 is the flydsl reader's own field, untouched by this change.
• Tier ordering is alias-before-fuzzy as described, and aliased_identity returns None on an empty implementation_signature, so a name-only match can never hand a record over.
• CHANGELOG entry and the read_reason list in docs/kernelforge/how-to/autonomous-loop.md are both updated and match the diff.
• No import cycle: experience_sink and experience_reader both import cleanly with identity_alias at module scope.
• Tests could not be run here (Windows host; kernelforge.tracker imports fcntl), so the assertions above come from reading the code plus the stubbed repros.

[chennyiiis]

@xiaofei-zheng thank you — issue 1 decided the design. The alias tier is gone; this PR now does something much smaller. Force-pushed 61a703bd7 → 0c2d261c6, so the line anchors on your comments are stale.

What your repro turned out to be a sample of. I ran a read-only scan of the production store (producer=forge-loop, 144 pages, 106 distinct names) looking for one implementation_signature spanning several kernel_names. 18 groups. One is a rename. The other 17 are your case — distinct operators sharing a source file and therefore a file-level signature:

group	names sharing one signature
vllm/_aiter_ops.py	6 operators (rocm_aiter_gemm_a8w8_blockscale, rocm_aiter_act_mul_and_fp8_group_quant, rocm_aiter_fused_allreduce_rmsnorm{,_decode,_prefill}, rocm_aiter_triton_gemm_a8w8_blockscale)
aiter/fused_moe.py	moe_flydsl_stage1, moe_flydsl_stage2
mxfp4 MoE	mxfp4_moe_2stage_t16 / t32 / t64 / t128
…	ck_moe_stage1/stage2, rmsnorm/add_rmsnorm, mxfp8_linear/mxfp8_quant, …

So the blast radius of signature-equality-as-an-alias-test is the majority structure of the store, not a corner case, and the benefit it was buying is one fork whose two pages hold the same 1.4439 champion — a wasted campaign budget, not a lost speedup. Issues 2 and 3 were real too; I'm not fixing them, because the code they're in is deleted.

What the defect actually was. The one genuine rename is kdapackeddecodekernel vs kda_packed_decode — a camelCase/snake_case pair. normalize_operator_name lowercased, replaced ., collapsed separators and stripped a trailing _kernel, but never split camel case. The task contract asks the agent for the camel spelling (aiter::fusedAddRmsNorm) so upstream PR search has term boundaries; that instruction was silently costing the operator its page. Splitting camel boundaries before lowercasing collapses both spellings onto one address, with no probing and no cross-operator reach.

The third design, also rejected on the data. Before landing this I costed "derive kernel_name from the file's kernel symbols instead of from prose". Re-deriving all 144 pages that way: 21 anchored, 49 drift, 74 have no recorded symbols. The drift is disqualifying — implementation_symbols holds GPU kernels while the operator a campaign targets is usually the host API, so unified_attention_with_output becomes apply_softcap, mxfp4_moe_2stage_t{16,32,64,128} collapse to one name and custom_all_reduce_tp2/tp4 to another. operator_name names the right thing; only its normalization was lossy.

Migration. Of 106 distinct names, 9 contain an unbroken 11+ letter run and 4 of those are camel-squashed (kdapackeddecodekernel, topkgatingsoftmax, hipbsolgemm, hipgraphlaunch_*); the rest (bpreshuffle, unquantized) are lowercase in source and normalize to themselves. kdapackeddecodekernel merges into the existing kda_packed_decode — the reported fork, closed at no cost. The other three strand one record each.

All four CodeQL cyclic-import alerts were on the deleted imports and are gone. Suite on crsuse2-m2m-069: 4351 passed; the 8 *_gpu.py failures reproduce identically on an unpatched tree (No module named 'torch').