This is an automated email from the ASF dual-hosted git repository. tqchen pushed a commit to branch script/canonical-parser-df in repository https://gitbox.apache.org/repos/asf/tvm.git
commit 8d3f02362a7f0bebcc52489bf8a85bbdd742a4fe Author: Tianqi Chen <[email protected]> AuthorDate: Wed Sep 23 18:35:16 2026 +0000 [Script] Propagate original errors and consolidate source inspection --- docs/arch/tvmscript.rst | 15 +- python/tvm/relax/script/__init__.py | 5 +- python/tvm/relax/script/builder/parser_protocol.py | 2 +- python/tvm/script/ir_builder/base.py | 33 +-- python/tvm/script/ir_builder/ir/parser_protocol.py | 2 +- python/tvm/script/parser/call_args_policy.py | 12 +- python/tvm/script/parser/diagnostics.py | 159 -------------- python/tvm/script/parser/entry.py | 31 ++- python/tvm/script/parser/inspect_source.py | 112 +++++++++- python/tvm/script/parser/source.py | 125 ----------- python/tvm/script/parser/transpile.py | 19 +- python/tvm/tirx/script/builder/parser_protocol.py | 2 +- .../test_frontend_relax_construction_validation.py | 3 +- .../test_relax_parser_comparison_integration.py | 4 +- tests/python/relax/test_op_view.py | 21 +- tests/python/relax/test_transform.py | 12 +- tests/python/relax/test_tvmscript_parser.py | 97 +++++---- .../s_tir/test_s_tir_axis_binding_validation.py | 3 +- .../python/script/test_capture_definition_scope.py | 9 +- tests/python/script/test_capture_frontend_scope.py | 2 +- tests/python/script/test_parser.py | 29 +-- .../python/script/test_parser_capture_lifetime.py | 33 ++- tests/python/script/test_parser_host_selection.py | 22 +- tests/python/script/test_parser_latest_design.py | 47 ++++- tests/python/script/test_parser_scope_helpers.py | 11 +- tests/python/script/test_parser_scoped_symbols.py | 7 +- .../python/tirx/codegen/test_codegen_blackwell.py | 10 +- tests/python/tirx/codegen/test_codegen_cuda.py | 2 +- tests/python/tirx/codegen/test_ptx_addr.py | 24 +-- tests/python/tirx/codegen/test_ptx_dialect.py | 132 ++++++------ .../test_frontend_construction_validation.py | 5 +- .../script/test_native_span_and_loop_contracts.py | 21 +- .../python/tirx/script/test_parser_buffer_views.py | 3 +- .../tirx/script/test_parser_jit_specialization.py | 6 +- .../tirx/script/test_parser_tirx_emission.py | 5 +- .../test_tirx_parser_comparison_integration.py | 4 +- tests/python/tirx/test_jit.py | 21 +- tests/python/tirx/test_parser_printer.py | 19 +- .../tvmscript/test_tvmscript_error_report.py | 232 +++++++++------------ tests/python/tvmscript/test_tvmscript_ops.py | 2 +- .../tvmscript/test_tvmscript_parser_source.py | 4 +- .../python/tvmscript/test_tvmscript_parser_tir.py | 8 +- .../tvmscript/test_tvmscript_syntax_sugar.py | 2 +- 43 files changed, 585 insertions(+), 732 deletions(-) diff --git a/docs/arch/tvmscript.rst b/docs/arch/tvmscript.rst index a3aae7c1e9..df5d045485 100644 --- a/docs/arch/tvmscript.rst +++ b/docs/arch/tvmscript.rst @@ -156,8 +156,10 @@ mechanism. Parser Architecture ------------------- -The parser lives in ``python/tvm/script/parser/``. Documented functions in ``entry.py`` -manage source acquisition, definition-context capture, callable composition and execution. +The parser lives in ``python/tvm/script/parser/``. ``inspect_source.py`` owns source +acquisition, the ``Source`` representation and its coordinate helpers, and definition-context +capture. Documented functions in ``entry.py`` manage the parse sequence, callable composition +and execution. The transpiler handles syntax; the builders interpret concrete values. TIRx JIT validation, defaults and specialization caching live in ``tvm.tirx.script.jit``; shared parser utilities remain independent of that policy. @@ -170,8 +172,8 @@ then lowers to operations supplied by that namespace: for example, unmarked ``if TIR control-flow frames in a primitive function and Relax control-flow frames in a Relax function. The transpiler emits these operations without implementing dialect IR semantics. -``protocol.py`` records callable syntax metadata. Argument policies such as ``expr_str`` -translate symbolic strings written directly in source expressions, while ``global_info`` +``ir_builder/ir/parser_protocol.py`` records callable syntax metadata. Argument policies such +as ``expr_str`` translate symbolic strings written directly in source expressions, while ``global_info`` preserves a module reference for builder-side lookup. Dtype and placement strings remain literal. Captured or computed symbolic shapes must already contain explicit IR variables; the parser does not interpret expression strings found inside captured values. @@ -181,6 +183,11 @@ and loops and scopes become builder contexts. Concrete binding, type checking, c construction, and frame finalization belong to the builders. Ordinary host calls and operator overloads execute as part of the generated Python program. +Parser syntax restrictions raise ``SyntaxError`` with the original filename and source +range. Python helper and builder exceptions propagate unchanged, preserving their identity, +type and source-mapped traceback. Temporary parser contexts are released on both successful +and exceptional exits. + Parse flow ~~~~~~~~~~ diff --git a/python/tvm/relax/script/__init__.py b/python/tvm/relax/script/__init__.py index d173ace380..58e59d4c45 100644 --- a/python/tvm/relax/script/__init__.py +++ b/python/tvm/relax/script/__init__.py @@ -98,8 +98,9 @@ def py_module( TypeError If the decorated value is not a class, or an unsupported parser option is supplied. - DiagnosticError - If shared parsing, construction or enabled validation fails. + SyntaxError + If the source violates a parser syntax restriction. Construction and + enabled validation propagate their original exceptions unchanged. OSError If source inspection cannot recover a Python function's definition. diff --git a/python/tvm/relax/script/builder/parser_protocol.py b/python/tvm/relax/script/builder/parser_protocol.py index 7a8d0f08ab..357c3de0cb 100644 --- a/python/tvm/relax/script/builder/parser_protocol.py +++ b/python/tvm/relax/script/builder/parser_protocol.py @@ -1425,7 +1425,7 @@ def check_well_formed_(function: _relax.Function) -> None: whole-module checks, s_tir verification, then TIRx verification on each non-s_tir primitive function, preserving cross-function checks. Dialect X hooks validate a standalone completed function under that dialect policy. Invalid IR raises ValueError - retaining native details; the parser converts it to source diagnostics. + retaining native details; the parser propagates the exception unchanged. check_well_formed=False omits the generated call entirely. .. code:: python diff --git a/python/tvm/script/ir_builder/base.py b/python/tvm/script/ir_builder/base.py index 83f7e84e75..ebae9be7bd 100644 --- a/python/tvm/script/ir_builder/base.py +++ b/python/tvm/script/ir_builder/base.py @@ -68,21 +68,13 @@ class IRBuilderFrame(_Object): """ def __enter__(self) -> "IRBuilderFrame": - try: - _ffi_api.IRBuilderFrameEnter(self) # type: ignore[attr-defined] # pylint: disable=no-member - except Exception as error: - _attach_diagnostic_span(error, self.source_span) - raise + _ffi_api.IRBuilderFrameEnter(self) # type: ignore[attr-defined] # pylint: disable=no-member return self def __exit__(self, exc_type, exc_value, trace) -> None: # pylint: disable=unused-argument if exc_type is None and exc_value is None: # Do not execute `FrameExit` if the with scope exits because of exceptions - try: - _ffi_api.IRBuilderFrameExit(self) # type: ignore[attr-defined] # pylint: disable=no-member - except Exception as error: - _attach_diagnostic_span(error, self.source_span) - raise + _ffi_api.IRBuilderFrameExit(self) # type: ignore[attr-defined] # pylint: disable=no-member def add_callback(self, callback: Callable[[], None]) -> None: """Add a callback method invoked when exiting the with-scope. @@ -261,19 +253,6 @@ def source_span( return ir.Span(source_name, line, end_line, column, end_column) -def _attach_diagnostic_span(error: Exception, span: ir.Span | None) -> None: - """Retain an operation's explicit location without opening a source scope.""" - if span is not None and not hasattr(error, "__tvm_script_location__"): - diagnostic_span = span.spans[-1] if isinstance(span, ir.SequentialSpan) else span - error.__tvm_script_location__ = ( - str(diagnostic_span.source_name.name), - diagnostic_span.line, - diagnostic_span.end_line, - diagnostic_span.column, - diagnostic_span.end_column, - ) - - _T = TypeVar("_T") @@ -321,12 +300,8 @@ def with_at_group_( if span is not None and IRBuilder.is_in_scope() else nullcontext() ) - try: - with context: - return at(location, thunk()) - except Exception as error: - _attach_diagnostic_span(error, span) - raise + with context: + return at(location, thunk()) at_ = at diff --git a/python/tvm/script/ir_builder/ir/parser_protocol.py b/python/tvm/script/ir_builder/ir/parser_protocol.py index 91982704e9..a680c9ca4d 100644 --- a/python/tvm/script/ir_builder/ir/parser_protocol.py +++ b/python/tvm/script/ir_builder/ir/parser_protocol.py @@ -510,7 +510,7 @@ def check_well_formed_(module: _ir.IRModule) -> None: whole-module checks, s_tir verification, then TIRx verification on each non-s_tir primitive function, preserving cross-function checks. Dialect X hooks validate a standalone completed function under that dialect policy. Invalid IR raises ValueError - retaining native details; the parser converts it to source diagnostics. + retaining native details; the parser propagates the exception unchanged. check_well_formed=False omits the generated call entirely. .. code:: python diff --git a/python/tvm/script/parser/call_args_policy.py b/python/tvm/script/parser/call_args_policy.py index fb1cdb5d4d..196b2d1aa6 100644 --- a/python/tvm/script/parser/call_args_policy.py +++ b/python/tvm/script/parser/call_args_policy.py @@ -43,7 +43,17 @@ class _LiteralParser: try: expression = ast.parse(node.value, mode="eval").body except SyntaxError as error: - raise SyntaxError(f"Invalid annotation expression: {error.msg}") from error + raise SyntaxError( + f"Invalid annotation expression: {error.msg}", + ( + self.filename, + node.lineno, + node.col_offset + 1, + None, + node.end_lineno, + node.end_col_offset + 1, + ), + ) from error # Example: in R.Tensor(("n + 1",), "float32"), the generated resolve # call for n retains the byte range of n inside the quoted literal, not # the whole constructor. A triple-quoted physical newline advances lineno; diff --git a/python/tvm/script/parser/diagnostics.py b/python/tvm/script/parser/diagnostics.py deleted file mode 100644 index 5605011ee4..0000000000 --- a/python/tvm/script/parser/diagnostics.py +++ /dev/null @@ -1,159 +0,0 @@ -# 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. -"""Source-located construction errors without replacing their original traceback.""" - -from __future__ import annotations - -import ast -import linecache -import traceback - -from tvm.error import DiagnosticError - - -def diagnostic_error(error: Exception, filename: str, tree: ast.Module) -> DiagnosticError: - """Return a DiagnosticError highlighting the original source operation. - - error is the caught Python/builder exception. filename and tree supply the source - filename and original AST for fallback coordinates. Builder errors carry - plain location tuples; otherwise Python traceback/SyntaxError coordinates - identify the innermost original operation. Python 3.11+ supplies precise - traceback columns, while earlier versions fall back to the source line. - No frame or IR object is inspected; the caller attaches error as __cause__. - """ - location = getattr(error, "__tvm_script_location__", None) - if location is not None: - filename, start, end, column, end_column = location - column, end_column = column - 1, end_column - 1 - elif isinstance(error, SyntaxError) and error.lineno: - start = error.lineno - end = error.end_lineno or start - column = max((error.offset or 1) - 1, 0) - end_column = max((error.end_offset or column + 2) - 1, column + 1) - else: - frames = [ - frame - for frame in traceback.extract_tb(error.__traceback__) - if frame.filename == filename - ] - if frames: - frame = frames[-1] - start, end = frame.lineno, getattr(frame, "end_lineno", None) or frame.lineno - column = getattr(frame, "colno", None) or 0 - end_column = getattr(frame, "end_colno", None) - if getattr(frame, "end_lineno", None) is None: - # Python before 3.11 supplies only a line. Recover the original - # statement range from source AST, without an IR span context or - # executing/reparsing a generated program. Prefer the narrowest - # statement beginning on that line (a nested body over its def). - candidates = [ - node - for node in ast.walk(tree) - if isinstance(node, ast.stmt) and node.lineno == start - ] - if candidates: - node = min(candidates, key=lambda item: item.end_lineno - item.lineno) - end, column, end_column = node.end_lineno, node.col_offset, node.end_col_offset - else: - node = tree.body[-1] - start, end = node.lineno, node.lineno - column, end_column = node.col_offset, None - lines = [f"{filename}:{start}: {type(error).__name__}: {error}"] - for number in range(start, end + 1): - source = linecache.getline(filename, number).rstrip("\n") - first = column if number == start else len(source) - len(source.lstrip()) - last = end_column if number == end and end_column is not None else len(source) - lines.extend( - ( - f" {number} | {source}", - " " * (len(str(number)) + 4 + first) + "^" * max(last - first, 1), - ) - ) - return DiagnosticError("\n".join(lines)) - - -def _format_source_snippet( - source_lines: list[str], - lineno: int, - col_offset: int, - end_lineno: int, - end_col_offset: int, -) -> str: - """Format a source code snippet with a column/span marker. - - Renders every source line spanned by the diagnostic (``lineno`` through - ``end_lineno``, inclusive) with a per-line gutter, followed by a caret - underline covering the offending span. For a single-line span - (``end_lineno == lineno``) only the columns ``col_offset`` (inclusive) - through ``end_col_offset`` (exclusive) are underlined; for a multi-line - span the underline covers the start column to end-of-line on the first - line, the full text of interior lines, and the start of the final line up - to ``end_col_offset``. - - Parameters - ---------- - source_lines : list of str - Lines of the source code. - - lineno : int - 1-based starting line number in the source. - - col_offset : int - 1-based starting column (inclusive) on ``lineno``. - - end_lineno : int - 1-based ending line number in the source (>= ``lineno``). - - end_col_offset : int - 1-based ending column (exclusive) on ``end_lineno``. - - Returns - ------- - snippet : str - Formatted source snippet with caret-marker line(s). - """ - if end_lineno < lineno: - end_lineno = lineno - - # Determine the gutter width so that all line numbers line up. - header_width = len(f" {end_lineno} ") - no_line_header = " " * header_width - - parts = [f"{no_line_header}| "] - for cur_lineno in range(lineno, end_lineno + 1): - idx = cur_lineno - 1 - if not 0 <= idx < len(source_lines): - continue - line_text = source_lines[idx].rstrip("\n") - line_header = f" {cur_lineno} ".rjust(header_width) - - # Compute the underline span [start_col, stop_col) for this line. - start_col = col_offset if cur_lineno == lineno else 1 - stop_col = end_col_offset if cur_lineno == end_lineno else len(line_text) + 1 - - marker = "" - for i in range(1, len(line_text) + 1): - if start_col <= i < stop_col: - marker += "^" - else: - marker += " " - parts.append(f"{line_header}| {line_text}") - parts.append(f"{no_line_header}| {marker}") - - if len(parts) == 1: - return "" - return "\n".join(parts) diff --git a/python/tvm/script/parser/entry.py b/python/tvm/script/parser/entry.py index 727b434725..7fef2e1973 100644 --- a/python/tvm/script/parser/entry.py +++ b/python/tvm/script/parser/entry.py @@ -30,14 +30,12 @@ from functools import wraps from types import FrameType, FunctionType from typing import TYPE_CHECKING, Any, TypeVar -from tvm.error import DiagnosticError from tvm.ir import SourceName from tvm.script.ir_builder import base from tvm.script.ir_builder import ir as builder_ir from tvm.script.ir_builder.ir import parser_protocol as syntax_protocol from . import jit_support -from .diagnostics import diagnostic_error from .inspect_source import ( acquire_source, capture_annotation_bindings, @@ -260,8 +258,8 @@ def make_decorator( ------ ValueError When the returned decorator receives a non-function positional value. - DiagnosticError - When standalone construction fails during `parse` execution. + SyntaxError + When standalone source violates a parser restriction. Notes ----- @@ -269,7 +267,8 @@ def make_decorator( Standalone functions immediately transpile and execute a builder program. Annotations must be safe to re-evaluate: eager MissingType placeholders are not cached, and source annotations execute in declaration frames. - Registration persists for the lifetime of the returned decorator. + Construction errors propagate unchanged through `parse`. Registration + persists for the lifetime of the returned decorator. """ mapping, default_options = dict(option_map or {}), dict(defaults or {}) @@ -707,22 +706,21 @@ def parse( Raises ------ - DiagnosticError - If transpilation or host/builder execution fails. An existing - DiagnosticError is preserved; other execution errors gain original - source ranges. OSError If source inspection cannot recover the supplied object's text. TypeError If the source object cannot be inspected. SyntaxError - If initial source parsing fails before program execution. + If source parsing fails or transpilation detects a syntax restriction. + Parser restrictions retain their source filename and range. Notes ----- Each call owns one AST copy and a fresh lexical environment. Declaration and definition frames are entered only during generated execution. Source - acquisition errors propagate directly, before diagnostic conversion. + acquisition, host and builder errors propagate with their original type, + identity and traceback. Temporary captures are released even when execution + fails. """ # - Recover source and explicit lexical/definition inputs. # - Copy the AST once, then collect source syntax facts. @@ -780,10 +778,6 @@ def parse( with jit_support.use_specialization(root_name, specialization): result = _builder() return result - except DiagnosticError: - raise - except Exception as error: - raise diagnostic_error(error, filename, tree) from error finally: # Release temporary captures on both successful and exceptional exits. if _builder is not None and definition_scope_name is not None: @@ -811,14 +805,15 @@ def ir_module(module: type | None = None, **options: Any) -> IRModule | Callable Raises ------ - DiagnosticError - If module transpilation or builder execution fails. + SyntaxError + If module source violates a parser restriction. Notes ----- Class host bindings are captured before transpilation. Generated execution declares all registered signatures before defining their bodies. Source - acquisition errors and frame lifetime follow `parse`. + acquisition and builder errors propagate unchanged; frame lifetime follows + `parse`. """ def apply(module: type) -> IRModule: diff --git a/python/tvm/script/parser/inspect_source.py b/python/tvm/script/parser/inspect_source.py index a235448764..443d982330 100644 --- a/python/tvm/script/parser/inspect_source.py +++ b/python/tvm/script/parser/inspect_source.py @@ -14,7 +14,13 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. -"""Recover source coordinates and capture the explicit definition context.""" +"""Inspect source text, AST coordinates and the explicit definition context. + +``Source`` exposes dedented text and relative AST coordinates with conversion to +absolute IR spans. ``acquire_source`` returns an AST whose original coordinates +are already restored for compilation. Both keep inspection in this module; +source ownership and execution remain with the parser entry point. +""" from __future__ import annotations import __future__ @@ -28,10 +34,112 @@ from collections.abc import Mapping from types import CodeType, FrameType, FunctionType from typing import Any +from tvm.ir import SourceName, Span + from .call_args_policy import parse_annotation from .prescan import collect_annotation_free_names +class Source: + """Source code class for TVMScript. + + It is constructed by source code str or doc AST tree. + + Parameters + ---------- + source_name : str + The filename of the file where the source code locates. + + start_line : int + The first line number of the source code. + + start_column : int + The first column number of the first line of the source code. + + source : str + The source code str of source code. + + full_source : str + The complete source code of the file where the source code locates. + """ + + source_name: str | None + start_line: int + start_column: int + source: str + full_source: str + + def __init__(self, program: str | FunctionType | type) -> None: + if isinstance(program, str): + self.source_name = "<str>" + self.start_line = 1 + self.start_column = 0 + self.source = program + self.full_source = program + return + + self.source_name = inspect.getsourcefile(program) # type: ignore + lines, self.start_line = inspect.getsourcelines(program) # type: ignore + if lines: + self.start_column = len(lines[0]) - len(lines[0].lstrip()) + else: + self.start_column = 0 + if self.start_column and lines: + self.source = "\n".join([line[self.start_column :].rstrip() for line in lines]) + else: + self.source = "".join(lines) + try: + # It will cause a problem when running in Jupyter Notebook. + # `mod` will be <module '__main__'>, which is a built-in module + # and `getsource` will throw a TypeError + mod = inspect.getmodule(program) + if mod: + self.full_source = inspect.getsource(mod) + else: + self.full_source = self.source + except TypeError: + # It's a work around for Jupyter problem. + # Since `findsource` is an internal API of inspect, we just use it + # as a fallback method. + src, _ = inspect.findsource(program) # type: ignore + self.full_source = "".join(src) + + def as_ast(self) -> ast.Module: + """Parse the source code into AST. + + Returns + ------- + res : ast.AST + The AST of source code. + """ + return ast.parse(self.source) + + def location(self, node: ast.AST) -> tuple[int, int, int, int]: + """Return the absolute 1-based source range of an AST node.""" + lineno = getattr(node, "lineno", 1) or 1 + col_offset = getattr(node, "col_offset", self.start_column) + col_offset = self.start_column if col_offset is None else col_offset + end_lineno = getattr(node, "end_lineno", lineno) or lineno + end_col_offset = getattr(node, "end_col_offset", col_offset) + end_col_offset = col_offset if end_col_offset is None else end_col_offset + lineno += self.start_line - 1 + end_lineno += self.start_line - 1 + col_offset += self.start_column + 1 + end_col_offset += self.start_column + 1 + return lineno, col_offset, end_lineno, end_col_offset + + def to_span(self, node: ast.AST) -> Span: + """Convert an AST node to the canonical IR source span.""" + lineno, col_offset, end_lineno, end_col_offset = self.location(node) + return Span( + SourceName(self.source_name or "<unknown>"), + lineno, + end_lineno, + col_offset, + end_col_offset, + ) + + def capture_lexical_bindings(function: FunctionType) -> dict[str, Any]: """Snapshot only actual body global/closure reads for a deferred callable.""" bindings: dict[str, Any] = {} @@ -130,7 +238,7 @@ def acquire_source( Text uses ``<str>`` unless a filename is supplied. Function/class source retains its original file, line and UTF-8 column offsets, including the decoration-site fallback used by gallery runners. Source inspection and - parsing errors propagate to the caller before diagnostic conversion. + parsing errors propagate unchanged to the caller. The returned AST is the source snapshot; entry copies it once for rewriting. """ members = vars(source).values() if inspect.isclass(source) else (source,) diff --git a/python/tvm/script/parser/source.py b/python/tvm/script/parser/source.py deleted file mode 100644 index b035a09851..0000000000 --- a/python/tvm/script/parser/source.py +++ /dev/null @@ -1,125 +0,0 @@ -# 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. -"""Source text and canonical one-based IR coordinates.""" - -from __future__ import annotations - -import ast -import inspect -from types import FunctionType - -from tvm.ir import SourceName, Span - - -class Source: - """Source code class for TVMScript. - - It is constructed by source code str or doc AST tree. - - Parameters - ---------- - source_name : str - The filename of the file where the source code locates. - - start_line : int - The first line number of the source code. - - start_column : int - The first column number of the first line of the source code. - - source : str - The source code str of source code. - - full_source : str - The complete source code of the file where the source code locates. - """ - - source_name: str | None - start_line: int - start_column: int - source: str - full_source: str - - def __init__(self, program: str | FunctionType | type) -> None: - if isinstance(program, str): - self.source_name = "<str>" - self.start_line = 1 - self.start_column = 0 - self.source = program - self.full_source = program - return - - self.source_name = inspect.getsourcefile(program) # type: ignore - lines, self.start_line = inspect.getsourcelines(program) # type: ignore - if lines: - self.start_column = len(lines[0]) - len(lines[0].lstrip()) - else: - self.start_column = 0 - if self.start_column and lines: - self.source = "\n".join([line[self.start_column :].rstrip() for line in lines]) - else: - self.source = "".join(lines) - try: - # It will cause a problem when running in Jupyter Notebook. - # `mod` will be <module '__main__'>, which is a built-in module - # and `getsource` will throw a TypeError - mod = inspect.getmodule(program) - if mod: - self.full_source = inspect.getsource(mod) - else: - self.full_source = self.source - except TypeError: - # It's a work around for Jupyter problem. - # Since `findsource` is an internal API of inspect, we just use it - # as a fallback method. - src, _ = inspect.findsource(program) # type: ignore - self.full_source = "".join(src) - - def as_ast(self) -> ast.Module: - """Parse the source code into AST. - - Returns - ------- - res : ast.AST - The AST of source code. - """ - return ast.parse(self.source) - - def location(self, node: ast.AST) -> tuple[int, int, int, int]: - """Return the absolute 1-based source range of an AST node.""" - lineno = getattr(node, "lineno", 1) or 1 - col_offset = getattr(node, "col_offset", self.start_column) - col_offset = self.start_column if col_offset is None else col_offset - end_lineno = getattr(node, "end_lineno", lineno) or lineno - end_col_offset = getattr(node, "end_col_offset", col_offset) - end_col_offset = col_offset if end_col_offset is None else end_col_offset - lineno += self.start_line - 1 - end_lineno += self.start_line - 1 - col_offset += self.start_column + 1 - end_col_offset += self.start_column + 1 - return lineno, col_offset, end_lineno, end_col_offset - - def to_span(self, node: ast.AST) -> Span: - """Convert an AST node to the canonical IR source span.""" - lineno, col_offset, end_lineno, end_col_offset = self.location(node) - return Span( - SourceName(self.source_name or "<unknown>"), - lineno, - end_lineno, - col_offset, - end_col_offset, - ) diff --git a/python/tvm/script/parser/transpile.py b/python/tvm/script/parser/transpile.py index c7ff50688d..53975dc20d 100644 --- a/python/tvm/script/parser/transpile.py +++ b/python/tvm/script/parser/transpile.py @@ -161,7 +161,17 @@ class IRBuilderTranspiler(ast.NodeTransformer): return ast.Name(name, ast.Load()) def _raise_error(self, node: ast.AST, message: str) -> NoReturn: - raise SyntaxError(message, (self.module.filename, node.lineno, node.col_offset + 1, None)) + raise SyntaxError( + message, + ( + self.module.filename, + node.lineno, + node.col_offset + 1, + None, + node.end_lineno, + node.end_col_offset + 1, + ), + ) def _call( self, namespace: str, member: str, args: list[ast.expr], node: ast.AST, **keywords: ast.expr @@ -1017,6 +1027,10 @@ class IRBuilderTranspiler(ast.NodeTransformer): # y = X.bind_(value, name="y") # ------------------------------------------------- value = self._call_dialect("bind_", [value], statement, **keywords) + if frame_value: + # Binding an entered frame originates at the source as-target; + # retain the existing native span argument while locating this call. + ast.copy_location(value, target) return [ast.copy_location(ast.Assign([target], value), statement)] if isinstance(target, ast.Attribute): # -------------------- Pattern -------------------- @@ -1417,6 +1431,9 @@ class IRBuilderTranspiler(ast.NodeTransformer): else: self._raise_error(node.target, "Loop targets must be names or a flat tuple of names") context = self._call_dialect("for_", [iterable], node, names=names) + # The generated iteration check originates at the source iterable, not + # the final body line. Its native frame span still covers the whole loop. + ast.copy_location(context, node.iter) body = self.transform_statements(node.body) return ast.copy_location( ast.With([ast.withitem(context, node.target)], body or [ast.Pass()]), node diff --git a/python/tvm/tirx/script/builder/parser_protocol.py b/python/tvm/tirx/script/builder/parser_protocol.py index 418a86c39c..6d74d26eb3 100644 --- a/python/tvm/tirx/script/builder/parser_protocol.py +++ b/python/tvm/tirx/script/builder/parser_protocol.py @@ -1686,7 +1686,7 @@ def check_well_formed_(function: _tir.PrimFunc) -> None: whole-module checks, s_tir verification, then TIRx verification on each non-s_tir primitive function, preserving cross-function checks. Dialect X hooks validate a standalone completed function under that dialect policy. Invalid IR raises ValueError - retaining native details; the parser converts it to source diagnostics. + retaining native details; the parser propagates the exception unchanged. check_well_formed=False omits the generated call entirely. .. code:: python diff --git a/tests/python/relax/script/test_frontend_relax_construction_validation.py b/tests/python/relax/script/test_frontend_relax_construction_validation.py index d547113619..f45fe9d0e7 100644 --- a/tests/python/relax/script/test_frontend_relax_construction_validation.py +++ b/tests/python/relax/script/test_frontend_relax_construction_validation.py @@ -21,7 +21,6 @@ import pytest import tvm from tvm import relax -from tvm.error import DiagnosticError from tvm.relax import BasePyModule from tvm.script import ir as I from tvm.script import relax as R @@ -50,7 +49,7 @@ def _invalid_inplace_module(**options): def test_default_relax_validation_rejects_repeated_inplace_input(): - with pytest.raises(DiagnosticError): + with pytest.raises(ValueError): _invalid_inplace_module() diff --git a/tests/python/relax/script/test_relax_parser_comparison_integration.py b/tests/python/relax/script/test_relax_parser_comparison_integration.py index 34c6532bd1..de0c351a5b 100644 --- a/tests/python/relax/script/test_relax_parser_comparison_integration.py +++ b/tests/python/relax/script/test_relax_parser_comparison_integration.py @@ -26,8 +26,6 @@ OPERATORS = [("<", "LT"), ("<=", "LE"), (">", "GT"), (">=", "GE"), ("==", "EQ"), @pytest.mark.parametrize("operator,kind", OPERATORS) def test_unmarked_host_comparison_never_calls_python_overload(operator, kind): - from tvm.error import DiagnosticError - calls = [] class Host: @@ -51,7 +49,7 @@ def test_unmarked_host_comparison_never_calls_python_overload(operator, kind): ) decorator = "R.function" statement = "return left " + operator + " right" - with pytest.raises(DiagnosticError, match="(PrimExpr|primitive|convert|type)"): + with pytest.raises(TypeError, match="(PrimExpr|primitive|convert|type)"): parser.parse( f"@{decorator}\ndef main():\n {statement}\n", extra_vars={"left": Host(), "right": Host()}, diff --git a/tests/python/relax/test_op_view.py b/tests/python/relax/test_op_view.py index dc7e658ab9..9679f92b25 100644 --- a/tests/python/relax/test_op_view.py +++ b/tests/python/relax/test_op_view.py @@ -54,7 +54,7 @@ def test_infer_shape_of_2d_static_view(): def test_error_if_shape_argument_is_not_shape(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(TypeError): @R.function def func(A: R.Tensor([16])): @@ -105,7 +105,7 @@ def test_infer_shape_of_2d_static_view_same_size_as_2d_source(): def test_error_if_1d_static_view_larger_than_1d_source(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def func(A: R.Tensor([16])): @@ -114,7 +114,7 @@ def test_error_if_1d_static_view_larger_than_1d_source(): def test_error_if_static_2d_view_larger_than_source(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def func(A: R.Tensor([16])): @@ -171,7 +171,7 @@ def test_infer_shape_of_2d_dynamic_view(): def test_error_if_1d_dynamic_view_larger_than_1d_source(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def func(A: R.Tensor(["N"])): @@ -180,9 +180,8 @@ def test_error_if_1d_dynamic_view_larger_than_1d_source(): return B [email protected](reason="See https://github.com/apache/tvm/pull/16877";) def test_error_if_1d_dynamic_view_provably_larger_than_1d_source(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def func(A: R.Tensor(["N"])): @@ -192,7 +191,7 @@ def test_error_if_1d_dynamic_view_provably_larger_than_1d_source(): def test_error_if_2d_dynamic_view_provably_larger_than_1d_source(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def func(A: R.Tensor(["N"])): @@ -245,7 +244,7 @@ def test_infer_dtype_of_float32_view(): def test_error_if_view_dtype_is_void(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(TypeError): @R.function def func(A: R.Tensor("float32")): @@ -338,7 +337,7 @@ def test_view_dtype_may_be_smaller_than_input_dtype(): def test_error_if_view_dtype_is_larger_than_input_dtype(): """A view may not exceed the bounds of the viewed array""" - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def func(A: R.Tensor([16], "uint8")): @@ -393,7 +392,7 @@ def test_error_if_number_of_bytes_of_view_is_larger_than_original(): byte/element). """ - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def func(A: R.Tensor([8], "float16")): @@ -410,7 +409,7 @@ def test_error_for_non_zero_relative_byte_offset(): """ - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def func(A: R.Tensor): diff --git a/tests/python/relax/test_transform.py b/tests/python/relax/test_transform.py index dd811855d8..a505cc5e26 100644 --- a/tests/python/relax/test_transform.py +++ b/tests/python/relax/test_transform.py @@ -589,7 +589,7 @@ def test_call_tir_inplace_some_new(): def test_call_tir_inplace_repeated_input(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @tvm.script.ir_module class Input: @@ -619,7 +619,7 @@ def test_call_tir_inplace_repeated_input(): def test_call_tir_inplace_all_new(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @tvm.script.ir_module class Input: @@ -648,7 +648,7 @@ def test_inplace_mutation_with_tuple_argument_raises_error(): triggered a segfault rather than raising an exception. """ - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @I.ir_module(s_tir=True) class Module: @@ -680,7 +680,7 @@ def test_inplace_mutation_with_non_tensor_argument_raises_error(): triggered a segfault rather than raising an exception. """ - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @I.ir_module(s_tir=True) class Module: @@ -709,7 +709,7 @@ def test_inplace_mutation_with_incompatible_tensor_shape_raises_error(): different than the output's shape (`[32]` as opposed to `[16]`). """ - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @I.ir_module(s_tir=True) class Module: @@ -738,7 +738,7 @@ def test_inplace_mutation_with_incompatible_tensor_dtype_raises_error(): different than the output's dtype (`int32` as opposed to `float32`). """ - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @I.ir_module(s_tir=True) class Module: diff --git a/tests/python/relax/test_tvmscript_parser.py b/tests/python/relax/test_tvmscript_parser.py index 875d2619e6..da4246f5ee 100644 --- a/tests/python/relax/test_tvmscript_parser.py +++ b/tests/python/relax/test_tvmscript_parser.py @@ -15,6 +15,7 @@ # specific language governing permissions and limitations # under the License. # ruff: noqa: F401, F821, F841 +import inspect import sys from typing import Optional, Union @@ -64,12 +65,12 @@ def test_simple_func(): _check(foo, bb.get()["foo"]) -def test_error_report(): - with pytest.raises(tvm.error.DiagnosticError): +def test_call_tir_requires_global_var(): + with pytest.raises(tvm.error.InternalError, match="first argument to be a GlobalVar"): @R.function def foo(x: R.Tensor((128, 128), "float32")) -> R.Tensor(None, "float32", ndim=2): - # error: a = b = c is not allowed. + # call_tir requires a GlobalVar rather than a packed-function name. gv0 = gv1 = R.call_tir("extern_func", x, R.Tensor((128, 128), dtype="float32")) return gv0 @@ -110,7 +111,7 @@ def test_unexpected_ndim_type(): def test_unexpected_tir_cast_args(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(TypeError): @R.function def f(x: R.Tensor(("m",), "float32")): @@ -120,7 +121,7 @@ def test_unexpected_tir_cast_args(): def test_unexpected_tir_args(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(TypeError): @tvm.script.ir_module class TestWellCallTIR: @@ -139,7 +140,7 @@ def test_unexpected_tir_args(): gv = R.call_tir(tir_addone, (x,), R.Tensor((T.max(16),), dtype="float32")) return gv - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(TypeError): @R.function def f(x: R.Tensor(("m", "n"), "float32")): @@ -149,7 +150,7 @@ def test_unexpected_tir_args(): def test_func_type_annotation_fail(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(SyntaxError, match="Parameter 'x' requires an annotation") as error: @R.function def f(x, y): # error: the parameter type annotation is missing @@ -157,9 +158,16 @@ def test_func_type_annotation_fail(): y = z return y + lines, start = inspect.getsourcelines(test_func_type_annotation_fail) + line_index = next(i for i, line in enumerate(lines) if "def f(x, y):" in line) + assert error.value.filename == __file__ + assert error.value.lineno == error.value.end_lineno == start + line_index + assert error.value.offset == lines[line_index].index("x, y") + 1 + assert error.value.end_offset == error.value.offset + 1 + def test_if_mismatch_var_fail(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(SyntaxError, match="same named output") as error: @R.function def f(cond: R.Tensor((), "bool"), x: R.Tensor((1,), "float32")): @@ -171,9 +179,18 @@ def test_if_mismatch_var_fail(): z = R.add(w, w) # error: The binding var is expected to `y` return z + lines, start = inspect.getsourcelines(test_if_mismatch_var_fail) + last = next(i for i, line in enumerate(lines) if "z = R.add(w, w)" in line) + assert error.value.filename == __file__ + assert error.value.lineno == error.value.end_lineno == start + last + assert error.value.offset == lines[last].index("z = R.add(w, w)") + 1 + assert ( + error.value.end_offset == lines[last].index("z = R.add(w, w)") + len("z = R.add(w, w)") + 1 + ) + def test_unassigned_call_fail(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def f(x: R.Tensor): @@ -182,7 +199,7 @@ def test_unassigned_call_fail(): def test_incorrect_tensor_shape(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(tvm.error.InternalError): @R.function def f(x: R.Tensor([16])): @@ -426,7 +443,7 @@ def test_symbolic_shape(): gv0 = R.call_dps_packed("extern_func", x, R.Tensor((m, n), dtype="float32")) return gv0 - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(tvm.error.InternalError): @R.function def mismatch_dtype(x: R.Tensor(("m", "n"), "float32")) -> R.Tensor(None, "float32", ndim=2): @@ -667,53 +684,57 @@ def test_dataflow_block_advanced(): def test_dataflow_binding_after_output(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError, match="New binding is not allowed after dataflow block output"): @R.function def foo(x: R.Tensor((128, 128), "float32")) -> R.Tensor(None, "float32", ndim=2): with R.dataflow(): - gv = R.call_tir("extern_func", x, R.Tensor((128, 128), dtype="float32")) + gv = R.call_dps_packed("extern_func", x, R.Tensor((128, 128), dtype="float32")) R.output(gv) - lv = R.call_tir("extern_func", gv, R.Tensor((128, 128), dtype="float32")) + lv = R.call_dps_packed("extern_func", gv, R.Tensor((128, 128), dtype="float32")) return gv def test_dataflow_output_global_var(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises( + ValueError, match="An output variable is not emitted by this dataflow block" + ): @R.function def foo(x: R.Tensor((128, 128), "float32")) -> R.Tensor(None, "float32", ndim=2): - gv0 = R.call_tir("extern_func", x, R.Tensor((128, 128), dtype="float32")) + gv0 = R.call_dps_packed("extern_func", x, R.Tensor((128, 128), dtype="float32")) with R.dataflow(): - gv1 = R.call_tir("extern_func", gv0, R.Tensor((128, 128), dtype="float32")) + gv1 = R.call_dps_packed("extern_func", gv0, R.Tensor((128, 128), dtype="float32")) R.output(gv0, gv1) return gv1 def test_dataflow_multiple_output(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises( + ValueError, match="not allowed for a dataflow block to have multiple output" + ): @R.function def foo(x: R.Tensor((128, 128), "float32")) -> R.Tensor(None, "float32", ndim=2): with R.dataflow(): - gv = R.call_tir("extern_func", x, R.Tensor((128, 128), dtype="float32")) + gv = R.call_dps_packed("extern_func", x, R.Tensor((128, 128), dtype="float32")) R.output(gv) R.output(gv) return gv def test_dataflow_output_outside_dataflow_block(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError, match="`R.output` should appear inside a dataflow block"): @R.function def foo(x: R.Tensor((128, 128), "float32")) -> R.Tensor(None, "float32", ndim=2): - gv = R.call_tir("extern_func", x, R.Tensor((128, 128), dtype="float32")) + gv = R.call_dps_packed("extern_func", x, R.Tensor((128, 128), dtype="float32")) R.output(gv) return gv def test_dataflow_scope_fail(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def f(x: R.Tensor(ndim=2)): @@ -740,7 +761,7 @@ def test_return_without_binding(): def test_multiple_return(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def foo(x: R.Tensor((128, 128), "float32")): @@ -749,11 +770,11 @@ def test_multiple_return(): def test_function_without_return(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError, match="A Relax function must have a return value"): @R.function def foo(x: R.Tensor((128, 128), "float32")): - gv0 = R.call_tir("extern_func", x, R.Tensor((128, 128), dtype="float32")) + gv0 = R.call_dps_packed("extern_func", x, R.Tensor((128, 128), dtype="float32")) def test_tensor_type_without_args(): @@ -931,7 +952,7 @@ def test_annotate_override(): assert isinstance(y_bind.var.ty, relax.TensorType) assert isinstance(z_bind.var.ty, relax.AnyType) - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(tvm.error.InternalError): @R.function def test(x: R.Tensor): @@ -1062,7 +1083,7 @@ def test_call_tir_inplace(): def test_call_tir_inplace_with_tuple_var_raises_error(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(TypeError): @tvm.script.ir_module class Module: @@ -1128,7 +1149,7 @@ def test_local_function(): def test_inline_prim_func(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(TypeError, match="unexpected keyword argument.*local"): @I.ir_module(s_tir=True) class TestModule: @@ -1270,7 +1291,7 @@ def test_if_branch_with_match_cast(): def test_if_inside_dataflow(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def foo(cond: R.Tensor((), "bool"), x: R.Tensor((1,), "float32")): @@ -1286,7 +1307,7 @@ def test_if_inside_dataflow(): def test_var_if_scoping_fail(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(NameError): @R.function def f(cond: R.Tensor((), "bool"), x: R.Tensor((1,), "float32")): @@ -1620,7 +1641,7 @@ class Module: def test_later_prim_param_not_adopted_by_usage_first_symbol(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): tvm.script.from_source( """ @R.function @@ -1631,7 +1652,7 @@ def main(x: R.Tensor(["n"], "float32"), n: R.Prim("int64")): def test_non_int64_prim_param_rejected_in_shape_annotation(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(tvm.error.InternalError): tvm.script.from_source( """ @R.function @@ -1698,7 +1719,7 @@ def test_undefined_symbolic_var_raises_error(): shape. That is, this test case raises an error, and will not attempt to define `m` as either `x.shape[0]-1` or `x.shape[1]//2`. """ - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function def foo(x: R.Tensor(("m + 1", "m * 2"), "float32")): # name 'm' is not defined @@ -2079,7 +2100,7 @@ def test_function_with_void_return_type_may_be_used_as_statements(): def test_function_with_non_void_return_type_must_be_assigned(): """Non-void results must be assigned to a variable""" - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function(pure=False) def func(x: R.Tensor((), "int32")) -> R.Tensor((), "int32"): @@ -2166,7 +2187,7 @@ def test_private_function(): def test_private_function_with_global_symbol_fail(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @I.ir_module(s_tir=True) class Addition: @@ -2183,7 +2204,7 @@ def test_private_function_with_global_symbol_fail(): def test_private_function_with_global_symbol_no_module_fail(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function(private=True) def func(x: R.Tensor((), "int32")) -> R.Tensor((), "int32"): @@ -2260,7 +2281,7 @@ def test_macro_non_hygienic(): def test_macro_no_variable_leak(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(NameError): @R.macro(hygienic=True) def add_two(value): @@ -2483,7 +2504,7 @@ def test_shared_meta_var_skips_relax_bindings(): assert "meta_var" not in source _check(func) - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @R.function(private=True) def mismatched_declaration(): diff --git a/tests/python/s_tir/test_s_tir_axis_binding_validation.py b/tests/python/s_tir/test_s_tir_axis_binding_validation.py index f82c77d040..4fda4b0e40 100644 --- a/tests/python/s_tir/test_s_tir_axis_binding_validation.py +++ b/tests/python/s_tir/test_s_tir_axis_binding_validation.py @@ -18,7 +18,6 @@ import pytest -import tvm from tvm.script import parser @@ -39,5 +38,5 @@ def main(): {axes} T.evaluate(vi) """ - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): parser.parse(source) diff --git a/tests/python/script/test_capture_definition_scope.py b/tests/python/script/test_capture_definition_scope.py index 2fb1fb5949..21613eeccb 100644 --- a/tests/python/script/test_capture_definition_scope.py +++ b/tests/python/script/test_capture_definition_scope.py @@ -23,7 +23,6 @@ from types import SimpleNamespace import pytest -from tvm.error import DiagnosticError from tvm.script import ir as I from tvm.script.ir_builder import IRBuilder from tvm.tirx.script.jit import make_jit @@ -234,7 +233,7 @@ def test_unrelated_same_file_caller_does_not_supply_annotation_locals(language): def test_inactive_lexical_ancestor_is_not_replaced_by_unrelated_caller(language): # Before: annotation-only extent has left scope while an unrelated caller has extent. - # Expected builder: raise DiagnosticError naming extent; do not capture the caller. + # Expected builder: raise NameError naming extent; do not capture the caller. X = language.X def outer(extent): @@ -249,7 +248,7 @@ def test_inactive_lexical_ancestor_is_not_replaced_by_unrelated_caller(language) middle = outer(8) extent = 99 # noqa: F841 - with pytest.raises(DiagnosticError, match="extent"): + with pytest.raises(NameError, match="extent"): middle() @@ -297,7 +296,7 @@ def test_jit_retains_active_lexical_ancestor_snapshot(factory, language): def test_unrelated_intervening_call_ends_annotation_capture(language): # Before: an unrelated invoke(callback) frame interrupts lexical ancestor capture. - # Expected builder: raise DiagnosticError naming unavailable extent. + # Expected builder: raise NameError naming unavailable extent. X = language.X def invoke(callback): @@ -313,7 +312,7 @@ def test_unrelated_intervening_call_ends_annotation_capture(language): return invoke(middle) - with pytest.raises(DiagnosticError, match="extent"): + with pytest.raises(NameError, match="extent"): outer(8) diff --git a/tests/python/script/test_capture_frontend_scope.py b/tests/python/script/test_capture_frontend_scope.py index 2f0d784332..a530af61f2 100644 --- a/tests/python/script/test_capture_frontend_scope.py +++ b/tests/python/script/test_capture_frontend_scope.py @@ -23,7 +23,7 @@ from dummy_builder import Function from tvm import ir from tvm.script import ir as I from tvm.script.parser import entry -from tvm.script.parser.source import Source +from tvm.script.parser.inspect_source import Source from tvm.tirx.script.jit import make_jit VALUE = 1 diff --git a/tests/python/script/test_parser.py b/tests/python/script/test_parser.py index 42e83d0ac5..9dd0212945 100644 --- a/tests/python/script/test_parser.py +++ b/tests/python/script/test_parser.py @@ -441,13 +441,13 @@ def test_quoted_symbols_share_identity_without_introducing_python_bindings(langu # Before: def main(x: X.tensor(("n", "n"))): X.record(n) # Expected builder program: X.tensor((X.resolve_type_var_("n"), X.resolve_type_var_("n"))) # The unquoted body name n raises NameError until explicitly declared. - with pytest.raises(Exception) as error: + with pytest.raises(NameError) as error: language.parse(""" @X.script def main(x: X.tensor(("n", "n"))): X.record(n) """) - assert isinstance(error.value.__cause__, NameError) + assert isinstance(error.value, NameError) def test_explicit_symbol_declaration_reuses_annotation_identity(language): @@ -469,7 +469,7 @@ def main(x: X.tensor(("n",))): def test_script_namespace_cannot_be_rebound(language, binding): # Before: X = 1 (also loop/with targets) # Expected builder program: reject a binding that shadows the fixed script namespace X. - with pytest.raises(Exception, match="namespace|shadow|rebind"): + with pytest.raises(SyntaxError, match="namespace|shadow|rebind"): language.parse(f"@X.script\ndef main():\n {binding}\n") @@ -616,7 +616,7 @@ def test_conditional_branches_require_matching_output_names(language, other): # Before: if condition: y = left; else: z = right # Expected builder program: reject branches without the same named terminal output. language.X.__tvm_value_if__ = True - with pytest.raises(Exception, match="same named output"): + with pytest.raises(SyntaxError, match="same named output"): language.parse( "@X.script\ndef main():\n if condition:\n y = left\n" f" else:\n {other}\n" @@ -641,7 +641,7 @@ def test_entry_keeps_reusable_input_tree_unchanged(language, monkeypatch): def test_bare_callable_alias_does_not_acquire_constexpr_syntax(language): # Before: marker = I.constexpr; if marker(True): ... # Expected builder program: an ordinary call to marker raises its runtime syntax-marker error. - with pytest.raises(Exception, match="syntax marker"): + with pytest.raises(TypeError, match="syntax marker"): language.parse( """ @X.script @@ -703,30 +703,33 @@ class Module: @pytest.mark.parametrize( - "body, line, column, message", + "body, line, column, end_column, message", [ - (" X = 1\n", 3, 5, "namespace"), - (" for X in range(2):\n pass\n", 3, 9, "namespace"), + (" X = 1\n", 3, 5, 6, "namespace"), + (" for X in range(2):\n pass\n", 3, 9, 10, "namespace"), ( " if condition:\n y = left\n else:\n z = right\n", 6, 9, + 18, "same named output", ), ], ) -def test_syntax_diagnostics_point_to_the_offending_binding(language, body, line, column, message): +def test_syntax_diagnostics_point_to_the_offending_binding( + language, body, line, column, end_column, message +): # Before: X = 1; or if condition: y = left; else: z = right # Expected builder program: reject the offending namespace/output binding # at its original filename, line and column, before running the builder. language.X.__tvm_value_if__ = True - with pytest.raises(Exception, match=message) as error: + with pytest.raises(SyntaxError, match=message) as error: language.parse("@X.script\ndef main():\n" + body) - cause = error.value.__cause__ + cause = error.value assert isinstance(cause, SyntaxError) assert (cause.filename, cause.lineno, cause.offset) == ("dummy.py", line, column) - assert f"dummy.py:{line}: SyntaxError:" in str(error.value) - assert f" {line} | {body.splitlines()[line - 3]}" in str(error.value) + assert cause.end_lineno == line + assert cause.end_offset == end_column assert not language.events diff --git a/tests/python/script/test_parser_capture_lifetime.py b/tests/python/script/test_parser_capture_lifetime.py index 23348bffa7..c62d421ebb 100644 --- a/tests/python/script/test_parser_capture_lifetime.py +++ b/tests/python/script/test_parser_capture_lifetime.py @@ -20,12 +20,12 @@ from __future__ import annotations import ast import gc +import inspect import weakref from contextlib import contextmanager import pytest -from tvm.error import DiagnosticError from tvm.script import ir as I from tvm.script.parser import entry @@ -115,7 +115,7 @@ def test_owned_ast_context_and_builder_are_temporary(language, monkeypatch, fail if failure: try: language.parse(source) - except DiagnosticError: + except NameError: pass else: pytest.fail("expected failed construction") @@ -130,27 +130,40 @@ def test_owned_ast_context_and_builder_are_temporary(language, monkeypatch, fail ) -def test_expression_exception_keeps_original_source_call_traceback(language): +class FixedFailure(ValueError): + def __setattr__(self, name, value): + if name == "__tvm_script_location__": + raise AttributeError("exception metadata is immutable") + super().__setattr__(name, value) + + [email protected]("error_type", [ValueError, FixedFailure]) +def test_expression_exception_keeps_original_source_call_traceback(spanned_language, error_type): # Before: a normal Python helper fails within a generated function body. # Expected builder program: its real call frame retains the source expression range. import traceback - original = ValueError("original expression failure") + # Exercise the production shared span boundary as well as the generated call. + language = spanned_language + original = error_type("original expression failure") def explode(): raise original source = "@X.script\ndef main():\n X.record(explode())\n" filename = "expression_traceback.py" - with pytest.raises(DiagnosticError) as caught: + with pytest.raises(error_type) as caught: entry.parse(source, extra_vars={"X": language.X, "explode": explode}, filename=filename) - cause = caught.value.__cause__ - assert cause is original - assert type(cause) is ValueError - frames = traceback.extract_tb(cause.__traceback__) + assert caught.value is original + assert type(caught.value) is error_type + frames = traceback.extract_tb(caught.value.__traceback__) calls = [frame for frame in frames if frame.filename == filename and frame.lineno == 3] assert calls - if calls[-1].colno is not None: + helper_frames = [frame for frame in frames if frame.name == "explode"] + assert len(helper_frames) == 1 + assert helper_frames[0].filename == inspect.getsourcefile(explode) + assert helper_frames[0].lineno == inspect.getsourcelines(explode)[1] + 1 + if getattr(calls[-1], "colno", None) is not None: assert (calls[-1].colno, calls[-1].end_lineno, calls[-1].end_colno) == (13, 3, 22) diff --git a/tests/python/script/test_parser_host_selection.py b/tests/python/script/test_parser_host_selection.py index d8ab5ce206..d133f7888e 100644 --- a/tests/python/script/test_parser_host_selection.py +++ b/tests/python/script/test_parser_host_selection.py @@ -113,7 +113,7 @@ def main(): @pytest.mark.parametrize("track_span", [True, False]) def test_constexpr_keeps_named_expression_unsupported(language, track_span): - with pytest.raises(Exception, match="Unsupported expression: NamedExpr"): + with pytest.raises(SyntaxError, match="Unsupported expression: NamedExpr"): parse( language, """ @@ -147,9 +147,7 @@ def main(condition: X.value): "expression", ["1 if I.constexpr(x) else 0", "I.constexpr(x) and 1", "I.constexpr(x) or 1"] ) def test_missing_host_binding_cannot_be_truth_tested(language, expression): - from tvm.error import DiagnosticError - - with pytest.raises(DiagnosticError) as caught: + with pytest.raises(NameError) as caught: parse( language, f""" @@ -160,13 +158,11 @@ def main(): X.record({expression}) """, ) - assert isinstance(caught.value.__cause__, NameError) + assert isinstance(caught.value, NameError) def test_missing_host_if_binding_raises_before_branch_assignments(language): - from tvm.error import DiagnosticError - - with pytest.raises(DiagnosticError) as caught: + with pytest.raises(NameError) as caught: parse( language, """ @@ -180,7 +176,7 @@ def main(): x = 3 """, ) - assert isinstance(caught.value.__cause__, NameError) + assert isinstance(caught.value, NameError) def test_host_lambda_local_does_not_capture_optional_binding(language): @@ -246,9 +242,7 @@ def main(a: X.value(), b: X.value(), c: X.value()): def test_ir_branch_incoming_read_and_rebinding_obeys_python_scope(language): - from tvm.error import DiagnosticError - - with pytest.raises(DiagnosticError) as error: + with pytest.raises(NameError) as error: language.parse(""" @X.script def main(condition: X.value(), x: X.value()): @@ -259,5 +253,5 @@ def main(condition: X.value(), x: X.value()): y = x return y """) - assert isinstance(error.value.__cause__, NameError) - assert "y" in str(error.value.__cause__) + assert isinstance(error.value, NameError) + assert "y" in str(error.value) diff --git a/tests/python/script/test_parser_latest_design.py b/tests/python/script/test_parser_latest_design.py index 59c5c4cf62..f8cea3847c 100644 --- a/tests/python/script/test_parser_latest_design.py +++ b/tests/python/script/test_parser_latest_design.py @@ -18,9 +18,10 @@ from __future__ import annotations +import ast + import pytest -from tvm.error import DiagnosticError from tvm.script import ir as I from tvm.script.ir_builder.ir import parser_protocol as protocol from tvm.script.parser import entry @@ -98,13 +99,13 @@ def test_symbol_reassignment_reports_introduction_and_exact_write( ): # Before: annotation-free/body-declared n, followed by an ordinary write. # Expected diagnostic: n, its introduction line, and the actual offending source target. - with pytest.raises(DiagnosticError) as caught: + with pytest.raises(SyntaxError) as caught: language.parse(source, n=4) message = str(caught.value) assert "Symbolic variable 'n' cannot be reassigned" in message assert f"introduced at line {introduction}" in message - assert f"dummy.py:{offending}:" in message - cause = caught.value.__cause__ + assert caught.value.filename == "dummy.py" + cause = caught.value assert isinstance(cause, SyntaxError) assert (cause.lineno, cause.end_lineno) == (offending, offending) assert cause.end_offset - cause.offset == 1 @@ -264,9 +265,45 @@ def test_scope_declaration_reuses_symbol_spelling_but_plain_write_is_rejected( source = "@X.script\ndef main():\n n = X.symbol()\n n = declared()\n" if ordinary_write: source += " n = 3\n" - with pytest.raises(DiagnosticError, match="Symbolic variable 'n' cannot be reassigned"): + with pytest.raises(SyntaxError, match="Symbolic variable 'n' cannot be reassigned"): language.parse(source, declared=declared) else: result = language.parse(source + " X.record(n)\n", declared=declared) assert result.body == [("emit", value)] assert not any(event[:2] == ("bind", "n") for event in language.events) + + [email protected]( + "source, node_type, message", + [ + ("@X.script\ndef main(value):\n pass\n", ast.arg, "requires an annotation"), + ( + '@X.script\ndef main(value: "invalid +"):\n pass\n', + ast.Constant, + "Invalid annotation expression", + ), + ( + "@X.script\ndef main():\n X.record((value :=\n 1))\n", + ast.NamedExpr, + "Unsupported expression: NamedExpr", + ), + ], +) +def test_transpilation_restrictions_keep_original_source_ranges( + language, source, node_type, message +): + # Before: invalid signature, quoted annotation or multiline expression syntax. + # Expected: direct SyntaxError points at the original node before builder execution. + expected = next(node for node in ast.walk(ast.parse(source)) if isinstance(node, node_type)) + with pytest.raises(SyntaxError, match=message) as caught: + entry.parse(source, extra_vars={"X": language.X}, filename="restriction.py") + error = caught.value + assert type(error) is SyntaxError + assert (error.filename, error.lineno, error.offset, error.end_lineno, error.end_offset) == ( + "restriction.py", + expected.lineno, + expected.col_offset + 1, + expected.end_lineno, + expected.end_col_offset + 1, + ) + assert not language.events diff --git a/tests/python/script/test_parser_scope_helpers.py b/tests/python/script/test_parser_scope_helpers.py index c22b32142e..1e951878d0 100644 --- a/tests/python/script/test_parser_scope_helpers.py +++ b/tests/python/script/test_parser_scope_helpers.py @@ -18,10 +18,10 @@ import ast import copy +import traceback import pytest -from tvm.error import DiagnosticError from tvm.script.ir_builder import IRBuilder from tvm.script.parser import entry @@ -141,13 +141,16 @@ def main(x: X.tensor((4,))): observe("outer_after", x, x, scope) """ if fail: - with pytest.raises(DiagnosticError) as caught: + with pytest.raises(ValueError) as caught: language.parse(source, observe=observe) - assert caught.value.__cause__ is failure + assert caught.value is failure source_line = next( index for index, line in enumerate(source.splitlines(), 1) if 'observe("inner"' in line ) - assert f"dummy.py:{source_line}:" in str(caught.value) + assert any( + frame.filename == "dummy.py" and frame.lineno == source_line + for frame in traceback.extract_tb(caught.value.__traceback__) + ) else: result = language.parse(source, observe=observe) assert result is language.functions["main"] diff --git a/tests/python/script/test_parser_scoped_symbols.py b/tests/python/script/test_parser_scoped_symbols.py index f7b7311e75..dfa219e889 100644 --- a/tests/python/script/test_parser_scoped_symbols.py +++ b/tests/python/script/test_parser_scoped_symbols.py @@ -18,7 +18,6 @@ import pytest -from tvm.error import DiagnosticError from tvm.script.parser import entry @@ -80,10 +79,10 @@ def test_signature_read_before_introduction_remains_unbound(language, annotation def main(x: {annotation}): return x """ - with pytest.raises(DiagnosticError) as error: + with pytest.raises(NameError) as error: parse(language, source) - assert isinstance(error.value.__cause__, NameError) - assert "n" in str(error.value.__cause__) + assert isinstance(error.value, NameError) + assert "n" in str(error.value) def test_signature_strings_do_not_replace_captured_python_names(language): diff --git a/tests/python/tirx/codegen/test_codegen_blackwell.py b/tests/python/tirx/codegen/test_codegen_blackwell.py index 26138ffd68..dadd71f435 100644 --- a/tests/python/tirx/codegen/test_codegen_blackwell.py +++ b/tests/python/tirx/codegen/test_codegen_blackwell.py @@ -245,7 +245,7 @@ def test_tma_mbarrier_remote_view_codegen(): def test_mbarrier_remote_view_rejects_invalid_operations(): from tvm.tirx.lang.pipeline import MBarrier, TMABar - with pytest.raises(tvm.error.DiagnosticError, match=r"remote_view\(\) cannot be initialized"): + with pytest.raises(ValueError, match=r"remote_view\(\) cannot be initialized"): # fmt: off @T.prim_func def invalid_init(): @@ -257,7 +257,7 @@ def test_mbarrier_remote_view_rejects_invalid_operations(): bar.remote_view(0).init(1) # fmt: on - with pytest.raises(tvm.error.DiagnosticError, match=r"remote_view\(\) cannot be waited on"): + with pytest.raises(ValueError, match=r"remote_view\(\) cannot be waited on"): # fmt: off @T.prim_func def invalid_wait(): @@ -269,7 +269,7 @@ def test_mbarrier_remote_view_rejects_invalid_operations(): bar.remote_view(0).wait(0, 0) # fmt: on - with pytest.raises(tvm.error.DiagnosticError, match="cannot also specify remote"): + with pytest.raises(ValueError, match="cannot also specify remote"): # fmt: off @T.prim_func def ambiguous_mbarrier_arrive(): @@ -281,7 +281,7 @@ def test_mbarrier_remote_view_rejects_invalid_operations(): bar.remote_view(0).arrive(0, remote=1) # fmt: on - with pytest.raises(tvm.error.DiagnosticError, match="cannot also specify remote"): + with pytest.raises(ValueError, match="cannot also specify remote"): # fmt: off @T.prim_func def ambiguous_tma_arrive(): @@ -294,7 +294,7 @@ def test_mbarrier_remote_view_rejects_invalid_operations(): # fmt: on with pytest.raises( - tvm.error.DiagnosticError, + ValueError, match=r"remote_view\(\) cannot be applied to a remote view", ): # fmt: off diff --git a/tests/python/tirx/codegen/test_codegen_cuda.py b/tests/python/tirx/codegen/test_codegen_cuda.py index 0500bee7c7..a352c19fe6 100644 --- a/tests/python/tirx/codegen/test_codegen_cuda.py +++ b/tests/python/tirx/codegen/test_codegen_cuda.py @@ -678,7 +678,7 @@ def test_cuda_ldg_vector_dtype_codegen(dtype, suffix, c_type, vec_base, vec): def test_cuda_ldg_vector_rejects_unsupported_dtype(): - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="Unsupported vector CUDA"): + with pytest.raises(ValueError, match="Unsupported vector CUDA"): _get_source(_cuda_ldg_vector_kernel("float16", "v2")) diff --git a/tests/python/tirx/codegen/test_ptx_addr.py b/tests/python/tirx/codegen/test_ptx_addr.py index b26750370d..e706a0304c 100644 --- a/tests/python/tirx/codegen/test_ptx_addr.py +++ b/tests/python/tirx/codegen/test_ptx_addr.py @@ -242,9 +242,7 @@ def test_ptx_addr_offset_type_and_range_rejections(): def test_ptx_addr_pointer_and_raw_address_validation(): - with pytest.raises( - (ValueError, tvm.error.DiagnosticError), match="uint32 address requires shared" - ): + with pytest.raises(ValueError, match="uint32 address requires shared"): @T.prim_func def global_u32(raw: T.uint32): @@ -252,9 +250,7 @@ def test_ptx_addr_pointer_and_raw_address_validation(): value = T.local_scalar("uint32") T.ptx.ld.global_.b32(value, T.ptx.addr(raw, 4)) - with pytest.raises( - (ValueError, tvm.error.DiagnosticError), match="does not support T.ptx.addr" - ): + with pytest.raises(ValueError, match="does not support T.ptx.addr"): @T.prim_func def ptr_operand(src: T.Buffer((8,), "uint32")): @@ -264,9 +260,7 @@ def test_ptx_addr_pointer_and_raw_address_validation(): def test_ptx_addr_tma_tmem_and_independent_immediate_rejections(): - with pytest.raises( - (ValueError, tvm.error.DiagnosticError), match="does not support T.ptx.addr" - ): + with pytest.raises(ValueError, match="does not support T.ptx.addr"): @T.prim_func def tma(tmap: T.Buffer((8,), "uint64")): @@ -277,9 +271,7 @@ def test_ptx_addr_tma_tmem_and_independent_immediate_rejections(): shared_buf.data, T.ptx.addr(tmap.data, 16), T.int32(0), barrier.data ) - with pytest.raises( - (ValueError, tvm.error.DiagnosticError), match="does not support T.ptx.addr" - ): + with pytest.raises(ValueError, match="does not support T.ptx.addr"): @T.prim_func def tmem(raw: T.uint32): @@ -314,7 +306,7 @@ def test_ptx_addr_printer_script_and_json_roundtrip(): def test_ptx_addr_legacy_positional_offsets_rejected(): - with pytest.raises((ValueError, tvm.error.DiagnosticError)): + with pytest.raises(ValueError): @T.prim_func def scalar_load(src: T.Buffer((8,), "uint32")): @@ -322,7 +314,7 @@ def test_ptx_addr_legacy_positional_offsets_rejected(): value = T.local_scalar("uint32") T.ptx.ld.global_.b32(value, src.data, 16) - with pytest.raises((ValueError, tvm.error.DiagnosticError)): + with pytest.raises(ValueError): @T.prim_func def vector_load(src: T.Buffer((8,), "uint32")): @@ -330,14 +322,14 @@ def test_ptx_addr_legacy_positional_offsets_rejected(): values = T.alloc_local((2,), "uint32") T.ptx.ld.global_.v2.b32(values[0], values[1], src.data, 16) - with pytest.raises((ValueError, tvm.error.DiagnosticError)): + with pytest.raises(ValueError): @T.prim_func def scalar_store(dst: T.Buffer((8,), "uint32")): T.device_entry() T.ptx.st.global_.b32(dst.data, 16, T.uint32(0)) - with pytest.raises((ValueError, tvm.error.DiagnosticError)): + with pytest.raises(ValueError): @T.prim_func def vector_store(dst: T.Buffer((8,), "uint32")): diff --git a/tests/python/tirx/codegen/test_ptx_dialect.py b/tests/python/tirx/codegen/test_ptx_dialect.py index 9db4666189..81b0d7ec29 100644 --- a/tests/python/tirx/codegen/test_ptx_dialect.py +++ b/tests/python/tirx/codegen/test_ptx_dialect.py @@ -234,7 +234,7 @@ def test_ptx_red_vector_codegen_and_roundtrip(): reparsed = tvm.script.from_source(kernel.script()) tvm.ir.assert_structural_equal(kernel, reparsed) - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="already a 32-bit pair"): + with pytest.raises(ValueError, match="already a 32-bit pair"): @T.prim_func def packed_v8(a_ptr: T.handle): @@ -243,7 +243,7 @@ def test_ptx_red_vector_codegen_and_roundtrip(): v = T.local_scalar("uint32") T.ptx.red.global_.add.noftz.v8.f16x2(A.ptr_to([0]), v, v, v, v, v, v, v, v) - with pytest.raises((AttributeError, tvm.error.DiagnosticError), match="not a valid modifier"): + with pytest.raises(AttributeError, match="not a valid modifier"): @T.prim_func def f32_max(a_ptr: T.handle): @@ -672,7 +672,7 @@ def test_ptx_92_cp_bulk_roundtrip(): def test_ptx_trace_time_errors(): # Global ld fed a raw uint32 address. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="shared state space"): + with pytest.raises(ValueError, match="shared state space"): @T.prim_func def bad_global_addr(out: T.Buffer((1,), "uint32")): @@ -680,7 +680,7 @@ def test_ptx_trace_time_errors(): T.ptx.ld.global_.b32(out[0], T.uint32(0)) # Bogus modifier token. - with pytest.raises((AttributeError, tvm.error.DiagnosticError), match="not a valid modifier"): + with pytest.raises(AttributeError, match="not a valid modifier"): @T.prim_func def bad_modifier(out: T.Buffer((1,), "uint32")): @@ -690,7 +690,7 @@ def test_ptx_trace_time_errors(): # A floating register is not compatible with an integer instruction type, # even when it is wider (the relaxed ld/st rule admits a bit carrier, not a # differently typed floating register). - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="must have dtype"): + with pytest.raises(ValueError, match="must have dtype"): @T.prim_func def bad_value_dtype(a_ptr: T.handle): @@ -699,7 +699,7 @@ def test_ptx_trace_time_errors(): T.ptx.st.global_.u32(A.ptr_to([0]), T.float64(1.0)) # Missing required modifier (no type token). - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="missing required modifier"): + with pytest.raises(ValueError, match="missing required modifier"): @T.prim_func def missing_type(a_ptr: T.handle): @@ -709,7 +709,7 @@ def test_ptx_trace_time_errors(): # Per-slot legal tokens whose combination is illegal PTX: acquire # requires a scope — rejected by the entry's check function. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="requires a scope"): + with pytest.raises(ValueError, match="requires a scope"): @T.prim_func def acquire_without_scope(out: T.Buffer((1,), "uint32"), a_ptr: T.handle): @@ -718,7 +718,7 @@ def test_ptx_trace_time_errors(): T.ptx.ld.global_.acquire.b32(out[0], A.ptr_to([0])) # A float is not an address in any state space. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="pointer or uint64 handle"): + with pytest.raises(ValueError, match="pointer or uint64 handle"): @T.prim_func def bad_addr_dtype(out: T.Buffer((1,), "uint32")): @@ -730,7 +730,7 @@ def test_ptx_destination_errors(): """A destination is a register the caller declared: it must be a writable lvalue.""" # A floating destination is incompatible with an integer instruction type; # relaxed widening does not numerically convert between the two. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="must have dtype"): + with pytest.raises(ValueError, match="must have dtype"): @T.prim_func def wrong_dst_dtype(out: T.Buffer((1,), "float64"), a_ptr: T.handle): @@ -740,7 +740,7 @@ def test_ptx_destination_errors(): # A T.let binding is immutable, so it cannot be written into. This is the # gate that keeps the analyzer from re-expanding one call into N. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="writable scalar"): + with pytest.raises(ValueError, match="writable scalar"): @T.prim_func def let_destination(out: T.Buffer((1,), "uint32"), a_ptr: T.handle): @@ -750,7 +750,7 @@ def test_ptx_destination_errors(): T.ptx.ld.global_.b32(bound, A.ptr_to([0])) # An rvalue is not a destination either. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="writable scalar"): + with pytest.raises(ValueError, match="writable scalar"): @T.prim_func def rvalue_destination(a_ptr: T.handle): @@ -792,7 +792,7 @@ def test_ptx_register_group_codegen(): def test_ptx_register_group_errors(): """A register group is one operand: its arity is fixed and its lanes agree.""" # No `mov` shape takes two operands, so nothing in the family matches. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match=r"expects \d+ operand"): + with pytest.raises(ValueError, match=r"expects \d+ operand"): @T.prim_func def wrong_arity(a_ptr: T.handle): @@ -804,7 +804,7 @@ def test_ptx_register_group_errors(): # Each lane of a destination group is its own register the caller declared, # so each has to be an lvalue -- not just the first. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="writable scalar"): + with pytest.raises(ValueError, match="writable scalar"): @T.prim_func def non_lvalue_lane(a_ptr: T.handle): @@ -818,7 +818,7 @@ def test_ptx_register_group_errors(): # Lanes disagreeing on dtype: legal for each lane alone (both are 32-bit), # but the group is one operand with one C parameter type, so binding the # odd lane to it would be a numeric conversion. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="must have one dtype"): + with pytest.raises(ValueError, match="must have one dtype"): @T.prim_func def mixed_lane_dtypes(a_ptr: T.handle): @@ -832,7 +832,7 @@ def test_ptx_register_group_errors(): # A bare float literal names no dtype: on a .b32 lane it could be the # float's bits or the number, and it used to silently become T.uint32(1). - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="is ambiguous"): + with pytest.raises(ValueError, match="is ambiguous"): @T.prim_func def bare_float_literal(a_ptr: T.handle): @@ -1052,7 +1052,7 @@ def test_ptx_relaxed_load_store_typing(): assert '"l"(__value)' in src or '"l"(__value0)' in src assert "cvt." not in src - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="must have dtype"): + with pytest.raises(ValueError, match="must have dtype"): @T.prim_func def floating_source_for_integer_type(a_ptr: T.handle): @@ -1245,7 +1245,7 @@ def test_ptx_integer_arithmetic_dispatch(): # A derived dtype is enforced at trace time like any other: `.wide.s32` # writes 64 bits, so a 32-bit destination is rejected before codegen. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="must have dtype int64"): + with pytest.raises(ValueError, match="must have dtype int64"): @T.prim_func def narrow_wide_dst(out: T.Buffer((1,), "int32")): @@ -1254,7 +1254,7 @@ def test_ptx_integer_arithmetic_dispatch(): # `.sat` is a syntax line of its own, not a free qualifier: the check # rejects it where the ISA does not spell it. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="hi.s32"): + with pytest.raises(ValueError, match="hi.s32"): @T.prim_func def sat_on_lo(out: T.Buffer((1,), "int32")): @@ -1265,21 +1265,21 @@ def test_ptx_integer_arithmetic_dispatch(): # integer lines on which `.sat` has never been legal. Keep both reasons # visible so an invalid spelling is not misreported as an architecture # gate. In particular, sub.sat.u32 is not the counterpart of add.sat.u32. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="sm_120f-only"): + with pytest.raises(ValueError, match="sm_120f-only"): @T.prim_func def add_sat_sm120(out: T.Buffer((1,), "uint32")): T.device_entry() T.ptx.add.sat.u32(out[0], T.uint32(2), T.uint32(3)) - with pytest.raises((ValueError, tvm.error.DiagnosticError), match=r"not on the add\.u64"): + with pytest.raises(ValueError, match=r"not on the add\.u64"): @T.prim_func def add_sat_no_syntax_line(out: T.Buffer((1,), "uint64")): T.device_entry() T.ptx.add.sat.u64(out[0], T.uint64(2), T.uint64(3)) - with pytest.raises((ValueError, tvm.error.DiagnosticError), match=r"not on the sub\.u32"): + with pytest.raises(ValueError, match=r"not on the sub\.u32"): @T.prim_func def sub_sat_no_syntax_line(out: T.Buffer((1,), "uint32")): @@ -1350,7 +1350,7 @@ def test_ptx_floating_point_dispatch(): tvm.ir.assert_structural_equal(kernel, reparsed) # sqrt has no f64 approximation at any spelling, unlike rcp (9.7.3.14). - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="only on the .f32 line"): + with pytest.raises(ValueError, match="only on the .f32 line"): @T.prim_func def sqrt_approx_f64(out: T.Buffer((1,), "float64")): @@ -1358,7 +1358,7 @@ def test_ptx_floating_point_dispatch(): T.ptx.sqrt.approx.f64(out[0], T.float64(2.0)) # ... and rcp's f64 approximation is unreachable without its mandatory .ftz. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="rcp.approx.ftz.f64"): + with pytest.raises(ValueError, match="rcp.approx.ftz.f64"): @T.prim_func def rcp_approx_f64(out: T.Buffer((1,), "float64")): @@ -1366,7 +1366,7 @@ def test_ptx_floating_point_dispatch(): T.ptx.rcp.approx.f64(out[0], T.float64(2.0)) # One of .approx/.full/.rnd is required: the bare mnemonic names no line. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="missing required modifier"): + with pytest.raises(ValueError, match="missing required modifier"): @T.prim_func def div_without_mode(out: T.Buffer((1,), "float32")): @@ -1426,7 +1426,7 @@ def test_ptx_half_precision_dispatch(): tvm.ir.assert_structural_equal(kernel, reparsed) # The `.oob` line spells neither .ftz nor .sat beside it. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="oob line spells no"): + with pytest.raises(ValueError, match="oob line spells no"): @T.prim_func def oob_with_ftz(out: T.Buffer((1,), "uint16")): @@ -1434,7 +1434,7 @@ def test_ptx_half_precision_dispatch(): T.ptx.fma.rn.oob.ftz.f16(out[0], T.uint16(0), T.uint16(0), T.uint16(0)) # .sat and .relu are two clampings on two different syntax lines. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="separate syntax lines"): + with pytest.raises(ValueError, match="separate syntax lines"): @T.prim_func def sat_and_relu(out: T.Buffer((1,), "uint16")): @@ -1442,7 +1442,7 @@ def test_ptx_half_precision_dispatch(): T.ptx.fma.rn.sat.relu.f16(out[0], T.uint16(0), T.uint16(0), T.uint16(0)) # ex2's bf16 line without its mandatory .ftz names no syntax line. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="mandatorily"): + with pytest.raises(ValueError, match="mandatorily"): @T.prim_func def ex2_bf16_no_ftz(out: T.Buffer((1,), "uint16")): @@ -1450,7 +1450,7 @@ def test_ptx_half_precision_dispatch(): T.ptx.ex2.approx.bf16(out[0], T.uint16(0)) # ... and the bf16 abs/neg lines spell no .ftz at all. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="takes no .ftz"): + with pytest.raises(ValueError, match="takes no .ftz"): @T.prim_func def abs_bf16_ftz(out: T.Buffer((1,), "uint16")): @@ -1497,7 +1497,7 @@ def test_ptx_mixed_precision_dispatch(): tvm.ir.assert_structural_equal(kernel, reparsed) # The mixed line spells no .ftz, unlike the .f32 line it shares an entry with. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="takes no .ftz"): + with pytest.raises(ValueError, match="takes no .ftz"): @T.prim_func def mixed_with_ftz(out: T.Buffer((1,), "float32")): @@ -1505,7 +1505,7 @@ def test_ptx_mixed_precision_dispatch(): T.ptx.add.rn.ftz.f32.f16(out[0], T.uint16(0), T.float32(0)) # A converted source exists only on the .f32 line -- there is no f64 form. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="only exists on the .f32"): + with pytest.raises(ValueError, match="only exists on the .f32"): @T.prim_func def mixed_f64(out: T.Buffer((1,), "float64")): @@ -1514,7 +1514,7 @@ def test_ptx_mixed_precision_dispatch(): # And `mul` is the one mnemonic of the four with no mixed line at all, so # its entry declares no srctype slot and the token does not resolve. - with pytest.raises((AttributeError, tvm.error.DiagnosticError), match="not a valid modifier"): + with pytest.raises(AttributeError, match="not a valid modifier"): @T.prim_func def mul_mixed(out: T.Buffer((1,), "float32")): @@ -1574,7 +1574,7 @@ def test_ptx_comparison_selection_dispatch(): # Each source type takes its own operator set, and ptxas agrees with every # rejection below (probed at sm_90 before they were written into the check). # lo/ls/hi/hs are unsigned-only alternates for lt/le/gt/ge. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="signed"): + with pytest.raises(ValueError, match="signed"): @T.prim_func def unsigned_op_on_signed(out: T.Buffer((1,), "uint32")): @@ -1582,7 +1582,7 @@ def test_ptx_comparison_selection_dispatch(): T.ptx.setp.lo.s32(out[0], T.int32(1), T.int32(2)) # The unordered comparisons and the NaN predicates are floating point only. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="floating-point comparison"): + with pytest.raises(ValueError, match="floating-point comparison"): @T.prim_func def float_op_on_integer(out: T.Buffer((1,), "uint32")): @@ -1590,7 +1590,7 @@ def test_ptx_comparison_selection_dispatch(): T.ptx.setp.nan.u32(out[0], T.uint32(1), T.uint32(2)) # A bit-size type has no ordering, only equality. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="only with eq/ne"): + with pytest.raises(ValueError, match="only with eq/ne"): @T.prim_func def ordered_on_bitsize(out: T.Buffer((1,), "uint32")): @@ -1598,7 +1598,7 @@ def test_ptx_comparison_selection_dispatch(): T.ptx.setp.lt.b32(out[0], T.uint32(1), T.uint32(2)) # .ftz applies only to .f32 comparisons ... - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="only to .f32 comparisons"): + with pytest.raises(ValueError, match="only to .f32 comparisons"): @T.prim_func def ftz_off_f32(out: T.Buffer((1,), "uint32")): @@ -1606,7 +1606,7 @@ def test_ptx_comparison_selection_dispatch(): T.ptx.setp.eq.ftz.s32(out[0], T.int32(1), T.int32(2)) # ... and on slct, only to the .f32 selector line. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="f32 selector line"): + with pytest.raises(ValueError, match="f32 selector line"): @T.prim_func def ftz_on_s32_selector(out: T.Buffer((1,), "uint32")): @@ -1616,9 +1616,7 @@ def test_ptx_comparison_selection_dispatch(): # Only d/a/b receive slct's relaxed bit-size typing. The selector must # still match the second instruction type exactly, even when another # same-constraint integer carrier would fit through inline asm. - with pytest.raises( - (ValueError, tvm.error.DiagnosticError), match=r"operand 'c'.*int32.*uint32" - ): + with pytest.raises(ValueError, match=r"operand 'c'.*int32.*uint32"): @T.prim_func def relaxed_selector(out: T.Buffer((1,), "uint32")): @@ -1716,7 +1714,7 @@ def test_ptx_half_comparison_dispatch(): # Every rejection below was probed against ptxas before being written into # the check. The (dtype, stype) grid: a packed source needs a 32-bit answer. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="no half-precision syntax"): + with pytest.raises(ValueError, match="no half-precision syntax"): @T.prim_func def bad_pair(out: T.Buffer((1,), "uint16")): @@ -1724,7 +1722,7 @@ def test_ptx_half_comparison_dispatch(): T.ptx.set.eq.u16.f16x2(out[0], T.uint32(0), T.uint32(0)) # `.ftz` follows the source precision: a bf16 source has none to flush. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="does not take it"): + with pytest.raises(ValueError, match="does not take it"): @T.prim_func def ftz_bf16_source(out: T.Buffer((1,), "uint32")): @@ -1732,14 +1730,14 @@ def test_ptx_half_comparison_dispatch(): T.ptx.set.eq.ftz.u32.bf16(out[0], T.uint16(0), T.uint16(0)) # ... and no bf16 destination takes it either. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="destination takes no"): + with pytest.raises(ValueError, match="destination takes no"): @T.prim_func def ftz_bf16_dest(out: T.Buffer((1,), "uint16")): T.device_entry() T.ptx.set.eq.ftz.bf16.f32(out[0], T.float32(0), T.float32(0)) - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="spells no .ftz"): + with pytest.raises(ValueError, match="spells no .ftz"): @T.prim_func def setp_ftz_bf16(out: T.Buffer((1,), "uint32")): @@ -1747,7 +1745,7 @@ def test_ptx_half_comparison_dispatch(): T.ptx.setp.eq.ftz.bf16(out[0], T.uint16(0), T.uint16(0)) # The unordered comparisons need a floating-point source. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="integer source"): + with pytest.raises(ValueError, match="integer source"): @T.prim_func def unordered_on_integer(out: T.Buffer((1,), "uint16")): @@ -1756,7 +1754,7 @@ def test_ptx_half_comparison_dispatch(): # The classification fix is narrow: a real .b16 source remains a bit-size # value and therefore still permits equality comparisons only. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="bit-size source"): + with pytest.raises(ValueError, match="bit-size source"): @T.prim_func def ordered_on_bitsize(out: T.Buffer((1,), "uint16")): @@ -1764,7 +1762,7 @@ def test_ptx_half_comparison_dispatch(): T.ptx.set.lt.f16.b16(out[0], T.uint16(0), T.uint16(0)) # 9.7.8 spells no unsigned alternates at all, so `lo` never resolves here. - with pytest.raises((AttributeError, tvm.error.DiagnosticError), match="not a valid modifier"): + with pytest.raises(AttributeError, match="not a valid modifier"): @T.prim_func def unsigned_alternate(out: T.Buffer((1,), "uint32")): @@ -1836,7 +1834,7 @@ def test_ptx_logic_shift_dispatch(): # cnot has no predicate line -- ptxas: "Unexpected instruction types # specified for 'cnot'" -- so the token is not in its slot at all. - with pytest.raises((AttributeError, tvm.error.DiagnosticError), match="not a valid modifier"): + with pytest.raises(AttributeError, match="not a valid modifier"): @T.prim_func def cnot_pred(out: T.Buffer((1,), "uint32")): @@ -1844,7 +1842,7 @@ def test_ptx_logic_shift_dispatch(): T.ptx.cnot.pred(out[0], T.uint32(0)) # lop3's BoolOp line stops at .or/.and; ptxas rejects .xor outright. - with pytest.raises((AttributeError, tvm.error.DiagnosticError), match="not a valid modifier"): + with pytest.raises(AttributeError, match="not a valid modifier"): @T.prim_func def lop3_xor(out: T.Buffer((1,), "uint32")): @@ -1853,7 +1851,7 @@ def test_ptx_logic_shift_dispatch(): # shl is untyped -- a left shift zero-fills whatever the bits mean, so # there is no signed line (unlike shr, which needs to know how to fill). - with pytest.raises((AttributeError, tvm.error.DiagnosticError), match="not a valid modifier"): + with pytest.raises(AttributeError, match="not a valid modifier"): @T.prim_func def shl_signed(out: T.Buffer((1,), "int32")): @@ -1903,9 +1901,7 @@ def test_ptx_logic_shift_dispatch(): assert "lop3.b32 %0, %1, %2, %3, 0;" in boundary_src assert "lop3.b32 %0, %1, %2, %3, 255;" in boundary_src - with pytest.raises( - (ValueError, tvm.error.DiagnosticError), match=r"inclusive range 0\.\.255, got -1" - ): + with pytest.raises(ValueError, match=r"inclusive range 0\.\.255, got -1"): @T.prim_func def lut_below_range(a_ptr: T.handle): @@ -1914,9 +1910,7 @@ def test_ptx_logic_shift_dispatch(): d = T.local_scalar("uint32") T.ptx.lop3.b32(d, A[0], A[0], A[0], -1) - with pytest.raises( - (ValueError, tvm.error.DiagnosticError), match=r"inclusive range 0\.\.255, got 256" - ): + with pytest.raises(ValueError, match=r"inclusive range 0\.\.255, got 256"): @T.prim_func def lut_above_range(a_ptr: T.handle): @@ -2018,7 +2012,7 @@ def test_ptx_data_movement_dispatch(): # Every rejection below was probed against ptxas before it was written into # a check. multimem pairs an ordering semantic with a scope ... - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="go together"): + with pytest.raises(ValueError, match="go together"): @T.prim_func def sem_without_scope(a_ptr: T.handle): @@ -2028,7 +2022,7 @@ def test_ptx_data_movement_dispatch(): T.ptx.multimem_ld_reduce.relaxed.add.u32(v, A.ptr_to([0])) # ... and `.weak` is the line that has neither. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="takes no scope"): + with pytest.raises(ValueError, match="takes no scope"): @T.prim_func def weak_with_scope(a_ptr: T.handle): @@ -2038,7 +2032,7 @@ def test_ptx_data_movement_dispatch(): T.ptx.multimem_ld_reduce.weak.gpu.add.u32(v, A.ptr_to([0])) # The op x type table: `.add` is the row that takes .s32 but not .s64. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match=r"\.add takes"): + with pytest.raises(ValueError, match=r"\.add takes"): @T.prim_func def add_s64(a_ptr: T.handle): @@ -2048,7 +2042,7 @@ def test_ptx_data_movement_dispatch(): T.ptx.multimem_ld_reduce.add.s64(v, A.ptr_to([0])) # The scalar float line has no lone half: a width has to reach 32 bits. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="the scalar line takes"): + with pytest.raises(ValueError, match="the scalar line takes"): @T.prim_func def scalar_f16(a_ptr: T.handle): @@ -2058,7 +2052,7 @@ def test_ptx_data_movement_dispatch(): T.ptx.multimem_ld_reduce.add.f16(v, A.ptr_to([0])) # `.acc::f32` raises an accumulation, so it needs something to accumulate. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="applies to .add"): + with pytest.raises(ValueError, match="applies to .add"): @T.prim_func def acc_on_min(a_ptr: T.handle): @@ -2068,7 +2062,7 @@ def test_ptx_data_movement_dispatch(): T.ptx.multimem_ld_reduce.min.acc__f32.f16x2(v, A.ptr_to([0])) # st.async's mmio line is system-scoped only. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="requires .sys"): + with pytest.raises(ValueError, match="requires .sys"): @T.prim_func def mmio_gpu(a_ptr: T.handle): @@ -2139,7 +2133,7 @@ def test_ptx_parallel_sync_dispatch(): # Every rejection below was probed against ptxas first. redux pairs the # arithmetic ops with a signed type and the bitwise ones with the untyped # word -- they are two syntax lines, so the token does not even resolve. - with pytest.raises((AttributeError, tvm.error.DiagnosticError), match="not a valid modifier"): + with pytest.raises(AttributeError, match="not a valid modifier"): @T.prim_func def redux_add_b32(out: T.Buffer((1,), "uint32")): @@ -2148,7 +2142,7 @@ def test_ptx_parallel_sync_dispatch(): # atom's vector lines bound the width by the element: a packed pair stops # at .v4, and only a lone half reaches .v8. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="already a 32-bit pair"): + with pytest.raises(ValueError, match="already a 32-bit pair"): @T.prim_func def packed_v8(a_ptr: T.handle): @@ -2160,7 +2154,7 @@ def test_ptx_parallel_sync_dispatch(): ) # red.async's op groups: `.add` reaches 64 bits, the bitwise ops do not. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match=r"\.and takes"): + with pytest.raises(ValueError, match=r"\.and takes"): @T.prim_func def red_async_and_u64(a_ptr: T.handle): @@ -2291,7 +2285,7 @@ def test_ptx_wgmma_scale_d_runtime_predicate_roundtrip(): # There is deliberately no compatibility overload for the old integer # immediate spelling. A bool already identifies a predicate; a runtime # integer must carry the explicit T.ptx.pred(...) register-class marker. - with pytest.raises((ValueError, tvm.error.DiagnosticError), match=r"T\.ptx\.pred"): + with pytest.raises(ValueError, match=r"T\.ptx\.pred"): @T.prim_func def bare_integer_scale_d(): @@ -2436,7 +2430,7 @@ def test_ptx_tcgen05_mma_block_size_form(): assert "tcgen05.mma.cta_group::1.kind::mxf4.block_scale.block32" in src _assert_ptxas_ok(src, arch="sm_100a") - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="mxf4.*block32"): + with pytest.raises(ValueError, match="mxf4.*block32"): @T.prim_func def invalid_mxf4_block16(): @@ -2603,7 +2597,7 @@ def test_ptx_pred_operand_rejects_untagged_integer(): tells the two `cp.async` optional-operand syntax lines apart. A bool needs no tag -- it already says what it is. """ - with pytest.raises((ValueError, tvm.error.DiagnosticError), match=r"T\.ptx\.pred"): + with pytest.raises(ValueError, match=r"T\.ptx\.pred"): @T.prim_func def untagged_integer(): @@ -2670,7 +2664,7 @@ def test_ptx_sink_rejected_where_the_isa_has_no_underscore(): Not from the family either: mov's *pack* shape has a scalar destination and a vector source, and the ISA gives neither of them a sink. """ - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="not sinkable here"): + with pytest.raises(ValueError, match="not sinkable here"): @T.prim_func def sink_a_source(): @@ -2682,7 +2676,7 @@ def test_ptx_sink_rejected_where_the_isa_has_no_underscore(): T.evaluate(hi) # ISA 9.7.10.4: "provided that at least one element is a scalar register". - with pytest.raises((ValueError, tvm.error.DiagnosticError), match="must be a real register"): + with pytest.raises(ValueError, match="must be a real register"): @T.prim_func def sink_every_lane(): diff --git a/tests/python/tirx/script/test_frontend_construction_validation.py b/tests/python/tirx/script/test_frontend_construction_validation.py index a219763551..ba31e70eea 100644 --- a/tests/python/tirx/script/test_frontend_construction_validation.py +++ b/tests/python/tirx/script/test_frontend_construction_validation.py @@ -22,7 +22,6 @@ import pytest import tvm from tvm import tirx -from tvm.error import DiagnosticError from tvm.script import ir as I from tvm.script import tirx as T @@ -38,7 +37,7 @@ def _build_lexical_global(): def test_default_tirx_validation_rejects_non_divisible_scope_extents(): - with pytest.raises(DiagnosticError): + with pytest.raises(ValueError): @T.prim_func def invalid(): @@ -66,7 +65,7 @@ def test_tirx_validation_can_be_explicitly_disabled(): def test_optional_annotation_is_restricted_to_jit(): - with pytest.raises(DiagnosticError, match="only supported by @T.jit"): + with pytest.raises(TypeError, match="only supported by @T.jit"): @T.prim_func(private=True) def invalid(value: T.Optional(T.handle)): diff --git a/tests/python/tirx/script/test_native_span_and_loop_contracts.py b/tests/python/tirx/script/test_native_span_and_loop_contracts.py index 42285ff38e..fed75c5e2c 100644 --- a/tests/python/tirx/script/test_native_span_and_loop_contracts.py +++ b/tests/python/tirx/script/test_native_span_and_loop_contracts.py @@ -17,6 +17,8 @@ """Deferred native results own spans and loop variables before frame entry.""" +import traceback + import pytest import tvm_ffi @@ -168,19 +170,34 @@ def test_source_call_preserves_cyclic_containers_and_native_member_spans(contain def test_native_diagnostics_keep_definition_location_and_restore_source_context(failure): with IRBuilder() as builder: if failure == "source_call": + original_error = ValueError("native helper failed") def fail(): - raise ValueError("native helper failed") + raise original_error with pytest.raises(ValueError, match="native helper failed") as error: _captured(fail) + assert error.value is original_error + assert type(error.value) is ValueError + frames = traceback.extract_tb(error.value.__traceback__) + assert frames[-1].filename == __file__ + assert frames[-1].name == "fail" + assert frames[-1].lineno == fail.__code__.co_firstlineno + 1 + assert any( + frame.filename == __file__ and frame.line.strip() == "_captured(fail)" + for frame in frames + ) else: frame = _captured(lambda: native.If(tirx.IntImm("bool", True))) + definition_span = frame.source_span + assert positions(definition_span) == [("caller.py", 3), ("frame.py", 7)] frame.__enter__() with builder.with_source_span(base.source_span(location(90, "unrelated.py"))): with pytest.raises(InternalError, match="then") as error: frame.__exit__(None, None, None) - assert error.value.__tvm_script_location__ == location(7) + assert type(error.value) is InternalError + assert frame.source_span.same_as(definition_span) + assert positions(frame.source_span) == [("caller.py", 3), ("frame.py", 7)] following = native.serial(2) assert following.source_span is None with following: diff --git a/tests/python/tirx/script/test_parser_buffer_views.py b/tests/python/tirx/script/test_parser_buffer_views.py index d3675a6da7..6c59cfb613 100644 --- a/tests/python/tirx/script/test_parser_buffer_views.py +++ b/tests/python/tirx/script/test_parser_buffer_views.py @@ -22,7 +22,6 @@ from types import MethodType import pytest from tvm import ir, tirx -from tvm.error import DiagnosticError from tvm.script import parser from tvm.script.ir_builder import IRBuilder from tvm.script.ir_builder import ir as I @@ -168,7 +167,7 @@ def test_native_roundtrip_preserves_buffer_identity_and_shared_methods(): ], ) def test_view_validation(view, error): - with pytest.raises(DiagnosticError, match=error): + with pytest.raises(ValueError, match=error): parse(f""" @T.prim_func def main(): diff --git a/tests/python/tirx/script/test_parser_jit_specialization.py b/tests/python/tirx/script/test_parser_jit_specialization.py index 7ead95acc7..3e13a0cd1c 100644 --- a/tests/python/tirx/script/test_parser_jit_specialization.py +++ b/tests/python/tirx/script/test_parser_jit_specialization.py @@ -123,7 +123,7 @@ def test_specialization_context_restored_after_failure(): missing_call() # noqa: F821 T.evaluate(N) - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(NameError): broken.specialize(N=1) assert len(broken.specialize(N=2).params) == 0 result = parser.parse("@T.prim_func\ndef normal(N: T.int32):\n T.evaluate(N)\n") @@ -232,7 +232,7 @@ def test_capture_transport_does_not_specialize_ordinary_captured_names(): def test_missing_constexpr_binding_has_explicit_diagnostic(): - with pytest.raises(tvm.error.DiagnosticError, match="requires a specialization binding"): + with pytest.raises(TypeError, match="requires a specialization binding"): parser.parse("@T.jit\ndef kernel(N: T.constexpr):\n T.evaluate(N)\n") @@ -250,7 +250,7 @@ def kernel(a: T.Optional(annotation())): missing_call() T.evaluate(1) """ - with pytest.raises(tvm.error.DiagnosticError, match="missing_call"): + with pytest.raises(NameError, match="missing_call"): parser.parse( source, extra_vars={"annotation": annotation}, _specialization_bindings={"a": None} ) diff --git a/tests/python/tirx/script/test_parser_tirx_emission.py b/tests/python/tirx/script/test_parser_tirx_emission.py index 8bbd845a31..ea32d5489c 100644 --- a/tests/python/tirx/script/test_parser_tirx_emission.py +++ b/tests/python/tirx/script/test_parser_tirx_emission.py @@ -18,7 +18,6 @@ import pytest -import tvm from tvm import tirx from tvm.relax.script import builder as R from tvm.script import parser @@ -47,9 +46,9 @@ def test_source_store_uses_public_builder_and_propagates_errors(monkeypatch, err monkeypatch.setattr(tir_builder, "buffer_store", reject_store) source = f"@T.prim_func\ndef main():\n {body}\n" - with pytest.raises(tvm.error.DiagnosticError, match="store rejected") as raised: + with pytest.raises(error_type, match="store rejected") as raised: parser.parse(source) - assert raised.value.__cause__ is error + assert raised.value is error assert len(calls) == 1 diff --git a/tests/python/tirx/script/test_tirx_parser_comparison_integration.py b/tests/python/tirx/script/test_tirx_parser_comparison_integration.py index fcc739a664..ae5918bfaf 100644 --- a/tests/python/tirx/script/test_tirx_parser_comparison_integration.py +++ b/tests/python/tirx/script/test_tirx_parser_comparison_integration.py @@ -105,8 +105,6 @@ def test_iterator_comparison_keeps_written_operand_order(operator, kind): @pytest.mark.parametrize("operator,kind", OPERATORS) def test_unmarked_host_comparison_never_calls_python_overload(operator, kind): - from tvm.error import DiagnosticError - calls = [] class Host: @@ -130,7 +128,7 @@ def test_unmarked_host_comparison_never_calls_python_overload(operator, kind): ) decorator = "T.prim_func" statement = "T.evaluate(left " + operator + " right)" - with pytest.raises(DiagnosticError, match="(PrimExpr|primitive|convert|type)"): + with pytest.raises(TypeError, match="(PrimExpr|primitive|convert|type)"): parser.parse( f"@{decorator}\ndef main():\n {statement}\n", extra_vars={"left": Host(), "right": Host()}, diff --git a/tests/python/tirx/test_jit.py b/tests/python/tirx/test_jit.py index 00af56df61..bf7d8b38fb 100644 --- a/tests/python/tirx/test_jit.py +++ b/tests/python/tirx/test_jit.py @@ -19,6 +19,7 @@ from __future__ import annotations +import traceback import typing import pytest @@ -364,7 +365,7 @@ def test_only_explicit_t_optional_is_specializable(): def test_t_optional_is_restricted_to_jit(): - with pytest.raises(tvm.error.DiagnosticError, match="only supported by @T.jit"): + with pytest.raises(TypeError, match="only supported by @T.jit"): @T.prim_func(private=True) def invalid(a: T.Optional(T.handle)): @@ -411,19 +412,19 @@ def test_runtime_tir_if_cannot_guard_absent_optional_param(): if flag != 0: T.match_buffer(a, (1,), "int32") - with pytest.raises(tvm.error.DiagnosticError, match="match_buffer"): + with pytest.raises(ValueError, match="match_buffer"): kernel.specialize(a=None) @pytest.mark.parametrize( - ("operation", "source_text"), + ("operation", "source_text", "error_type"), [ - ("subscript", "a[10]"), - ("attribute", "a.ptr_to"), - ("match_buffer", "T.match_buffer"), + ("subscript", "a[10]", TypeError), + ("attribute", "a.ptr_to", AttributeError), + ("match_buffer", "T.match_buffer", tvm.error.InternalError), ], ) -def test_unguarded_absent_optional_param_reports_source(operation, source_text): +def test_unguarded_absent_optional_param_reports_source(operation, source_text, error_type): @T.jit(private=True) def kernel(a: T.Optional(T.handle)): if T.constexpr(operation == "subscript"): @@ -433,9 +434,11 @@ def test_unguarded_absent_optional_param_reports_source(operation, source_text): else: T.match_buffer(a, (1,), "int32") - with pytest.raises(tvm.error.DiagnosticError) as exc_info: + with pytest.raises(error_type) as exc_info: kernel.specialize(a=None) - assert source_text in str(exc_info.value) + assert type(exc_info.value) is error_type + frames = traceback.extract_tb(exc_info.value.__traceback__) + assert any(frame.filename == __file__ and source_text in frame.line for frame in frames) def test_present_optional_param_still_rejects_ffi_none(): diff --git a/tests/python/tirx/test_parser_printer.py b/tests/python/tirx/test_parser_printer.py index d98cb2c36d..603d004857 100644 --- a/tests/python/tirx/test_parser_printer.py +++ b/tests/python/tirx/test_parser_printer.py @@ -757,7 +757,7 @@ def test_meta_class_constructor_rejects_unowned_resource(): def __init__(self): tmp = T.alloc_buffer((1,), "int32", scope="local") - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @T.prim_func def test(): @@ -1088,11 +1088,10 @@ def func(): v: T.int32 v = v + T.int32(1) """ - # The ValueError propagates through the parser framework which wraps it - # into a DiagnosticError. Before the fix the broad ``except Exception`` - # would silently swallow it and fall through to eval_assign. + # The ValueError propagates unchanged. A broad ``except Exception`` here + # previously swallowed it and fell through to eval_assign. with patch("tvm.tirx.script.builder.buffer_store", side_effect=bomb): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError, match="boom"): from_source(src) @@ -1251,7 +1250,7 @@ from tvm.script import tirx as T def func(): x: T.handle = T.int64(0) """ - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(tvm.error.InternalError): from_source(src_handle) # 3. Banned: non-PrimType annotation without T.let @@ -1262,7 +1261,7 @@ from tvm.ir import PointerType, PrimType def func(): x: T.Var(name="x", ty=PointerType(PrimType("float16"))) = T.int64(0) """ - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(tvm.error.InternalError): from_source(src_ptr) # 4. An explicit mutable scalar declaration retains updates — round-trip @@ -1503,7 +1502,7 @@ def test_buffer_local_compose_layout_printer_roundtrip(): def test_buffer_local_inference_without_parent_layout_has_clear_diagnostic(): """Shape inference requires a parent storage layout.""" - with pytest.raises(tvm.error.DiagnosticError, match="parent buffer has layout=None"): + with pytest.raises(ValueError, match="parent buffer has layout=None"): # fmt: off @T.prim_func def func() -> None: @@ -1632,7 +1631,7 @@ def test_buffer_local_printer_preserves_inherited_metadata(): def test_buffer_local_rejects_shape_that_does_not_match_physical_span(): """An explicit local shape product must preserve the physical span.""" - with pytest.raises(tvm.error.DiagnosticError, match="physical storage span 6 per thread"): + with pytest.raises(ValueError, match="physical storage span 6 per thread"): # fmt: off @T.prim_func def func() -> None: @@ -2628,7 +2627,7 @@ def test_buffer_sub_tmem_rejects_partial_column_offset(): return func - with pytest.raises(tvm.error.DiagnosticError, match="aligned to a physical 32-bit column"): + with pytest.raises(ValueError, match="aligned to a physical 32-bit column"): build() diff --git a/tests/python/tvmscript/test_tvmscript_error_report.py b/tests/python/tvmscript/test_tvmscript_error_report.py index aec053fbe7..a59c6f83b6 100644 --- a/tests/python/tvmscript/test_tvmscript_error_report.py +++ b/tests/python/tvmscript/test_tvmscript_error_report.py @@ -15,8 +15,10 @@ # specific language governing permissions and limitations # under the License. # ruff: noqa: E741, F821, F841, RUF005 +import ast import inspect import re +import traceback import pytest @@ -27,39 +29,42 @@ from tvm.script import from_source from tvm.script import tirx as T -def check_error(func, rel_lineno): - """check if TIR script throws error""" - check_error_re = re.compile(r"^.*# check_error: (.+)$") +def check_error(func, rel_lineno, error_type): + """Check the original exception class and its real source location.""" source_code = inspect.getsource(func) indent = len(re.match(r"^\s*", source_code).group(0)) source_code = "@T.prim_func(s_tir=True)\n" + "\n".join( line[indent:] for line in source_code.splitlines() ) - # Parse errors now raise DiagnosticError with formatted source location. - with pytest.raises(tvm.error.DiagnosticError) as execinfo: + with pytest.raises(error_type) as caught: from_source(source_code) - err_str = str(execinfo.value) - if rel_lineno is None: - return - # The error message contains " --> <source>:<lineno>:<col>" formatted by Diagnostics.error(). - # Accept either rel_lineno or rel_lineno+1 to match old tolerance. - assert f":{rel_lineno}:" in err_str or f":{rel_lineno + 1}:" in err_str, ( - f"Expected error message to contain line {rel_lineno}, got:\n{err_str}" - ) - error_line = source_code.split("\n")[rel_lineno] - m = check_error_re.match(error_line) - if m: - expected_error_text = m.group(1) - assert expected_error_text in err_str, ( - f'check_error expects "{expected_error_text}" in error: {err_str}' - ) + assert type(caught.value) is error_type + if isinstance(caught.value, SyntaxError): + assert caught.value.filename == "<str>" + assert caught.value.lineno == rel_lineno + 1 + assert caught.value.offset is not None + assert caught.value.end_lineno >= caught.value.lineno + assert caught.value.end_offset is not None + else: + frames = [ + frame + for frame in traceback.extract_tb(caught.value.__traceback__) + if frame.filename == "<str>" + ] + assert frames + if rel_lineno is not None: + assert frames[-1].lineno == rel_lineno + 1 + if rel_lineno is not None: + match = re.match(r"^.*# check_error: (.+)$", source_code.splitlines()[rel_lineno]) + if match: + assert match.group(1) in str(caught.value) def test_buffer_bind(): def buffer_bind_missing_args(a: T.handle) -> None: A = T.match_buffer((16, 16), "float32") # error - check_error(buffer_bind_missing_args, 2) + check_error(buffer_bind_missing_args, 2, tvm.error.InternalError) def test_undefined_buffer(): @@ -70,7 +75,7 @@ def test_undefined_buffer(): for j in T.serial(0, 16): C[i, j] = 0.0 # error - check_error(undefined_buffer, 6) + check_error(undefined_buffer, 6, NameError) def test_unsupported_function_call(): @@ -81,14 +86,14 @@ def test_unsupported_function_call(): for j in T.serial(0, 16): A[i, j] = 0.0 - check_error(unsupported_function_call, 4) + check_error(unsupported_function_call, 4, AttributeError) def test_missing_type_annotation(): def missing_type_annotation(a) -> None: # error T.evaluate(0.0) - check_error(missing_type_annotation, 1) + check_error(missing_type_annotation, 1, SyntaxError) def test_invalid_for_function(): @@ -98,7 +103,7 @@ def test_invalid_for_function(): for j in T.serial(0, 16): A[i, j] = 0.0 - check_error(invalid_for_function, 4) + check_error(invalid_for_function, 3, TypeError) def test_invalid_block_function(): @@ -108,14 +113,14 @@ def test_invalid_block_function(): with T.evaluate(0.0): # error T.evaluate(1.0) - check_error(invalid_block_function, 4) + check_error(invalid_block_function, 4, TypeError) def test_return_not_allowed(): def return_not_allowed(a: T.handle) -> None: return T.evaluate(0) # error - check_error(return_not_allowed, 2) + check_error(return_not_allowed, 2, NotImplementedError) def test_no_body(): @@ -123,7 +128,7 @@ def test_no_body(): A = T.match_buffer(a, (16, 16), "float32") T.realize(A, "") # error - check_error(no_body, 3) + check_error(no_body, 3, AttributeError) def test_inconsistent_binding(): @@ -137,8 +142,8 @@ def test_inconsistent_binding(): vi, vj = T.axis.remap("S", [i, j]) # error T.evaluate(1.0) - check_error(inconsistent_binding_value, 3) - check_error(inconsistent_binding_type, 3) + check_error(inconsistent_binding_value, 3, tvm.error.InternalError) + check_error(inconsistent_binding_type, 3, tvm.error.InternalError) def test_error_remap_args(): @@ -154,8 +159,8 @@ def test_error_remap_args(): vi, vj = T.axis.remap("SS", [i + j, j]) # error T.evaluate(1.0) - check_error(error_remap_type, 4) - check_error(error_remap_value, 4) + check_error(error_remap_type, 4, tvm.error.InternalError) + check_error(error_remap_value, 4, tvm.error.InternalError) def test_invalid_block_axes(): @@ -166,7 +171,7 @@ def test_invalid_block_axes(): vi = T.axis.S(i, A) # error T.evaluate(1.0) - check_error(invalid_block_axes, 5) + check_error(invalid_block_axes, 5, TypeError) def test_duplicate_block_axes(): @@ -183,8 +188,8 @@ def test_duplicate_block_axes(): vi, vi = T.axis.remap("SS", [i, j]) # error T.evaluate(1.0) - check_error(duplicate_block_axes, 5) - check_error(duplicate_block_axes_remap, 4) + check_error(duplicate_block_axes, 5, ValueError) + check_error(duplicate_block_axes_remap, 4, ValueError) def test_miss_block_bind(): @@ -194,7 +199,7 @@ def test_miss_block_bind(): vi = T.axis.S(i) # error T.evaluate(1.0) - check_error(miss_block_bind_value, 4) + check_error(miss_block_bind_value, 4, TypeError) def test_invalid_loop_var(): @@ -202,15 +207,15 @@ def test_invalid_loop_var(): for i, j in range(0, 16): # error T.evaluate(1.0) - check_error(invalid_loop_var, 2) + check_error(invalid_loop_var, 2, ValueError) def test_inconsistent_grid(): def inconsistent_grid(A: T.Buffer(16)) -> None: - for i in T.grid(16, 16): # valid, i is a tuple (iter0, iter1) - T.evaluate(A[i]) # error + for i in T.grid(16, 16): # error: one target cannot unpack two loop variables + T.evaluate(A[i]) - check_error(inconsistent_grid, 3) + check_error(inconsistent_grid, 2, ValueError) def test_invalid_match_buffer_region(): @@ -221,7 +226,7 @@ def test_invalid_match_buffer_region(): A = T.match_buffer(vi) # error T.evaluate(1.0) - check_error(invalid_match_buffer_region, 5) + check_error(invalid_match_buffer_region, 5, ValueError) def test_buffer_rebinding_preserves_distinct_allocations(): @@ -292,12 +297,12 @@ def test_duplicate_block_signature(): T.sblock_attr({"key1": "block2"}) # error T.evaluate(1.0) - check_error(duplicate_reads, 7) - check_error(duplicate_writes, 7) - check_error(duplicate_predicate, 6) - check_error(duplicate_init, 7) - check_error(duplicate_axes, 5) - check_error(duplicate_sblock_attrs_with_same_key_diff_value, 6) + check_error(duplicate_reads, 7, tvm.error.InternalError) + check_error(duplicate_writes, 7, tvm.error.InternalError) + check_error(duplicate_predicate, 6, tvm.error.InternalError) + check_error(duplicate_init, 7, ValueError) + check_error(duplicate_axes, 5, ValueError) + check_error(duplicate_sblock_attrs_with_same_key_diff_value, 6, tvm.error.InternalError) def test_opaque_access_during_complete(): @@ -307,7 +312,7 @@ def test_opaque_access_during_complete(): with T.sblock(): T.evaluate(T.call_extern("dummy_extern_function", A.data, dtype="int32")) - check_error(opaque_access_during_complete, None) + check_error(opaque_access_during_complete, None, ValueError) def test_convert_slice_to_bufferload(): @@ -318,7 +323,7 @@ def test_convert_slice_to_bufferload(): vi, vj = T.axis.remap("SS", [i, j]) A[vi, vj] = A[vi : vi + 2, vj] + 1 # error - check_error(convert_slice_to_bufferload, 6) + check_error(convert_slice_to_bufferload, 6, TypeError) def test_tvm_exception_catch_from_special_stmt(): @@ -326,7 +331,7 @@ def test_tvm_exception_catch_from_special_stmt(): A = T.sblock_alloc_buffer("(128, 128)", "float32") # error T.evaluate(1.0) - check_error(special_stmt_except, 2) + check_error(special_stmt_except, 2, TypeError) def test_tvm_exception_catch_from_scope_handler(): @@ -334,7 +339,7 @@ def test_tvm_exception_catch_from_scope_handler(): for i in T.serial("1", "1"): # error T.evaluate(1) - check_error(scope_handler_except, 2) + check_error(scope_handler_except, 2, TypeError) def test_tvm_exception_catch_from_bare_intrin(): @@ -342,7 +347,7 @@ def test_tvm_exception_catch_from_bare_intrin(): A = T.match_buffer(a, (16, 16), "float32") T.evaluate(A) # error - check_error(intrin_except_unassign, 3) + check_error(intrin_except_unassign, 3, tvm.error.InternalError) def test_tvm_exception_catch_from_assigned_intrin(): @@ -350,7 +355,7 @@ def test_tvm_exception_catch_from_assigned_intrin(): A = T.match_buffer(a, (16, 16), "float32") A[0, 0] = A[A] # error - check_error(intrin_except_assign, 3) + check_error(intrin_except_assign, 3, tvm.error.InternalError) def test_match_buffer_shape_mismatch(): @@ -366,7 +371,7 @@ def test_match_buffer_shape_mismatch(): for jj in range(0, 4): sub_A[i, j * 4 + jj] = 1 - check_error(buffer_shape_mismatch, 7) + check_error(buffer_shape_mismatch, 7, tvm.error.InternalError) def test_high_dim_store(): @@ -376,7 +381,7 @@ def test_high_dim_store(): for i, j in T.grid(16, 16): B[i, j] = 1.0 # error: Store is only allowed with one index - check_error(high_dim_store, 5) + check_error(high_dim_store, 5, tvm.error.InternalError) def test_block_has_option_vars(): @@ -384,7 +389,7 @@ def test_block_has_option_vars(): with T.sblock("root") as x: # error: block does not support option_vars T.evaluate(0.0) - check_error(block_has_option_vars, 2) + check_error(block_has_option_vars, 2, TypeError) def test_implicit_root_has_attrs(): @@ -408,11 +413,11 @@ def test_implicit_root_has_attrs(): v = T.axis.S(0, 0) # error: implicit root does not support axis define T.evaluate(0.0) - check_error(implicit_root_has_read, 2) - check_error(implicit_root_has_write, 2) - check_error(implicit_root_has_attrs, 2) - check_error(implicit_root_has_predicate, 2) - check_error(implicit_root_has_axes, 2) + check_error(implicit_root_has_read, 2, ValueError) + check_error(implicit_root_has_write, 2, ValueError) + check_error(implicit_root_has_attrs, 2, tvm.error.InternalError) + check_error(implicit_root_has_predicate, 2, ValueError) + check_error(implicit_root_has_axes, 2, tvm.error.InternalError) @T.prim_func(s_tir=True) @@ -502,7 +507,7 @@ def test_load_var(): d = T.float32() d[2] = d[2, 1] # error cannot provide two indices to load - check_error(load_var_multiple, 3) + check_error(load_var_multiple, 3, TypeError) def test_store_var(): @@ -510,7 +515,7 @@ def test_store_var(): d = T.float32() d[2, 1] = d[1] # error cannot provide two indices to store - check_error(store_var_multiple, 3) + check_error(store_var_multiple, 3, TypeError) def test_load_handle(): @@ -518,7 +523,7 @@ def test_load_handle(): h_ = T.match_buffer(h, [1]) h_[0] = h[0] # error cannot load from handle - check_error(load_handle, 3) + check_error(load_handle, 3, TypeError) def test_store_handle(): @@ -526,7 +531,7 @@ def test_store_handle(): h_ = T.match_buffer(h, [1]) h[0] = h_[0] # error cannot store to handle - check_error(store_handle, 3) + check_error(store_handle, 3, TypeError) def test_binop_bad_ast_type(): @@ -534,7 +539,7 @@ def test_binop_bad_ast_type(): h_ = T.match_buffer(h, [1]) h_[0] = h + [2] # error rhs should be a primexpr - check_error(binop_bad_ast_type, 3) + check_error(binop_bad_ast_type, 3, TypeError) def test_binop_bad_type(): @@ -542,7 +547,7 @@ def test_binop_bad_type(): h_ = T.match_buffer(h, [1]) h_[0] = h + 2 # error lhs and rhs should be the same type - check_error(binop_bad_type, 3) + check_error(binop_bad_type, 3, TypeError) def test_non_integer_typed_block_iter(): @@ -550,7 +555,7 @@ def test_non_integer_typed_block_iter(): with T.sblock(): i = T.axis.S(0.1, 0.1) # error IterVar requires an integer dtype - check_error(non_integer_typed_block_iter, 3) + check_error(non_integer_typed_block_iter, 3, tvm.error.InternalError) def test_illegal_buffer_slice(): @@ -573,9 +578,9 @@ def test_illegal_buffer_slice(): for i in range(4): T.evaluate(A[0:i:1]) # error - check_error(strided_buffer_region, 3) - check_error(access_reversed_slice, 3) - check_error(access_non_const_slice_length, 3) + check_error(strided_buffer_region, 3, tvm.error.InternalError) + check_error(access_reversed_slice, 3, tvm.error.InternalError) + check_error(access_non_const_slice_length, 3, tvm.error.InternalError) def test_syntax_sugar_fail(): @@ -584,13 +589,11 @@ def test_syntax_sugar_fail(): for i in T.thread_binding(128, 128): A[i] = A[i] * 2.0 - check_error(loop_syntax_sugar_fail, 3) + check_error(loop_syntax_sugar_fail, 3, ValueError) def test_multi_line_error_report(): - """A parse error whose offending AST node spans several physical source - lines must render ALL spanned lines (each with its own gutter line number - and an underline covering the span), not just the first line.""" + """An original builder failure retains the full source-call traceback range.""" # The offending call (`T.axis.remap(...)`) is deliberately split across # four physical lines so its AST node spans lineno..end_lineno > lineno. @@ -607,69 +610,26 @@ def test_multi_line_error_report(): ] ) - with pytest.raises(tvm.error.DiagnosticError) as execinfo: + with pytest.raises(tvm.error.InternalError) as caught: from_source(source_code) - err_str = str(execinfo.value) - - # All four spanned source lines must appear in the rendered snippet. - assert "T.axis.remap(" in err_str, err_str - assert '"S",' in err_str, err_str - assert "[i, j]," in err_str, err_str - # The trailing `)` closing line is also part of the span. - rendered_lines = err_str.splitlines() - assert any(" 7 " in line and ")" in line for line in rendered_lines), err_str - # The underline carets must be present on more than one line (multi-line). - marker_lines = [line for line in rendered_lines if "^" in line] - assert len(marker_lines) >= 2, err_str - # The gutter must show distinct line numbers for the spanned lines. - assert " 4 " in err_str and " 5 " in err_str and " 6 " in err_str, err_str - - -def test_format_source_snippet_multi_line(): - """Unit-level check that _format_source_snippet renders every line in a - multi-line span, with the underline covering start-col..EOL on the first - line, full interior lines, and col-1..end-col on the last line.""" - from tvm.script.parser.diagnostics import _format_source_snippet - - source_lines = [ - "first ignored line\n", - " foo(bar,\n", - " baz,\n", - " qux)\n", - "last ignored line\n", + frames = [ + frame + for frame in traceback.extract_tb(caught.value.__traceback__) + if frame.filename == "<str>" ] - # Span lines 2..4 (1-based), starting at col 5 ('foo'), ending at col 13 - # (exclusive) on line 4. - snippet = _format_source_snippet( - source_lines, lineno=2, col_offset=5, end_lineno=4, end_col_offset=13 - ) - lines = snippet.splitlines() - # All three spanned source lines must be present. - assert any("foo(bar," in line for line in lines), snippet - assert any("baz," in line for line in lines), snippet - assert any("qux)" in line for line in lines), snippet - # Underline carets present on the first line under 'foo(bar,'. - assert "^" in snippet, snippet - # The line numbers 2, 3, 4 appear in the gutter. - assert " 2 |" in snippet and " 3 |" in snippet and " 4 |" in snippet, snippet - - -def test_format_source_snippet_single_line_unchanged(): - """A single-line span (end_lineno == lineno) underlines only the - [col_offset, end_col_offset) columns on that one line.""" - from tvm.script.parser.diagnostics import _format_source_snippet - - source_lines = ["ignored\n", " abc + def\n", "ignored\n"] - # Underline just 'abc' (cols 5..8 exclusive) on line 2. - snippet = _format_source_snippet( - source_lines, lineno=2, col_offset=5, end_lineno=2, end_col_offset=8 - ) - lines = snippet.splitlines() - # Exactly one source-text line and one marker line (plus the leading gutter). - text_lines = [line for line in lines if "abc + def" in line] - assert len(text_lines) == 1, snippet - marker_line = next(line for line in lines if "^" in line) - assert marker_line.count("^") == 3, snippet + assert frames[-1].lineno == 4 + if getattr(frames[-1], "colno", None) is not None: + call = next( + node + for node in ast.walk(ast.parse(source_code)) + if isinstance(node, ast.Call) and node.lineno == 4 + ) + assert ( + frames[-1].lineno, + frames[-1].colno, + frames[-1].end_lineno, + frames[-1].end_colno, + ) == (call.lineno, call.col_offset, call.end_lineno, call.end_col_offset) if __name__ == "__main__": diff --git a/tests/python/tvmscript/test_tvmscript_ops.py b/tests/python/tvmscript/test_tvmscript_ops.py index 7793fb72c6..5184043a7e 100644 --- a/tests/python/tvmscript/test_tvmscript_ops.py +++ b/tests/python/tvmscript/test_tvmscript_ops.py @@ -247,7 +247,7 @@ try: T.bitwise_or(A[0:5] > T.Broadcast(T.float32(0), 5), B[0:5] > T.Broadcast(T.float32(0), 5)) T.bitwise_and(A[0:5] < T.Broadcast(T.float32(0), 5), T.Broadcast(T.bool(1), 5)) T.bitwise_or(A[0:5] > T.Broadcast(T.float32(0), 5), T.Broadcast(T.bool(1), 5)) -except tvm.error.DiagnosticError: +except TypeError: slice_op_test = None slice_op_test_ref = None diff --git a/tests/python/tvmscript/test_tvmscript_parser_source.py b/tests/python/tvmscript/test_tvmscript_parser_source.py index 6ab4dc552b..3939a76cfb 100644 --- a/tests/python/tvmscript/test_tvmscript_parser_source.py +++ b/tests/python/tvmscript/test_tvmscript_parser_source.py @@ -17,6 +17,7 @@ # ruff: noqa: F401 """Source and span tests for the canonical parser""" +import ast as doc import inspect import pytest @@ -27,8 +28,7 @@ import tvm import tvm.testing from tvm.ir import Call, SequentialSpan, TensorLoad, assert_structural_equal from tvm.script import tirx as T -import ast as doc -from tvm.script.parser.source import Source +from tvm.script.parser.inspect_source import Source from tvm.script.tirx import tile as Tx from tvm.tirx.stmt import TilePrimitiveCall diff --git a/tests/python/tvmscript/test_tvmscript_parser_tir.py b/tests/python/tvmscript/test_tvmscript_parser_tir.py index 0321bedc3e..1e3837025e 100644 --- a/tests/python/tvmscript/test_tvmscript_parser_tir.py +++ b/tests/python/tvmscript/test_tvmscript_parser_tir.py @@ -134,7 +134,7 @@ def main(A: T.Buffer(("n",), "float32")): def test_tir_direct_use_before_string_definition_is_undefined(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(NameError): tvm.script.from_source( """ @T.prim_func @@ -162,12 +162,12 @@ class Module: ], ) def test_tir_direct_later_prim_param_is_undefined(source): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(NameError): tvm.script.from_source(source) def test_tir_return_annotation_does_not_define_symbolic_var(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): tvm.script.from_source( """ @T.prim_func @@ -228,7 +228,7 @@ def test_tir_func_private_attrs(): def test_tir_func_private_manual_global_symbol_fail(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(tvm.error.InternalError): @T.prim_func(private=True, s_tir=True) def matmul(a: T.handle, b: T.handle, c: T.handle) -> None: diff --git a/tests/python/tvmscript/test_tvmscript_syntax_sugar.py b/tests/python/tvmscript/test_tvmscript_syntax_sugar.py index 006e415012..2184f9a09f 100644 --- a/tests/python/tvmscript/test_tvmscript_syntax_sugar.py +++ b/tests/python/tvmscript/test_tvmscript_syntax_sugar.py @@ -233,7 +233,7 @@ def test_match_buffer_region_has_implicit_shape_dtype(): def test_match_buffer_input_requires_shape_arg(): - with pytest.raises(tvm.error.DiagnosticError): + with pytest.raises(ValueError): @T.prim_func(s_tir=True) def func(a: T.handle):
