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/07/06 17:08:43 UTC
[GitHub] [tvm] wrongtest opened a new pull request, #12019: [TIR] Revert #11428 and move loop dependent alloc extent check after region union
wrongtest opened a new pull request, #12019:
URL: https://github.com/apache/tvm/pull/12019
cc @junrushao1994
--
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 #12019: [TIR] Revert #11428 and move loop dependent alloc extent check after region union
Posted by GitBox <gi...@apache.org>.
wrongtest commented on PR #12019:
URL: https://github.com/apache/tvm/pull/12019#issuecomment-1177599118
Trim https://gist.github.com/junrushao1994/943eb0e9232a21a2748eb35af0fac4ff as a testcase.
--
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 #12019: [TIR] Revert #11428 and move loop dependent alloc extent check after region union
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on PR #12019:
URL: https://github.com/apache/tvm/pull/12019#issuecomment-1176795281
Looks like we are now calculating the upper bound of the union bound, which I think is correct. Would love to have confirmation from both of you guys @wrongtest @ArmageddonKnight
--
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 #12019: [TIR] Revert #11428 and move loop dependent alloc extent check after region union
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on PR #12019:
URL: https://github.com/apache/tvm/pull/12019#issuecomment-1176764243
Moving our discussion thread from #11428. Thanks @ArmageddonKnight for spotting this nasty corner case, and @comaniac for offline discussion :-)
@junrushao1994:
> 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)).
@junrushao1994
> A minimal reproducible example: https://gist.github.com/junrushao1994/943eb0e9232a21a2748eb35af0fac4ff
>
> The region of `A_shared` should be [128, 4] instead of [192, 4]
@junrushao1994
> 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)
@wrongtest
> > 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 a diff in pull request #12019: [TIR] Revert #11428 and move loop dependent alloc extent check after region union
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on code in PR #12019:
URL: https://github.com/apache/tvm/pull/12019#discussion_r915288059
##########
src/tir/transforms/compact_buffer_region.cc:
##########
@@ -45,25 +45,43 @@ using support::NDIntSet;
* \brief simplify and return the region collected by NDIntSet. return the original
* buffer shape if the int_set is empty.
*/
-Region SimplifyAndNarrowBufferRegionFromNDIntSet(const NDIntSet& nd_int_set,
- const Array<PrimExpr>& original_shape,
- arith::Analyzer* analyzer) {
+Region SimplifyAndNarrowBufferRegionFromNDIntSet(
+ const NDIntSet& nd_int_set, const Array<PrimExpr>& original_shape, arith::Analyzer* analyzer,
+ const std::vector<const ForNode*>& ancestor_loops) {
Array<Range> result;
result.reserve(nd_int_set.size());
for (size_t i = 0; i < nd_int_set.size(); ++i) {
const arith::IntSet& int_set = nd_int_set[i];
Range range = int_set.CoverRange(Range(/*begin=*/0, /*end=*/original_shape[i]));
- result.push_back(
- Range::FromMinExtent(analyzer->Simplify(max(0, range->min)),
- analyzer->Simplify(min(original_shape[i], range->extent))));
+ PrimExpr min = analyzer->Simplify(tvm::max(0, range->min));
+ PrimExpr extent = analyzer->Simplify(tvm::min(original_shape[i], range->extent));
+
+ // Check the buffer region is not loop dependent, since loop dependent
+ // allocation is not supported yet.
+ auto is_loop_var = [&ancestor_loops](const VarNode* v) {
+ return std::any_of(ancestor_loops.begin(), ancestor_loops.end(),
+ [v](const ForNode* n) { return n->loop_var.get() == v; });
+ };
+ if (UsesVar(extent, is_loop_var)) {
+ // try estimate a constant upperbound on region's extent
+ int64_t upperbound = analyzer->const_int_bound(extent)->max_value;
+ if (upperbound != arith::ConstIntBound::kPosInf) {
+ extent = make_const(extent->dtype, upperbound);
+ } else {
+ // or else we have to fallback to full region
+ min = 0;
Review Comment:
use `make_zero(original_shape[i]->dtype)`
--
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 merged pull request #12019: [TIR] Revert #11428 and move loop dependent alloc extent check after region union
Posted by GitBox <gi...@apache.org>.
junrushao1994 merged PR #12019:
URL: https://github.com/apache/tvm/pull/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 #12019: [TIR] Revert #11428 and move loop dependent alloc extent check after region union
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on PR #12019:
URL: https://github.com/apache/tvm/pull/12019#issuecomment-1176795611
Let's also add my minimal reproducible example as a regression unittest in case it happens in the future
--
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] ArmageddonKnight commented on pull request #12019: [TIR] Revert #11428 and move loop dependent alloc extent check after region union
Posted by GitBox <gi...@apache.org>.
ArmageddonKnight commented on PR #12019:
URL: https://github.com/apache/tvm/pull/12019#issuecomment-1176779913
I think this implementation makes sense. Please let me know when it is merged so that I can rebase onto the latest main branch.
--
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