roastduck commented on a change in pull request #5382:
URL: https://github.com/apache/incubator-tvm/pull/5382#discussion_r420230912
##########
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:
As for the reason of making the message a warning rather than an error,
consider 3 situations:
1. One uses `threadIdx.x` only. No `threadIdx.y` or `threadIdx.z`, no
warning is shown.
2. One uses `threadIdx.x` and `threadIdx.y`, and the extent of `threadIdx.x`
is 32 (warp size), which is a common use case. Now a warning is shown, but the
code is correct.
3. One uses `threadIdx.x` and `threadIdx.y`, but the extent of `threadIdx.x`
< 32, now a warning is shown, and the code is wrong.
We can still proceed in Situation 2.
----------------------------------------------------------------
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]
