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/04/11 23:20:16 UTC
[GitHub] [tvm] altanh opened a new pull request, #10969: [Bug] hotfix edge case for broadcast_to const shape int64
altanh opened a new pull request, #10969:
URL: https://github.com/apache/tvm/pull/10969
When the target shape of `broadcast_to` is an `int64` Constant, something breaks in lowering for CUDA (the thread extent ends up being `int64` and causes a mismatch error with the iteration variable). I could only get the error to show up when following up with a `sum` op.
This is a hotfix to get around the problem for now while we dig into the root cause.
cc @Lunderberg @mbrookhart
--
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] ganler commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
ganler commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097322856
@altanh I see. According to previous fixes (https://github.com/apache/tvm/pull/10172 https://github.com/apache/tvm/pull/9582 https://github.com/apache/tvm/pull/10519 https://github.com/apache/tvm/pull/10571 https://github.com/apache/tvm/pull/10584), it seems to be more compatible if we just change the data type when constructing a new IR node. And https://github.com/apache/tvm/pull/10983 did pass the CI (but I made further edits to adapt your suggestion about `vi_dtype.bits() < var.dtype().bits()`).
--
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] altanh commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
altanh commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097232343
Yep I think we're solving the same problem haha
--
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] altanh commented on a diff in pull request #10969: [Bug] hotfix edge case for broadcast_to const shape int64
Posted by GitBox <gi...@apache.org>.
altanh commented on code in PR #10969:
URL: https://github.com/apache/tvm/pull/10969#discussion_r848709195
##########
python/tvm/relay/op/transform.py:
##########
@@ -836,7 +836,7 @@ def broadcast_to(data, shape):
The resulting tensor.
"""
if isinstance(shape, Constant):
- shape = list(shape.data.numpy())
+ shape = [int(i) for i in shape.data.numpy()]
Review Comment:
It's possible, although the default handler for a python list of ints would hit the same problem. In either case, I'll push a proper fix for the reduction schedule on CUDA (forcibly cast the extent to `int32`, since that's needed for CUDA anyway if I understand correctly).
Do you think I should wrap ints using numpy `int64` to avoid this problem? I'm slightly worried it will break assumptions elsewhere.
--
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] ganler commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
ganler commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097218004
@masahi @altanh I just tried your test case and https://github.com/apache/tvm/pull/10983 can fix it.
--
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] masahi commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
masahi commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097205281
There is also a PR https://github.com/apache/tvm/pull/10983 touching `narrow_datatype.cc`. Can you discuss which patch to merge? cc @ganler
--
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] ganler commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
ganler commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097232811
> interesting, I'm confused how it does the cast when the IterVar is int32 but the extent is int64 (since then `vi_dtype.bits() < var.dtype().bits()` is false), there must be some interaction with the `DataTypeVisitor` that I don't understand. I'll defer to @Lunderberg, I'm fine with either fix
Very great point and this is where I think your PR is better than mine as I simply assume that extend.dtype.bits <= var.dtype.bits.
--
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] ganler commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
ganler commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097231489
@altanh I just tried your PR, which also fix the issue from a PyTorch-generated ONNX model in https://github.com/apache/tvm/pull/10983. So it seems the main issue is from the thread extend?
--
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] Lunderberg commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
Lunderberg commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1098135030
I'm fine with either fix as well, and thank you for tracking it down!
--
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] altanh commented on a diff in pull request #10969: [Bug] hotfix edge case for broadcast_to const shape int64
Posted by GitBox <gi...@apache.org>.
altanh commented on code in PR #10969:
URL: https://github.com/apache/tvm/pull/10969#discussion_r848710144
##########
python/tvm/relay/op/transform.py:
##########
@@ -836,7 +836,7 @@ def broadcast_to(data, shape):
The resulting tensor.
"""
if isinstance(shape, Constant):
- shape = list(shape.data.numpy())
+ shape = [int(i) for i in shape.data.numpy()]
Review Comment:
I guess one possible option is to change the behavior `tvm.runtime.convert` to check for overflow and use `int64` when necessary
--
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] altanh commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
altanh commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097240002
yeah I see those tests failures now... my operating assumption is that for the `thread_extent` attr, it's only for GPUs and currently GPUs only support int32 IterVar (e.g. threadIdx.x) and extents. So maybe I should just go ahead and force everything to be int32 here? I also assumed that the extents will be concrete, but maybe that's too strict too?
--
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] ganler commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
ganler commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097242207
If that's the case, probably we should just force the thread index to be int32. i.e., removing CanProveLess.
--
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] altanh commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
altanh commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097230170
interesting, I'm confused how it does the cast when the IterVar is int32 but the extent is int64 (since then `vi_dtype.bits() < var.dtype().bits()` is false), there must be some interaction with the `DataTypeVisitor` that I don't understand. I'll defer to @Lunderberg, I'm fine with either fix
--
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] Lunderberg commented on a diff in pull request #10969: [Bug] hotfix edge case for broadcast_to const shape int64
Posted by GitBox <gi...@apache.org>.
Lunderberg commented on code in PR #10969:
URL: https://github.com/apache/tvm/pull/10969#discussion_r848520342
##########
python/tvm/relay/op/transform.py:
##########
@@ -836,7 +836,7 @@ def broadcast_to(data, shape):
The resulting tensor.
"""
if isinstance(shape, Constant):
- shape = list(shape.data.numpy())
+ shape = [int(i) for i in shape.data.numpy()]
Review Comment:
Because tvm.runtime.convert [handles all integer types as int32](https://github.com/apache/tvm/blob/main/python/tvm/runtime/object_generic.py#L121), would this cause issues with arrays larger than 4 GB? I don't think we have many of those in practice, but I think it could then cause a similar issue.
--
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] ganler commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
ganler commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097235787
But I am not sure if `CanProveLess` is able to consider the integer bits. That said, if the `var` of `IterVal` is `int64` and `extend` is `int32`. They are both symbolic that `CanProveLess` seems to return `false` but we can actually cast `extend` to `int64` right?
--
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] masahi closed pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
masahi closed pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
URL: https://github.com/apache/tvm/pull/10969
--
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] masahi commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
masahi commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097839425
The other PR https://github.com/apache/tvm/pull/10983 has been merged
--
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] altanh commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
altanh commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097211893
> There is also a PR (looks related) #10983 touching `narrow_datatype.cc`. Can you discuss which patch to merge? cc @ganler
haha great timing, @ganler do you mind trying the test case I added here on your branch? I suspect your fix might miss the case of mismatched int32 var with int64 extent.
--
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] ganler commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
ganler commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097224456
OK, IMHO, the difference between this fix and https://github.com/apache/tvm/pull/10983 is that:
- this one fixes the bug in Visitor by doing a casting for thread extend;
- https://github.com/apache/tvm/pull/10983 fixes a bigger scope of similar bugs as it does casting in the Rewrite phase and does not specifically target `attr::thread_extent` (can be any other incompatible integer attributes).
--
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] altanh commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
altanh commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097248929
let me try just removing the CanProveLess and see if it can pass CI
--
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] altanh commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
altanh commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097318502
@ganler it seems like I'm breaking some assumptions elsewhere with this change- some unit tests seem to want the thread extents to explicitly be int64. Maybe we should just go with your change if you're happy with it
--
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] altanh commented on pull request #10969: [Bug] narrow thread extents to 32 bits for GPU lowering
Posted by GitBox <gi...@apache.org>.
altanh commented on PR #10969:
URL: https://github.com/apache/tvm/pull/10969#issuecomment-1097328029
I'll close this PR once yours passes CI, thx for tracking down the other relevant PRs!
--
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