MasterJH5574 commented on a change in pull request #8716:
URL: https://github.com/apache/tvm/pull/8716#discussion_r688905123
##########
File path: src/tir/transforms/flatten_buffer.cc
##########
@@ -140,7 +140,10 @@ class BufferFlattener : public StmtExprMutator {
/*var=*/std::move(var),
/*iter_type=*/IterVarType::kThreadIndex,
/*thread_tag=*/thread_tag);
- String attr_key = thread_tag == "vthread" ? attr::virtual_thread :
attr::thread_extent;
+ String attr_key = (thread_tag == "vthread" || thread_tag == "vthread.x" ||
+ thread_tag == "vthread.y" || thread_tag == "vthread.z")
+ ? attr::virtual_thread
+ : attr::thread_extent;
Review comment:
I find myself not understanding it...
Take `threadIdx.x` as an example: when we bind two different loops to
`threadIdx.x` (both in TE and TIR) and then lower the schedule, there will be
two `IterVar`s and two `AttrStmt`, such as
```
attr [IterVar(threadIdx.x: int32, (nullptr), "ThreadIndex", "threadIdx.x")]
"thread_extent" = 32;
attr [IterVar(threadIdx.x_1: int32, (nullptr), "ThreadIndex",
"threadIdx.x")] "thread_extent" = 32;
```
And the occurrences of the two different loops were substituted to
`threadIdx.x` and `threadIdx.x_1` respectively. Then it is the code generation
phase which turns the two `IterVar`s to the real `threadIdx.x` in CUDA kernel.
So if we bind two loops to `vthread.x`, we want the two loop vars to always
take the same value. Maintaining a map and using a same `IterVar` is
inconsistent with the way we deal with `threadIdx.x`. So do you think we should
insist on making changes here, instead of changing the codegen part 🤔.
--
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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
