junrushao1994 commented on a change in pull request #8716:
URL: https://github.com/apache/tvm/pull/8716#discussion_r688988611
##########
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:
The particular case in my mind is [cooperative
fetching](https://tvm.apache.org/docs/tutorials/optimize/opt_conv_cuda.html#cooperative-fetching).
In this example, `threadIdx.x` is bound twice in nested loops: The outer loop
specifies the kernel launch condition, while the inner one specifies how data
are fetched cooperatively (i.e. which thread is responsible for which data)
You may play with the tutorial by printing out the TIR:
```python
print(tvm.script.asscript(tvm.lower(s, [A, W, B], simple_mode=True)))
```
We will see there is only one threadIdx.x created
--
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]
