tqchen opened a new pull request, #20079:
URL: https://github.com/apache/tvm/pull/20079
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.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
