Don't replace reduction init axis with new axis if bound to a thread. #3408
Conversation
|
Thanks for the PR, sorry for the delayed review as many of folks I know in the community are traveling to FCRC recently. Can you please add a regression test case to prevent it from happening again? |
|
Sure, I should be able to base it off the repro in the issue. I'll put it on my todo list but may not get to it for a while. I've been prototyping some kernels in TE, which is why the few small bug fixes. When I get a break I'll come back and write some tests. |
|
@csarofeen can you add a test case? |
|
ping @csarofeen |
|
Sorry this was so delayed, was actually preparing a talk on TE for CUDA code generation. Simple test added @tqchen |
|
I can't tell if the linter error is from my change or not. I tried to get it running locally but I got a docker build error. |
|
|
||
| mo, _ = s[B].split(B.op.axis[0], 32) | ||
| s[B].bind(mo, tvm.thread_axis("blockIdx.x")) | ||
| fcuda = tvm.build(s, [A, B], "cuda") |
There was a problem hiding this comment.
Add guard to this test (see Line 143). Because some test environments do not have a gpu.
csarofeen
force-pushed
the
reduce_bind_fix
branch
from
8484a5e
to
a077503
Compare
|
@merrymercy @tqchen I rebased at top of tree and added 2 tests to cover: #3444 and #3382 hope that's okay. |
|
Ping. If this looks good I can rebase so it can be merged. |
|
Thanks @csarofeen please rebase |
csarofeen
force-pushed
the
reduce_bind_fix
branch
from
a077503
to
852f3db
Compare
|
Looks like this got stuck, is there a way to retry the tests? |
|
Thanks @csarofeen , this PR is now merged! |
tqchen
added
status: accepted
and removed
status: need test case
…apache#3408) * Don't replace reduction init axis with new axis if bound to a thread. * Linter. * Reduce bind test case. * Guard test on CUDA support. * [CUDA TE TESTS] Add rfactor predicate test, add global bx and tx. * [CUDA TE TESTS] Add loop partition test for simple rfactor case.
…apache#3408) * Don't replace reduction init axis with new axis if bound to a thread. * Linter. * Reduce bind test case. * Guard test on CUDA support. * [CUDA TE TESTS] Add rfactor predicate test, add global bx and tx. * [CUDA TE TESTS] Add loop partition test for simple rfactor case.
…apache#3408) * Don't replace reduction init axis with new axis if bound to a thread. * Linter. * Reduce bind test case. * Guard test on CUDA support. * [CUDA TE TESTS] Add rfactor predicate test, add global bx and tx. * [CUDA TE TESTS] Add loop partition test for simple rfactor case.
Fixes #3407