yongfeng-nv commented on a change in pull request #5382:
URL: https://github.com/apache/incubator-tvm/pull/5382#discussion_r420543138
##########
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:
threadIdx.x means a lower level of hierarchy than threadIdx.y and
threadIdx.z for "warp" scope -- "warp" memory is local to threadIdx.y and
threadIdx.z, but shared by threadIdx.x. One idea for you to consider is to
change threadIdx.x's rank to 2 and keep threadIdx.y and threadIdx.z's rank 1.
It will avoid special casing "warp" here, but may cause other regression.
> In fact, things are more complicated because threadIdx can be less than
the warp size. To determine whether an index is inside or outside a warp, we
need to know the extent of that index. But this information is not available in
MakeLoopNest.
Does `dom` and/or `bind_iv->dom` have the extent for the checks?
> Here is an example for the 2nd situation: Consider extent of threadIdx.x =
32, extent of threadIdx.y = 4. Now there are 4 warps, each consisting of 32
threads.
This example clarifies a lot. Is it possible to make it a test to cover the
"warp" specific code here?
----------------------------------------------------------------
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]
