You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/05/24 13:53:25 UTC
[GitHub] [tvm] wrongtest opened a new pull request, #11428: [TIR] Avoid loop dependent allocation in buffer compaction
wrongtest opened a new pull request, #11428:
URL: https://github.com/apache/tvm/pull/11428
`EstimateRegionLowerBound` sometimes gives "too accurate" region extents (dependent by outer loop vars). However, the lowering (`StorageRewrite`) would lift buffer allocation to the global scope, bring the loop vars out of it's def scope illegally.
For the padded and tiled pooling case in the test, the h-dimension's region extent could be infered as `h_o * 8 + 8 - max(h_o * 8 - 1, 0)`, it is correct since with padding, the edge tiles do not need same buffer size as others. (h_o = 0 => extent = 8, h_o > 0 => extent = 9). But the dynamic allocation for this extent expression can not get properly handled afterwards...
Here is just one work-around, very glad to see any robust solutions. @Hzfengsy @spectrometerHBH
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao1994 commented on pull request #11428: [TIR] Avoid loop dependent allocation in buffer compaction
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on PR #11428:
URL: https://github.com/apache/tvm/pull/11428#issuecomment-1175634426
This PR introduces a regression with those lines: https://github.com/apache/tvm/blob/d4be49aec62299275565066b56a0555bafc2ccac/src/tir/transforms/compact_buffer_region.cc#L77-L85
Consider a case below:
```python
extent = 960 - max(bx // 18 * 128, 832)
where bx \in [0, 144)
```
Even if `bx` exists in `extent`, the upper bound of the allocation is not affected (which is 128). However, the check there introduces a false positive which falls back to a much more conservative estimate ([full region](https://github.com/apache/tvm/blob/d4be49aec62299275565066b56a0555bafc2ccac/src/tir/transforms/compact_buffer_region.cc#L77-L85)).
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] wrongtest commented on pull request #11428: [TIR] Avoid loop dependent allocation in buffer compaction
Posted by GitBox <gi...@apache.org>.
wrongtest commented on PR #11428:
URL: https://github.com/apache/tvm/pull/11428#issuecomment-1176478434
> calculating the region union bound first
Many thanks! It indeed should check loop dependency for union region.
https://gist.github.com/junrushao1994/943eb0e9232a21a2748eb35af0fac4ff
In the case it seems we get extent=128 for `A_shared[128, 4]` after fix, even without const bound estimation... Perhaps the loop dependent region parts get covered during region unions?
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao1994 commented on pull request #11428: [TIR] Avoid loop dependent allocation in buffer compaction
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on PR #11428:
URL: https://github.com/apache/tvm/pull/11428#issuecomment-1176761434
Hey let's move our discussion to #12019
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao1994 commented on pull request #11428: [TIR] Avoid loop dependent allocation in buffer compaction
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on PR #11428:
URL: https://github.com/apache/tvm/pull/11428#issuecomment-1175750797
so let me briefly introduce the workflow of Compute-Buffer-Region first, and then let’s think about how to make it perfectly fit in our usecase.
On each use-site of a buffer, we calculate the region it is touched by the loops above it, and then by unionizing those regions, it’s made possible to trim out those regions that are not touched.
As a simplest example, if the original size of a buffer in shared memory is `[1024, 1024]`, but only `[bx : bx + 10, by : by + 20]` is touched in a threadblock, then the buffer size could shrink to `[10, 20]` which is significantly smaller.
Then to provide a slightly more complicated example, if the buffer is touched in two different places, where the region are `[bx : bx + 10 : by : by + 20]`, `[bx + 10: bx + 20, by : by + 30]`, then we calculate the union of the two regions, i.e. `[bx : bx + 20, by : by + 30]`, and then do shrinking so that its size of `[20, 30]`.
The problem this PR aims to address is that some buffer’s size actually depends on irrelevant outer variables, for example, the shared memory size to be `[0, 960 - max(bx // 18 * 128, 832)]`, where `bx` is the extent of `blockIdx.x`. However, it takes an approach to over-relax the region when seeing any undesired variables (in our case, it’s `bx`), which is less optimal.
The actual solution should be: calculating the region union bound first, and then remove the undesired variables with something similar to analyzer's `const_int_bound` (note that we don’t want the bound to be strictly constant)
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] spectrometerHBH merged pull request #11428: [TIR] Avoid loop dependent allocation in buffer compaction
Posted by GitBox <gi...@apache.org>.
spectrometerHBH merged PR #11428:
URL: https://github.com/apache/tvm/pull/11428
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao1994 commented on pull request #11428: [TIR] Avoid loop dependent allocation in buffer compaction
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on PR #11428:
URL: https://github.com/apache/tvm/pull/11428#issuecomment-1175634709
A minimal reproducible example: https://gist.github.com/junrushao1994/943eb0e9232a21a2748eb35af0fac4ff
The region of `A_shared` should be [128, 4] instead of [192, 4]
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao1994 commented on pull request #11428: [TIR] Avoid loop dependent allocation in buffer compaction
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on PR #11428:
URL: https://github.com/apache/tvm/pull/11428#issuecomment-1136761189
Indeed. It is a rare bug that we noticed but haven't got priority to fix yet. Thanks for fixing this!
Would you like to verify if it works on this conv2d winograd example: https://gist.github.com/junrushao1994/e287cfb8091f18eb9fad90f535c70218?
Thanks a lot!
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org