You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/08/20 15:37:08 UTC
[GitHub] [tvm] wrongtest-intellif opened a new pull request, #12515: [TIR][Arith] Add more strict checking in imm construction and folding.
wrongtest-intellif opened a new pull request, #12515:
URL: https://github.com/apache/tvm/pull/12515
Refer to issues https://github.com/apache/tvm/issues/12377 and https://github.com/apache/tvm/issues/12378. The behavior of compile time constant folding and general path computation differs.
The changes are:
- Add more strict value range checking when build IntImm object.
- For constant folding on integers
- Use intN arithmetic on intN folding.
- For constant folding on floats
- Only support float32 and float64, use float32 arithmetic on float32 folding.
- Do not allow fold float16 since we do not have implemented fp16 arithmetics yet and there may exist compatibility issues on different targets.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ganler commented on a diff in pull request #12515: [TIR][Arith] Add more strict checking in imm construction and folding.
Posted by GitBox <gi...@apache.org>.
ganler commented on code in PR #12515:
URL: https://github.com/apache/tvm/pull/12515#discussion_r950715982
##########
include/tvm/tir/op.h:
##########
@@ -911,7 +911,9 @@ inline PrimExpr MakeConstScalar(DataType t, ValueType value, Span span = Span())
if (t.is_uint()) {
// Use IntImm if it is a small integer
uint64_t uval = static_cast<uint64_t>(value);
- if (uval <= static_cast<uint64_t>(std::numeric_limits<int64_t>::max())) {
+ if (static_cast<int64_t>(value) < 0) {
Review Comment:
Maybe here we should do `value < static_cast<ValueType>(0)`. Because if `value` is of `uint64_t`, it will be regarded as a negative number if it is greater than `numeric_limits<int64_t>::max()`.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ganler commented on a diff in pull request #12515: [TIR][Arith] Add more strict checking in imm construction and folding.
Posted by GitBox <gi...@apache.org>.
ganler commented on code in PR #12515:
URL: https://github.com/apache/tvm/pull/12515#discussion_r950711151
##########
tests/python/unittest/test_tir_imm_values.py:
##########
@@ -0,0 +1,155 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+import math
+import numpy as np
+import tvm
+import tvm.testing
+import pytest
+from tvm import tir
+from tvm.script import tir as T
+import pytest
+
+
+@pytest.mark.parametrize(
+ "dtype, literals",
+ [
+ ["int8", [-128, 0, 127]],
+ ["uint8", [0, 255]],
+ ["int32", [-2147483648, 2147483647]],
+ ["uint32", [0, 4294967295]],
+ ["int64", [-9223372036854775808, 9223372036854775807]],
+ ["uint64", [0, 9223372036854775807]],
+ ],
+)
+def test_tir_make_intimm(dtype, literals):
+ for l in literals:
+ imm = tir.const(l, dtype)
+ assert imm.value == l, imm
+
+
+@pytest.mark.parametrize(
+ "dtype, literals",
+ [
+ ["int8", [-129, 128]],
+ ["uint8", [-1, 256]],
+ ["int32", [-2147483650, 2147483648]],
+ ["uint32", [-1, 4294967296]],
+ ["uint64", [-1, 18446744073709551616]],
+ ],
+)
+def test_tir_invalid_intimm(dtype, literals):
+ for l in literals:
+ with pytest.raises(tvm.TVMError):
+ tir.const(l, dtype)
+
+
+@pytest.mark.parametrize(
+ "dtype, literals",
+ [
+ [
+ "int64",
+ {
+ -9223372036854775810: 9223372036854775806,
+ 9223372036854775808: -9223372036854775808,
+ },
+ ],
+ [
+ "uint64",
+ {
+ 9223372036854775807: 9223372036854775807,
+ 18446744073709551615: 18446744073709551615,
+ },
+ ],
+ ],
+)
+def test_tir_large_py_int_literals(dtype, literals):
+ """
+ For large uint value, use LargeUIntImm intrin,
+ For large int value exceed int64_t value range, the value is wrapped back.
+ """
+ for l in literals:
+ x = tir.const(l, dtype)
+ if isinstance(x, (tir.IntImm, tir.FloatImm)):
+ assert x.value == literals[l]
+ else:
+ # LargeUIntImm(low32, hi32)
+ assert (int(x.args[1]) << 32) + int(x.args[0]) == literals[l]
+
+
+def test_tir_intimm_overflow():
+ assert int(tir.const(127, "int8") + tir.const(1, "int8")) == -128
+ assert int(tir.const(127, "int8") + tir.const(2, "int8")) == -127
+ assert int(tir.const(255, "uint8") + tir.const(1, "uint8")) == 0
+ assert int(tir.const(2 ** 31 - 1, "int32") + tir.const(1, "int32")) == -(2 ** 31)
+ assert int(tir.const(2 ** 32 - 1, "uint32") + tir.const(1, "uint32")) == 0
+ assert int(tir.const(2 ** 63 - 1, "int64") + tir.const(1, "int64")) == -(2 ** 63)
+ assert int(tir.const(2 ** 32, "uint64") * tir.const(2 ** 32, "uint64")) == 0
+
+
+def compare_float_value(value, expect):
+ if math.isfinite(value):
+ assert value == expect
+ elif math.isnan(value):
+ assert math.isnan(expect)
+ elif math.isinf(value):
+ assert math.isinf(expect)
+
+
+@pytest.mark.parametrize("dtype", ["float16", "float32", "float64"])
+@pytest.mark.parametrize("literal", [3.14, np.nan, np.inf])
+def test_tir_special_floatimms(dtype, literal):
+ x = tir.const(literal, dtype)
+ compare_float_value(x.value, literal)
+
+
+@tvm.testing.requires_llvm()
+def test_tir_floatimm_overflow():
+ # Behavior check: if literal value is out of dtype range, the
+ # object is still constructed, and eval to infinity.
+ @T.prim_func
+ def imm_overflow_fp16() -> T.float16:
+ T.evaluate(T.ret(T.float16(65536), dtype="float16"))
Review Comment:
Just curious, should we just raise an exception or warning to tell the user that this is actually an ill-formed behavior, i.e., constructing an out-of-bound constant? For example, gcc 12 will warn such behaviors.
##########
include/tvm/tir/op.h:
##########
@@ -911,7 +911,9 @@ inline PrimExpr MakeConstScalar(DataType t, ValueType value, Span span = Span())
if (t.is_uint()) {
// Use IntImm if it is a small integer
uint64_t uval = static_cast<uint64_t>(value);
- if (uval <= static_cast<uint64_t>(std::numeric_limits<int64_t>::max())) {
+ if (value < 0) {
Review Comment:
What if `value` is greater than `numeric_limit<int64_t>::max()`? Is it possible to keep the sign information. If it is from Python end, `int` in Python is of unlimited length so catching the sign is easy. The sign information in C++ probably can be catched from the type information, i.e., ValueType.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] wrongtest-intellif commented on pull request #12515: [TIR][Arith] Add more strict checking in imm construction and folding.
Posted by GitBox <gi...@apache.org>.
wrongtest-intellif commented on PR #12515:
URL: https://github.com/apache/tvm/pull/12515#issuecomment-1230223040
cc @ganler Hi, could you kindly take another look at the change~?
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] wrongtest-intellif commented on a diff in pull request #12515: [TIR][Arith] Add more strict checking in imm construction and folding.
Posted by GitBox <gi...@apache.org>.
wrongtest-intellif commented on code in PR #12515:
URL: https://github.com/apache/tvm/pull/12515#discussion_r950713439
##########
tests/python/unittest/test_tir_imm_values.py:
##########
@@ -0,0 +1,155 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+import math
+import numpy as np
+import tvm
+import tvm.testing
+import pytest
+from tvm import tir
+from tvm.script import tir as T
+import pytest
+
+
+@pytest.mark.parametrize(
+ "dtype, literals",
+ [
+ ["int8", [-128, 0, 127]],
+ ["uint8", [0, 255]],
+ ["int32", [-2147483648, 2147483647]],
+ ["uint32", [0, 4294967295]],
+ ["int64", [-9223372036854775808, 9223372036854775807]],
+ ["uint64", [0, 9223372036854775807]],
+ ],
+)
+def test_tir_make_intimm(dtype, literals):
+ for l in literals:
+ imm = tir.const(l, dtype)
+ assert imm.value == l, imm
+
+
+@pytest.mark.parametrize(
+ "dtype, literals",
+ [
+ ["int8", [-129, 128]],
+ ["uint8", [-1, 256]],
+ ["int32", [-2147483650, 2147483648]],
+ ["uint32", [-1, 4294967296]],
+ ["uint64", [-1, 18446744073709551616]],
+ ],
+)
+def test_tir_invalid_intimm(dtype, literals):
+ for l in literals:
+ with pytest.raises(tvm.TVMError):
+ tir.const(l, dtype)
+
+
+@pytest.mark.parametrize(
+ "dtype, literals",
+ [
+ [
+ "int64",
+ {
+ -9223372036854775810: 9223372036854775806,
+ 9223372036854775808: -9223372036854775808,
+ },
+ ],
+ [
+ "uint64",
+ {
+ 9223372036854775807: 9223372036854775807,
+ 18446744073709551615: 18446744073709551615,
+ },
+ ],
+ ],
+)
+def test_tir_large_py_int_literals(dtype, literals):
+ """
+ For large uint value, use LargeUIntImm intrin,
+ For large int value exceed int64_t value range, the value is wrapped back.
+ """
+ for l in literals:
+ x = tir.const(l, dtype)
+ if isinstance(x, (tir.IntImm, tir.FloatImm)):
+ assert x.value == literals[l]
+ else:
+ # LargeUIntImm(low32, hi32)
+ assert (int(x.args[1]) << 32) + int(x.args[0]) == literals[l]
+
+
+def test_tir_intimm_overflow():
+ assert int(tir.const(127, "int8") + tir.const(1, "int8")) == -128
+ assert int(tir.const(127, "int8") + tir.const(2, "int8")) == -127
+ assert int(tir.const(255, "uint8") + tir.const(1, "uint8")) == 0
+ assert int(tir.const(2 ** 31 - 1, "int32") + tir.const(1, "int32")) == -(2 ** 31)
+ assert int(tir.const(2 ** 32 - 1, "uint32") + tir.const(1, "uint32")) == 0
+ assert int(tir.const(2 ** 63 - 1, "int64") + tir.const(1, "int64")) == -(2 ** 63)
+ assert int(tir.const(2 ** 32, "uint64") * tir.const(2 ** 32, "uint64")) == 0
+
+
+def compare_float_value(value, expect):
+ if math.isfinite(value):
+ assert value == expect
+ elif math.isnan(value):
+ assert math.isnan(expect)
+ elif math.isinf(value):
+ assert math.isinf(expect)
+
+
+@pytest.mark.parametrize("dtype", ["float16", "float32", "float64"])
+@pytest.mark.parametrize("literal", [3.14, np.nan, np.inf])
+def test_tir_special_floatimms(dtype, literal):
+ x = tir.const(literal, dtype)
+ compare_float_value(x.value, literal)
+
+
+@tvm.testing.requires_llvm()
+def test_tir_floatimm_overflow():
+ # Behavior check: if literal value is out of dtype range, the
+ # object is still constructed, and eval to infinity.
+ @T.prim_func
+ def imm_overflow_fp16() -> T.float16:
+ T.evaluate(T.ret(T.float16(65536), dtype="float16"))
Review Comment:
Well I agree~ For fp16, fp32, fp64 since we know the range actually. I could try for that.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 merged pull request #12515: [TIR][Arith] Add more strict checking in imm construction and folding.
Posted by GitBox <gi...@apache.org>.
vinx13 merged PR #12515:
URL: https://github.com/apache/tvm/pull/12515
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org