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 + + [email protected]( + "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 + + [email protected]( + "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) + + [email protected]( + "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) + + [email protected]("dtype", ["float16", "float32", "float64"]) [email protected]("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) + + [email protected]_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: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
