SiriusNEO opened a new pull request, #14928:
URL: https://github.com/apache/tvm/pull/14928
The `DefaultGPUSchedule` pass aims to provide a simplest schedule to
generate a valid GPU kernel. Its logic is just to fuse all data parallel loops
together and then to split them according to `max_thread_per_block` and
`max_threadblocks`. However, currently this pass can not handle the scalar
program well such as:
```
@I.ir_module
class Add:
@R.function
def main(x: R.Tensor((), "int64")) -> R.Tensor((), "int64"):
gv: R.Tensor((), "int64") = R.add(x, R.const(1, "int64"))
return gv
```
When we lower `Add`, it will create a prim func like:
```
@tvm.script.ir_module
class Before:
@T.prim_func
def add(rxplaceholder: T.Buffer((), "int64"), T_add: T.Buffer((),
"int64")):
T.func_attr({"tir.noalias": T.bool(True)})
with T.block("T_add"):
vi = T.axis.spatial(1, T.int64(0))
T.reads(rxplaceholder[()])
T.writes(T_add[()])
T_add[()] = rxplaceholder[()] + T.int64(1)
```
Note that `te/operation/create_primfunc.cc:L321` will create a dummy iter
var for no loops case. So when running `DefaultGPUSchedule`, it will report
`Check failed: loops.size() == iters.size() (1 vs. 0)`.
This PR fixes this case by adding a unit loop.
--
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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
