Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[TIR] Output DeclBuffer in FlattenBuffer #15757

Open
wants to merge 3 commits into
base: main
Choose a base branch
from

Conversation

Lunderberg
Copy link
Contributor

If a flattened buffer is produced for use in BufferLoad and BufferStore statements, generate a DeclBuffer. This PR updates the behavior for both TE-derived schedules (tir.StorageFlatten transform) and TIR-derived schedules (tir.FlattenBuffer transform).

This is a subset of the changes made in #14778, broken out for ease of testing and review.

@Lunderberg Lunderberg force-pushed the declbuffer_output_from_buffer_flattening branch from f97dbd9 to 9875bd3 Compare December 19, 2023 18:53
If a flattened buffer is produced for use in `BufferLoad` and
`BufferStore` statements, generate a `DeclBuffer`.

This is a subset of the changes made in
apache#14778, broken out for ease of
testing and review.
When producing a flattened buffer for use in `BufferLoad` and
`BufferStore` nodes, generate a `DeclBuffer` for the flattened buffer.

This is a subset of the changes made in
apache#14778, broken out for ease of
testing and review.
Prior to this commit, the functions in
`python/tvm/relay/backend/contrib/ethosu` tracked buffers based on
`tir.Var`, typically determined from `buffer_load.buffer.data`.  This
commit updates these funcitons to instead track based on `tir.Buffer`,
now determined as `buffer_load.buffer`.

This change allows for tracking of buffer objects, to determine which
`DeclBuffer` statements should be removed.
@Lunderberg
Copy link
Contributor Author

This is blocked on the ethos-u and ethos-m codegen, which does not currently support DeclBuffer nodes.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

1 participant