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 2021/02/01 19:26:27 UTC
[GitHub] [tvm] masahi opened a new pull request #7385: [TIR] Add additional termination condition to For node to enable While loop like feature
masahi opened a new pull request #7385:
URL: https://github.com/apache/tvm/pull/7385
This is my proposed solution to add While loop like feature to TIR, in the simplest, the least invasive way. It generalizes the For node termination condition from
```
loop_var < extent
```
to
```
loop_var < extent && test
```
Using this, we can write binary search as follows (see the complete test case, which implements numpy `searchsorted` function, [here](https://github.com/masahi/tvm/blob/a1b2c4a57e136186165e54bb1148faa44ad0a899/tests/python/unittest/test_tir_ir_builder.py#L176)).
```
lo[0] = 0
hi[0] = n
v = Bptr[i]
num_loop = int(np.log2(n)) + 1
with ib.for_range(0, num_loop, test=(lo[0] < hi[0])) as _:
mid = lo[0] + tvm.tir.floordiv(hi[0] - lo[0], 2).astype("int32")
with ib.if_scope(Aptr[mid] < v):
lo[0] = mid + 1
with ib.else_scope():
hi[0] = mid
Cptr[i] = lo[0]
```
My motivation was to improve GPU NMS performance using while loop, and it indeed did:
NMS workload from PyTorch MaskRCNN:
```
Without while loop (current main): 4.11 milli sec
With while loop (my branch): 3.66 milli sec
```
And a crazy 120000 box + 100 max_out_size NMS workload from TF MaskRCNN. The difference is huge because the # of iterations changed from 120000 to 100
```
Without while loop (current main): 51.31 milli sec
With while loop (my branch): 17.63 milli sec
```
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen commented on pull request #7385: [TIR] Add additional termination condition to For node to enable While loop like feature
Posted by GitBox <gi...@apache.org>.
tqchen commented on pull request #7385:
URL: https://github.com/apache/tvm/pull/7385#issuecomment-771192686
Thanks @masahi . I think it is great to enable support for some form of While loop.
It would be great to have an RFC thread discussing the alternatives in the IR node design. Since the IR node design can impact the general ability to do analysis. For example, I can see two possible variants
- V0: Put the condition into the for loop as it is (the current approach)
- V1: Introduce a separate While node for while loops
V0 means the for loop is somewhat overloaded for both while and For. On one hand it brings the benefit of richer semantics and the minimum set of changes to enable such feature.
This does mean that the visitors to For would need to handle the semantics of break. Given that the for loop is used for regular interval analysis, it could be beneficial to distinguish between "regular structured loop" vs "un-structured loop", unless the condition testing also offers some analysis benefits
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7385: [TIR] Add additional termination condition to For node to enable While loop like feature
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7385:
URL: https://github.com/apache/tvm/pull/7385#issuecomment-771926977
The rfc posted https://discuss.tvm.apache.org/t/rfc-add-while-loop-node-to-tir/9028
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7385: [TIR] Add additional termination condition to For node to enable While loop like feature
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7385:
URL: https://github.com/apache/tvm/pull/7385#issuecomment-771210096
@tqchen Yes, what you said totally makes sense to me. As you said, this is a minimal-change solution, but I think ideally we want a separate `While` node. Since while loop only makes sense for sequential loop (I think), I think it is better to decouple a simpler `While` node that doesn't need any analysis from heavy-duty `For` node.
I'll send a RFC, sure.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi closed pull request #7385: [TIR] Add additional termination condition to For node to enable While loop like feature
Posted by GitBox <gi...@apache.org>.
masahi closed pull request #7385:
URL: https://github.com/apache/tvm/pull/7385
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen edited a comment on pull request #7385: [TIR] Add additional termination condition to For node to enable While loop like feature
Posted by GitBox <gi...@apache.org>.
tqchen edited a comment on pull request #7385:
URL: https://github.com/apache/tvm/pull/7385#issuecomment-771192686
Thanks @masahi . I think it is great to enable support for some form of While loop.
It would be great to have an RFC thread discussing the alternatives in the IR node design. Since the IR node design can impact the general ability to do analysis and will impact how would we engineer future transformation passes.
For example, I can see two possible variants:
- V0: Put the condition into the for loop as it is (the current approach)
- V1: Introduce a separate While node for while loops
V0 means the for loop is somewhat overloaded for both while and For. On one hand it brings the benefit of richer semantics and the minimum set of changes to enable such feature.
This does mean that the visitors to For would need to handle the semantics of break. Given that the for loop is used for regular interval analysis, it could be beneficial to distinguish between "regular structured loop" vs "un-structured loop", unless the condition testing also offers some analysis benefits.
I think this PoC is a great start and would be great to have a design discussion over the RFC
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen commented on a change in pull request #7385: [TIR] Add additional termination condition to For node to enable While loop like feature
Posted by GitBox <gi...@apache.org>.
tqchen commented on a change in pull request #7385:
URL: https://github.com/apache/tvm/pull/7385#discussion_r568169870
##########
File path: include/tvm/tir/stmt.h
##########
@@ -802,6 +802,9 @@ class ForNode : public StmtNode {
ForKind kind;
/*! \brief The body of the for loop. */
Stmt body;
+ /*! \brief The additional termination condition of the for loop. */
+ Optional<PrimExpr> test;
Review comment:
It would be helpful to have a RFC discussion, since different strategies changes to the IR can have different implications
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7385: [TIR] Add additional termination condition to For node to enable While loop like feature
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7385:
URL: https://github.com/apache/tvm/pull/7385#issuecomment-771926977
The rfc posted https://discuss.tvm.apache.org/t/rfc-add-while-loop-node-to-tir/9028
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7385: [TIR] Add additional termination condition to For node to enable While loop like feature
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7385:
URL: https://github.com/apache/tvm/pull/7385#issuecomment-775922122
TIR While node added in https://github.com/apache/tvm/pull/7425
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org