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