junrushao1994 commented on code in PR #11177:
URL: https://github.com/apache/tvm/pull/11177#discussion_r862148840
##########
src/meta_schedule/postproc/rewrite_unbound_block.cc:
##########
@@ -194,45 +121,9 @@ bool RewriteUnboundBlockNode::Apply(const tir::Schedule&
sch) {
for (const auto& kv : unbound_blocks) {
tir::StmtSRef block_sref = kv.first;
String global_var_name = kv.second;
- int fuse_first_num = 0;
- tir::BindType bind_type = tir::GetBindType(block_sref, &fuse_first_num);
- if (bind_type == tir::BindType::kNoBind) {
- continue;
- }
BlockRV block_rv = GetRVFromSRef(sch, block_sref, global_var_name);
- Array<LoopRV> loop_rvs = sch->GetLoops(block_rv);
- LoopRV fused = sch->Fuse({loop_rvs.begin(), loop_rvs.begin() +
fuse_first_num});
- if (bind_type == tir::BindType::kBindBlock) {
- sch->Bind(fused, "blockIdx.x");
- } else if (bind_type == tir::BindType::kBindBlockThread) {
- int64_t extent_size = 0;
- Array<LoopRV> splits;
- if (const int64_t* extent_ptr =
tir::GetLoopIntExtent(sch->Get(fused).get())) {
- extent_size = *extent_ptr;
- if (extent_size > max_threadblock_ * max_num_threads_) {
- splits =
- sch->Split(fused, {NullOpt, Integer(max_threadblock_),
Integer(max_num_threads_)});
- ICHECK_EQ(splits.size(), 3);
- sch->Reorder({splits[1], splits[2], splits[0]});
- sch->Bind(splits[1], "blockIdx.x");
- sch->Bind(splits[2], "threadIdx.x");
- } else {
- ICHECK_NE(extent_size, 0);
- splits = sch->Split(
- fused,
- {NullOpt,
Integer(std::min(static_cast<int64_t>(max_num_threads_), extent_size))});
- ICHECK_EQ(splits.size(), 2);
- sch->Bind(splits[0], "blockIdx.x");
- sch->Bind(splits[1], "threadIdx.x");
- }
- } else {
- // loop is dynamic, returns nullptr
- splits = sch->Split(fused, {NullOpt, Integer(max_num_threads_)});
- ICHECK_EQ(splits.size(), 2);
- sch->Bind(splits[0], "blockIdx.x");
- sch->Bind(splits[1], "threadIdx.x");
- }
- }
+ BindThreadsForUnboundBlock(sch, block_rv, max_threadblock_,
max_num_threads_,
Review Comment:
Looks like we are introducing randomness in postproc, which violates the
system's assumption where postproc should be deterministic. Any ideas how to
fix this?
--
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]
