vinx13 commented on a change in pull request #9432: URL: https://github.com/apache/tvm/pull/9432#discussion_r743148098
########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... +def max_value(dtype): ... +def floordiv(x: PrimExpr, y: PrimExpr): ... +def floormod(x: PrimExpr, y: PrimExpr): ... +def abs(x): ... +def load(dtype, var, index, predicate=None): ... +def cast(value, dtype): ... +def ramp(base, stride, lanes): ... +def broadcast(value, lanes): ... +def iter_var(var, dom, iter_type, thread_tag): ... +def max(a, b): ... +def min(a, b): ... +def get_axis(begin, end, iter_type): ... +def Select(cond, if_body, else_body): ... +def evaluate(value): ... +def store(var, index, value, predicate=True): ... +def comm_reducer(lambda_io, identities): ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype, storage_scope) -> IterVar: ... Review comment: ```suggestion def buffer_var(dtype: str, storage_scope: str) -> IterVar: ... ``` ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... Review comment: can we type these functions? ########## File path: tests/python/unittest/test_tvmscript_type.py ########## @@ -0,0 +1,88 @@ +# 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. +# pylint: disable=missing-function-docstring,missing-module-docstring +import sys Review comment: not needed ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... +def max_value(dtype): ... +def floordiv(x: PrimExpr, y: PrimExpr): ... +def floormod(x: PrimExpr, y: PrimExpr): ... +def abs(x): ... +def load(dtype, var, index, predicate=None): ... +def cast(value, dtype): ... +def ramp(base, stride, lanes): ... +def broadcast(value, lanes): ... +def iter_var(var, dom, iter_type, thread_tag): ... +def max(a, b): ... +def min(a, b): ... +def get_axis(begin, end, iter_type): ... +def Select(cond, if_body, else_body): ... +def evaluate(value): ... +def store(var, index, value, predicate=True): ... +def comm_reducer(lambda_io, identities): ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype, storage_scope) -> IterVar: ... +def func_attr(attrs: Dict) -> None: ... +def prim_func(input_func: Callable) -> PrimFunc: ... + +""" +special_stmt - Threads and Bindings +""" + +def env_thread(env_name: str) -> IterVar: ... +def bind(iter_var: IterVar, expr: PrimExpr) -> None: ... + +""" +Scope handler +""" + +class block(ContextManager): + def __init__(self, name_hint: str = "") -> None: ... + def __enter__(self) -> Sequence[IterVar]: ... + +class init(ContextManager): + def __init__(self) -> None: ... + +class let(ContextManager): + def __init__(self, var: Var, value: PrimExpr) -> None: ... + +def where(cond: PrimExpr) -> None: ... +def allocate( + extents, dtype, scope: str, condition: builtins.bool = True, annotations=None +) -> None: ... +def launch_thread(env_var, extent) -> None: ... +def realize(buffer_slice: BufferSlice, scope: str, condition: builtins.bool = True) -> None: ... Review comment: same as above ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... +def max_value(dtype): ... +def floordiv(x: PrimExpr, y: PrimExpr): ... +def floormod(x: PrimExpr, y: PrimExpr): ... +def abs(x): ... +def load(dtype, var, index, predicate=None): ... +def cast(value, dtype): ... +def ramp(base, stride, lanes): ... +def broadcast(value, lanes): ... +def iter_var(var, dom, iter_type, thread_tag): ... +def max(a, b): ... +def min(a, b): ... +def get_axis(begin, end, iter_type): ... +def Select(cond, if_body, else_body): ... +def evaluate(value): ... +def store(var, index, value, predicate=True): ... +def comm_reducer(lambda_io, identities): ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype, storage_scope) -> IterVar: ... +def func_attr(attrs: Dict) -> None: ... +def prim_func(input_func: Callable) -> PrimFunc: ... + +""" +special_stmt - Threads and Bindings +""" + +def env_thread(env_name: str) -> IterVar: ... +def bind(iter_var: IterVar, expr: PrimExpr) -> None: ... + +""" +Scope handler +""" + +class block(ContextManager): + def __init__(self, name_hint: str = "") -> None: ... + def __enter__(self) -> Sequence[IterVar]: ... + +class init(ContextManager): + def __init__(self) -> None: ... + +class let(ContextManager): + def __init__(self, var: Var, value: PrimExpr) -> None: ... + +def where(cond: PrimExpr) -> None: ... +def allocate( + extents, dtype, scope: str, condition: builtins.bool = True, annotations=None Review comment: ```suggestion extents, dtype, scope: str, condition: Union[PrimExpr, builtins.bool] = True, annotations=None ``` if possible, please also type other args ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... +def max_value(dtype): ... +def floordiv(x: PrimExpr, y: PrimExpr): ... +def floormod(x: PrimExpr, y: PrimExpr): ... +def abs(x): ... +def load(dtype, var, index, predicate=None): ... +def cast(value, dtype): ... +def ramp(base, stride, lanes): ... +def broadcast(value, lanes): ... +def iter_var(var, dom, iter_type, thread_tag): ... +def max(a, b): ... +def min(a, b): ... +def get_axis(begin, end, iter_type): ... +def Select(cond, if_body, else_body): ... +def evaluate(value): ... +def store(var, index, value, predicate=True): ... +def comm_reducer(lambda_io, identities): ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype, storage_scope) -> IterVar: ... Review comment: ```suggestion def buffer_var(dtype: str, storage_scope: str) -> Var: ... ``` ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... Review comment: can we type these functions? If there are some args that can't be typed, please explicitly add `Any` as its type ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,278 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__(self: Buffer, *pos: Union[int, PrimExpr, slice, Any]) -> Buffer: ... Review comment: isn't the return value `PrimExpr`? ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,278 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): Review comment: `Buffer` shouldn't subclass `Var` ########## File path: tests/python/unittest/test_tvmscript_type.py ########## @@ -0,0 +1,87 @@ +# 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. +# pylint: disable=missing-function-docstring,missing-module-docstring +from tvm.script import tir as T + +""" +This prim_func tests the type of Review comment: I think rather than list these types here, it would be clearer to just generally describe the intention of these cases. ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,278 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): Review comment: is it because `__setitem__` is missing? ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,285 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__(self: Buffer, *pos: Union[int, PrimExpr, slice, Any]) -> PrimExpr: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype: str) -> PrimExpr: ... +def max_value(dtype: str) -> PrimExpr: ... +def floordiv(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def floormod(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def truncmod(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def truncdiv(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def abs(x: PrimExpr) -> PrimExpr: ... +def load(dtype: str, var: Var, index: PrimExpr, predicate: PrimExpr = None) -> PrimExpr: ... +def cast(value: PrimExpr, dtype: str) -> PrimExpr: ... +def ramp(base: PrimExpr, stride: Any, lanes: int) -> PrimExpr: ... +def broadcast(value: PrimExpr, lanes: int) -> PrimExpr: ... +def iter_var(var: Union[Var, str], dom: Range, iter_type: int, thread_tag: str) -> IterVar: ... +def max(a: PrimExpr, b: PrimExpr) -> PrimExpr: ... +def min(a: PrimExpr, b: PrimExpr) -> PrimExpr: ... +def Select(cond: PrimExpr, if_body: PrimExpr, else_body: PrimExpr) -> PrimExpr: ... +def if_then_else(cond: PrimExpr, t: PrimExpr, f: PrimExpr, dtype: str) -> PrimExpr: ... +def evaluate(value: PrimExpr) -> PrimExpr: ... +def reinterpret(value: PrimExpr, dtype: str) -> PrimExpr: ... +def store(var: Var, index: PrimExpr, value: PrimExpr, predicate: PrimExpr = True) -> None: ... +def comm_reducer(lambda_io: Tuple[List[PrimExpr]], identities: List[PrimExpr]) -> PrimExpr: ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +def get_axis(begin: PrimExpr, end: PrimExpr, iter_type: int) -> IterVar: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype: str, storage_scope: str) -> Var: ... +def func_attr(attrs: Mapping[str, Object]) -> None: ... +def prim_func(input_func: Callable) -> PrimFunc: ... + +""" +special_stmt - Threads and Bindings +""" + +def env_thread(env_name: str) -> IterVar: ... +def bind(iter_var: IterVar, expr: PrimExpr) -> None: ... + +""" +Scope handler +""" + +class block(ContextManager): + def __init__(self, name_hint: str = "") -> None: ... + def __enter__(self) -> Sequence[IterVar]: ... + +class init(ContextManager): + def __init__(self) -> None: ... + +class let(ContextManager): + def __init__(self, var: Var, value: PrimExpr) -> None: ... + +def where(cond: PrimExpr) -> None: ... +def allocate( + extents: List[PrimExpr], + dtype: str, + scope: str, + condition: Union[PrimExpr, builtins.bool] = True, + annotations: Optional[Mapping[str, Object]] = None, +) -> Var: ... +def launch_thread(env_var: Var, extent: PrimExpr) -> Var: ... +def realize( + buffer_slice: BufferSlice, scope: str, condition: Union[PrimExpr, builtins.bool] = True +) -> None: ... +def attr(node: PrimExpr, attr_key: str, value: PrimExpr) -> None: ... +def Assert(condition: Union[PrimExpr, builtins.bool], message: str) -> PrimExpr: ... + +""" +Scope handler - Loops +""" + +def serial( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def parallel( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def vectorized( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def unroll( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def thread_binding( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + thread: str, + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def for_range( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int] = None, + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def grid(*extents: Union[PrimExpr, int]) -> Iterable[Tuple[IterVar]]: ... + +""" +ty - redefine types +""" + +class boolean: ... + +class handle: + def __getitem__(self: handle, pos: Tuple[Union[Number, PrimExpr, slice]]) -> Buffer: ... + @property + def data(self: handle) -> Ptr: ... Review comment: are they needed? do you have a use case? ########## File path: python/tvm/script/tir/ty.py ########## @@ -30,8 +30,11 @@ def evaluate(self): """Return an actual ir.Type Object that this Generic class wraps""" raise TypeError("Cannot get tvm.Type from a generic type") + def __call__(self): + raise NotImplementedError Review comment: also add a comment explain why this is needed ```suggestion raise NotImplementedError() ``` ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,285 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__(self: Buffer, *pos: Union[int, PrimExpr, slice, Any]) -> PrimExpr: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype: str) -> PrimExpr: ... +def max_value(dtype: str) -> PrimExpr: ... +def floordiv(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def floormod(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def truncmod(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def truncdiv(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def abs(x: PrimExpr) -> PrimExpr: ... +def load(dtype: str, var: Var, index: PrimExpr, predicate: PrimExpr = None) -> PrimExpr: ... +def cast(value: PrimExpr, dtype: str) -> PrimExpr: ... +def ramp(base: PrimExpr, stride: Any, lanes: int) -> PrimExpr: ... +def broadcast(value: PrimExpr, lanes: int) -> PrimExpr: ... +def iter_var(var: Union[Var, str], dom: Range, iter_type: int, thread_tag: str) -> IterVar: ... +def max(a: PrimExpr, b: PrimExpr) -> PrimExpr: ... +def min(a: PrimExpr, b: PrimExpr) -> PrimExpr: ... +def Select(cond: PrimExpr, if_body: PrimExpr, else_body: PrimExpr) -> PrimExpr: ... +def if_then_else(cond: PrimExpr, t: PrimExpr, f: PrimExpr, dtype: str) -> PrimExpr: ... +def evaluate(value: PrimExpr) -> PrimExpr: ... +def reinterpret(value: PrimExpr, dtype: str) -> PrimExpr: ... +def store(var: Var, index: PrimExpr, value: PrimExpr, predicate: PrimExpr = True) -> None: ... +def comm_reducer(lambda_io: Tuple[List[PrimExpr]], identities: List[PrimExpr]) -> PrimExpr: ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +def get_axis(begin: PrimExpr, end: PrimExpr, iter_type: int) -> IterVar: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype: str, storage_scope: str) -> Var: ... +def func_attr(attrs: Mapping[str, Object]) -> None: ... +def prim_func(input_func: Callable) -> PrimFunc: ... + +""" +special_stmt - Threads and Bindings +""" + +def env_thread(env_name: str) -> IterVar: ... +def bind(iter_var: IterVar, expr: PrimExpr) -> None: ... + +""" +Scope handler +""" + +class block(ContextManager): + def __init__(self, name_hint: str = "") -> None: ... + def __enter__(self) -> Sequence[IterVar]: ... + +class init(ContextManager): + def __init__(self) -> None: ... + +class let(ContextManager): + def __init__(self, var: Var, value: PrimExpr) -> None: ... + +def where(cond: PrimExpr) -> None: ... +def allocate( + extents: List[PrimExpr], + dtype: str, + scope: str, + condition: Union[PrimExpr, builtins.bool] = True, + annotations: Optional[Mapping[str, Object]] = None, +) -> Var: ... +def launch_thread(env_var: Var, extent: PrimExpr) -> Var: ... +def realize( + buffer_slice: BufferSlice, scope: str, condition: Union[PrimExpr, builtins.bool] = True +) -> None: ... +def attr(node: PrimExpr, attr_key: str, value: PrimExpr) -> None: ... +def Assert(condition: Union[PrimExpr, builtins.bool], message: str) -> PrimExpr: ... + +""" +Scope handler - Loops +""" + +def serial( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def parallel( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def vectorized( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def unroll( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def thread_binding( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + thread: str, + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def for_range( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int] = None, + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def grid(*extents: Union[PrimExpr, int]) -> Iterable[Tuple[IterVar]]: ... + +""" +ty - redefine types +""" + +class boolean: ... + +class handle: + def __getitem__(self: handle, pos: Tuple[Union[Number, PrimExpr, slice]]) -> Buffer: ... + @property + def data(self: handle) -> Ptr: ... + +# class float32: +# def __new__(self, imm: Union[PrimExpr, Number]) -> PrimExpr: ... +# def __init__(self, imm: Union[PrimExpr, Number]) -> None: ... Review comment: remove if not needed ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... +def max_value(dtype): ... +def floordiv(x: PrimExpr, y: PrimExpr): ... +def floormod(x: PrimExpr, y: PrimExpr): ... +def abs(x): ... +def load(dtype, var, index, predicate=None): ... +def cast(value, dtype): ... +def ramp(base, stride, lanes): ... +def broadcast(value, lanes): ... +def iter_var(var, dom, iter_type, thread_tag): ... +def max(a, b): ... +def min(a, b): ... +def get_axis(begin, end, iter_type): ... +def Select(cond, if_body, else_body): ... +def evaluate(value): ... +def store(var, index, value, predicate=True): ... +def comm_reducer(lambda_io, identities): ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype, storage_scope) -> IterVar: ... Review comment: ```suggestion def buffer_var(dtype: str, storage_scope: str) -> IterVar: ... ``` ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... Review comment: can we type these functions? ########## File path: tests/python/unittest/test_tvmscript_type.py ########## @@ -0,0 +1,88 @@ +# 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. +# pylint: disable=missing-function-docstring,missing-module-docstring +import sys Review comment: not needed ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... +def max_value(dtype): ... +def floordiv(x: PrimExpr, y: PrimExpr): ... +def floormod(x: PrimExpr, y: PrimExpr): ... +def abs(x): ... +def load(dtype, var, index, predicate=None): ... +def cast(value, dtype): ... +def ramp(base, stride, lanes): ... +def broadcast(value, lanes): ... +def iter_var(var, dom, iter_type, thread_tag): ... +def max(a, b): ... +def min(a, b): ... +def get_axis(begin, end, iter_type): ... +def Select(cond, if_body, else_body): ... +def evaluate(value): ... +def store(var, index, value, predicate=True): ... +def comm_reducer(lambda_io, identities): ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype, storage_scope) -> IterVar: ... +def func_attr(attrs: Dict) -> None: ... +def prim_func(input_func: Callable) -> PrimFunc: ... + +""" +special_stmt - Threads and Bindings +""" + +def env_thread(env_name: str) -> IterVar: ... +def bind(iter_var: IterVar, expr: PrimExpr) -> None: ... + +""" +Scope handler +""" + +class block(ContextManager): + def __init__(self, name_hint: str = "") -> None: ... + def __enter__(self) -> Sequence[IterVar]: ... + +class init(ContextManager): + def __init__(self) -> None: ... + +class let(ContextManager): + def __init__(self, var: Var, value: PrimExpr) -> None: ... + +def where(cond: PrimExpr) -> None: ... +def allocate( + extents, dtype, scope: str, condition: builtins.bool = True, annotations=None +) -> None: ... +def launch_thread(env_var, extent) -> None: ... +def realize(buffer_slice: BufferSlice, scope: str, condition: builtins.bool = True) -> None: ... Review comment: same as above ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... +def max_value(dtype): ... +def floordiv(x: PrimExpr, y: PrimExpr): ... +def floormod(x: PrimExpr, y: PrimExpr): ... +def abs(x): ... +def load(dtype, var, index, predicate=None): ... +def cast(value, dtype): ... +def ramp(base, stride, lanes): ... +def broadcast(value, lanes): ... +def iter_var(var, dom, iter_type, thread_tag): ... +def max(a, b): ... +def min(a, b): ... +def get_axis(begin, end, iter_type): ... +def Select(cond, if_body, else_body): ... +def evaluate(value): ... +def store(var, index, value, predicate=True): ... +def comm_reducer(lambda_io, identities): ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype, storage_scope) -> IterVar: ... +def func_attr(attrs: Dict) -> None: ... +def prim_func(input_func: Callable) -> PrimFunc: ... + +""" +special_stmt - Threads and Bindings +""" + +def env_thread(env_name: str) -> IterVar: ... +def bind(iter_var: IterVar, expr: PrimExpr) -> None: ... + +""" +Scope handler +""" + +class block(ContextManager): + def __init__(self, name_hint: str = "") -> None: ... + def __enter__(self) -> Sequence[IterVar]: ... + +class init(ContextManager): + def __init__(self) -> None: ... + +class let(ContextManager): + def __init__(self, var: Var, value: PrimExpr) -> None: ... + +def where(cond: PrimExpr) -> None: ... +def allocate( + extents, dtype, scope: str, condition: builtins.bool = True, annotations=None Review comment: ```suggestion extents, dtype, scope: str, condition: Union[PrimExpr, builtins.bool] = True, annotations=None ``` if possible, please also type other args ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... +def max_value(dtype): ... +def floordiv(x: PrimExpr, y: PrimExpr): ... +def floormod(x: PrimExpr, y: PrimExpr): ... +def abs(x): ... +def load(dtype, var, index, predicate=None): ... +def cast(value, dtype): ... +def ramp(base, stride, lanes): ... +def broadcast(value, lanes): ... +def iter_var(var, dom, iter_type, thread_tag): ... +def max(a, b): ... +def min(a, b): ... +def get_axis(begin, end, iter_type): ... +def Select(cond, if_body, else_body): ... +def evaluate(value): ... +def store(var, index, value, predicate=True): ... +def comm_reducer(lambda_io, identities): ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data=None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype, storage_scope) -> IterVar: ... Review comment: ```suggestion def buffer_var(dtype: str, storage_scope: str) -> Var: ... ``` ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,267 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__( + self: Buffer, pos: Tuple[Union[int, PrimExpr], Union[int, PrimExpr]] + ) -> Buffer: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype): ... Review comment: can we type these functions? If there are some args that can't be typed, please explicitly add `Any` as its type ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,278 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__(self: Buffer, *pos: Union[int, PrimExpr, slice, Any]) -> Buffer: ... Review comment: isn't the return value `PrimExpr`? ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,278 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): Review comment: `Buffer` shouldn't subclass `Var` ########## File path: tests/python/unittest/test_tvmscript_type.py ########## @@ -0,0 +1,87 @@ +# 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. +# pylint: disable=missing-function-docstring,missing-module-docstring +from tvm.script import tir as T + +""" +This prim_func tests the type of Review comment: I think rather than list these types here, it would be clearer to just generally describe the intention of these cases. ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,278 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): Review comment: is it because `__setitem__` is missing? ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,285 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__(self: Buffer, *pos: Union[int, PrimExpr, slice, Any]) -> PrimExpr: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype: str) -> PrimExpr: ... +def max_value(dtype: str) -> PrimExpr: ... +def floordiv(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def floormod(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def truncmod(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def truncdiv(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def abs(x: PrimExpr) -> PrimExpr: ... +def load(dtype: str, var: Var, index: PrimExpr, predicate: PrimExpr = None) -> PrimExpr: ... +def cast(value: PrimExpr, dtype: str) -> PrimExpr: ... +def ramp(base: PrimExpr, stride: Any, lanes: int) -> PrimExpr: ... +def broadcast(value: PrimExpr, lanes: int) -> PrimExpr: ... +def iter_var(var: Union[Var, str], dom: Range, iter_type: int, thread_tag: str) -> IterVar: ... +def max(a: PrimExpr, b: PrimExpr) -> PrimExpr: ... +def min(a: PrimExpr, b: PrimExpr) -> PrimExpr: ... +def Select(cond: PrimExpr, if_body: PrimExpr, else_body: PrimExpr) -> PrimExpr: ... +def if_then_else(cond: PrimExpr, t: PrimExpr, f: PrimExpr, dtype: str) -> PrimExpr: ... +def evaluate(value: PrimExpr) -> PrimExpr: ... +def reinterpret(value: PrimExpr, dtype: str) -> PrimExpr: ... +def store(var: Var, index: PrimExpr, value: PrimExpr, predicate: PrimExpr = True) -> None: ... +def comm_reducer(lambda_io: Tuple[List[PrimExpr]], identities: List[PrimExpr]) -> PrimExpr: ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +def get_axis(begin: PrimExpr, end: PrimExpr, iter_type: int) -> IterVar: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype: str, storage_scope: str) -> Var: ... +def func_attr(attrs: Mapping[str, Object]) -> None: ... +def prim_func(input_func: Callable) -> PrimFunc: ... + +""" +special_stmt - Threads and Bindings +""" + +def env_thread(env_name: str) -> IterVar: ... +def bind(iter_var: IterVar, expr: PrimExpr) -> None: ... + +""" +Scope handler +""" + +class block(ContextManager): + def __init__(self, name_hint: str = "") -> None: ... + def __enter__(self) -> Sequence[IterVar]: ... + +class init(ContextManager): + def __init__(self) -> None: ... + +class let(ContextManager): + def __init__(self, var: Var, value: PrimExpr) -> None: ... + +def where(cond: PrimExpr) -> None: ... +def allocate( + extents: List[PrimExpr], + dtype: str, + scope: str, + condition: Union[PrimExpr, builtins.bool] = True, + annotations: Optional[Mapping[str, Object]] = None, +) -> Var: ... +def launch_thread(env_var: Var, extent: PrimExpr) -> Var: ... +def realize( + buffer_slice: BufferSlice, scope: str, condition: Union[PrimExpr, builtins.bool] = True +) -> None: ... +def attr(node: PrimExpr, attr_key: str, value: PrimExpr) -> None: ... +def Assert(condition: Union[PrimExpr, builtins.bool], message: str) -> PrimExpr: ... + +""" +Scope handler - Loops +""" + +def serial( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def parallel( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def vectorized( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def unroll( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def thread_binding( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + thread: str, + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def for_range( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int] = None, + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def grid(*extents: Union[PrimExpr, int]) -> Iterable[Tuple[IterVar]]: ... + +""" +ty - redefine types +""" + +class boolean: ... + +class handle: + def __getitem__(self: handle, pos: Tuple[Union[Number, PrimExpr, slice]]) -> Buffer: ... + @property + def data(self: handle) -> Ptr: ... Review comment: are they needed? do you have a use case? ########## File path: python/tvm/script/tir/ty.py ########## @@ -30,8 +30,11 @@ def evaluate(self): """Return an actual ir.Type Object that this Generic class wraps""" raise TypeError("Cannot get tvm.Type from a generic type") + def __call__(self): + raise NotImplementedError Review comment: also add a comment explain why this is needed ```suggestion raise NotImplementedError() ``` ########## File path: python/tvm/script/tir/__init__.pyi ########## @@ -0,0 +1,285 @@ +# 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. +# pylint: disable=redefined-builtin +from typing import ( + Any, + Callable, + ContextManager, + Dict, + Iterable, + Optional, + Tuple, + Union, + Sequence, + List, + Mapping, + overload, +) +from numbers import Number +import builtins + +from tvm.tir.function import PrimFunc +from tvm.tir import PrimExpr, Range, IterVar, Var +from tvm.runtime import Object +from .node import BufferSlice + +""" +redefine types +""" + +class Buffer(Var): + def __getitem__(self: Buffer, *pos: Union[int, PrimExpr, slice, Any]) -> PrimExpr: ... + @property + def data(self: Buffer) -> Ptr: ... + +""" +Variables and constants +""" + +def bool(imm: Union[PrimExpr, builtins.bool, Number]) -> PrimExpr: ... +def int8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def int64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def uint64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float8(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float16(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float32(imm: Union[PrimExpr, Number]) -> PrimExpr: ... +def float64(imm: Union[PrimExpr, Number]) -> PrimExpr: ... + +""" +Intrinsic +""" + +def min_value(dtype: str) -> PrimExpr: ... +def max_value(dtype: str) -> PrimExpr: ... +def floordiv(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def floormod(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def truncmod(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def truncdiv(x: PrimExpr, y: PrimExpr) -> PrimExpr: ... +def abs(x: PrimExpr) -> PrimExpr: ... +def load(dtype: str, var: Var, index: PrimExpr, predicate: PrimExpr = None) -> PrimExpr: ... +def cast(value: PrimExpr, dtype: str) -> PrimExpr: ... +def ramp(base: PrimExpr, stride: Any, lanes: int) -> PrimExpr: ... +def broadcast(value: PrimExpr, lanes: int) -> PrimExpr: ... +def iter_var(var: Union[Var, str], dom: Range, iter_type: int, thread_tag: str) -> IterVar: ... +def max(a: PrimExpr, b: PrimExpr) -> PrimExpr: ... +def min(a: PrimExpr, b: PrimExpr) -> PrimExpr: ... +def Select(cond: PrimExpr, if_body: PrimExpr, else_body: PrimExpr) -> PrimExpr: ... +def if_then_else(cond: PrimExpr, t: PrimExpr, f: PrimExpr, dtype: str) -> PrimExpr: ... +def evaluate(value: PrimExpr) -> PrimExpr: ... +def reinterpret(value: PrimExpr, dtype: str) -> PrimExpr: ... +def store(var: Var, index: PrimExpr, value: PrimExpr, predicate: PrimExpr = True) -> None: ... +def comm_reducer(lambda_io: Tuple[List[PrimExpr]], identities: List[PrimExpr]) -> PrimExpr: ... + +""" +Unary operator +""" + +def exp2(x: PrimExpr) -> PrimExpr: ... +def exp10(x: PrimExpr) -> PrimExpr: ... +def erf(x: PrimExpr) -> PrimExpr: ... +def tanh(x: PrimExpr) -> PrimExpr: ... +def sigmoid(x: PrimExpr) -> PrimExpr: ... +def log(x: PrimExpr) -> PrimExpr: ... +def log2(x: PrimExpr) -> PrimExpr: ... +def log10(x: PrimExpr) -> PrimExpr: ... +def log1p(x: PrimExpr) -> PrimExpr: ... +def tan(x: PrimExpr) -> PrimExpr: ... +def cos(x: PrimExpr) -> PrimExpr: ... +def cosh(x: PrimExpr) -> PrimExpr: ... +def acos(x: PrimExpr) -> PrimExpr: ... +def acosh(x: PrimExpr) -> PrimExpr: ... +def sin(x: PrimExpr) -> PrimExpr: ... +def sinh(x: PrimExpr) -> PrimExpr: ... +def asin(x: PrimExpr) -> PrimExpr: ... +def asinh(x: PrimExpr) -> PrimExpr: ... +def atan(x: PrimExpr) -> PrimExpr: ... +def atanh(x: PrimExpr) -> PrimExpr: ... +def atan2(x: PrimExpr) -> PrimExpr: ... +def sqrt(x: PrimExpr) -> PrimExpr: ... +def rsqrt(x: PrimExpr) -> PrimExpr: ... + +""" +special_stmt - Buffers +""" + +def match_buffer( + param: Union[Var, BufferSlice], + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def buffer_decl( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... +def alloc_buffer( + shape: Sequence[Union[PrimExpr, int]], + dtype: str = "float32", + data: Var = None, + strides: Optional[Sequence[int]] = None, + elem_offset: Optional[int] = None, + scope: str = "global", + align: int = -1, + offset_factor: int = 0, + buffer_type: str = "default", +) -> Buffer: ... + +""" +special_stmt - Reads/Writes +""" + +def reads(read_regions: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def writes(write_region: Union[BufferSlice, List[BufferSlice]]) -> None: ... +def block_attr(attrs: Mapping[str, Object]) -> None: ... + +""" +special_stmt - Axis +""" + +class axis: + @staticmethod + def spatial(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def S(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def reduce(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def R(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def scan(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def opaque(dom: Union[PrimExpr, Tuple[PrimExpr, PrimExpr]], value: PrimExpr) -> IterVar: ... + @staticmethod + def remap(iter_types: str, loop_vars: List[Var]) -> List[IterVar]: ... + +def get_axis(begin: PrimExpr, end: PrimExpr, iter_type: int) -> IterVar: ... + +""" +special_stmt - Annotations +""" + +def buffer_var(dtype: str, storage_scope: str) -> Var: ... +def func_attr(attrs: Mapping[str, Object]) -> None: ... +def prim_func(input_func: Callable) -> PrimFunc: ... + +""" +special_stmt - Threads and Bindings +""" + +def env_thread(env_name: str) -> IterVar: ... +def bind(iter_var: IterVar, expr: PrimExpr) -> None: ... + +""" +Scope handler +""" + +class block(ContextManager): + def __init__(self, name_hint: str = "") -> None: ... + def __enter__(self) -> Sequence[IterVar]: ... + +class init(ContextManager): + def __init__(self) -> None: ... + +class let(ContextManager): + def __init__(self, var: Var, value: PrimExpr) -> None: ... + +def where(cond: PrimExpr) -> None: ... +def allocate( + extents: List[PrimExpr], + dtype: str, + scope: str, + condition: Union[PrimExpr, builtins.bool] = True, + annotations: Optional[Mapping[str, Object]] = None, +) -> Var: ... +def launch_thread(env_var: Var, extent: PrimExpr) -> Var: ... +def realize( + buffer_slice: BufferSlice, scope: str, condition: Union[PrimExpr, builtins.bool] = True +) -> None: ... +def attr(node: PrimExpr, attr_key: str, value: PrimExpr) -> None: ... +def Assert(condition: Union[PrimExpr, builtins.bool], message: str) -> PrimExpr: ... + +""" +Scope handler - Loops +""" + +def serial( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def parallel( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def vectorized( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def unroll( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def thread_binding( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int], + thread: str, + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def for_range( + begin: Union[PrimExpr, int], + end: Union[PrimExpr, int] = None, + annotations: Optional[Mapping[str, Object]] = None, +) -> Iterable[IterVar]: ... +def grid(*extents: Union[PrimExpr, int]) -> Iterable[Tuple[IterVar]]: ... + +""" +ty - redefine types +""" + +class boolean: ... + +class handle: + def __getitem__(self: handle, pos: Tuple[Union[Number, PrimExpr, slice]]) -> Buffer: ... + @property + def data(self: handle) -> Ptr: ... + +# class float32: +# def __new__(self, imm: Union[PrimExpr, Number]) -> PrimExpr: ... +# def __init__(self, imm: Union[PrimExpr, Number]) -> None: ... Review comment: remove if not needed -- 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]
