[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]

Rebase validation is complete on exact base 771193a5c26188713b2ec01f5bcd937bf7966fae and exact head 90460d0575efeaa7ddaded591a036a65cf7778f6. The three commits are patch-identical to the previously validated series by range-diff. The only newly merged overlap is TMEM datapath B in python/tvm/tirx/script/builder/ir.py; its changes are disjoint from the typed-buffer edits and the combined focused suite passed 102 tests with 256 compute>=10 hardware skips on H200. The Debug LLVM/CUDA rebuild, matched CUDA 13.1 NVRTC+NVCC suites (376 passed, 159 skipped), H200 MLA runtime (4 passed), exact CI=true optimize-LLM tutorial, repository-wide pre-commit, and per-commit/diff audits all passed. The broad TIRx/S-TIR run completed with 2,438 passed, 1,642 skipped, 10 xfailed, and 2 xpassed; its sole failure is the preserved local build omitting optional runtime.RPCTimeEvaluator (USE_RPC), not a source regression. No merge was performed.