You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by lm...@apache.org on 2020/03/10 15:18:12 UTC
[incubator-tvm] branch
revert-4931-split-node-min-range-only-at-final-pass created (now 13588f0)
This is an automated email from the ASF dual-hosted git repository.
lmzheng pushed a change to branch revert-4931-split-node-min-range-only-at-final-pass
in repository https://gitbox.apache.org/repos/asf/incubator-tvm.git.
at 13588f0 Revert "Tighten split's extent (#4931)"
This branch includes the following new commits:
new 13588f0 Revert "Tighten split's extent (#4931)"
The 1 revisions listed above as "new" are entirely new to this
repository and will be described in separate emails. The revisions
listed as "add" were already present in the repository and have only
been added to this reference.
[incubator-tvm] 01/01: Revert "Tighten split's extent (#4931)"
Posted by lm...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
lmzheng pushed a commit to branch revert-4931-split-node-min-range-only-at-final-pass
in repository https://gitbox.apache.org/repos/asf/incubator-tvm.git
commit 13588f0ac709165467472f261f351bd86a17dfe1
Author: Lianmin Zheng <li...@gmail.com>
AuthorDate: Tue Mar 10 08:17:59 2020 -0700
Revert "Tighten split's extent (#4931)"
This reverts commit 585f9ce6e7bef7d0e8902b1c1e55dcb3bbe84eed.
---
src/te/schedule/message_passing.cc | 76 +---------------------
.../unittest/test_schedule_bound_inference.py | 26 --------
2 files changed, 3 insertions(+), 99 deletions(-)
diff --git a/src/te/schedule/message_passing.cc b/src/te/schedule/message_passing.cc
index a7b2482..5b6fa86 100644
--- a/src/te/schedule/message_passing.cc
+++ b/src/te/schedule/message_passing.cc
@@ -51,66 +51,17 @@ void Update(std::unordered_map<IterVar, Range>* p_state,
}
}
-/*!
- * \param Upward propagating whether an IterVar derives at least one leaf IterVar that binds to
- * a thread.
- *
- * \param stage The stage to operate on.
- * \param p_state The propagation result of each IterVar.
- */
-void PassUpThreadBinding(const Stage& stage, std::unordered_map<IterVar, bool>* p_state) {
- auto bound_to_thread = [&stage](const IterVar& iv) {
- bool bound = false;
- auto it = stage->iter_var_attrs.find(iv);
- if (it != stage->iter_var_attrs.end()) {
- bound = (*it).second->bind_thread.defined();
- }
- return bound;
- };
-
- auto& state = *p_state;
- // Fill p_state with leaf itervars
- for (const IterVar& iv : stage->leaf_iter_vars) {
- state[iv] = bound_to_thread(iv);
- }
- // Traverse the graph bottom-up to propagate thread binding information
- for (size_t i = stage->relations.size(); i != 0; --i) {
- IterVarRelation rel = stage->relations[i - 1];
- if (const SplitNode* s = rel.as<SplitNode>()) {
- state[s->parent] = state[s->inner] || state[s->outer];
- } else if (const FuseNode* s = rel.as<FuseNode>()) {
- state[s->inner] = state[s->fused];
- state[s->outer] = state[s->fused];
- } else if (const RebaseNode* s = rel.as<RebaseNode>()) {
- state[s->parent] = state[s->rebased];
- } else if (rel.as<SingletonNode>()) {
- } else {
- LOG(FATAL) << "unknown relation type";
- }
- }
-}
-
void PassDownDomain(const Stage& stage,
std::unordered_map<IterVar, Range>* p_state,
arith::Analyzer* actx,
bool allow_missing) {
- auto ceil_div = [actx](const PrimExpr& a, const PrimExpr& b) {
+ auto ceil_div = [actx](PrimExpr a, PrimExpr b) {
if (actx->CanProve(indexmod(a, b) == 0)) {
return actx->Simplify(indexdiv(a, b));
}
return actx->Simplify(indexdiv(a + (b - 1), b));
};
- auto minimum_or_later = [actx](const PrimExpr& a, const PrimExpr& b) {
- if (actx->CanProve(a < b)) {
- return actx->Simplify(a);
- }
- return actx->Simplify(b);
- };
-
- std::unordered_map<IterVar, bool> dominating_thread;
- PassUpThreadBinding(stage, &dominating_thread);
-
auto& state = *p_state;
// forwar iteration on relations
for (IterVarRelation rel : stage->relations) {
@@ -121,35 +72,14 @@ void PassDownDomain(const Stage& stage,
}
CHECK(!state.count(r->inner));
const Range& range_parent = state.at(r->parent);
- // Tighten iv's extent to min(parent_extent, factor_or_nparts), only if all of the
- // following conditions are met:
- // 1. No leaf IterVar derived from iv binds to any thread. People may use split
- // to force an IterVar extent to match the number of allocated threads to fuse stages
- // that require different number of threads. We don't want to change these extents.
- // 2. allow_missing is false, i.e. that PassDownDomain is called by the final InferBound,
- // rather than by an early compiler phase, such as rfactor(). We don't want to tighten an
- // IterVar in an early phase allowing missing IterVars, because it may bind to a thread later.
- // 3. range_parent's extent is not 0. At lest one Topi test has a case where a tensor has one
- // zero-sized dimension. Split creates iv with a positive extent to avoid zero-extent
- // IterVar. We don't touch it.
- auto resolve_min_extent_for_split = [&](const IterVar& iv, const PrimExpr& factor_or_nparts) {
- return dominating_thread[iv] || allow_missing || is_zero(range_parent->extent)
- ? factor_or_nparts
- : minimum_or_later(range_parent->extent, factor_or_nparts);
- };
if (r->factor.defined()) {
Update(p_state, r->inner,
- Range::make_by_min_extent(
- 0, resolve_min_extent_for_split(r->inner, r->factor)),
- actx);
+ Range::make_by_min_extent(0, r->factor), actx);
Update(p_state, r->outer,
Range::make_by_min_extent(
0, ceil_div(range_parent->extent, r->factor)), actx);
} else {
- Update(p_state, r->outer,
- Range::make_by_min_extent(
- 0, resolve_min_extent_for_split(r->outer, r->nparts)),
- actx);
+ Update(p_state, r->outer, Range::make_by_min_extent(0, r->nparts), actx);
Update(p_state, r->inner,
Range::make_by_min_extent(
0, ceil_div(range_parent->extent, r->nparts)), actx);
diff --git a/tests/python/unittest/test_schedule_bound_inference.py b/tests/python/unittest/test_schedule_bound_inference.py
index edae527..484aa50 100644
--- a/tests/python/unittest/test_schedule_bound_inference.py
+++ b/tests/python/unittest/test_schedule_bound_inference.py
@@ -70,32 +70,6 @@ def test_bound3():
assert(bounds[A1.op.axis[0]].extent.value==32)
assert(bounds[A1.op.axis[1]].extent.value==16)
-def test_bound_split_ext_less_than_factor():
- m = 8
- I = te.placeholder((m,), name='I')
- EF = te.compute((m,), lambda i: I[i] * 2, name = "EF")
- E = te.compute((m,), lambda i: EF[i] * 2, name = "E")
- s = te.create_schedule([E.op])
- xo, xi = s[E].split(s[E].op.axis[0], factor = 32)
- s[EF].compute_at(s[E], xo)
-
- bounds = tvm.te.schedule.InferBound(s)
- assert isinstance(bounds, tvm.container.Map)
- assert bounds[xi].extent.value == m
-
-def test_bound_split_ext_less_than_naprts():
- m = 8
- I = te.placeholder((m,), name='I')
- EF = te.compute((m,), lambda i: I[i] * 2, name = "EF")
- E = te.compute((m,), lambda i: EF[i] * 2, name = "E")
- s = te.create_schedule([E.op])
- xo, xi = s[E].split(s[E].op.axis[0], nparts = 32)
- s[EF].compute_at(s[E], xo)
-
- bounds = tvm.te.schedule.InferBound(s)
- assert isinstance(bounds, tvm.container.Map)
- assert bounds[xo].extent.value == m
-
def test_bound_split_divisible():
m = te.var('m')
l = te.var('l')