From 2b2fceb22dce6094c546578979ad354446e7f373 Mon Sep 17 00:00:00 2001 From: Haozheng Fan Date: Tue, 27 Oct 2020 17:49:42 +0800 Subject: [PATCH 1/4] fine-grained estimation for floormod --- src/arith/const_int_bound.cc | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/src/arith/const_int_bound.cc b/src/arith/const_int_bound.cc index f39ce4b05643..7b9aa1fa5c50 100644 --- a/src/arith/const_int_bound.cc +++ b/src/arith/const_int_bound.cc @@ -261,7 +261,10 @@ class ConstIntBoundAnalyzer::Impl ICHECK(!b.is_const(0)) << "floormod by zero"; // mod by negative value is rare, // and we just use the simpliest rule. - return Everything(op->dtype); + int64_t b_min_cap = InfAwareAdd(b.min_value, 1); + int64_t b_max_cap = InfAwareAdd(b.max_value, -1); + return Intersect(MakeBound(std::min(0ll, b_min_cap), std::max(0ll, b_max_cap)), + Everything(op->dtype)); } } From e90cc29aeca7b54ed5ac20a8daaf92823a279e87 Mon Sep 17 00:00:00 2001 From: Haozheng Fan Date: Tue, 27 Oct 2020 17:53:43 +0800 Subject: [PATCH 2/4] fix --- src/arith/const_int_bound.cc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/arith/const_int_bound.cc b/src/arith/const_int_bound.cc index 7b9aa1fa5c50..2ab262563fc9 100644 --- a/src/arith/const_int_bound.cc +++ b/src/arith/const_int_bound.cc @@ -263,7 +263,8 @@ class ConstIntBoundAnalyzer::Impl // and we just use the simpliest rule. int64_t b_min_cap = InfAwareAdd(b.min_value, 1); int64_t b_max_cap = InfAwareAdd(b.max_value, -1); - return Intersect(MakeBound(std::min(0ll, b_min_cap), std::max(0ll, b_max_cap)), + return Intersect(MakeBound(std::min(static_cast(0), b_min_cap), + std::max(static_cast(0), b_max_cap)), Everything(op->dtype)); } } From 60b479d61f007673750bbcf0d027af786381b772 Mon Sep 17 00:00:00 2001 From: Haozheng Fan Date: Tue, 27 Oct 2020 18:07:25 +0800 Subject: [PATCH 3/4] comments --- src/arith/const_int_bound.cc | 19 +++++++++++++++++-- 1 file changed, 17 insertions(+), 2 deletions(-) diff --git a/src/arith/const_int_bound.cc b/src/arith/const_int_bound.cc index 2ab262563fc9..1a5cce4ee4b1 100644 --- a/src/arith/const_int_bound.cc +++ b/src/arith/const_int_bound.cc @@ -259,8 +259,23 @@ class ConstIntBoundAnalyzer::Impl } } else { ICHECK(!b.is_const(0)) << "floormod by zero"; - // mod by negative value is rare, - // and we just use the simpliest rule. + /* let a / b = x + y, where x is integer, y \in [0, 1) + * floormod(a, b) = a - floordiv(a, b) * b + * floordiv(a, b) = x + * floormod(a, b) = a - floordiv(a, b) * b + * = a - x * b + * = a - (a / b - y) * b + * = a - a + y * b + * = y * b + * note that 0 <= y < 1 + * when b > 0, 0 <= b * y < b + * 0 <= b * y <= b - 1 + * when b < 0, b < b * y <= 0 + * b + 1 <= b * y <= 0 + * In all cases, min(0, b + 1) <= b * y <= max(0, b - 1) + * min(0, b_min + 1) <= b * y <= max(0, b_max - 1) + * That is, min(0, b_min + 1) <= floormod(a, b) <= max(0, b_max - 1) + */ int64_t b_min_cap = InfAwareAdd(b.min_value, 1); int64_t b_max_cap = InfAwareAdd(b.max_value, -1); return Intersect(MakeBound(std::min(static_cast(0), b_min_cap), From 69f8fb23d72b8c14e1d4bc63bcc7b8233f00e75d Mon Sep 17 00:00:00 2001 From: Haozheng Fan Date: Wed, 28 Oct 2020 01:53:17 +0000 Subject: [PATCH 4/4] negative divisor test --- src/arith/const_int_bound.cc | 34 +++++++++---------- .../unittest/test_arith_const_int_bound.py | 12 +++++++ 2 files changed, 29 insertions(+), 17 deletions(-) diff --git a/src/arith/const_int_bound.cc b/src/arith/const_int_bound.cc index 1a5cce4ee4b1..2c01b9143155 100644 --- a/src/arith/const_int_bound.cc +++ b/src/arith/const_int_bound.cc @@ -245,6 +245,23 @@ class ConstIntBoundAnalyzer::Impl } Entry VisitExpr_(const FloorModNode* op) final { + /* let a / b = x + y, where x is integer, y \in [0, 1) + * floormod(a, b) = a - floordiv(a, b) * b + * floordiv(a, b) = x + * floormod(a, b) = a - floordiv(a, b) * b + * = a - x * b + * = a - (a / b - y) * b + * = a - a + y * b + * = y * b + * note that 0 <= y < 1 + * when b > 0, 0 <= b * y < b + * 0 <= b * y <= b - 1 + * when b < 0, b < b * y <= 0 + * b + 1 <= b * y <= 0 + * In all cases, min(0, b + 1) <= b * y <= max(0, b - 1) + * min(0, b_min + 1) <= b * y <= max(0, b_max - 1) + * That is, min(0, b_min + 1) <= floormod(a, b) <= max(0, b_max - 1) + */ Entry a = VisitExpr(op->a); Entry b = VisitExpr(op->b); if (b.min_value > 0) { @@ -259,23 +276,6 @@ class ConstIntBoundAnalyzer::Impl } } else { ICHECK(!b.is_const(0)) << "floormod by zero"; - /* let a / b = x + y, where x is integer, y \in [0, 1) - * floormod(a, b) = a - floordiv(a, b) * b - * floordiv(a, b) = x - * floormod(a, b) = a - floordiv(a, b) * b - * = a - x * b - * = a - (a / b - y) * b - * = a - a + y * b - * = y * b - * note that 0 <= y < 1 - * when b > 0, 0 <= b * y < b - * 0 <= b * y <= b - 1 - * when b < 0, b < b * y <= 0 - * b + 1 <= b * y <= 0 - * In all cases, min(0, b + 1) <= b * y <= max(0, b - 1) - * min(0, b_min + 1) <= b * y <= max(0, b_max - 1) - * That is, min(0, b_min + 1) <= floormod(a, b) <= max(0, b_max - 1) - */ int64_t b_min_cap = InfAwareAdd(b.min_value, 1); int64_t b_max_cap = InfAwareAdd(b.max_value, -1); return Intersect(MakeBound(std::min(static_cast(0), b_min_cap), diff --git a/tests/python/unittest/test_arith_const_int_bound.py b/tests/python/unittest/test_arith_const_int_bound.py index badbcbcf1bb3..84fc7fd64614 100644 --- a/tests/python/unittest/test_arith_const_int_bound.py +++ b/tests/python/unittest/test_arith_const_int_bound.py @@ -303,6 +303,17 @@ def test_let_bound(): assert bd.max_value == 2 +def test_floormod_negative_divisor(): + analyzer = tvm.arith.Analyzer() + flm, fld = tvm.te.floormod, tvm.te.floordiv + a, b = te.var("a"), te.var("b") + analyzer.update(a, tvm.arith.ConstIntBound(0, 6)) + analyzer.update(b, tvm.arith.ConstIntBound(-5, 7)) + bd = analyzer.const_int_bound(flm(a, b)) + assert bd.min_value == -4 + assert bd.max_value == 6 + + if __name__ == "__main__": test_let_bound() test_dtype_bound() @@ -318,3 +329,4 @@ def test_let_bound(): test_shift_and_bound() test_mix_index_bound() test_size_var_bound() + test_floormod_negative_divisor()