[TIR] Require buffer declaration (DeclBuffer) before use#14778
Closed
Lunderberg wants to merge 30 commits intoapache:mainfrom
Closed
[TIR] Require buffer declaration (DeclBuffer) before use#14778Lunderberg wants to merge 30 commits intoapache:mainfrom
Lunderberg wants to merge 30 commits intoapache:mainfrom
Conversation
Collaborator
|
Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from Reviewers by @-ing them in a comment.
Generated by tvm-bot |
Member
|
Super excited about this feature! |
Lunderberg
changed the title
Contributor
Author
|
@tvm-bot rerun |
Contributor
|
Failed to re-run CI in https://github.com/apache/tvm/actions/runs/4996801890 Details |
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
May 25, 2023
Prior to this commit, any trivial let binding of `var1 = var2` is
inlined. However, buffer definitions are not updated, so this can
result in dangling `tir::Var` instances. This commit updates the
`tir.Simplify` pass to keep trivial let bindings if they are used as
part of a buffer definition.
Ideally, the trivial `LetStmt` variable would be inlined into the
buffer definition as well as other expressions. However, because a
buffer may be implicitly declared, the first usage may be within a
constrained context. If that happens, the simplified shape/strides
expression cannot be used to update the buffer definition, as that
simplification is not valid at all possible usage points of the
buffer.
```python
for i in range(n):
elem_offset = i
view = T.Buffer(1, data=buf, elem_offset = elem_offset)
if i == 0:
# First occurrence in TIR is here, where elem_offset would
# simplify to zero.
view[0] = 1
else:
# But the same buffer is used here, where elem_offset doesn't
# simplify to zero.
view[0] = 2
```
This will be resolvable after apache#14778
lands, requiring all buffers to be declared with `DeclBuffer` prior to
usage.
```python
for i in range(n):
elem_offset = i
# All variables used by the DeclBuffer are valid across the entire
# body of the DeclBuffer.
view = T.decl_buffer(1, data=buf, elem_offset = elem_offset)
if i == 0:
view[0] = 1
else:
view[0] = 2
```
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
May 30, 2023
Prior to this commit, any trivial let binding of `var1 = var2` is
inlined. However, buffer definitions are not updated, so this can
result in dangling `tir::Var` instances. This commit updates the
`tir.Simplify` pass to keep trivial let bindings if they are used as
part of a buffer definition.
Ideally, the trivial `LetStmt` variable would be inlined into the
buffer definition as well as other expressions. However, because a
buffer may be implicitly declared, the first usage may be within a
constrained context. If that happens, the simplified shape/strides
expression cannot be used to update the buffer definition, as that
simplification is not valid at all possible usage points of the
buffer.
```python
for i in range(n):
elem_offset = i
view = T.Buffer(1, data=buf, elem_offset = elem_offset)
if i == 0:
# First occurrence in TIR is here, where elem_offset would
# simplify to zero.
view[0] = 1
else:
# But the same buffer is used here, where elem_offset doesn't
# simplify to zero.
view[0] = 2
```
This will be resolvable after apache#14778
lands, requiring all buffers to be declared with `DeclBuffer` prior to
usage.
```python
for i in range(n):
elem_offset = i
# All variables used by the DeclBuffer are valid across the entire
# body of the DeclBuffer.
view = T.decl_buffer(1, data=buf, elem_offset = elem_offset)
if i == 0:
view[0] = 1
else:
view[0] = 2
```
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
May 30, 2023
Prior to this commit, any trivial let binding of `var1 = var2` is
inlined. However, buffer definitions are not updated, so this can
result in dangling `tir::Var` instances. This commit updates the
`tir.Simplify` pass to keep trivial let bindings if they are used as
part of a buffer definition.
Ideally, the trivial `LetStmt` variable would be inlined into the
buffer definition as well as other expressions. However, because a
buffer may be implicitly declared, the first usage may be within a
constrained context. If that happens, the simplified shape/strides
expression cannot be used to update the buffer definition, as that
simplification is not valid at all possible usage points of the
buffer.
```python
for i in range(n):
elem_offset = i
view = T.Buffer(1, data=buf, elem_offset = elem_offset)
if i == 0:
# First occurrence in TIR is here, where elem_offset would
# simplify to zero.
view[0] = 1
else:
# But the same buffer is used here, where elem_offset doesn't
# simplify to zero.
view[0] = 2
```
This will be resolvable after apache#14778
lands, requiring all buffers to be declared with `DeclBuffer` prior to
usage.
```python
for i in range(n):
elem_offset = i
# All variables used by the DeclBuffer are valid across the entire
# body of the DeclBuffer.
view = T.decl_buffer(1, data=buf, elem_offset = elem_offset)
if i == 0:
view[0] = 1
else:
view[0] = 2
```
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
Jun 3, 2023
Prior to this commit, any trivial let binding of `var1 = var2` is
inlined. However, buffer definitions are not updated, so this can
result in dangling `tir::Var` instances. This commit updates the
`tir.Simplify` pass to keep trivial let bindings if they are used as
part of a buffer definition.
Ideally, the trivial `LetStmt` variable would be inlined into the
buffer definition as well as other expressions. However, because a
buffer may be implicitly declared, the first usage may be within a
constrained context. If that happens, the simplified shape/strides
expression cannot be used to update the buffer definition, as that
simplification is not valid at all possible usage points of the
buffer.
```python
for i in range(n):
elem_offset = i
view = T.Buffer(1, data=buf, elem_offset = elem_offset)
if i == 0:
# First occurrence in TIR is here, where elem_offset would
# simplify to zero.
view[0] = 1
else:
# But the same buffer is used here, where elem_offset doesn't
# simplify to zero.
view[0] = 2
```
This will be resolvable after apache#14778
lands, requiring all buffers to be declared with `DeclBuffer` prior to
usage.
```python
for i in range(n):
elem_offset = i
# All variables used by the DeclBuffer are valid across the entire
# body of the DeclBuffer.
view = T.decl_buffer(1, data=buf, elem_offset = elem_offset)
if i == 0:
view[0] = 1
else:
view[0] = 2
```
masahi
pushed a commit
that referenced
this pull request
Jun 4, 2023
…14951) Prior to this commit, any trivial let binding of `var1 = var2` is inlined. However, buffer definitions are not updated, so this can result in dangling `tir::Var` instances. This commit updates the `tir.Simplify` pass to keep trivial let bindings if they are used as part of a buffer definition. Ideally, the trivial `LetStmt` variable would be inlined into the buffer definition as well as other expressions. However, because a buffer may be implicitly declared, the first usage may be within a constrained context. If that happens, the simplified shape/strides expression cannot be used to update the buffer definition, as that simplification is not valid at all possible usage points of the buffer. ```python for i in range(n): elem_offset = i view = T.Buffer(1, data=buf, elem_offset = elem_offset) if i == 0: # First occurrence in TIR is here, where elem_offset would # simplify to zero. view[0] = 1 else: # But the same buffer is used here, where elem_offset doesn't # simplify to zero. view[0] = 2 ``` This will be resolvable after #14778 lands, requiring all buffers to be declared with `DeclBuffer` prior to usage. ```python for i in range(n): elem_offset = i # All variables used by the DeclBuffer are valid across the entire # body of the DeclBuffer. view = T.decl_buffer(1, data=buf, elem_offset = elem_offset) if i == 0: view[0] = 1 else: view[0] = 2 ```
Lunderberg
force-pushed
the
require_tir_buffer_declaration
branch
from
f8e0b4d to
72ed1d4
Compare
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
Jun 6, 2023
Part of changes being split out from apache#14778 into independent portions. This commit allows DeclBuffer to occur in the lowered TIR passed to CodeGenSPIRV.
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
Jun 6, 2023
Part of changes being split out from apache#14778 into independent portions. This commit allows DeclBuffer to occur in the lowered TIR passed to CodeGenStackVM.
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
Jun 6, 2023
Part of changes being split out from apache#14778 into independent portions. This commit allows TIR `cache_read` and `cache_write` schedule primitives to preserve `DeclBuffer` nodes.
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
Jun 6, 2023
Part of changes being split out from apache#14778 into independent portions. This commit allows TIR `compute_inline`, `compute_at`, and `reverse_compute_at` schedule primitives to preserve `DeclBuffer` nodes.
This was referenced Jun 6, 2023
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
Jun 6, 2023
Preserve DeclBuffer node when transforming with `LowerCustomDatatypes` This is a subset of changes, being split out from apache#14778 into independent portions.
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
Jun 6, 2023
This is a subset of changes, being split out from apache#14778 into independent portions.
Lunderberg
added a commit
to Lunderberg/tvm
that referenced
this pull request
Jun 6, 2023
Previously, `PoolAllocationToOffsetConverter` did not remap buffer objects occurring in `DeclBuffer` nodes. This commit updates `PoolAllocationToOffsetConverter` to handle `DeclBuffer` nodes. This is a subset of changes, being split out from apache#14778 into independent portions.
For any local buffers allocated as part of `LowerThreadAllreduce`, generate a `DeclBuffer`. This is a subset of the changes made in apache#14778, broken out for ease of testing and review.
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.
Expect the DeclBuffer that is now generated in `ir_builder.allocate(...)`
For the `res_buffer` and `normal_res_buffer`, output a `DeclBuffer` node. This is a subset of the changes made in apache#14778, broken out for ease of testing and review.
Previously, each bound constant resulted in a `AllocateConst` node. Now, each constant results in a `AllocateConst` node containing a `DeclBuffer` node.
This updates the older `tvm.tir.ir_builder.IRBuilder`, not the newer one that parses TVMScript. To maintain backwards compatibility with the older IRBuilder, the `ib.allocate` method also inserts a `DeclBuffer`. This is a subset of the changes made in apache#14778, broken out for ease of testing and review.
Implemented as part of VerifyWellFormed. Currently, fails after FlattenStorage.
Since it is no longer valid to omit the `DeclBuffer`, these test cases are no longer required.
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.
Previously, the `StorageRewrite` would hoist a single `DeclBuffer` node that was directly underneath a hoisted `Allocate` node. This commit updates `StorageRewrite` to hoist `DeclBuffer` nodes belonging to each allocation begin merged, and to merge the buffers themselves if several identical allocations are being merged.
In MergeConstants, hoist all `DeclBuffer` instances that use a merged allocation.
3 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Part of RFC#70, requiring buffers to be declared prior to their usage. This is currently a draft PR, to see what errors are found by the CI.