Tianqi Chen
1e1920bcbd
In the past we have been using `DataType` in PrimExpr.dtype field to check type information for PrimExpr while still having BaseExpr.ty for richer type information. DataType is also used both in runtime and compiler. This PR streamlines the boundary: - PrimExpr.ty now carries PrimType that replaces original use of `DataType` - Runtime use will now favor DLPack DLDataType, removing one layer of indirection. - Constants attributes where values are usually runtime values, will use `DLDataType` - DataType will be phased out after this PR We also brings up helper functions in PrimType, but also limits them to a more concise set so the functions do not grow with the data type codes in DLPack. This is a major refactor that changes the IR primitive. It helps to bring possible future benefits: - Unified type mechanism through Expr.ty - Possibility of carry future Type nodes Migration Guide: - Use `PrimType` when code reasons about compiler expression types, tensor element compiler types, or constructs a `PrimExpr`/compiler type. - Use existing source types such as `expr.ty()`, `ExprOp.expr_ty()`, or TE tensor element `dtype` where possible instead of rebuilding a type from dtype text. - Use raw `DLDataType` for runtime constants, ABI paths, dtype-valued attrs, and storage/runtime helper logic. - Prefer direct `PrimType` equality, `MatchesCode(...)`, `MatchesElementType(...)`, and `WithCode(...)` over local wrappers or string dtype checks. Performance: Using Object type instead of DLDataType would indeed bring some performance impact to the IR. We have done the following performance optimizations: - Make sure most of the outputs reuse one of the PrimType from inputs - Cache a thread local PrimType based on input so we don't repeatly realloc We did benchmarks show that rewrite simplify operation stays within +-10% overhead of original one. Which merits the refactor given the benefit the unfication brings
95 lines
2.9 KiB
Python
95 lines
2.9 KiB
Python
# 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.
|
|
# ruff: noqa: F841
|
|
import numpy
|
|
|
|
import tvm
|
|
import tvm.testing
|
|
from tvm.script import tirx as T
|
|
|
|
# This numpy array is used to test the comparison between the global objects and the
|
|
# `tvm.script.tirx` submodule.
|
|
np_array = numpy.array([0, 1, 2, 3])
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def matmul(a: T.handle, b: T.handle, c: T.handle) -> None:
|
|
A = T.match_buffer(a, [128, 128])
|
|
B = T.match_buffer(b, [128, 128])
|
|
C = T.match_buffer(c, [128, 128])
|
|
|
|
for i, j, k in T.grid(128, 128, 128):
|
|
with T.sblock("update"):
|
|
vi, vj, vk = T.axis.remap("SSR", [i, j, k])
|
|
with T.init():
|
|
C[vi, vj] = T.float32(0)
|
|
C[vi, vj] = C[vi, vj] + A[vi, vk] * B[vj, vk]
|
|
|
|
|
|
def test_multi_element_array_in_outmost_namespace():
|
|
func = matmul
|
|
rt_func = tvm.script.from_source(func.script())
|
|
tvm.ir.assert_structural_equal(func, rt_func)
|
|
|
|
|
|
def test_different_dtype_assignment_to_var():
|
|
@T.prim_func(s_tir=True)
|
|
def test_case():
|
|
a = T.sblock_alloc_buffer((10, 10), dtype="int8")
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def func_ref():
|
|
a = T.sblock_alloc_buffer([10, 10], dtype="int8")
|
|
T.evaluate(0)
|
|
|
|
tvm.ir.assert_structural_equal(
|
|
test_case.with_attr("global_symbol", "main"), func_ref.with_attr("global_symbol", "main")
|
|
)
|
|
|
|
|
|
def test_var_capturing_order():
|
|
b = 2
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def test_case():
|
|
k: T.let[T.int32] = b
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def func_ref():
|
|
k: T.let[T.int32] = 2
|
|
T.evaluate(0)
|
|
|
|
tvm.ir.assert_structural_equal(
|
|
test_case.with_attr("global_symbol", "main"), func_ref.with_attr("global_symbol", "main")
|
|
)
|
|
|
|
|
|
def test_tir_buffer_region_extent_correct_dtype():
|
|
@T.prim_func(s_tir=True)
|
|
def func(A: T.Buffer((T.int64(16), T.int64(1)), "float32")):
|
|
for i in T.grid(T.int64(16)):
|
|
with T.sblock("block"):
|
|
vi = T.axis.remap("S", [i])
|
|
T.reads(A[vi, T.int64(0) : T.int64(1)])
|
|
T.evaluate(0)
|
|
|
|
assert func.body.block.body.body.block.reads[0].region[0].extent.ty.dtype == "int64"
|
|
|
|
|
|
if __name__ == "__main__":
|
|
tvm.testing.main()
|