.. code:: cpp
- rmap[split->inner] = Range::make_by_min_extent(0, split->factor)
- rmap[split->outer] = Range::make_by_min_extent(0, DivCeil(rmap[split->parent]->extent, split->factor))
+ rmap[split->inner] = Range::FromMinExtent(0, split->factor)
+ rmap[split->outer] = Range::FromMinExtent(0, DivCeil(rmap[split->parent]->extent, split->factor))
There is an opportunity here to tighten the bounds produced by InferBound, when ``split->factor`` does not evenly divide the parent's extent. Suppose the parent's extent is 20, and the split factor is 16. Then on the second iteration of the outer loop, the inner loop only needs to perform 4 iterations, not 16. If PassDownDomain could set the extent of ``split->inner`` to ``min(split->factor, rmap[split->parent]->extent - (split->outer * split->factor))``, then the extent of the inner variable would properly adapt, based on which iteration of the outer loop is being executed.
.. code:: cpp
- rmap[fuse->fused] = Range::make_by_min_extent(0, rmap[fuse->outer]->extent * rmap[fuse->inner]->extent)
+ rmap[fuse->fused] = Range::FromMinExtent(0, rmap[fuse->outer]->extent * rmap[fuse->inner]->extent)
InferRootBound
*/
class IntSet : public ObjectRef {
public:
- /*! \brief constructor */
- IntSet() {}
- // constructor from not container.
- explicit IntSet(ObjectPtr<Object> n) : ObjectRef(n) {}
- /*!
- * \brief access the internal node container
- * \return the pointer to the internal node container
- */
- const IntSetNode* operator->() const { return static_cast<const IntSetNode*>(get()); }
/*!
* \brief Find a range that covers the region.
* \param max_range The range to be covered.
* \return The covering range.
*/
- Range cover_range(Range max_range) const;
+ Range CoverRange(Range max_range) const;
/*! \return Lower bound of the set */
PrimExpr min() const;
/*! \return upper bound of the set */
PrimExpr max() const;
+ /*! \return The sign of the elements in the integer set */
+ SignType GetSignType() const;
/*! \return Whether the set represent nothing */
- bool is_nothing() const;
+ bool IsNothing() const;
/*! \return Whether the set represent everything */
- bool is_everything() const;
+ bool IsEverything() const;
/*! \return Whether the set is a single point */
- bool is_single_point() const;
+ bool IsSinglePoint() const;
/*! \return Whether the set is proved to be bigger than 0 */
- bool can_prove_positive() const;
+ bool CanProvePositive() const;
/*! \return Whether the set is proved to be smaller than 0 */
- bool can_prove_negative() const;
+ bool CanProveNegative() const;
/*! \return Whether the set is proved to be smaller than or equal to 0 */
- bool can_prove_non_positive() const;
+ bool CanProveNonPositive() const;
/*! \return Whether the set is proved to be larger than or equal to 0 */
- bool can_prove_non_negative() const;
- /*! \return The sign of the elements in the integer set */
- SignType sign_type() const;
+ bool CanProveNonNegative() const;
/*!
- * \brief The single point value, call only if is_single_point is true
+ * \brief The single point value, call only if IsSinglePoint is true
* \return The point value.
*/
- PrimExpr point_value() const;
+ PrimExpr PointValue() const;
/*!
* \brief Try to match IntSet with range r.
*
- * \note It is guanrateed that IntSet::range(r).match_range(r) == true
+ * \note It is guanrateed that IntSet::FromRange(r).MatchRange(r) == true
* \return true if we can prove they are the same.
*/
- bool match_range(const Range& r) const;
+ bool MatchRange(const tvm::Range& r) const;
/*! \return The set contains nothing */
- static IntSet nothing();
+ static IntSet Nothing();
/*! \return The set contains everything */
- static IntSet everything();
+ static IntSet Everything();
/*!
* \brief construct a point set.
* \param point The point in the set.
* \return construct a single point set
*/
- static IntSet single_point(PrimExpr point);
+ static IntSet SinglePoint(PrimExpr point);
/*!
* \brief construct a integer set from vector expression.
* \param vec The vector expression, can also be single point.
* \return The result set containing the indices in the vector.
*/
- static IntSet vector(PrimExpr vec);
+ static IntSet Vector(PrimExpr vec);
/*!
* \brief Construct a set representing a range.
* \param r The range
* \return constructed set.
*/
- static IntSet range(Range r);
+ static IntSet FromRange(tvm::Range r);
/*!
* \brief Construct a set representing a interval.
* \param min The minimum value of the interval.
* \param max The maximum value of the interval.
* \return constructed set.
*/
- static IntSet interval(PrimExpr min, PrimExpr max);
+ static IntSet Interval(PrimExpr min, PrimExpr max);
+
+ TVM_DEFINE_OBJECT_REF_METHODS(IntSet, ObjectRef, IntSetNode);
};
//-----------------------------------------------
* \param min The minimum range.
* \param extent The extent of the range.
*/
- static Range make_by_min_extent(PrimExpr min, PrimExpr extent);
+ static Range FromMinExtent(PrimExpr min, PrimExpr extent);
// declare range.
TVM_DEFINE_OBJECT_REF_METHODS(Range, ObjectRef, RangeNode);
};
_ffi_api.Range, begin, end)
@staticmethod
- def make_by_min_extent(min_value, extent):
+ def from_min_extent(min_value, extent):
"""Construct a Range by min and extent.
This constructs a range in [min_value, min_value + extent)
rng : Range
The constructed range.
"""
- return _ffi_api.range_by_min_extent(min_value, extent)
+ return _ffi_api.Range_from_min_extent(min_value, extent)
if _scope == 'global':
body = self.wrap_up_binds(body)
- _domain = [Range.make_by_min_extent(0, i) for i in _buf.shape]
+ _domain = [Range.from_min_extent(0, i) for i in _buf.shape]
_dtype = _buf.dtype
_true = tvm.runtime.convert(True)
body = tvm.tir.ProducerRealize(_buf, _domain, _true, body)
begin = idx.var
else:
begin = idx
- region.append(Range.make_by_min_extent(begin, 1))
+ region.append(Range.from_min_extent(begin, 1))
return region
if (operand.dtype().is_uint()) {
sign_operand = kPositive;
} else {
- sign_operand = expr_map_[operand].sign_type();
+ sign_operand = expr_map_[operand].GetSignType();
}
if (sign_operand == SignType::kNegative) {
void BoundDeducer::Relax() {
IntSet a = EvalSet(expr_, relax_map_);
IntSet b = EvalSet(result_, relax_map_);
- if (a.is_everything() || b.is_everything()) {
+ if (a.IsEverything() || b.IsEverything()) {
success_ = false;
return;
}
const std::unordered_map<const VarNode*, IntSet>& relax_map) {
BoundDeducer d(v, e, hint_map, relax_map);
d.Deduce();
- if (!d.success_) return IntSet::nothing();
+ if (!d.success_) return IntSet::Nothing();
PrimExpr min = neg_inf(), max = pos_inf();
if (d.comp_op == kEqual) {
min = d.result_;
} else {
max = d.result_;
}
- return IntSet::interval(min, max);
+ return IntSet::Interval(min, max);
}
// assuming e >= 0, deduce the bound of variable from it.
Region ret;
Range none;
for (size_t i = 0; i < bounds_.size(); ++i) {
- ret.push_back(arith::Union(bounds_[i]).cover_range(none));
+ ret.push_back(arith::Union(bounds_[i]).CoverRange(none));
}
return ret;
}
void VisitStmt_(const ForNode* op) final {
const VarNode* var = op->loop_var.get();
- dom_map_[var] = IntSet::range(Range::make_by_min_extent(op->min, op->extent));
+ dom_map_[var] = IntSet::FromRange(Range::FromMinExtent(op->min, op->extent));
StmtExprVisitor::VisitStmt_(op);
dom_map_.erase(var);
}
const IterVarNode* thread_axis = op->node.as<IterVarNode>();
CHECK(thread_axis);
const VarNode* var = thread_axis->var.get();
- dom_map_[var] = IntSet::range(Range(make_zero(op->value.dtype()), op->value));
+ dom_map_[var] = IntSet::FromRange(Range(make_zero(op->value.dtype()), op->value));
StmtExprVisitor::VisitStmt_(op);
dom_map_.erase(var);
} else {
// Quickly adapt to IntSet interface
// TODO(tqchen): revisit IntSet interface as well.
-Range IntSet::cover_range(Range max_range) const {
+Range IntSet::CoverRange(Range max_range) const {
IntSet temp;
Analyzer analyzer;
const IntervalSetNode* s_int = (*this).as<IntervalSetNode>();
CHECK(s_int != nullptr);
if (s_int->HasUpperBound() && s_int->HasLowerBound()) {
- return Range::make_by_min_extent(s_int->min_value,
- analyzer.Simplify(s_int->max_value + 1 - s_int->min_value));
+ return Range::FromMinExtent(s_int->min_value,
+ analyzer.Simplify(s_int->max_value + 1 - s_int->min_value));
}
return max_range;
}
return s_int->max_value;
}
-bool IntSet::is_nothing() const {
+bool IntSet::IsNothing() const {
const IntervalSetNode* s_int = (*this).as<IntervalSetNode>();
return (s_int && s_int->IsEmpty());
}
-bool IntSet::is_everything() const {
+bool IntSet::IsEverything() const {
const IntervalSetNode* s_int = (*this).as<IntervalSetNode>();
return (s_int && s_int->IsEverything());
}
-bool IntSet::is_single_point() const {
+bool IntSet::IsSinglePoint() const {
const IntervalSetNode* s_int = (*this).as<IntervalSetNode>();
return (s_int && s_int->IsSinglePoint());
}
-bool IntSet::can_prove_positive() const {
+bool IntSet::CanProvePositive() const {
Analyzer analyzer;
const IntervalSetNode* s_int = (*this).as<IntervalSetNode>();
return (s_int && is_positive_const(analyzer.Simplify(s_int->min_value)));
}
-bool IntSet::can_prove_negative() const {
+bool IntSet::CanProveNegative() const {
Analyzer analyzer;
const IntervalSetNode* s_int = (*this).as<IntervalSetNode>();
return (s_int && is_negative_const(analyzer.Simplify(s_int->max_value)));
}
-bool IntSet::can_prove_non_positive() const {
+bool IntSet::CanProveNonPositive() const {
Analyzer analyzer;
if (const auto* s_int = (*this).as<IntervalSetNode>()) {
auto max = analyzer.Simplify(s_int->max_value);
return false;
}
-bool IntSet::can_prove_non_negative() const {
+bool IntSet::CanProveNonNegative() const {
Analyzer analyzer;
if (const IntervalSetNode* s_int = (*this).as<IntervalSetNode>()) {
auto min = analyzer.Simplify(s_int->min_value);
return false;
}
-SignType IntSet::sign_type() const {
- if (can_prove_positive()) {
+SignType IntSet::GetSignType() const {
+ if (CanProvePositive()) {
return kPositive;
- } else if (can_prove_negative()) {
+ } else if (CanProveNegative()) {
return kNegative;
- } else if (is_single_point() && is_zero(point_value())) {
+ } else if (IsSinglePoint() && is_zero(PointValue())) {
return kZero;
} else {
return kUnknown;
}
}
-PrimExpr IntSet::point_value() const {
+PrimExpr IntSet::PointValue() const {
const IntervalSetNode* s_int = (*this).as<IntervalSetNode>();
CHECK(s_int && s_int->IsSinglePoint());
return s_int->min_value;
}
-IntSet IntSet::nothing() { return IntervalSet::Empty(); }
+IntSet IntSet::Nothing() { return IntervalSet::Empty(); }
-IntSet IntSet::everything() { return IntervalSet::Everything(); }
+IntSet IntSet::Everything() { return IntervalSet::Everything(); }
-IntSet IntSet::single_point(PrimExpr x) { return IntervalSet::SinglePoint(x); }
+IntSet IntSet::SinglePoint(PrimExpr x) { return IntervalSet::SinglePoint(x); }
-IntSet IntSet::interval(PrimExpr min, PrimExpr max) {
+IntSet IntSet::Interval(PrimExpr min, PrimExpr max) {
if (min.same_as(max)) {
- return IntSet::single_point(min);
+ return IntSet::SinglePoint(min);
}
return IntervalSet(min, max);
}
return is_zero(analyzer->Simplify(lhs - rhs));
}
-IntSet IntSet::range(Range r) {
+IntSet IntSet::FromRange(Range r) {
// must make sure it can be matched back by MatchRange.
if (is_one(r->extent)) {
- return IntSet::single_point(r->min);
+ return IntSet::SinglePoint(r->min);
}
return IntervalSet(r->min, r->extent + r->min - 1);
}
-bool IntSet::match_range(const Range& b) const {
+bool IntSet::MatchRange(const Range& b) const {
const IntSet& a = *this;
const IntervalSetNode* a_int = a.as<IntervalSetNode>();
if (!a_int) return false;
}
IntSet Union(const Array<IntSet>& sets) {
- if (sets.size() == 0) return IntSet::nothing();
+ if (sets.size() == 0) return IntSet::Nothing();
if (sets.size() == 1) return sets[0];
Analyzer ana;
IntervalSet x = ToIntervalSet(sets[0]);
}
IntSet Intersect(const Array<IntSet>& sets) {
- if (sets.size() == 0) return IntSet::nothing();
+ if (sets.size() == 0) return IntSet::Nothing();
if (sets.size() == 1) return sets[0];
Analyzer ana;
IntervalSet x = ToIntervalSet(sets[0]);
return IntervalSetEvaluator(&ana, dom_map, false).Eval(e);
}
-IntSet IntSet::vector(PrimExpr x) {
+IntSet IntSet::Vector(PrimExpr x) {
Analyzer ana;
Map<Var, IntSet> dmap;
return IntervalSetEvaluator(&ana, dmap, true).Eval(x);
<< "[" << op->min_value << ", " << op->max_value << ']';
});
-TVM_REGISTER_GLOBAL("arith.intset_single_point").set_body_typed(IntSet::single_point);
+TVM_REGISTER_GLOBAL("arith.intset_single_point").set_body_typed(IntSet::SinglePoint);
-TVM_REGISTER_GLOBAL("arith.intset_vector").set_body_typed(IntSet::vector);
+TVM_REGISTER_GLOBAL("arith.intset_vector").set_body_typed(IntSet::Vector);
-TVM_REGISTER_GLOBAL("arith.intset_interval").set_body_typed(IntSet::interval);
+TVM_REGISTER_GLOBAL("arith.intset_interval").set_body_typed(IntSet::Interval);
TVM_REGISTER_GLOBAL("arith.IntervalSetGetMin").set_body_method(&IntSet::min);
TVM_REGISTER_GLOBAL("arith.IntervalSetGetMax").set_body_method(&IntSet::max);
-TVM_REGISTER_GLOBAL("arith.IntSetIsNothing").set_body_method(&IntSet::is_nothing);
+TVM_REGISTER_GLOBAL("arith.IntSetIsNothing").set_body_method(&IntSet::IsNothing);
-TVM_REGISTER_GLOBAL("arith.IntSetIsEverything").set_body_method(&IntSet::is_everything);
+TVM_REGISTER_GLOBAL("arith.IntSetIsEverything").set_body_method(&IntSet::IsEverything);
} // namespace arith
} // namespace tvm
using namespace tir;
Stmt IRMutatorWithAnalyzer::VisitStmt_(const ForNode* op) {
- analyzer_->Bind(op->loop_var, Range::make_by_min_extent(op->min, op->extent));
+ analyzer_->Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
return StmtExprMutator::VisitStmt_(op);
}
if (op->attr_key == tir::attr::thread_extent || op->attr_key == tir::attr::virtual_thread) {
IterVar iv = Downcast<IterVar>(op->node);
CHECK_NE(iv->thread_tag.length(), 0U);
- analyzer_->Bind(iv->var, Range::make_by_min_extent(0, op->value));
+ analyzer_->Bind(iv->var, Range::FromMinExtent(0, op->value));
Stmt stmt = StmtExprMutator::VisitStmt_(op);
return stmt;
} else {
PrimExpr Simplify(const PrimExpr& expr) { return analyzer_.Simplify(expr); }
void VisitStmt_(const ForNode* op) {
- analyzer_.Bind(op->loop_var, Range::make_by_min_extent(op->min, op->extent));
+ analyzer_.Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
return StmtExprVisitor::VisitStmt_(op);
}
if (op->attr_key == attr::thread_extent || op->attr_key == attr::virtual_thread) {
IterVar iv = Downcast<IterVar>(op->node);
CHECK_NE(iv->thread_tag.length(), 0U);
- analyzer_.Bind(iv->var, Range::make_by_min_extent(0, op->value));
+ analyzer_.Bind(iv->var, Range::FromMinExtent(0, op->value));
StmtExprVisitor::VisitStmt_(op);
} else {
StmtExprVisitor::VisitStmt_(op);
new_ranges.Set(p.first, p.second);
}
// Convert original ranges to IntSets
- var_intsets[p.first.get()] = IntSet::range(p.second);
+ var_intsets[p.first.get()] = IntSet::FromRange(p.second);
}
// Infer ranges for the new variables and add them to the resulting ranges
for (const auto& p : vars_to_infer) {
const auto& var = p.first;
const auto& expr = p.second;
- Range range = EvalSet(expr, var_intsets).cover_range(Range());
+ Range range = EvalSet(expr, var_intsets).CoverRange(Range());
if (range.defined()) {
new_ranges.Set(var, range);
}
Range::Range(PrimExpr begin, PrimExpr end)
: Range(make_object<RangeNode>(begin, tir::is_zero(begin) ? end : (end - begin))) {}
-Range Range::make_by_min_extent(PrimExpr min, PrimExpr extent) {
+Range Range::FromMinExtent(PrimExpr min, PrimExpr extent) {
return Range(make_object<RangeNode>(min, extent));
}
-TVM_REGISTER_GLOBAL("ir.range_by_min_extent").set_body_typed(Range::make_by_min_extent);
+TVM_REGISTER_GLOBAL("ir.Range_from_min_extent").set_body_typed(Range::FromMinExtent);
TVM_REGISTER_GLOBAL("ir.Range").set_body([](TVMArgs args, TVMRetValue* ret) {
*ret = Range(args[0], args[1]);
void CodeGenLLVM::VisitStmt_(const ForNode* op) {
CHECK(is_zero(op->min));
- analyzer_->Bind(op->loop_var, Range::make_by_min_extent(op->min, op->extent));
+ analyzer_->Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
if (op->for_type == ForType::Unrolled) {
LOG(WARNING) << "Unroll hint get ignore at CodeGenLLVM backend, "
<< " consider set unroll_explicit=True";
if (iv->thread_tag.length() != 0) {
if (!var_map_.count(iv->var.get())) {
var_map_[iv->var.get()] = GetThreadIndex(iv);
- analyzer_->Bind(iv->var, Range::make_by_min_extent(0, op->value));
+ analyzer_->Bind(iv->var, Range::FromMinExtent(0, op->value));
}
}
} else if (op->attr_key == tir::attr::storage_scope) {
void CodeGenSPIRV::VisitStmt_(const ForNode* op) {
CHECK(is_zero(op->min));
- analyzer_->Bind(op->loop_var, Range::make_by_min_extent(op->min, op->extent));
+ analyzer_->Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
spirv::Value init_value = MakeValue(op->min);
spirv::Value extent_value = MakeValue(op->extent);
// Must get init label after making value(to make sure they are correct)
if (iv->thread_tag.length() != 0) {
if (!var_map_.count(iv->var.get())) {
var_map_[iv->var.get()] = GetThreadIndex(iv, op->value);
- analyzer_->Bind(iv->var, Range::make_by_min_extent(0, op->value));
+ analyzer_->Bind(iv->var, Range::FromMinExtent(0, op->value));
}
}
} else if (op->attr_key == tir::attr::storage_scope) {
min_value = shape_i_min_value;
max_value = shape_i_max_value;
}
- dom.data[i].push_back(IntSet::interval(min_value, max_value));
+ dom.data[i].push_back(IntSet::Interval(min_value, max_value));
} else {
dom.data[i].push_back(arg_intset);
}
CHECK_EQ(self.operator->(), this);
const TensorDom& tdom = tensor_dom.at(self.output(0));
for (size_t i = 0; i < this->axis.size(); ++i) {
- Range r = arith::Union(tdom.data.at(i)).cover_range(this->axis[i]->dom);
+ Range r = arith::Union(tdom.data.at(i)).CoverRange(this->axis[i]->dom);
CHECK(!out_dom_map->count(this->axis[i]));
(*out_dom_map)[this->axis[i]] = r;
}
if (it == out_dom_map->end()) continue;
TensorDom& dom = it->second;
for (size_t i = 0; i < t->shape.size(); ++i) {
- dom.data[i].emplace_back(IntSet::range(
- Range::make_by_min_extent(make_const(t->shape[i].dtype(), 0), t->shape[i])));
+ dom.data[i].emplace_back(
+ IntSet::FromRange(Range::FromMinExtent(make_const(t->shape[i].dtype(), 0), t->shape[i])));
}
}
}
Tensor t = stage->op.output(k);
Region bounds;
for (size_t i = 0; i < t->shape.size(); ++i) {
- bounds.push_back(Range::make_by_min_extent(make_const(t->shape[i].dtype(), 0), t->shape[i]));
+ bounds.push_back(Range::FromMinExtent(make_const(t->shape[i].dtype(), 0), t->shape[i]));
}
realize_body = tir::ProducerRealize(t, bounds, const_true(), realize_body);
}
if (it == out_dom_map->end()) continue;
TensorDom& dom = it->second;
for (size_t i = 0; i < t->shape.size(); ++i) {
- dom.data[i].emplace_back(IntSet::range(
- Range::make_by_min_extent(make_const(t->shape[i].dtype(), 0), t->shape[i])));
+ dom.data[i].emplace_back(
+ IntSet::FromRange(Range::FromMinExtent(make_const(t->shape[i].dtype(), 0), t->shape[i])));
}
}
}
Tensor t = stage->op.output(k);
Region bounds;
for (size_t i = 0; i < t->shape.size(); ++i) {
- bounds.push_back(Range::make_by_min_extent(make_const(t->shape[i].dtype(), 0), t->shape[i]));
+ bounds.push_back(Range::FromMinExtent(make_const(t->shape[i].dtype(), 0), t->shape[i]));
}
realize_body = tir::ProducerRealize(t, bounds, const_true(), realize_body);
}
PostOrderVisit(stmt, [&res_](const ObjectRef& node) {
if (const ForNode* op = node.as<ForNode>()) {
Var loop_var(op->loop_var);
- Range dom = Range::make_by_min_extent(op->min, op->extent);
+ Range dom = Range::FromMinExtent(op->min, op->extent);
res_.push_back(IterVar(dom, loop_var, ForTypeToIterVarType(op->for_type)));
}
});
// setup spatial axis
std::ostringstream spatial_name;
spatial_name << name << ".out" << i << ".i" << k;
- n->spatial_axis_.push_back(IterVar(Range::make_by_min_extent(0, update[i]->shape[k]),
+ n->spatial_axis_.push_back(IterVar(Range::FromMinExtent(0, update[i]->shape[k]),
Var(spatial_name.str()), kOpaque));
}
}
Array<Tensor> inputs, std::string name, std::string tag,
Map<String, ObjectRef> attrs) {
IterVar scan_axis =
- IterVar(Range::make_by_min_extent(init[0]->shape[0], update[0]->shape[0] - init[0]->shape[0]),
+ IterVar(Range::FromMinExtent(init[0]->shape[0], update[0]->shape[0] - init[0]->shape[0]),
Var(name + ".idx"), kOrdered);
Operation op = ScanOp(name, tag, attrs, scan_axis, init, update, state_placeholder, inputs);
Array<Tensor> res;
// first dimension, always needed.
if (init_dom) {
init_dom->data[0].push_back(
- IntSet::range(Range::make_by_min_extent(0, this->init[i]->shape[0])));
+ IntSet::FromRange(Range::FromMinExtent(0, this->init[i]->shape[0])));
}
if (update_dom) {
update_dom->data[0].push_back(dom_map.at(this->scan_axis->var.get()));
CHECK(!out_dom_map->count(this->scan_axis));
arith::Analyzer analyzer;
Range sdom = this->scan_axis->dom;
- Range r = arith::Union(time_dom).cover_range(sdom);
+ Range r = arith::Union(time_dom).CoverRange(sdom);
(*out_dom_map)[this->scan_axis] =
- Range::make_by_min_extent(sdom->min, analyzer.Simplify(r->extent + r->min - sdom->min));
+ Range::FromMinExtent(sdom->min, analyzer.Simplify(r->extent + r->min - sdom->min));
Map<IterVar, PrimExpr> fix_pt = ScanFixPointAnalysis(self);
// Update for spatial axis.
size_t sp_idx = 0;
CHECK(fix_pt.count(sp_ax));
if (fix_pt[sp_ax].as<tir::IntImmNode>()->value) {
// fix point, we can slice it.
- (*out_dom_map)[sp_ax] = arith::Union(d.data[k]).cover_range(sp_ax->dom);
+ (*out_dom_map)[sp_ax] = arith::Union(d.data[k]).CoverRange(sp_ax->dom);
} else {
// not a fix point, need to include everything.
(*out_dom_map)[sp_ax] = sp_ax->dom;
arith::Analyzer analyzer;
CHECK_EQ(stage->op.get(), this);
Range sdom = dom_map.at(this->scan_axis);
- Range tdom = Range::make_by_min_extent(0, analyzer.Simplify(sdom->extent + sdom->min));
+ Range tdom = Range::FromMinExtent(0, analyzer.Simplify(sdom->extent + sdom->min));
Stmt ret = body;
size_t sp_idx = 0;
for (size_t i = 0; i < update.size(); ++i) {
CHECK(vit != dom_map.end());
const Range& vrange = vit->second;
if (is_one(vrange->extent)) {
- up_state[iv] = IntSet::single_point(vrange->min);
+ up_state[iv] = IntSet::SinglePoint(vrange->min);
} else if (found_point) {
CHECK(is_zero(vrange->min));
- up_state[iv] = IntSet::single_point(iv->var);
+ up_state[iv] = IntSet::SinglePoint(iv->var);
} else {
- up_state[iv] = IntSet::range(vrange);
+ up_state[iv] = IntSet::FromRange(vrange);
}
auto iit = stage->iter_var_attrs.find(iv);
if (iit != stage->iter_var_attrs.end()) {
}
for (IterVar iv : self->root_iter_vars()) {
IntSet iset = up_state.at(iv);
- Range iv_range = iset.cover_range(dom_map.at(iv));
+ Range iv_range = iset.CoverRange(dom_map.at(iv));
(*out_dom)[iv] = iv_range;
analyzer.Bind(iv->var, iv_range);
temp_dmap[iv->var.get()] = iset;
Array<Range> vec;
const Tensor& t = kv.first;
for (size_t i = 0; i < t.ndim(); ++i) {
- Range r = arith::Union(kv.second.data.at(i)).cover_range(none);
+ Range r = arith::Union(kv.second.data.at(i)).CoverRange(none);
CHECK(r.defined()) << "cannot deduce region of tensorized scope for input " << t;
vec.push_back(std::move(r));
}
CHECK(it != rmap->end());
const Range& vrange = it->second;
if (is_one(vrange->extent)) {
- up_state[iv] = IntSet::single_point(vrange->min);
+ up_state[iv] = IntSet::SinglePoint(vrange->min);
} else if (!NeedRelax(iv, found_attach, ctx.bind_map, scope)) {
CHECK(is_zero(vrange->min)) << "InferBound requires every leaf iter var's min equals 0, "
<< " call schedule.normalize to achieve this. ";
if (ctx.bind_map.count(iv)) {
- up_state[iv] = IntSet::single_point(ctx.bind_map.at(iv)->var);
+ up_state[iv] = IntSet::SinglePoint(ctx.bind_map.at(iv)->var);
} else {
- up_state[iv] = IntSet::single_point(iv->var);
+ up_state[iv] = IntSet::SinglePoint(iv->var);
}
} else {
- up_state[iv] = IntSet::range(vrange);
+ up_state[iv] = IntSet::FromRange(vrange);
}
}
// Consumer's attach nest
CHECK(is_zero(vrange->min)) << "InferBound requires every leaf iter var's min equals 0, "
<< "call schedule.normalize to achieve this.";
if (NeedRelax(iv, found_attach, ctx.bind_map, scope)) {
- relax_set.Set(iv->var, IntSet::range(vrange));
+ relax_set.Set(iv->var, IntSet::FromRange(vrange));
if (ctx.bind_map.count(iv)) {
- relax_set.Set(ctx.bind_map.at(iv)->var, IntSet::range(vrange));
+ relax_set.Set(ctx.bind_map.at(iv)->var, IntSet::FromRange(vrange));
}
}
}
for (auto iv : op->root_iter_vars()) {
Range r;
if (up_state.count(iv)) {
- r = up_state.at(iv).cover_range(iv->dom);
+ r = up_state.at(iv).CoverRange(iv->dom);
} else {
r = iv->dom;
}
if (relax_set.size() != 0) {
dom_map[iv->var.get()] =
- IntSet::interval(analyzer.int_set(r->min, relax_set).min(),
+ IntSet::Interval(analyzer.int_set(r->min, relax_set).min(),
analyzer.int_set(r->min + r->extent - 1, relax_set).max());
} else {
- dom_map[iv->var.get()] = IntSet::range(r);
+ dom_map[iv->var.get()] = IntSet::FromRange(r);
}
analyzer.Bind(iv->var, r, true);
}
}
}
for (auto& p : ret) {
- ret[p.first] = Range::make_by_min_extent(analyzer.Simplify(p.second->min),
- analyzer.Simplify(p.second->extent));
+ ret[p.first] =
+ Range::FromMinExtent(analyzer.Simplify(p.second->min), analyzer.Simplify(p.second->extent));
}
return Map<IterVar, Range>(ret.begin(), ret.end());
}
};
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::FromMinExtent(0, resolve_min_extent_for_split(r->inner, r->factor)), actx);
Update(p_state, r->outer,
- Range::make_by_min_extent(0, ceil_div(range_parent->extent, r->factor)), actx);
+ Range::FromMinExtent(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);
+ Range::FromMinExtent(0, resolve_min_extent_for_split(r->outer, r->nparts)), actx);
Update(p_state, r->inner,
- Range::make_by_min_extent(0, ceil_div(range_parent->extent, r->nparts)), actx);
+ Range::FromMinExtent(0, ceil_div(range_parent->extent, r->nparts)), actx);
}
} else if (const FuseNode* r = rel.as<FuseNode>()) {
if (!state.count(r->outer) || !state.count(r->inner)) {
}
const Range& range_outer = state.at(r->outer);
const Range& range_inner = state.at(r->inner);
- state[r->fused] = Range::make_by_min_extent(0, range_outer->extent * range_inner->extent);
+ state[r->fused] = Range::FromMinExtent(0, range_outer->extent * range_inner->extent);
} else if (const RebaseNode* r = rel.as<RebaseNode>()) {
if (!state.count(r->parent)) {
CHECK(allow_missing);
continue;
}
- Update(p_state, r->rebased, Range::make_by_min_extent(0, state.at(r->parent)->extent), actx);
+ Update(p_state, r->rebased, Range::FromMinExtent(0, state.at(r->parent)->extent), actx);
} else if (const SingletonNode* s = rel.as<SingletonNode>()) {
- Update(p_state, s->iter, Range::make_by_min_extent(0, 1), actx);
+ Update(p_state, s->iter, Range::FromMinExtent(0, 1), actx);
} else {
LOG(FATAL) << "unknown relation type";
}
void PassUpDomain(const SplitNode* s, const std::unordered_map<IterVar, Range>& dom_map,
const IntSet& outer, const IntSet& inner, IntSet* parent) {
if (dom_map.count(s->outer) && dom_map.count(s->inner) && dom_map.count(s->parent) &&
- outer.match_range(dom_map.at(s->outer)) && inner.match_range(dom_map.at(s->inner))) {
- *parent = IntSet::range(dom_map.at(s->parent));
+ outer.MatchRange(dom_map.at(s->outer)) && inner.MatchRange(dom_map.at(s->inner))) {
+ *parent = IntSet::FromRange(dom_map.at(s->parent));
return;
}
PrimExpr factor = dom_map.at(s->inner)->extent;
CHECK(dom_map.count(s->fused));
arith::Analyzer ana;
- if (fused.match_range(dom_map.at(s->fused))) {
- *outer = IntSet::range(dom_map.at(s->outer));
- *inner = IntSet::range(dom_map.at(s->inner));
+ if (fused.MatchRange(dom_map.at(s->fused))) {
+ *outer = IntSet::FromRange(dom_map.at(s->outer));
+ *inner = IntSet::FromRange(dom_map.at(s->inner));
return;
}
PrimExpr outer_min = dom_map.at(s->outer)->min;
PrimExpr inner_min = dom_map.at(s->inner)->min;
- if (fused.is_single_point()) {
- PrimExpr value = fused.point_value();
+ if (fused.IsSinglePoint()) {
+ PrimExpr value = fused.PointValue();
PrimExpr factor = dom_map.at(s->inner)->extent;
PrimExpr v_outer = indexdiv(value, factor);
PrimExpr v_inner = indexmod(value, factor);
if (!is_zero(outer_min)) v_outer = v_outer + outer_min;
if (!is_zero(inner_min)) v_inner = v_inner + inner_min;
- *outer = IntSet::single_point(v_outer);
- *inner = IntSet::single_point(v_inner);
+ *outer = IntSet::SinglePoint(v_outer);
+ *inner = IntSet::SinglePoint(v_inner);
} else {
PrimExpr fused_extent = (fused.max() - fused.min() + 1);
PrimExpr inner_extent = dom_map.at(s->inner)->extent;
- *outer = IntSet::interval(outer_min + indexdiv(fused.min(), inner_extent),
+ *outer = IntSet::Interval(outer_min + indexdiv(fused.min(), inner_extent),
outer_min + indexdiv(fused.max(), inner_extent));
if (is_zero(ana.Simplify(indexmod(inner_extent, fused_extent))) &&
is_zero(ana.Simplify(indexmod(fused.min(), fused_extent)))) {
// fused never spans multiple rows, make a tight bounding box
// there may be other cases when bounding box could be tightened
- *inner = IntSet::interval(inner_min + indexmod(fused.min(), inner_extent),
+ *inner = IntSet::Interval(inner_min + indexmod(fused.min(), inner_extent),
inner_min + indexmod(fused.max(), inner_extent));
} else { // fused may span multiple rows, use full row widths
if (!is_zero(ana.Simplify(indexmod(fused_extent, inner_extent))) ||
LOG(WARNING)
<< "fused and original axes are not aligned, this may cause redundant computations";
}
- *inner = IntSet::range(dom_map.at(s->inner));
+ *inner = IntSet::FromRange(dom_map.at(s->inner));
}
return;
}
void PassUpDomain(const RebaseNode* s, const std::unordered_map<IterVar, Range>& dom_map,
const IntSet& rebased, IntSet* parent) {
CHECK(dom_map.count(s->parent));
- if (rebased.match_range(dom_map.at(s->rebased))) {
- *parent = IntSet::range(dom_map.at(s->parent));
+ if (rebased.MatchRange(dom_map.at(s->rebased))) {
+ *parent = IntSet::FromRange(dom_map.at(s->parent));
return;
}
PrimExpr parent_min = dom_map.at(s->parent)->min;
// setup domain map for set analysis
for (const auto& kv : dom_map) {
- iset_dmap.Set(kv.first->var, IntSet::range(kv.second));
+ iset_dmap.Set(kv.first->var, IntSet::FromRange(kv.second));
}
for (auto entry : dom_map) {
for (Range r : old_region) {
PrimExpr min = VarReplacer(vsub2newvar)(r->min);
PrimExpr extent = VarReplacer(vsub2newvar)(r->extent);
- region.push_back(Range::make_by_min_extent(min, extent));
+ region.push_back(Range::FromMinExtent(min, extent));
}
new_regions.push_back(region);
}
// special handle fuse empty array.
// insert at the outer most loop
IterVar singleton =
- IterVar(Range::make_by_min_extent(0, 1), Var("singleton", DataType::Int(32)), kDataPar);
+ IterVar(Range::FromMinExtent(0, 1), Var("singleton", DataType::Int(32)), kDataPar);
self->relations.push_back(Singleton(singleton));
Array<IterVar>& all_vars = self->all_iter_vars;
Array<IterVar>& leaf_vars = self->leaf_iter_vars;
}
CHECK_GE(op->bounds.size(), 2) << "Less than 2 dimensions for matrix " << key->GetNameHint();
new_bounds.push_back(
- Range::make_by_min_extent(op->bounds[op->bounds.size() - 2]->min, new_extents[0]));
+ Range::FromMinExtent(op->bounds[op->bounds.size() - 2]->min, new_extents[0]));
new_bounds.push_back(
- Range::make_by_min_extent(op->bounds[op->bounds.size() - 1]->min, new_extents[1]));
+ Range::FromMinExtent(op->bounds[op->bounds.size() - 1]->min, new_extents[1]));
return ProducerRealize(op->producer, new_bounds, op->condition, op->body);
}
if (min.same_as(r->min) && extent.same_as(r->extent)) {
return v;
} else {
- return IterVar(Range::make_by_min_extent(min, extent), v->var, v->iter_type, v->thread_tag);
+ return IterVar(Range::FromMinExtent(min, extent), v->var, v->iter_type, v->thread_tag);
}
};
Array<IterVar> axis = MutateArray(op->axis, fitervar);
if (min.same_as(r->min) && extent.same_as(r->extent)) {
return r;
} else {
- return Range::make_by_min_extent(min, extent);
+ return Range::FromMinExtent(min, extent);
}
};
return MutateArray(arr, fmutate, self->allow_copy_on_write_);
wset.push_back(acc.touched);
}
Range none;
- Range r = arith::Union(wset).cover_range(none);
+ Range r = arith::Union(wset).CoverRange(none);
CHECK(r.defined()) << "Cannot deduce write range of " << wvec[0].buffer;
PrimExpr min = r->min;
PrimExpr extent = r->extent;
Region region;
auto iter_var = loop_nest_.back().get();
- vectorized_[iter_var] = IntSet::single_point(loop_nest_.back() + op->value);
+ vectorized_[iter_var] = IntSet::SinglePoint(loop_nest_.back() + op->value);
for (Range r : domain) {
if (!r.defined()) {
LOG(WARNING) << "Cannot decide prefetch region for " << buffer;
return op->body;
}
- Range res(EvalSet(r, vectorized_).cover_range(none));
- region.push_back(Range::make_by_min_extent(res->min, res->extent));
+ Range res(EvalSet(r, vectorized_).CoverRange(none));
+ region.push_back(Range::FromMinExtent(res->min, res->extent));
}
vectorized_.erase(iter_var);
auto& var = op->loop_var;
loop_nest_.push_back(var);
if (op->for_type == ForType::Vectorized) {
- vectorized_[var.get()] = IntSet::interval(op->min, (op->min + op->extent) - 1);
+ vectorized_[var.get()] = IntSet::Interval(op->min, (op->min + op->extent) - 1);
}
Stmt ret = StmtMutator::VisitStmt_(op);
if (op->for_type == ForType::Vectorized) {
if (ExprUseVars(op->min, out_vars_) || ExprUseVars(op->extent, out_vars_)) return;
const VarNode* var = op->loop_var.get();
- hint_map_.insert({var, IntSet::interval(op->min, op->min + op->extent - 1)});
- relax_map_.insert({var, IntSet::interval(op->min, op->min + op->extent - 1)});
+ hint_map_.insert({var, IntSet::Interval(op->min, op->min + op->extent - 1)});
+ relax_map_.insert({var, IntSet::Interval(op->min, op->min + op->extent - 1)});
StmtExprVisitor::VisitStmt_(op);
relax_map_.erase(var);
hint_map_.erase(var);
const IterVarNode* thread_axis = op->node.as<IterVarNode>();
CHECK(thread_axis);
const VarNode* var = thread_axis->var.get();
- IntSet dom = IntSet::range(Range(make_zero(op->value.dtype()), op->value));
+ IntSet dom = IntSet::FromRange(Range(make_zero(op->value.dtype()), op->value));
hint_map_.insert({var, dom});
relax_map_.insert({var, dom});
StmtExprVisitor::VisitStmt_(op);
// true. Also find the interval, if exists, in which we can prove that cond is
// false.
IntSet interval = DeduceBound(current_var_, cond, hint_map_, relax_map_);
- if (!interval.is_nothing()) {
+ if (!interval.IsNothing()) {
// cond is true within interval
partitions[{cond.get(), true}] = interval;
}
PrimExpr inverse_cond = InverseCond(cond);
if (inverse_cond.defined()) {
IntSet interval = DeduceBound(current_var_, inverse_cond, hint_map_, relax_map_);
- if (!interval.is_nothing()) {
+ if (!interval.IsNothing()) {
// cond is false within interval
partitions[{cond.get(), false}] = interval;
}
// normal path when loop partition fails
// normal loop variable can be put into hint map.
- hint_map_.insert({op->loop_var.get(), IntSet::interval(op->min, op->min + op->extent - 1)});
+ hint_map_.insert({op->loop_var.get(), IntSet::Interval(op->min, op->min + op->extent - 1)});
Stmt res = StmtMutator::VisitStmt_(op);
hint_map_.erase(op->loop_var.get());
return res;
Stmt res;
if (scope.rank == 1) {
// threadIdx should be put into relax map, in case of divergence.
- relax_map_.insert({var.get(), IntSet::interval(make_zero(var.dtype()), op->value - 1)});
+ relax_map_.insert({var.get(), IntSet::Interval(make_zero(var.dtype()), op->value - 1)});
res = StmtMutator::VisitStmt_(op);
relax_map_.erase(var.get());
} else {
- hint_map_.insert({var.get(), IntSet::interval(make_zero(var.dtype()), op->value - 1)});
+ hint_map_.insert({var.get(), IntSet::Interval(make_zero(var.dtype()), op->value - 1)});
res = StmtMutator::VisitStmt_(op);
hint_map_.erase(var.get());
}
}
}
}
- IntSet interval = sets.empty() ? IntSet::nothing() : Intersect(sets);
+ IntSet interval = sets.empty() ? IntSet::Nothing() : Intersect(sets);
return std::make_pair(interval, cond_set);
}
PrimExpr max, Stmt body, bool partition_thread_scope) {
using namespace arith;
// include hint of var.
- hint_map_.insert({var.get(), IntSet::interval(min, max)});
+ hint_map_.insert({var.get(), IntSet::Interval(min, max)});
PartitionFinder finder(var, hint_map_, relax_map_);
finder(body);
// find an interval in which all conditions on var are true
std::tie(middle_interval, cond_set) =
GetIntervalAndCondset(finder.partitions, for_interval, true);
- if (middle_interval.is_nothing()) {
+ if (middle_interval.IsNothing()) {
// if such interval doesn't exist, find an interval in which all
// conditions on var are false
std::tie(middle_interval, cond_set) =
GetIntervalAndCondset(finder.partitions, for_interval, false);
- if (middle_interval.is_nothing())
+ if (middle_interval.IsNothing())
// we couldn't find an interval in which the conditions are provably true or false
// Therefore, we can't partition the loop based on those conds
return Stmt();
void VisitStmt_(const ForNode* op) final {
const Var& loop_var = op->loop_var;
- analyzer_->Bind(loop_var, Range::make_by_min_extent(op->min, op->extent));
+ analyzer_->Bind(loop_var, Range::FromMinExtent(op->min, op->extent));
StmtVisitor::VisitStmt_(op);
}
IterVar iv = Downcast<IterVar>(op->node);
CHECK_NE(iv->thread_tag.length(), 0U);
if (!var_dom_.count(iv->var.get())) {
- Range dom = Range::make_by_min_extent(0, op->value);
+ Range dom = Range::FromMinExtent(0, op->value);
var_dom_[iv->var.get()] = dom;
analyzer_->Bind(iv->var, dom);
}
}
void VisitStmt_(const ForNode* op) {
- analyzer_.Bind(op->loop_var, Range::make_by_min_extent(op->min, op->extent));
+ analyzer_.Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
vextent_[op->loop_var.as<VarNode>()] = op->extent.dtype();
return StmtExprVisitor::VisitStmt_(op);
}
if (op->attr_key == attr::thread_extent || op->attr_key == attr::virtual_thread) {
IterVar iv = Downcast<IterVar>(op->node);
CHECK_NE(iv->thread_tag.length(), 0U);
- analyzer_.Bind(iv->var, Range::make_by_min_extent(0, op->value));
+ analyzer_.Bind(iv->var, Range::FromMinExtent(0, op->value));
vextent_[iv->var.as<VarNode>()] = op->value.dtype();
StmtExprVisitor::VisitStmt_(op);
} else {
Stmt Simplify(Stmt stmt) { return operator()(std::move(stmt)); }
Stmt VisitStmt_(const ForNode* op) final {
- analyzer_->Bind(op->loop_var, Range::make_by_min_extent(op->min, op->extent));
+ analyzer_->Bind(op->loop_var, Range::FromMinExtent(op->min, op->extent));
With<ConstraintContext> ctx1(analyzer_, op->loop_var >= op->min);
With<ConstraintContext> ctx2(analyzer_, op->loop_var < op->min + op->extent);
return Parent::VisitStmt_(op);
e.threads = env_threads();
e.buffer = op->buffer_var;
e.dtype = op->dtype.element_of();
- e.touched = arith::IntSet::vector(op->index);
+ e.touched = arith::IntSet::Vector(op->index);
e.type = kRead;
e.scope = scope;
curr_stmt_.access.emplace_back(std::move(e));
e.threads = env_threads();
e.buffer = op->buffer_var;
e.dtype = op->value.dtype().element_of();
- e.touched = arith::IntSet::vector(op->index);
+ e.touched = arith::IntSet::Vector(op->index);
e.type = kWrite;
e.scope = scope;
curr_stmt_.access.emplace_back(std::move(e));
// relax the touched set to contain all ranges in the loop.
std::unordered_map<const VarNode*, arith::IntSet> relax_map;
relax_map[op->loop_var.get()] =
- arith::IntSet::range(Range::make_by_min_extent(op->min, op->extent));
+ arith::IntSet::FromRange(Range::FromMinExtent(op->min, op->extent));
for (AccessEntry& e : s.access) {
if (e.buffer.defined()) {
CHECK(e.touched.defined());
e.threads = env_threads();
e.dtype = dtype;
e.buffer = Downcast<Var>(op->args[1]);
- e.touched = arith::IntSet::range(Range::make_by_min_extent(offset, extent));
+ e.touched = arith::IntSet::FromRange(Range::FromMinExtent(offset, extent));
e.scope = scope;
if (flag->value & 1) {
e.type = kRead;
// Assumes no race between threads
// Same index value means no conflicts
// TODO(tqchen) more standard set based testing.
- if (e.touched.is_single_point() && x.touched.is_single_point()) {
- if (ExprDeepEqual()(e.touched.point_value(), x.touched.point_value())) continue;
+ if (e.touched.IsSinglePoint() && x.touched.IsSinglePoint()) {
+ if (ExprDeepEqual()(e.touched.PointValue(), x.touched.PointValue())) continue;
}
if (x.double_buffer_write && e.type == kRead && !loop_carry) continue;
return true;
floordiv = tvm.te.floordiv
z = te.var("z")
- ck.analyzer.bind(x, tvm.ir.Range.make_by_min_extent(0, 3))
+ ck.analyzer.bind(x, tvm.ir.Range.from_min_extent(0, 3))
ck.verify(flm(y, 8), {y : tvm.arith.IntervalSet(z*8+x*4, z*8+x*4+3)},
(0, 7))
ck1 = IntSetChecker()
- ck1.analyzer.bind(x, tvm.ir.Range.make_by_min_extent(0, 2))
+ ck1.analyzer.bind(x, tvm.ir.Range.from_min_extent(0, 2))
ck1.verify(flm(y, 8), {y : tvm.arith.IntervalSet(z*8+x*4, z*8+x*4+3)}, (x*4, x*4+3))
def test_infer_range():
x, y = te.var("x"), te.var("y")
ranges = {
- x: tvm.ir.Range.make_by_min_extent(-5, 10),
- y: tvm.ir.Range.make_by_min_extent(0, 10),
+ x: tvm.ir.Range.from_min_extent(-5, 10),
+ y: tvm.ir.Range.from_min_extent(0, 10),
}
solution = arith.solve_linear_equations([
with ib.for_range(0, n, name="i") as i:
ib.emit(
tvm.tir.Prefetch(A,
- [tvm.ir.Range.make_by_min_extent(i+1, 2),
- tvm.ir.Range.make_by_min_extent(0, 20)]))
+ [tvm.ir.Range.from_min_extent(i+1, 2),
+ tvm.ir.Range.from_min_extent(0, 20)]))
body = ib.get()
assert body.body.bounds[0].extent.value == 2
_A= tvm.tir.decl_buffer(A.shape, A.dtype, name = 'A');
i = te.size_var('i')
j = te.size_var('j')
- region = [tvm.ir.Range.make_by_min_extent(i[0], i[1]) for i in [(i, 2), (j, 8), (0, 4)]]
+ region = [tvm.ir.Range.from_min_extent(i[0], i[1]) for i in [(i, 2), (j, 8), (0, 4)]]
stmt = tvm.tir.Prefetch(_A, region)
func = tvm.te.schedule.SchedulePostProcToPrimFunc(