[FIX][TIRx] Traverse pointer expressions in tile calls

[jinhongyii]

Motivation

TilePrimitiveCall arguments and config values can contain any tvm.ir.Expr, not only PrimExpr. In particular, T.address_of(...) produces a pointer expression.

The Python StmtVisitor and StmtMutator only traversed values accepted by tvm.ir.is_prim_expr. Pointer-valued operands were therefore skipped entirely. Analyses could miss buffer uses below the pointer expression, and rewrites could leave stale buffer references behind.

Example

A tile call can pass an mbarrier address through its config:

@T.prim_func
def copy_async(
    A: T.Buffer((8,), "float16"),
    B: T.Buffer((8,), "float16"),
    mbar: T.Buffer((1,), "uint64"),
):
    Tx.copy_async(
        B[:],
        A[:],
        dispatch="tma_auto",
        mbar=T.address_of(mbar[0]),
    )

Before this PR, a StmtExprVisitor did not reach the BufferLoad for mbar[0], and a StmtExprMutator could not replace it. After this PR, both visitor and mutator traverse the pointer expression normally.

What changed

• Treat every tvm.ir.Expr in TilePrimitiveCall arguments and config values as traversable.
• Add a regression test that verifies both visiting and replacing the buffer load beneath T.address_of.

Tests

• python -m pytest tests/python/tirx/transform/test_stmt_functor.py -q (20 passed)
• pre-commit run --files python/tvm/tirx/stmt_functor.py tests/python/tirx/transform/test_stmt_functor.py

Ported from mlc-ai/tvm@98b7dec.