[TIR] Affine utility support iter lowerbound and diagnostics #9699

wrongtest-intellif · 2021-12-10T07:21:57Z

Hi, there. The PR originate from the issue detected in #9527, where a tiled block with both lowerbound and upperbound predicates fail to infer the region_cover property. Some tracing show that the DetectIterMap fail on such cases.

for hh_0, ww_0 in T.grid(28, 28):
    for ax0, ax1 in T.grid(10, 10):
        with T.block("cache"):
            h = T.axis.spatial(224, hh_0 * 8 - 1 + ax0)
            w = T.axis.spatial(224, ww_0 * 8 - 1 + ax1)
            T.where(1 <= hh_0 * 8 + ax0 and hh_0 * 8 + ax0 < 225 and 1 <= ww_0 * 8 + ax1 and ww_0 * 8 + ax1 < 225)
            cache[h, w] = X[h, w]

The PR modify affine utility in aspects below :

Free vars
For DetectIterMap, if the expression do not contain vars in domain map, it should be safe to not analyze affine form into it. eg, x*x + 8y + z could pass affine analysis if x is not an var of iter domains. We can just put x*x into base part of IterSumExpr. It benefit the above case since hh_0 * 8 + ax0 is not an affine form, but outer var might be free in certain analysis procedures.
Add min field to IterMark class
It is zero almost all time. But when lowerbound predicate exists, it seems much hard to represent the lowerbound with only source and extent.
More flexible predicate support
Support resolve all kinds of >, >=, <, <= integer comparisions, for the expr of form like i * 8 < 10 - j, try refactor the expr into domain var related part and domain var free parts.
Introduce DiagnosticContext for debug purpose
Replace each fail point (++unresolved_cnt_) with an error message recorded into diagnostic context. Or else it would be too hard to detect what happend for beginners. However, I do not find a proper point to call Render() of diagnostics for now.

Hzfengsy · 2021-12-10T08:59:26Z

cc @spectrometerHBH @junrushao1994

junrushao · 2021-12-14T22:41:23Z

Ping @spectrometerHBH

spectrometerHBH · 2021-12-16T23:32:43Z

Here are some comments.

Free vars
I think currently the DetectIterMap already supports free var, as long as you don't include the var in the domain map
Add min field to IterMark class
Iters with lower bound constraints can be substituted with a zero-starting iter plus offset, I think. If we want to allow non-zero lower bound constraints in predicates, then we may want to loose the check.
3 & 4 look good to me.

BTW, the case shown in your description is not the same as the case you added in the testcase. I thought you'd like to support the binding after #9527, so you may want to modify the test case, IIUC?

wrongtest-intellif · 2021-12-17T04:43:45Z

Many thanks to your comments!

Free vars
I think currently the DetectIterMap already supports free var, as long as you don't include the var in the domain map

Yes, I found that is the case. It should be my mistake on reason of failed cases.

2. Iters with lower bound constraints can be substituted with a zero-starting iter plus offset

Previous trials for me seems it require TryFuseIters be aware of those introduced offsets during rewrites of actual iter expr. So a mark of non-zero start is not a good idea? (for simplicity and correctness). I will try update to follow the suggestion.

the case shown in your description is not the same as the case you added in the testcase

I will change the testcase to match #9527,

wrongtest-intellif · 2021-12-20T07:34:41Z

2. Iters with lower bound constraints can be substituted with a zero-starting iter plus offset

Hi~ @spectrometerHBH Could you kindly take another round of review for lowerbound feature?

spectrometerHBH · 2021-12-21T03:43:29Z

tests/python/unittest/test_arith_intset.py

+    # h3 == 0: region is [1, 9]
+    # 0 < h3 <= 26: region is [h3 * 8, h3 * 8 + 9]
+    # h3 > 26: region is [h3 * 8, 223]
+    do_test_point_access(


Would be great if we can add 3 more test cases for each case of h3 :)

spectrometerHBH · 2021-12-21T05:02:37Z

src/arith/iter_affine_map.cc

+  // Then,    flattened form = IterSum(IterSplit(i, scale=9),
+  //                                   IterSplit(j, scale=2),
+  //                                   IterSplit(k, scale=1))
+  //          normal form    = IterSum(IterSplit(i, scale=9),


I don't think it's a good idea to rewrite i*9 + j*2 + k into such a normal form (i*9 + (j*2 + k - 1) + 1).
The reason is that j*2 + k - 1 has extent 8, it's weird that the scale of i is 9. Actually i*9 + j*2 + k with predicate 1 <= j*2 + k < 9 doesn't correspond to an iter, since the value of it is not continuous. Its value is [1, 8] \union [10, 17] \union [19, 26], ....

cc @wrongtest

What if we change this example to i*8 + j*2 + k , j*2 + k >= 1?

Then it looks good to me.

Then I think you might need to add more test cases to test the sanity.
Such as what if we try to fuse several iters all with lower bound constraints, and what if we split an iter with lower bound constraint.

wrongtest-intellif · 2021-12-21T05:32:36Z

BTW, I notice that negative scale is possble #9776, would we add some cases like -8 * x + y into testcases? @spectrometerHBH

spectrometerHBH · 2021-12-21T06:00:18Z

Try to add these test cases, which I think might be helpful to test the sanity.

i*8 + j*2 + k, with 1 <= j*2 + k < 9, i < 11, j < 5, i < 2
i*8 + j*2 + k, with 1 <= j*2 + k < 9, 3 <= i*8 + j*2 + k < 25, i < 11, j < 5, i < 2
i0*45 + i1*45 + i2*9 + i3*4 + i4, with 3<= i1*5 + i2 < 8, 1 <= i3*4 + i4 < 10, i0 < 3, i1 < 2, i2 < 5, i3 < 3, i4 < 4

These all should be detected successfully (I think). Correct me if anything wrong :)

…e constraints on iter

wrongtest-intellif · 2021-12-22T05:00:52Z

Rebase, add more testcases and try to fix three detected issues:

My previous code is buggy on duplicate lowerbounds on same iteration, I should not subtract min directly because it will break the assumption that mark->source be a structural form IterSumExpr. Instead, set -min to base as suggested.
Previously sum with single iter would not go into TryFuseIters thus if we have the i in [0, 48) and predicate is i < 10, it seems ignore the condition totally, result to an iteration of extent=48. A new testcase with single var depict this.

res = tvm.arith.detect_iter_map([i], var_dom([(i, 48)]), tvm.tir.all(i < 10))
assert_iter_sum_pattern(res[0], 10, 0)

If we fix (2), than the case below and several general unittests on tir will fail because TryFuseIters can not fully match the predicate constraint even all iters are visisted.

res = tvm.arith.detect_iter_map(
        [i, j, k],
        var_dom([(i, 128), (j, 128), (k, 128)]),
        tvm.tir.all(i * 16384 + j * 128 + k < 100),
    )

The workaround currently is that if we fail to fully match the predicate constraint, but a suffix of predicate is matched successfully, step into fallback behaviors and skip this constraint. For example, if target iter expr is k or j*128+k or i*128+j(fix me, this case will just not start a constraint match), the predicate i * 16384 + j * 128 + k < 100 will be skipped.

Unfortunately, the originated region_cover issue now seems can not be fully solved by this PR. I think the issue (2) make the corresponding case incorrectly passed previously. There is only one iteration except free vars in hh_0 * 8 - 1 + ax0 so the produced region is inferred regardless of predicate. A new case is added to intset test (previous case has two iters).

cc @spectrometerHBH

wrongtest-intellif · 2021-12-22T05:12:18Z

For region cover problem, there are two notes:

the consumer block has a conditional read accesses, so the T.reads() annotation of consumed region is larger than actual consumed.
the region_cover check in state.cc actually check produce^buffer_shape can cover consume^buffer_shape, even though, the analyzer can not prove it. I can manually prove it but it depend on a partition of free variable's range.

spectrometerHBH · 2021-12-23T02:16:19Z

Would be great if you can fix the CI :)

spectrometerHBH · 2021-12-23T04:26:31Z

src/arith/iter_affine_map.cc

+            if (i == 0 && num_visited == visited.size()) {
+              // if match failed because of iterations are used out instead of scale mismatch,
+              // and all used iters are visited during current match round, fallback to skip the
+              // constraint. Example: exprs = [i * 2 + j, k], i in [0, 3), j in [0, 2), k in [0, 4)


I think we should fail on such a case since there exists an unexpected predicate.

cc @wrongtest

spectrometerHBH · 2021-12-23T04:30:10Z

the region_cover check in state.cc actually check produce^buffer_shape can cover consume^buffer_shape, even though, the analyzer can not prove it. I can manually prove it but it depend on a partition of free variable's range.

Would you mind providing the produce^buffer_shape and consume^buffer_shape calculated by the region cover check? That would be a great help.

spectrometerHBH

Can be merged if we fail on the predicate unmatch problem.

spectrometerHBH · 2021-12-24T17:14:34Z

cc @Hzfengsy @junrushao1994

Hzfengsy · 2021-12-25T04:13:35Z

Thanks a lot @wrongtest and @spectrometerHBH

…9699) * Enable freevars, iter lowerbound and diagnostics in affine utility * fix lint issues and compare bug * update to use iter shift instead of itermark min for lowerbound * add testcase of fused iters sum with multiple lowerbounds * add more affine check testcases, fix bug for single iter and duplicate constraints on iter * add a newline to comment * forbidden predicate unmatch Co-authored-by: baoxinqi <[email protected]>

wrongtest-intellif requested review from areusch, comaniac, Hzfengsy, icemelon, jroesch, junrushao, kparzysz-quic, masahi, merrymercy, tqchen, vinx13, yzhliu and ZihengJiang as code owners December 10, 2021 07:21

wrongtest-intellif force-pushed the enforce_affine_utility branch 2 times, most recently from 1ef2e7d to 1b1852e Compare December 10, 2021 09:33

wrongtest-intellif changed the title ~~[TIR] Affine utility support freevars, iter lowerbound and diagnostics~~ [TIR] Affine utility support iter lowerbound and diagnostics Dec 17, 2021

wrongtest-intellif force-pushed the enforce_affine_utility branch from e642a84 to 21ccbf8 Compare December 17, 2021 12:12

spectrometerHBH requested changes Dec 21, 2021

View reviewed changes

baoxinqi and others added 5 commits December 22, 2021 12:04

Enable freevars, iter lowerbound and diagnostics in affine utility

df002d2

fix lint issues and compare bug

d86997d

update to use iter shift instead of itermark min for lowerbound

a4625a6

add testcase of fused iters sum with multiple lowerbounds

899e2eb

add more affine check testcases, fix bug for single iter and duplicat…

79b7277

…e constraints on iter

wrongtest-intellif force-pushed the enforce_affine_utility branch from 7ea21ef to 79b7277 Compare December 22, 2021 04:16

spectrometerHBH requested changes Dec 23, 2021

View reviewed changes

add a newline to comment

d71bd24

wrongtest-intellif force-pushed the enforce_affine_utility branch from fd18d84 to d71bd24 Compare December 23, 2021 10:33

spectrometerHBH approved these changes Dec 24, 2021

View reviewed changes

forbidden predicate unmatch

0337ffa

spectrometerHBH approved these changes Dec 24, 2021

View reviewed changes

Hzfengsy approved these changes Dec 25, 2021

View reviewed changes

Hzfengsy merged commit e2dcba2 into apache:main Dec 25, 2021

driazati mentioned this pull request Jul 14, 2022

TVM v0.9.0.rc0 Release Candidate Notes #12102

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[TIR] Affine utility support iter lowerbound and diagnostics #9699

[TIR] Affine utility support iter lowerbound and diagnostics #9699

wrongtest-intellif commented Dec 10, 2021

Hzfengsy commented Dec 10, 2021

junrushao commented Dec 14, 2021

spectrometerHBH commented Dec 16, 2021 •

edited

Loading

wrongtest-intellif commented Dec 17, 2021

wrongtest-intellif commented Dec 20, 2021

spectrometerHBH Dec 21, 2021

spectrometerHBH Dec 21, 2021 •

edited

Loading

spectrometerHBH Dec 21, 2021

wrongtest-intellif Dec 21, 2021

spectrometerHBH Dec 21, 2021

spectrometerHBH Dec 21, 2021 •

edited

Loading

wrongtest-intellif commented Dec 21, 2021

spectrometerHBH commented Dec 21, 2021 •

edited

Loading

wrongtest-intellif commented Dec 22, 2021

wrongtest-intellif commented Dec 22, 2021

spectrometerHBH commented Dec 23, 2021

spectrometerHBH Dec 23, 2021

spectrometerHBH Dec 23, 2021

spectrometerHBH commented Dec 23, 2021

spectrometerHBH left a comment

spectrometerHBH commented Dec 24, 2021

Hzfengsy commented Dec 25, 2021

[TIR] Affine utility support iter lowerbound and diagnostics #9699

[TIR] Affine utility support iter lowerbound and diagnostics #9699

Conversation

wrongtest-intellif commented Dec 10, 2021

Hzfengsy commented Dec 10, 2021

junrushao commented Dec 14, 2021

spectrometerHBH commented Dec 16, 2021 • edited Loading

wrongtest-intellif commented Dec 17, 2021

wrongtest-intellif commented Dec 20, 2021

spectrometerHBH Dec 21, 2021

Choose a reason for hiding this comment

spectrometerHBH Dec 21, 2021 • edited Loading

Choose a reason for hiding this comment

spectrometerHBH Dec 21, 2021

Choose a reason for hiding this comment

wrongtest-intellif Dec 21, 2021

Choose a reason for hiding this comment

spectrometerHBH Dec 21, 2021

Choose a reason for hiding this comment

spectrometerHBH Dec 21, 2021 • edited Loading

Choose a reason for hiding this comment

wrongtest-intellif commented Dec 21, 2021

spectrometerHBH commented Dec 21, 2021 • edited Loading

wrongtest-intellif commented Dec 22, 2021

wrongtest-intellif commented Dec 22, 2021

spectrometerHBH commented Dec 23, 2021

spectrometerHBH Dec 23, 2021

Choose a reason for hiding this comment

spectrometerHBH Dec 23, 2021

Choose a reason for hiding this comment

spectrometerHBH commented Dec 23, 2021

spectrometerHBH left a comment

Choose a reason for hiding this comment

spectrometerHBH commented Dec 24, 2021

Hzfengsy commented Dec 25, 2021

spectrometerHBH commented Dec 16, 2021 •

edited

Loading

spectrometerHBH Dec 21, 2021 •

edited

Loading

spectrometerHBH Dec 21, 2021 •

edited

Loading

spectrometerHBH commented Dec 21, 2021 •

edited

Loading