[S-TIR][Tests] Fix transform test failures after TIRx bringup

[tlopex]

This PR fixes 11 test failures in tests/python/s_tir/transform/ introduced as side effects of the TIRx bringup (#19581 / 859498d), 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.

[gemini-code-assist]

Code Review

This pull request introduces a helper method IsOp in inject_permuted_layout.cc to match device intrinsics by both their flat builtin name and canonical dotted name, updating the visitor to handle legacy forms of ptx_ldmatrix and mma_store. Additionally, test files are updated to use T.prim_func(s_tir=True) and replace T.decl_buffer with T.alloc_buffer containing explicit annotations. The reviewer suggests optimizing the IsOp helper by caching the canonical Ops as static local variables using Op::Get to avoid costly dynamic casts and string comparisons on every visited node.

Important

The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.

[tlopex]

@tvm-bot rerun