Skip to content

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