tlopex opened a new pull request, #19735:
URL: https://github.com/apache/tvm/pull/19735
This PR fixes 11 test failures in `tests/python/s_tir/transform/` introduced
as side effects of the TIRx bringup (#19581 / 859498dc01), in three independent
commits.
### 1. LowerOpaqueBlock: update expected IR for buffer metadata annotations
`LowerOpaqueBlock` now emits `buffer_allocated_addr` and
`buffer_data_alignment` annotations on lowered allocations (intentional in
#19581: the annotations are consumed downstream by `codegen_cuda.cc` /
`codegen_trn.cc`; the alignment value 64 comes from `kAllocAlignment`). The
tests' expected IR predates this, so `assert_structural_equal` failed on the
missing annotations.
Fix: update the expected IR in `test_s_tir_transform_lower_opaque_block.py`
to carry the annotations (`T.decl_buffer(...)` → `T.alloc_buffer(...,
annotations={...})`). Fixes 6 tests.
### 2. DefaultGPUSchedule: parse scalar-block test in s_tir mode
#19581 added a well-formedness rule rejecting `SBlockRealize` in `tirx=True`
mode, which is correct — sblocks are s_tir-mode constructs. The hand-written
`Before`/`Expected` modules in `test_scalar_block_no_loops` were the only ones
in the file still using plain `@T.prim_func`, so they failed at parse time
before the pass under test even ran.
Fix: parse both modules with `@T.prim_func(s_tir=True)`, consistent with
every other test in the file. Fixes 1 test.
### 3. InjectPermutedLayout: match legacy PTX intrinsics by canonical name
#19581 registers device intrinsics under two Op identities: a flat builtin
name (returned by `builtin::xxx()` in C++) and a canonical dotted name (e.g.
`tirx.ptx.ldmatrix_legacy`, produced when TVMScript / tensor intrinsics are
parsed). `InjectPermutedLayout` only compared with `same_as(builtin::...)`, so
it silently skipped rewriting the swizzled shared-memory offsets of parsed
legacy-form calls, leaving the expected swizzle index expressions unmatched.
Fix: match `ptx_ldmatrix_legacy` / `mma_store_legacy` by both the builtin Op
and the canonical name via an `IsOp` helper, following the existing pattern in
`lower_warp_memory.cc` and `codegen_cuda.cc`. Only the legacy intrinsic forms
fold shared-memory access into `tvm_access_ptr` + offset; non-legacy forms
address shared memory through `BufferLoad` and are already handled by the
BufferLoad visitor, so the unreachable `InternalError` throw is replaced by a
pass-through. (`mma_store_legacy` has no dotted alias, hence the asymmetric
name strings.) Fixes 4 tests.
### Validation
- `test_s_tir_transform_lower_opaque_block.py`,
`test_s_tir_transform_default_gpu_schedule.py`,
`test_s_tir_transform_inject_permuted_layout.py`: 27 passed.
- Full `tests/python/s_tir/transform/`: 243 passed; the remaining failures
are pre-existing on main and owned by other in-flight fixes (#19729 et al.),
none introduced here.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
