[TIRX] Represent buffers as typed variables

[tqchen]

This refactors TIRx buffers into ordinary ir.Var identities carrying an
immutable BufferType.

Rationale

The old representation kept a standalone Buffer object, its logical metadata,
and a separate pointer-typed data Var. Generic Var substitution and
buffer-specific rebuilding could therefore let identity, metadata, and pointer
type drift apart. In this change, BufferType is the single source of the
logical contract and buffer_data(buffer) is the explicit physical-pointer
projection; DeclBuffer records a view's binding.

Simplification

• Generic Var scoping, substitution, undefined-variable analysis, and identity
  maps now apply directly to buffers.
• The pointer type is derived from BufferType, so the physical pointer and
  buffer metadata cannot disagree.
• The common IR path no longer needs a second Buffer object, duplicate
  constructor state, or a parallel buffer-only substitution mechanism.
• Metadata-changing transforms create a fresh typed Var and bind it explicitly
  at the replacing AllocBuffer or DeclBuffer.

Physical-root tracking remains deliberately local to the passes whose
correctness depends on allocation ownership:

1. Storage liveness/reuse attributes alias accesses to the owning allocation.
2. Pointer vector-type inference and rewriting aggregate evidence at that root.
3. LowerTIRx cleanup projects aliases to the remapped physical allocation.
4. Thread-allreduce's deferred remapper replaces or removes the owning
   allocation.

Each is a plain pass-local Map<Var, Var> flattened when a DeclBuffer is
visited. Parameters, buffer_map, and AllocBuffer seed roots; an alias to an
unseen typed source is rejected rather than silently treated as a new root.
Simpler passes decode the one-step buffer_data projection locally.

Migration guide

C++

• Treat BufferVar as a checked view over an ordinary Var.
• Read logical metadata through buffer->...; use buffer.var() for identity.
• Use buffer.data() or builtin::buffer_data() only when a physical pointer
  expression is required.
• To change metadata, rebuild BufferType, create a fresh Var, and bind it at
  the replacement definition site.
• A transformed DeclBuffer replaces its superseded declaration consistently;
  it must not leave a stale declaration beside the new binding.

Python

• Buffer constructors return tvm.ir.Var; .ty contains BufferType.
• Use tvm.tirx.is_buffer_var(value) for runtime discrimination.
• tvm.tirx.Buffer remains an import and annotation compatibility alias for
  tvm.ir.Var, so isinstance(value, Buffer) matches every Var.
• Importing tvm.tirx installs the legacy metadata properties on Var.
  Non-buffer Vars raise AttributeError; buffer .dtype keeps its historical
  DataType result.

Behavior preservation

The behavior-preserving follow-up restores established LowerTIRx,
FlattenBuffer, BF16, alias-lifetime, and declaration-replacement behavior while
keeping ownership logic pass-local.

The final focused correction completes typed-buffer integration at structural
comparison, retyping, specialization, annotation, external-output, access
detection, projection lowering, printer, and target-codegen boundaries. It
also preserves the physical source handle type for declared aliases and
retains allocation-backed versus parameter-backed identity where lowering
requires that distinction.

Validation

• Debug build with LLVM 18.1.8 and CUDA: passed.
• LowerIntrin, C address-generation, padded-vector LLVM, packed CUDA, raw
  pointer, and rank-0/rank-2 access-pointer regressions: 11 passed.
• Coherent NumPy 2.3.5 reproduced both missed PrimType conversion
  boundaries; the corrected runtime-trace and minimal-LLVM files pass
  8 tests.
• Broader affected transform, C-host, runtime-trace, minimal-LLVM, and
  commit-3 regression suite: 467 passed, 8 xfailed, 1 xpassed.
• Matched CUDA 13.1 validation asserted nvrtcVersion() == 13.1 in the test
  process and loaded libnvrtc.so.13.1.115. After the representation-safe
  TensorMap address_of migration, the two complete CUDA-codegen files passed
  376 tests with 159 skipped.
• Complete TIR and Relax specialization files: 12 passed.
• Isolated schedule-state regression: 1 passed.
• H200 paged-attention float32 prefill/decode nvcc regression: 1 passed.
• Repository-pinned formatting, all pre-commit hooks, and git diff --check:
  passed.

[gemini-code-assist]

Caution

The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased.

[tqchen]

Rebased this PR's three-commit series onto current main@115029b3ab44c177bfca56c2c1bcf65566450d26; the new exact head is 8c82011301213c409d961b8fff1c732215d09ae5.

Integration audit:

• preserved the dependency-aware ConvertSSA buffer-remap invalidation added upstream and adapted it to BufferVar;
• preserved the new logical element-pointer handling and now derives pointer dtype/storage scope from BufferType;
• ported the three affected upstream tests to the typed-buffer representation.

Validation completed on the rebased tree:

• LLVM/CUDA compiler build and exact-head no-op rebuild;
• exact CI=true H200 docs/how_to/tutorials/optimize_llm.py run with CUDA/NVRTC 13.1;
• target-codegen suites: 804 passed, 32 skipped;
• H200 TIRx CUDA + MLA: 16 passed, 155 skipped;
• broad TIRx/S-TIR: 2,440 passed, 1,624 skipped, 10 xfailed, 2 xpassed. The sole failure was isolated to this build omitting the optional runtime.RPCTimeEvaluator RPC global, rather than a source regression;
• repository-wide all-files pre-commit and git diff --check.

The PR remains open and unmerged for fresh CI on the new exact head.

[tqchen]

@tvm-bot rerun

[github-actions]

Failed to re-run CI in https://github.com/apache/tvm/actions/runs/30517774442

Details

Traceback (most recent call last):
  File "/home/runner/work/tvm/tvm/ci/scripts/github/github_tvmbot.py", line 588, in comment_failure
    raise item
  File "/home/runner/work/tvm/tvm/ci/scripts/github/github_tvmbot.py", line 694, in run
    pr.rerun_jenkins_ci()
  File "/home/runner/work/tvm/tvm/ci/scripts/github/github_tvmbot.py", line 547, in rerun_jenkins_ci
    post(url, auth=("tvm-bot", TVM_BOT_JENKINS_TOKEN))
  File "/home/runner/work/tvm/tvm/ci/scripts/jenkins/git_utils.py", line 54, in post
    with request.urlopen(req, data) as response:
         ^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/usr/lib/python3.12/urllib/request.py", line 215, in urlopen
    return opener.open(url, data, timeout)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/usr/lib/python3.12/urllib/request.py", line 515, in open
    response = self._open(req, data)
               ^^^^^^^^^^^^^^^^^^^^^
  File "/usr/lib/python3.12/urllib/request.py", line 532, in _open
    result = self._call_chain(self.handle_open, protocol, protocol +
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/usr/lib/python3.12/urllib/request.py", line 492, in _call_chain
    result = func(*args)
             ^^^^^^^^^^^
  File "/usr/lib/python3.12/urllib/request.py", line 1392, in https_open
    return self.do_open(http.client.HTTPSConnection, req,
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/usr/lib/python3.12/urllib/request.py", line 1348, in do_open
    r = h.getresponse()
        ^^^^^^^^^^^^^^^
  File "/usr/lib/python3.12/http/client.py", line 1457, in getresponse
    response.begin()
  File "/usr/lib/python3.12/http/client.py", line 336, in begin
    version, status, reason = self._read_status()
                              ^^^^^^^^^^^^^^^^^^^
  File "/usr/lib/python3.12/http/client.py", line 305, in _read_status
    raise RemoteDisconnected("Remote end closed connection without"
http.client.RemoteDisconnected: Remote end closed connection without response