yongfeng-nv commented on a change in pull request #5382:
URL: https://github.com/apache/incubator-tvm/pull/5382#discussion_r421003668
##########
File path: src/te/operation/op_util.cc
##########
@@ -164,9 +164,21 @@ MakeLoopNest(const Stage& stage,
value_map[iv] = dom->min;
} else {
runtime::ThreadScope ts =
runtime::ThreadScope::make(bind_iv->thread_tag);
- if (stage->scope == "" || stage->scope == "warp" ||
+ if (stage->scope == "" ||
static_cast<int>(runtime::StorageScope::make(stage->scope).rank)
<= ts.rank) {
value_map[iv] = var;
+ } else if (stage->scope == "warp" && ts.rank == 1) {
+ // To determine whether a thread index is inside or outside a warp,
we need
+ // to know the thread extent. We leave a warning for now.
+ if (ts.dim_index == 0) {
+ value_map[iv] = var;
+ } else {
Review comment:
> Yes, dom does. But where can we get the warp size? Warp size is not a
constant value, it is 32 for NVIDIA GPU and 64 for AMD GPU. The warp size
information is stored in Target class. Since MakeLoopNest is designed to be a
target-irrelevant function, I don't think it a good idea to add an argument to
specify the target.
I agree with you about `MakeLoopNest ` being target independent. My first
thought is to defer the check and `threadIdx.x` vs. `threadIdx.y/z` handling
till LowerWarpMemory(), such as adjusting indices in WarpAccessRewriter(),
since it does target specific "warp" lowering. Maybe substitute
`threadIdx.y/z` with `0` in the group indices for the supported cases and give
error otherwise? However, it seems not the case, as your 3rd situation ends
with incorrect code instead of an error from LowerWarpMemory(). But I don't
know the reason.
> Actually I did try to run this example, but TVM somehow generated a
correct code. It only led to a wrong boundary checking bug in a very complex
program.
I see. Does your simplified case trigger the warning? If so, checking for
the warning can guard your changes from being accidentally deleted or skipped.
----------------------------------------------------------------
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:
[email protected]
