feat: unify TFILLPAD modes and FP operand forms

[HecreReed]

Summary

• Collapse pto.tfillpad_inplace and pto.tfillpad_expand into pto.tfillpad; PTOAS infers Normal, InPlace, or Expand after physical-shape analysis and memory planning. The textual op does not accept a mode attribute.
• Make the VPTO Normal fallback alias-safe by copying the complete valid source region before padding, matching the EmitC/PTO-ISA ordering when equality cannot be proven.
• Remove the textual IR ops pto.textract_fp, pto.tinsert_fp, pto.tmov.fp, and pto.tstore_fp; represent FP variants through optional fp operands on the unified ops.
• Keep one-to-one PTO-ISA lowering: no-mode FP forms select TEXTRACT_FP, TINSERT_FP, TMOV_FP, and TSTORE_FP, while explicit A5 acc-to-vec modes select the unified PTO-ISA overloads.
• Preserve PTO-BC v0 cross-build readability: historical FP opcodes and tfillpad opcodes 0x1024 / 0x1025 remain wire-only aliases that decode into the unified IR ops.
• Classify ACC-to-VEC pto.textract as PIPE_FIX, with a lit test pinning the required PIPE_FIX -> PIPE_V synchronization.
• Migrate Python bindings, PTODSL, TileLib, tests, samples, documentation, and PTOBC metadata; add the PTO-BC suite to CI.
• Update all seven managed PTO-ISA pin locations to 27386d906e8fdcbd93aec84197939bc0b2c6caea; extend update_pto_isa_pin.py and its workflow to cover ci_sim.yml and the compile-only guide.

Compatibility

• Existing pto.tfillpad source remains valid, but explicit mode attributes and the removed mode-specific textual op names must be dropped; lowering mode is compiler-inferred.
• Removed FP textual op names require migration to the corresponding unified op with an optional fp operand.
• PTO-BC v0 remains a cross-build exchange format. Legacy FP and tfillpad bytecode stays readable and is normalized to unified operations during decode.

Validation

• PTOASPythonPackage and ptobc builds passed.
• Full lit suite after merging current main: 1690 passed, 1 unsupported.
• PTO-BC suite: 21/21 passed, including legacy tfillpad opcode decode and unified FP operand round trips.
• Fillpad samples passed for Normal, Expand, and InPlace; invalid cases produced the expected XFAIL results.
• Focused tfillpad/ACC-to-VEC sync lit tests: 3/3 passed.
• git diff --check, shell syntax, ShellCheck, and PTO-BC schema checks passed.

[chatgpt-codex-connector]

You have reached your Codex usage limits for code reviews. You can see your limits in the Codex usage dashboard.
To continue using code reviews, you can upgrade your account or add credits to your account and enable them for code reviews in your settings.

[zhangstevenunity]

Deep correctness review. I built PR head (ef95f38) against LLVM 21 and ran the full lit suite: 1635/1635 pass (1 unsupported). I also ran the ptobc ctest scripts and the Fillpad samples locally, and checked the emitted code against the pinned pto-isa (e948507 = GitCode 27386d9).

The refactor itself is careful and the ODS/verifier/lowering/bytecode/docs layers are kept in sync. Four things below; the first is a definite blocker, and the second is a semantics question I would like answered before merge.

---

P1 - the new PTO-BC gate is dead on arrival

tools/ptobc/tests/fp_operand_forms_v0_encode.sh:49

[[ $(grep -Fc " fp " "${ROUNDTRIP}") -eq 3 ]]

Only textract, tinsert and tstore print the fp keyword. pto.tmov keeps its pre-existing positional spelling ((,$fp^: qualified(type($fp)))?), so it round-trips as , %1 : !pto.tile_buf<scaling, ...> with no fp token. The count is therefore always 2 and the assertion can never hold.

Reproduced on a clean build of PR head (inputs LF-normalised to rule out CRLF):

+ /home/rdp/.../ptobc encode  ... (ok)
+ /home/rdp/.../ptobc decode  ... (ok)
+ grep -F 'pto.textract ins('  ... (ok)
+ grep -F 'pto.tinsert ins('   ... (ok)
+ grep -F 'pto.tmov ins('      ... (ok)
++ grep -Fc ' fp ' .../fp_operand_forms_v0_roundtrip.roundtrip.pto
+ [[ 2 -eq 3 ]]
rc=1

The round-trip itself is correct - all three fp operands survive - only the assertion is wrong. ctest -R ptobc_fp_operand_forms_v0_encode fails. Note the PR body claims "PTOBC suite: 20/20 passed"; that does not hold for this new test. It also will not show up as CI red, because ci.yml only runs ctest ... -L PTODSL and never runs the ptobc *_v0_encode tests - so please do not read a green CI as validation here.

Fix: either drop the count assertion, or count the two spellings separately (e.g. 2 x fp plus one pto.tmov ins(%0 : ..., %1 :).

---

P2 - tfillpad in-place now has no escape hatch, and the VPTO normal template is not alias-safe

pto.tfillpad_inplace is removed and pto.tfillpad explicitly rejects a mode attribute, so in-place lowering is reachable only when inferTFillPadLoweringKindAfterMemoryPlanning can prove address equality. When it cannot, it silently returns Normal (lib/PTO/Transforms/Utils.cpp:119) - even though the op's own diagnostic says the inference "require[s] ... resolved planned addresses". There is no path left for a frontend to assert an aliasing relationship the compiler cannot see.

That fallback is safe on the EmitC path: pinned pto-isa TFILLPAD_GENERIC_IMPL does CopyValidElementsVec first and pads afterwards, so Normal on exactly-aliased storage degenerates to an identity copy plus the same pad writes.

It is not safe on the VPTO/TileLib path. lib/TileOps/a5/_fillpad.py interleaves the phases the other way for the normal/expand kinds:

_copy_region(src, dst, src_valid_rows, 0, aligned_cols)                # cols [0, align_down(v_col))
_fill(dst, 0, fill_row_stop, aligned_cols, dst_valid_cols)             # pads cols [align_down(v_col), dst.v_col)
_copy_region(src, dst, src_valid_rows, aligned_cols, src_valid_cols)   # copies cols [align_down(v_col), v_col)

The pad store covers the source tail before the tail is copied. Emitted VPTO IR for a normal tfillpad with src.valid = 8x7, dst.valid = 16x16 f32 (aligned_cols == 0, lanes 64) at PR head:

%cst = arith.constant 3.40282347E+38 : f32
%1   = pto.vdup %cst, %mask
scf.for %arg0 = %c0_i16 to %c8_i16 ... { pto.vsts %1, %5[%c0], %mask }        <-- pad rows 0..8, cols 0..16
scf.for %arg0 = %c0_i16 to %c8_i16 ... { pto.vlds ...; pto.vsts %result, ... } <-- then copy cols 0..7
scf.for %arg0 = %c8_i16 to %c16_i16 ... { pto.vsts %1, ... }

If src and dst are the same UB storage, the first loop overwrites the whole valid region with FLT_MAX before the second loop reads it, so the "copy" restores pad values. Before this PR that combination was unreachable because in-place had its own op and its own copy=False template.

Two related data points from the build:

• A private helper taking two !pto.tile_buf arguments (which a caller may alias) compiles to plain TFILLPAD(v2, v1) with no diagnostic. That spelling used to be pto.tfillpad_inplace.
• Conversely, the same-SSA spelling tfillpad ins(%t) outs(%t) now correctly infers InPlace, which fixes a pre-existing hazard (tfillpad_same_ssa_lowers_to_tfillpad.pto used to expect plain TFILLPAD(), so the change is a net improvement in that direction.

VPTO currently rejects tile-typed block arguments ("must be defined by ... pto.alloc_tile or pto.treshape"), which narrows the exposure to level3 allocs whose addr operands are equal at runtime but not provably so. Still, silently picking a non-alias-safe expansion is a bad failure mode for something that used to be explicit. Any one of these would close it:

1. Make the normal template alias-safe by padding only [src.v_col, dst.v_col) (the masked pxor idiom _fill_inplace already uses) instead of [align_down(src.v_col), dst.v_col), so Normal matches the pto-isa reference ordering.
2. Keep an explicit optional override attribute so a frontend can assert in-place when the compiler cannot prove it.
3. Make the inference hard-fail for a VEC tfillpad whose operand addresses are unresolved, matching what the diagnostic already claims.

---

P2 - v0 opcodes 0x1024 / 0x1025 deleted instead of aliased

tools/ptobc/generated/ptobc_opcodes_v0.h: every other removed op keeps its wire opcode as a legacy alias (0x1022 -> pto.textract, 0x102E -> pto.tinsert, 0x1039 -> pto.tmov, 0x1066 -> pto.tstore), but 0x1024 (pto.tfillpad_expand) and 0x1025 (pto.tfillpad_inplace) are dropped from kOpTable entirely. lookupByOpcode then returns null and buildKnownOpFromReader throws "missing opcode schema", so any previously emitted v0 file containing them stops decoding.

Both take the same two operands as pto.tfillpad, so they alias cleanly with no operand reordering - exactly the treatment the four FP opcodes got. As written this contradicts the compatibility contract this PR adds in tools/ptobc/MAINTENANCE.md ("PTO-BC v0 files are expected to remain readable across PTOAS builds"), and v0_fp_schema_compatibility_check.py does not cover the tfillpad opcodes so nothing catches it.

---

P2 - unannounced sync change: acc->vec pto.textract moves to PIPE_FIX

include/PTO/IR/PTOOps.td:4591 extends the PIPE_FIX arm from ACC->MAT to ACC->MAT | ACC->VEC. That is not scoped to the fp form - it changes the pipe of every plain pto.textract with loc=acc src and loc=vec dst.

A/B on the same input (acc->vec textract followed by pto.tabs on the vec tile, --enable-insert-sync --emit-pto-ir):

• three pre-PR builds: no flag at all (both ops classified PIPE_V)
• PR head: pto.set_flag[<PIPE_FIX>, <PIPE_V>, <EVENT_ID0>] / pto.wait_flag[...] inserted

I think the new value is the correct one: in the pinned pto-isa, an acc-source TEXTRACT with a vec destination lowers to pto_copy_matrix_cc_to_ub (include/pto/npu/a5/TExtract.hpp), the same L0C fixpipe family as the acc->mat pto_copy_matrix_cc_to_cbuf path, and TMovOp::getPipe already returns PIPE_FIX for ACC->VEC. So this most likely closes a latent missing-sync hole. (Correction to the original text of this review: I first cited a TEXTRACT_A2V opcode name - that name does not exist in pto-isa or in PTOAS; the TEXTRACT_A2M / TEXTRACT_V2M / TEXTRACT_M2LR spellings are PTOAS-side shorthand in the PTOOps.td comment, not pto-isa symbols. The conclusion is unchanged.) But it is a cross-cutting codegen change that is unrelated to the stated scope, is not mentioned in the PR body, has no test, and adds one more in-flight event-id pair per acc->vec textract (which matters near the 8-pair-per-pipe-pair limit before the pipe_barrier(PIPE_ALL) fallback). Please call it out and add a lit test pinning the flag pair.

---

Nits

• test/samples/runop.sh:991: the fillpad check was relaxed from TFILLPAD( + not-TFILLPAD_EXPAND( to a bare grep -Fq "TFILLPAD", which now also matches TFILLPAD<...InPlace> and TFILLPAD<...Expand>. I ran the sample - it still emits plain TFILLPAD(v12, v9) - so the strict form still passes and the relaxation just loses coverage.
• test/lit/pto/tfillpad_non_normal_mat_invalid.pto is a new file with no PR386 OAT.3 license header (.claude/CLAUDE.md requires it on new files, and on touched files that lack it).
• The PR description is stale: it says tfillpad is collapsed "with a mode attribute defaulting to normal", but the final state rejects mode outright and infers the lowering. Same for "no-mode FP forms select TEXTRACT_FP ..." which is accurate, but the tfillpad paragraph is not.
• ExpandTileOp::appendOpContextAttrs adds lowering_kind for tfillpad but InsertTemplateAttributes::appendOpContextAttrs does not, so the legality query and the instantiation query use different context attrs. Harmless today (one tfillpad template, and its constraints do not read the attr), but the two functions are otherwise kept mirrored - worth a comment if intentional.

---

Verified sound

• Unified TEXTRACT/TINSERT fp lowering now matches the pinned pto-isa signatures (TEXTRACT_FP<Dst,Src,Fp,relu>, TINSERT_FP<Dst,Src,Fp,relu>, TSTORE_FP<Tile,Global,Fp,atomic,relu>). Note this also repairs tinsert_a5_extended_modes.pto, which previously asserted TINSERT<Dst,Src,Fp>(...) - a spelling with no matching overload in pto-isa, since the AccToVecMode mode parameter has no default.
• Operand-order change on TExtractOp (dst moved ahead of the optional operands) keeps the v0 generic payload byte-identical for the 4-operand form, and the new omitsDerivedOperandSegmentsInV0 strip keeps operandSegmentSizes out of the wire dict now that the trait was added.
• The decoder now initialises the operandSegmentSizes property directly, which fixes a latent decode bug for pto.tinsert / pto.tmov (their v0 records carried the attribute in the generic dict, which OperationState never routed into properties).
• relu_pre_mode now reaches the A5 tinsert TileLib templates via both appendOpContextAttrs implementations; previously it was silently dropped on the VPTO path.
• TFillPadOp padValue is MAT-only per the verifier, and MAT can only infer Normal, so the else if in PTOFillPadToEmitC cannot drop a non-normal mode token.
• verifyTFillPadLike checks rank 2 before zip_equal, so no assert on rank mismatch.
• Fillpad samples end to end: fillpad.py -> TFILLPAD(v12, v9), fillpad_inplace.py -> TFILLPAD<pto::TFillPadMode::InPlace>(v11, v11), fillpad_expand.py -> TFILLPAD<pto::TFillPadMode::Expand>(v15, v12).

[zhangstevenunity]

This assertion can never hold, so ctest -R ptobc_fp_operand_forms_v0_encode always fails.

pto.textract and pto.tinsert print fp %x : <type> (keyword form), but pto.tmov keeps the older positional spelling from its assembly format:

`ins` `(` $src `:` qualified(type($src))
(`,` $fp^ `:` qualified(type($fp)))?

so the decoded tmov line is pto.tmov ins(%0 : !pto.tile_buf<acc, ...>, %1 : !pto.tile_buf<scaling, ...>) outs(...) with no fp token. grep -Fc " fp " therefore returns 2.

Reproduced on a clean build of PR head with LF-normalised inputs:

++ grep -Fc ' fp ' .../fp_operand_forms_v0_roundtrip.roundtrip.pto
+ [[ 2 -eq 3 ]]
rc=1

The round-trip itself is fine - all three fp operands survive and ptoas --emit-pto-ir on the result succeeds - only this count is wrong. Suggest -eq 2 plus a separate grep -F "pto.tmov ins(" | grep -F ", %" style check for the positional form, so the tmov fp operand is still asserted.

[zhangstevenunity]

This is the silent fallback. When the addresses cannot be compared, the op takes the Normal expansion with no diagnostic, even though PTOFillPadToEmitC / ExpandTileOp describe the inference as requiring "resolved planned addresses".

On the EmitC path that is benign: pinned pto-isa TFILLPAD_GENERIC_IMPL copies the valid region first and pads afterwards, so Normal on exactly-aliased storage is an identity copy plus the same pad writes.

On the VPTO path it is not. lib/TileOps/a5/_fillpad.py pads [align_down(src.v_col), dst.v_col) before copying the source tail [align_down(src.v_col), src.v_col). Emitted VPTO IR at PR head for src.valid = 8x7, dst.valid = 16x16 f32:

%cst = arith.constant 3.40282347E+38 : f32
%1   = pto.vdup %cst, %mask
scf.for %arg0 = %c0_i16 to %c8_i16 ... { pto.vsts %1, %5[%c0], %mask }         // pad rows 0..8, cols 0..16
scf.for %arg0 = %c0_i16 to %c8_i16 ... { pto.vlds ... ; pto.vsts %result, ... } // then copy cols 0..7

With aliased storage the first loop destroys the data the second loop reads. Since pto.tfillpad_inplace is gone and mode is rejected, there is no way to tell the compiler about an alias it cannot prove. Options: make the normal template pad only [src.v_col, dst.v_col) using the masked pxor idiom _fill_inplace already uses; keep an explicit override attribute; or hard-fail here for VEC operands with unresolved addresses.

[zhangstevenunity]

0x1022 / 0x102E / 0x1039 / 0x1066 are correctly retained as legacy wire aliases here, but 0x1024 (pto.tfillpad_expand) and 0x1025 (pto.tfillpad_inplace) were removed from kOpTable entirely rather than aliased to pto.tfillpad.

lookupByOpcode(0x1024) now returns null and buildKnownOpFromReader throws "missing opcode schema", so any previously emitted v0 file containing those opcodes stops decoding. Both ops had the same two operands (src, dst) as pto.tfillpad, so they alias with no reordering and no normalizeLegacyFpOperandOrder entry needed.

This contradicts the contract added in tools/ptobc/MAINTENANCE.md in this same PR, and v0_fp_schema_compatibility_check.py does not cover the tfillpad opcodes, so nothing catches the regression.

[zhangstevenunity]

This widens PIPE_FIX from ACC->MAT to ACC->MAT | ACC->VEC for all pto.textract, not just the new fp form, so it changes sync insertion for existing non-fp A5 acc->vec extracts.

A/B on the same input (acc->vec textract feeding a pto.tabs on the vec tile, --enable-insert-sync --emit-pto-ir):

• three pre-PR builds: no flag emitted (both ops were PIPE_V)
• PR head: pto.set_flag[<PIPE_FIX>, <PIPE_V>, <EVENT_ID0>] / pto.wait_flag[<PIPE_FIX>, <PIPE_V>, <EVENT_ID0>]

I believe the new classification is the correct one: in the pinned pto-isa, an acc-source TEXTRACT with a vec destination lowers to pto_copy_matrix_cc_to_ub (include/pto/npu/a5/TExtract.hpp), the same L0C fixpipe family as the acc->mat pto_copy_matrix_cc_to_cbuf path used two branches above it, and TMovOp::getPipe already returns PIPE_FIX for ACC->VEC. So this probably closes a latent missing-sync hole rather than creating one. (Correction to my original text: I first cited a TEXTRACT_A2V opcode name - that name exists neither in pto-isa nor in PTOAS; TEXTRACT_A2M / TEXTRACT_V2M / TEXTRACT_M2LR are PTOAS-side shorthand in the comment right above this hunk, not pto-isa symbols. The conclusion is unchanged.) But it is out of the PR's stated scope, is not mentioned in the description, has no test, and consumes one more in-flight event-id pair per acc->vec textract (relevant near the 8-pair limit before the pipe_barrier(PIPE_ALL) fallback). Please mention it in the description and add a lit test pinning the flag pair.

[zhangstevenunity]

This check no longer distinguishes the inferred kinds: bare grep -Fq "TFILLPAD" also matches TFILLPAD<pto::TFillPadMode::InPlace> and TFILLPAD<pto::TFillPadMode::Expand>, so the fillpad sample would pass even if the compiler picked the wrong mode - which is exactly the property this PR makes inference-dependent and therefore most worth pinning.

I ran the sample against a PR-head build: fillpad.py still emits TFILLPAD(v12, v9), so keeping the original strict pair (grep -Fq "TFILLPAD(" plus a negative grep -Fq "TFillPadMode::") still passes and preserves the coverage.

[zhangstevenunity]

New file is missing the PR386 OAT.3 license header. .claude/CLAUDE.md requires it on all new source/script files (the sibling new tests tfillpad_plan_memory_inference.pto and expand_tile_op_tilelang_tinsert_fp_acc2mat.pto do carry it).