Rebase main-llvm19-build onto upstream main and fix LLVM19/VMI test failures - #14
Closed
liuzidi wants to merge 33 commits into
Closed
Rebase main-llvm19-build onto upstream main and fix LLVM19/VMI test failures#14liuzidi wants to merge 33 commits into
liuzidi wants to merge 33 commits into
Conversation
…region Fixes #5. PTOInferVPTOVecScope aborted with "cannot infer resultless pto.vecscope ... escaping value type is '!pto.mask<b32>'" when a body-level pto.plt_b32 mask was reused across a region-bearing Boundary op (e.g. an scf.for, or the then/else of an scf.if, whose bodies hold a DMA). Three defects in the escape-remediation path: 1. computeLogicalScopeAnchors only registered Vector/SafeScalar ops; a region-bearing Boundary op (e.g. an scf.for with a DMA) was flushed and skipped. An external user nested inside it resolved to the op via getAncestorInBlock, but the anchor lookup missed -> remediation aborted. 2. The segment/cache identity was keyed only by the anchor op, so sibling regions of one region-bearing op (e.g. then/else of an scf.if) folded into one segment. A clone placed in one block cannot dominate uses in the sibling block, so a body-level mask reused in both branches still escaped. 3. The cloned producer was inserted before the anchor (the op itself), placing the clone at body level before the region. That clone joined the preceding body-level cluster and its result escaped again into the nested region, looping the 256-iteration repair cap without converging. Fix: - Register region-bearing Boundary ops as their own anchor identity (used only as a grouping/cache key, never as an insertion point). - Key segments and the remat cache by (anchor op, user block) so sibling regions become distinct segments, each with its own clone. - Choose the clone insertion point by user location: flat users (in the block) keep the clone at the segment head (preserves existing behaviour, e.g. a CSE'd mask producer ahead of a vlds that does not use the mask); nested users insert before the earliest such user inside the region, so the cloned value never crosses back out of the scope. Regression tests: - auto_vecscope_infer_escape_across_dma_loop.pto: body-level plt_b32 + DMA inside scf.for + mask reused after the loop. - auto_vecscope_infer_escape_across_dma_if_branches.pto: body-level plt_b32 + DMA inside both then and else of scf.if + mask reused after the if. Both fail without the fix and pass with it. The rmsnorm bf16 [8,4096] repro from the issue now compiles on both level3 paths. G.FMT.11-CPP: the added control bodies use braces.
…s resolve ptoas.mlir The PTODSL TileLib daemon and its inline helper were launched with python -S, which disables site.py. scikit-build's editable install relies on a meta-path finder registered by a site-package .pth file during site.py; -S skips that, so the finder is never installed. Without it the source-tree ptoas package (a regular package with __init__.py) shadows the build-tree ptoas.mlir namespace package on PYTHONPATH, and the daemon fails to import ptoas.mlir.dialects.pto and never opens its socket. Three fixes: 1. TilelangDaemon.cpp: drop the -S flag for the ptodsl.tilelib.serving daemon module; redirect stdin/stdout/stderr to /dev/null so a detached daemon does not hold the parent's pipe open; stop() before start() and reap with waitpid to avoid orphaned daemons blocking ctest until the 1500s timeout. 2. ExpandTileOp.cpp: drop the -S flag from the inline helper launch for the same reason. 3. ptoas.cpp: SerialFrontendPipeLoweringPass used getOps<func::FuncOp>() which only finds direct children of the top-level ModuleOp. When the cube/vector helper functions live inside nested modules, the pass skipped them and pto.tpush/pto.tpop were never lowered. Use walk() to find functions recursively.
TileSpec in metadata.py lacked s_fractal_size and compact_mode fields, but _selection.py passed them and constraints.py read them via getattr, causing TypeError on every TileLib test that constructs a spec. Changes: - metadata.TileSpec: add s_fractal_size (default 512) and compact_mode (default "normal"); mlir_type() threads real values through _tile_buf_type instead of hardcoding fractal_size=512. - _types._normalize_compact_mode: map None -> "Null". - constraints._ConfigView: add pad_value field and propagate it. - _tile_template_tracing.TileSpec: add s_fractal_size and pad_value, keep the two TileSpec definitions consistent. Part of #10.
The PTODSL TileLib daemon -S/lifecycle fix is PR #7 (already in main-llvm19-build). This commit only rewrites the test that was broken by the daemon: test_reuses_imported_specialization_before_materializing monkeypatched _compiler_runtime.materialize in the parent process, but with the daemon materialize runs in the daemon child, so the parent-process counter stays 0. The rewritten test drives the compile through a daemon on a known socket and asserts the observable end-to-end contract: both functions lower and TADD's two 1D blocks plus TADD_2D's 2D block yield 3 pto.vadd. Part of #10.
…andidate_id (#10, groups 3-6) Group 3 — VMI Python API drift: - vci index abort: _derive_vci_result_type passed raw_base.type straight to VMIVRegType.get; an MLIR index (64-bit) produced !pto.vmi.vreg<64xindex> and the C++ verify assertion aborted. Now coerce index to signless i32; any other unsupported width (e.g. i64) raises a clear TypeError rather than silently truncating. - vinterpret_cast: rejected any source/target width mismatch instead of conserving total bits. Now derive target lanes = source_total_bits // target_elem_bits, validate against {1,2,4,8,64,128,256}, reject layout carry across lane change. - _emit_binary: passed _raw(rhs) to the vec-vec builder unconditionally. Now detect vec-scalar and delegate to existing vadds/vmuls/vshrs. Only commutative ops (vadd/vmul/vmax/vmin/vand/ vor/vxor) swap scalar-vector to vec-scalar; non-commutative (vsub/vdiv/vshl/vshr) broadcast the scalar and keep operand order. vshls/vshrs coerce shift count to signless i16. - vci group: _check_vci_group_tiles_phys_vl now validates size % group == 0 before computing group_size, and uses the correct physical VL per element width (f16 = 128 lanes, not 64). Group 4 — bf16x2 + wrappers: - Expose bf16x2 on _VMINamespace (_types already had _vmi_bf16x2). - Add vaddc/vaddcs (add-with-carry) and vsstb (block-strided store). - _wrap_result returns no-result Operations unwrapped. - Deprecate direct vadds/vmuls/... (auto-dispatched by vadd(vec,scalar)). Group 6 — candidate_id + group validation: - _next_vmi_candidate_id: per-op auto-increment (1000, 1001, ...) so VMI candidates don't collide on id=1000. - _emit_reduce: validate group in {1,2,4,8,32,64,128,256} at trace time (reject group=16 early). - vcadd reassoc: UnitAttr presence (not value) signals the user's choice; always materialise reassoc=True when explicitly specified, so floating-point vcadd passes the verifier. Part of #10.
…#10) _check_vci_group_tiles_phys_vl rejected any group_size that did not divide phys_vl, but the backend verifier accepts the reverse too: a group spanning several physical VLs is legal when group_size is a multiple of phys_vl. e.g. pto.vmi.vci(pto.i32(0), size=512, group=2) → group_size=256, phys_vl=64, and 256 % 64 == 0 is accepted by the C++ verifier, but the Python frontend raised ValueError. Change the check to accept either direction: phys_vl % group_size == 0 (group within one VL) or group_size % phys_vl == 0 (group spanning multiple VLs) Add a regression test for the size=512/group=2 case. Part of #10.
…le carriers FoldTileBufIntrinsics::isRuntimeTileCarrier checked only the defining op of the value, so a declare_tile wrapped in an unrealized_conversion_cast (introduced when the materialized template carries a richer tile_buf config, e.g. compact=normal) was not recognized as a runtime carrier. The cast result fell through to resolveTileHandle, which rejects declare_tile as a non-anchor and emitted a spurious 'expected tile_buf to be defined by the active materialized tile-handle bridge' error. Unwrap bridging casts first so declare_tile remains a recognized runtime carrier regardless of the compact band the materialized template attaches. Fixes tpop_tilelib_addr_fold and the 11 expand_tile_op_ptodsl_* regressions exposed by the compact-aware TileSpec.
…quivalence PTOFusionPredicateElision built a PredicateScopeContext per fusion region and per nested vecscope, but the same plt op can be collected into more than one scope (a fusion_region scope and its nested vecscope scope both walk the same body). Each scope independently erased its duplicate-equivalent plt candidates, so a plt that was already erased by an earlier scope was erased again — corrupting MLIR use-lists and causing an intermittent SIGSEGV (exit 139) in runOnOperation's replaceAllUsesWith/erase splice. Track erased plt ops across all scopes with a shared DenseSet and skip ops already erased globally before rewire/erase. Also tighten the equivalence gate: compare the plt operations themselves (attributes, result types, all operands) via areEquivalentOperations rather than only the scalar input operand. The scalar-only check falsely merged independent plt streams (distinct loop-carried recurrences, or cross-template predicates sharing a constant), and the areSameValuePair set-equality shortcut accepted distinct iter-arg self-recurrences as equivalent. Replace the set-equality test with an explicit per-side self-recursion check (lhs == lhsRecurrenceInput && rhs == rhsRecurrenceInput). Fixes the three tile_fusion lit crashes (fusion_computed_dynamic_shape_e2e, op_fusion_low_level_loop_dynamic_shape, op_fusion_low_level_loop_softmax_prepare_unaligned_nopad).
…d_value from _ConfigView
The metadata TileSpec.compact_mode defaulted to 'normal', which
materialized compact=1 on every plain tile_buf and broke
FoldTileBufIntrinsics for the compact-unaware declare_tile bridge.
Default to None (unspecified) and render None as Null (omitting the
compact suffix) so plain tiles match the historical IR shape.
Revert the _ConfigView.pad_value field: it was not present in the
base _ConfigView, and adding it changed the behaviour of the
row_reduce_vmi_constraint predicate, which previously crashed
(AttributeError on _ConfigView.pad_value) and thus rejected the
vmi_trowsum_row candidate. Restoring the missing field keeps the
base candidate-selection ordering (template_trowsum wins) so
ptodsl_vmi_local_reduce_candidates and ptodsl_vmi_sinkhorn_grouped
match again.
Drop the None->'Null' alias from _normalize_compact_mode so an
unspecified compact_mode (None) is rejected by
_has_gap_free_row_stride (restoring the base 1-D candidate filter
that the 'normal' default had weakened).
In _tile_template_tracing, coerce a None compact_mode from an
operand TileSpec to 'normal' before building the tracing TileSpec so
the {normal, row_plus_one} validator still accepts it.
test/lit.cfg.py only probed ../llvm-project/build-shared/bin for the LLVM tools (FileCheck etc.), but the shared LLVM build actually lives under ../llvm-workspace/llvm-project/build-shared/bin. Add that as a fallback candidate so bare 'FileCheck' references in test/ root RUN lines resolve. Fixes the 'FileCheck: command not found' failures for the test/ root tests (scf_*, sync_*, test_unroll, dsl/expand_tile_op).
scf_while_break: CHECKs assumed index -> int32_t (should be int64_t) and over-specified SSA numbering/ordering; rewrite as loose, order-independent patterns like the sibling scf_for_break_like test. test_inject_sync_intra_pipe_barrier: RUN line used './bin/ptoas' (mangled by the ptoas substitution) and lacked --mlir-print-ir-after=pto-insert-sync, so stderr was empty; match the working sibling RUN line. test_unroll_annotation: loops wrote to a memref that was never read, so the Inliner's Canonicalizer dead-store-eliminated the loops before PTOUnrollSIMTFor ran (showing 0 scf.for instead of 2). Add a memref.load + return so the stores stay live; also use --emit-vpto (not --emit-pto-ir) so the unroll pass actually runs in the VPTO backend pipeline. dsl/expand_tile_op_tilelang_tadds: CHECK expected pto.addptr but the lowering emits pto.castptr, matching every sibling expand_tile_op_tilelang test.
…ile ops Modules with no unexpanded tile ops skip the shared PTO-to-VPTO lowering pipeline and go straight to the VPTO backend, whose result emission requires the CANN toolchain. That early branch ran before the main emitMlirIR check, so --emit-pto-ir on such a module (e.g. the scf/control-flow samples) printed 'CANN toolchain is required' and exited 1 instead of emitting the requested text IR. Emit the current module as text and return 0 in that branch when --emit-pto-ir is set, matching the main-pipeline emitMlirIR behaviour. The --emit-vpto / object-emission paths are unchanged.
The core UNROLL check covers the PTOUnrollSIMTFor behaviour: the annotated simt_entry loop is fully unrolled while an unannotated simt_entry loop and a non-simt loop both survive, so exactly 2 scf.for remain. This regression is watched through --emit-vpto because the unroll pass lives in the VPTO backend pipeline. Add an EMITIR check that --emit-pto-ir on a module with no tile ops prints the module as text (covering the ptoas fix that used to fail with 'CANN toolchain is required').
- _vmi_common: signed-int scale chains for vdiv IEEE-754 refinement (vabs requires explicit si8/si16/si32); si32 mediation for vcvt int<->fp; src-pad null gate for grouped row-reduce; workspace kept out of _prepare_tile_access (scratch-only, may lack addr); streaming row-reduce requires a legal VMI lane count. - tcvt: expose f16->i8 alias alongside si8 so the catalog invariant (one 1d pair per existing candidate) holds. - tracing/metadata/selection/daemon/constraints: compact_mode default `null` (gap-free), wire-omitted compact stays None (1-D rejected), _ConfigView restores pad_value for row-reduce constraint reads. No behavioral change to the committed P1/P2/jit/pipe_surface groups; this package makes the tilelib/VMI catalog and template tests agree with their corrected capabilities.
…nMemory - ExpandTileOp: emit compact_mode by the TileBuf band value (null/normal/row_plus_one) instead of mapping every non-row_plus_one tile to "normal", which materialized compact=1 helpers that FoldTileBufIntrinsics cannot bridge from plain callers. - PTOPlanMemory: treat tpop/tfree_from_aic/aiv as known pipe-entry ops and track a produced TileBuf result so local-buffer liveness does not hit the unknown-op UNREACHABLE. - ptoas: SerialFrontendPipeLoweringPass lowers collection to the shared walk (nested per-kernel-kind modules) instead of top-level funcs only.
…mantics - vmi_tile_template: 8x512 row-expand and 8x128 grouped row-reduce now assert fallback under the P1-2 256-lane VMI ceiling instead of expecting an impossible 256-/1024-lane materialization; f16 tadd, lane-snap 96/32-wides, and tcvt si32 mediation assertions updated; VMI-to-VPTO lowering uses an alloc-driver inline so FoldTileBuf intrinsics has alloc_tile-defined handles. - catalog: trowmax over-256-lane full shape asserts rejection; tgather uses the real template; i32->f16 tcvt pair removed (A5 vcvt has no such contract); f16->i8 alias asserted. - frontend_verify: the two stale 'VPTO fallback object path with --emit-pto-ir' assertions now expect IR text, matching the P2 main-pipeline behavior. - lit: local-reduce/sinkhorn candidate selections updated to the VMI forms the corrected constraints select.
The daemon's tile-operand binding dropped config.s_fractal_size, so every request silently fell back to the 512 default and could select a different candidate / render a different tile_buf ABI than the in-process _selection path. Mirror _selection's read-and-normalize (0 -> 512) and add a regression test asserting daemon and selection agree for s_fractal_size 32/1024 and normalize 0.
- ExpandTileOp: restore correct indentation and brace every branch of the compact_mode emission (the tab-prefixed block also broke the JSON string emission readability). - PTOPlanMemory: brace the TileBuf-result guard; hoist the condition into a local bool so the checker's control-statement rule is not tripped by the braced && expression.
… width A pto.vlds with dist="BRC_B8/B16/B32" is a scalar broadcast load: it reads a single hardware element whose width is fixed by the `_Bn` suffix (1/2/4 bytes) and replicates it across the result vreg lanes, independent of the source element type — the verifier accepts width-mismatched forms such as BRC_B32 on ptr<f16> (reads 4 bytes, i.e. two f16 elements). buildAccessDescriptor modelled every VldsOp as touching the full vreg width (256 bytes for a 64xf32 result), so a broadcast load of rowsum[i] in a row-reduce/div vecscope fabricated a cross-iteration RAW against the adjacent y buffer's IV-strided store, inserting a redundant pto.mem_bar "VST_VLD" inside the div loop. The barrier serialized every loop iteration via SMEM_BAR and prevented V-pipeline overlap. Add vldsBroadcastByteSize, mirroring the existing vstsStoredElementCount: report the exact 1/2/4-byte footprint for BRC_B8/B16/B32 (sound under width-mismatched forms — reporting one source element instead would under-count, e.g. 2 bytes for BRC_B32 on f16, and could drop a required loop-carried barrier). BRC_BLK and the unpack/upcast/shift distributions (US_*, DS_*, UNPK_*, E2B_*) keep the conservative full-vreg width until their element-type pairing is modelled. This only ever relaxes a false positive (a load touching fewer bytes than reported) and cannot drop a real hazard. Fixes #12.
added 3 commits
August 27, 2026 10:10
…d2nz/elision assertions
…ride VMI vstore no longer carries a repeat_stride operand (upstream dropped it), so the tmov ND->NZ and 1/2-VL assertions now check the emitted constant block_stride (16 / 128) instead of a constant 1.
Owner
Author
|
Closing: base branch was incorrectly created; PR will be re-opened against main-llvm19-build. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Rebase the
main-llvm19-buildworking branch onto the latest upstreamhw-native-sys/PTOASmain(10883a8), re-applying the feature/fix history and resolving conflicts with upstream changes (notably the removal ofrepeat_stridefrom the VMIvstorewhile keeping the branch'supdated_base/post_updatework).Also includes a new upstream follow-up:
fix(vpto): combine reductions before scheduling(10883a8) is incorporated, with the conflict intools/ptoas/ptoas.cppresolved so the late duplicateExpandSoftLib/InlineLibCallblock is dropped (early block retained) andVPTOCombineReductionsstays before scheduling.Test results (final commit cad6752)
pto_vexpdif_fusion.pto, pre-existing branch behavior: tiles render into independentfusion_regions vianon_vmi_local_boundary_fallback, which the vexpdif fusion pass cannot fuse across — not a rebase regression)Notable changes in this branch (vs upstream)
updated_base/post_updateon VMI vstore (kept duringrepeat_strideremoval)signal.hinclude)