This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm-tirx.git
The following commit(s) were added to refs/heads/main by this push:
new 2ffcce69c [CI] Tighten Ruff checks and preserve parser exceptions (#46)
2ffcce69c is described below
commit 2ffcce69ca11a68b87aab1cc8fd95e46f3af7742
Author: Tianqi Chen <[email protected]>
AuthorDate: Sat Oct 24 18:32:01 2026 -0400
[CI] Tighten Ruff checks and preserve parser exceptions (#46)
Enforce the existing Ruff families by removing global style ignores and
narrowing package and test exemptions to imports and source-bound DSL
expressions that require them. Update the managed Ruff hook and CUDA
generator dependency together to 0.17.2, where the regex and
unused-unpacking checks are stable, while retaining the pre-commit
mechanism and CI invocation.
Clean up ordinary Python unused bindings, imports, callbacks, shared
class data, formatting and exception handling while preserving
parser-generated IR and public namespace exports.
---
.pre-commit-config.yaml | 4 +--
ci/jenkins/scripts/determine_docker_images.py | 4 +--
ci/jenkins/scripts/git_skip_ci_globs.py | 6 ++--
pyproject.toml | 13 +-------
.../tvm_tirx/backend/trn/tile/binary/__init__.py | 3 ++
.../backend/trn/tile/compose_op/__init__.py | 3 ++
python/tvm_tirx/backend/trn/tile/copy/__init__.py | 2 +-
python/tvm_tirx/backend/trn/tile/gemm/__init__.py | 2 +-
python/tvm_tirx/backend/trn/tile/private_alloc.py | 2 +-
.../tvm_tirx/backend/trn/tile/select/__init__.py | 2 +-
python/tvm_tirx/backend/trn/tile/unary/__init__.py | 6 ++--
python/tvm_tirx/runtime/_tensor.py | 10 ++----
python/tvm_tirx/script/__init__.py | 4 +--
python/tvm_tirx/script/printer/highlight.py | 12 ++++---
python/tvm_tirx/support/popen_pool.py | 13 +++++---
python/tvm_tirx/support/popen_worker.py | 4 +--
python/tvm_tirx/support/utils.py | 7 ++--
python/tvm_tirx/support/xcode.py | 5 ++-
python/tvm_tirx/testing/__init__.py | 2 +-
python/tvm_tirx/testing/bench.py | 10 +++---
python/tvm_tirx/tirx/__init__.py | 2 ++
python/tvm_tirx/tirx/analysis/__init__.py | 2 +-
python/tvm_tirx/tirx/expr.py | 7 ++--
python/tvm_tirx/tirx/script/__init__.py | 3 +-
python/tvm_tirx/tirx/script/ir_builder/__init__.py | 9 +++---
.../tirx/script/ir_builder/external_kernel.py | 4 +--
python/tvm_tirx/tirx/script/ir_builder/triton.py | 9 +++---
python/tvm_tirx/tirx/transform/__init__.py | 2 +-
tests/lint/check_asf_header.py | 13 ++++----
tests/python/backend/cuda/codegen/test_basic.py | 2 +-
tests/python/backend/cuda/codegen/test_fp8.py | 9 +++---
tests/python/backend/rocm/codegen/test_basic.py | 3 ++
tests/python/codegen/llvm/test_basic.py | 37 +++++++++++++---------
tests/python/codegen/llvm/test_riscv.py | 2 +-
tests/python/codegen/test_extern.py | 3 +-
tests/python/ir/test_attrs.py | 6 ++--
tests/python/ir/test_roundtrip_runtime_module.py | 4 +--
tests/python/script/test_basic_usage.py | 4 ++-
tests/python/script/test_dynamic_shape.py | 5 ++-
tests/python/script/test_error_handling.py | 12 +++----
tests/python/script/test_meta_programming.py | 4 ++-
tests/python/support/test_popen_pool.py | 7 ++--
tests/python/sym/test_canonical_simplify.py | 3 +-
tests/python/sym/test_const_int_bound.py | 9 +++---
tests/python/sym/test_deduce_bound.py | 5 +--
tests/python/sym/test_intset.py | 4 +--
tests/python/sym/test_iter_affine_map.py | 6 ++--
tests/python/sym/test_modular_set.py | 8 ++---
tests/python/sym/test_rewrite_simplify.py | 5 +--
tests/python/target/test_parser_mprofile.py | 2 +-
tests/python/target/test_riscv_features.py | 5 ++-
tests/python/target/test_x86_features.py | 3 +-
.../tirx/analysis/test_check_well_defined.py | 2 +-
tests/python/tirx/base/test_index_map.py | 4 +--
tests/python/tirx/base/test_intrin.py | 5 +--
tests/python/tirx/base/test_nodes.py | 9 +++---
tests/python/tirx/codegen/test_ampere.py | 3 ++
tests/python/tirx/codegen/test_blackwell.py | 3 ++
tests/python/tirx/codegen/test_cuda.py | 11 ++++---
tests/python/tirx/codegen/test_cuda_cta_reduce.py | 3 ++
tests/python/tirx/codegen/test_cuda_warp_reduce.py | 3 ++
tests/python/tirx/codegen/test_dsmem.py | 3 ++
tests/python/tirx/codegen/test_hopper.py | 21 ++++++------
tests/python/tirx/codegen/test_ptx_addr.py | 8 ++---
tests/python/tirx/codegen/test_ptx_dialect.py | 28 ++++++++--------
tests/python/tirx/codegen/test_ptx_ld_st_ops.py | 1 -
tests/python/tirx/iket/test_orchestration.py | 2 +-
tests/python/tirx/iket/test_profiler.py | 6 ++--
tests/python/tirx/script/test_basic_usage.py | 11 +++++--
tests/python/tirx/script/test_dynamic_shape.py | 4 +--
tests/python/tirx/script/test_error_handling.py | 2 +-
tests/python/tirx/script/test_gpu.py | 3 ++
tests/python/tirx/script/test_ir_builder.py | 18 +++++------
tests/python/tirx/script/test_meta_programming.py | 5 ++-
tests/python/tirx/script/test_printer.py | 4 ---
tests/python/tirx/script/test_printer_function.py | 2 +-
.../tirx/script/test_printer_structural_equal.py | 3 ++
tests/python/tirx/script/test_source_locations.py | 11 +++++--
tests/python/tirx/script/test_tensors.py | 6 ++--
tests/python/tirx/test_base.py | 2 +-
tests/python/tirx/test_control_flow.py | 3 ++
tests/python/tirx/test_jit.py | 33 ++++++++++---------
tests/python/tirx/test_ops.py | 4 +--
tests/python/tirx/test_verifier.py | 5 ++-
.../python/tirx/tile/cuda/test_add_sub_mul_div.py | 7 ++--
tests/python/tirx/tile/cuda/test_cp_async.py | 3 ++
.../tirx/tile/cuda/test_cp_async_bulk_tensor.py | 6 ++--
tests/python/tirx/tile/cuda/test_cvt.py | 7 ++--
tests/python/tirx/tile/cuda/test_ewise.py | 9 ++++--
.../tirx/tile/cuda/test_explicit_patterns.py | 4 +--
tests/python/tirx/tile/cuda/test_fma.py | 3 ++
.../tirx/tile/cuda/test_ld_st_cache_policy.py | 2 +-
.../python/tirx/tile/cuda/test_ld_st_gmem_smem.py | 2 +-
.../tirx/tile/cuda/test_ldmatrix_stmatrix.py | 3 --
tests/python/tirx/tile/cuda/test_mma_sync.py | 2 +-
tests/python/tirx/tile/cuda/test_tcgen05_cp.py | 14 ++++----
tests/python/tirx/tile/cuda/test_tcgen05_ld_st.py | 7 ++--
tests/python/tirx/tile/cuda/test_tcgen05_mma.py | 15 ++++-----
.../tirx/tile/cuda/test_tcgen05_mma_codegen.py | 8 ++---
tests/python/tirx/tile/test_dispatcher.py | 4 +--
tests/python/tirx/tile/trn/test_copy.py | 2 +-
tests/python/tirx/transform/test_convert_ssa.py | 6 ++--
.../tirx/transform/test_lower_tensor_layouts.py | 3 ++
.../tirx/transform/test_lower_tvm_builtin.py | 3 ++
tests/python/tirx/transform/test_remove_no_op.py | 3 ++
tests/python/tirx/transform/test_simplify.py | 5 ++-
.../tirx/transform/test_split_host_device.py | 3 ++
.../python/tirx/transform/test_storage_rewrite.py | 3 ++
tests/python/tirx/transform/test_tile_dispatch.py | 3 ++
tests/python/tirx/transform/test_unroll.py | 3 ++
tests/python/tirx/transform/test_vectorize.py | 3 ++
111 files changed, 379 insertions(+), 295 deletions(-)
diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index 3d0139548..ff7532635 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -25,7 +25,7 @@ repos:
name: check CUDA launch generated files
entry: python python/tvm_tirx/backend/cuda/launch/generate.py --check
language: python
- additional_dependencies: [ruff==0.12.3, clang-format==20.1.8]
+ additional_dependencies: [ruff==0.17.2, clang-format==20.1.8]
pass_filenames: false
files: ^(python/tvm_tirx/backend/cuda/launch/|src/backend/cuda/launch/)
- id: check-asf-header
@@ -78,7 +78,7 @@ repos:
# ---------- Python: ruff ----------
- repo: https://github.com/astral-sh/ruff-pre-commit
- rev: v0.12.3
+ rev: v0.17.2
hooks:
- id: ruff-check
types_or: [python, pyi, jupyter]
diff --git a/ci/jenkins/scripts/determine_docker_images.py
b/ci/jenkins/scripts/determine_docker_images.py
index 9802ea784..3fe8c87fb 100755
--- a/ci/jenkins/scripts/determine_docker_images.py
+++ b/ci/jenkins/scripts/determine_docker_images.py
@@ -15,7 +15,6 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: E501
import argparse
import configparser
import datetime
@@ -54,7 +53,8 @@ def docker_api(url: str, use_pagination: bool = False) ->
dict[str, Any]:
reset = datetime.datetime.fromtimestamp(int(reset))
reset = reset.isoformat()
logging.info(
- f"Docker API Rate Limit: {headers.get('x-ratelimit-remaining')} /
{headers.get('x-ratelimit-limit')} (reset at {reset})"
+ f"Docker API Rate Limit: {headers.get('x-ratelimit-remaining')} / "
+ f"{headers.get('x-ratelimit-limit')} (reset at {reset})"
)
return r
diff --git a/ci/jenkins/scripts/git_skip_ci_globs.py
b/ci/jenkins/scripts/git_skip_ci_globs.py
index f7fc0860e..816217d94 100755
--- a/ci/jenkins/scripts/git_skip_ci_globs.py
+++ b/ci/jenkins/scripts/git_skip_ci_globs.py
@@ -15,7 +15,6 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: E501
import argparse
import fnmatch
@@ -43,7 +42,10 @@ def match_any(f: str) -> str | None:
if __name__ == "__main__":
- help = "Exits with code 1 if a change only touched files, indicating that
CI could be skipped for this changeset"
+ help = (
+ "Exits with code 1 if a change only touched files, "
+ "indicating that CI could be skipped for this changeset"
+ )
parser = argparse.ArgumentParser(description=help)
parser.add_argument("--files", help="(testing only) comma separated list
of files to check")
args = parser.parse_args()
diff --git a/pyproject.toml b/pyproject.toml
index 3be000989..8441b4349 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -194,22 +194,11 @@ select = [
"UP", # pyupgrade, https://docs.astral.sh/ruff/rules/#pyupgrade-up
"RUF", # ruff, https://docs.astral.sh/ruff/rules/#ruff-specific-rules-ruf
]
-ignore = [
- # pyupgrade (style preference or low-value)
- "UP028", # yield-in-for-loop — sometimes clearer to keep
- "UP030", # format-literals — minor
- "UP031", # printf-string-formatting — style preference
- # ruff (preview rules, not yet stable)
- "RUF043", # pytest-raises-ambiguous-pattern
- "RUF059", # unused-unpacked-variable
-]
fixable = ["ALL"]
unfixable = []
[tool.ruff.lint.per-file-ignores]
-"__init__.py" = ["E402", "F401", "F403", "F405"]
-# TIRx parser and lowering tests intentionally retain unused IR bindings.
-"tests/python/tirx/**/*.py" = ["F841"]
+"__init__.py" = ["F401"]
[tool.ruff.lint.isort]
known-first-party = ["tvm_tirx"]
diff --git a/python/tvm_tirx/backend/trn/tile/binary/__init__.py
b/python/tvm_tirx/backend/trn/tile/binary/__init__.py
index ed01927cf..4243d54a7 100644
--- a/python/tvm_tirx/backend/trn/tile/binary/__init__.py
+++ b/python/tvm_tirx/backend/trn/tile/binary/__init__.py
@@ -15,5 +15,8 @@
# specific language governing permissions and limitations
# under the License.
+# Re-export the public namespace.
+# ruff: noqa: F403
+
from .default import *
from .utils import *
diff --git a/python/tvm_tirx/backend/trn/tile/compose_op/__init__.py
b/python/tvm_tirx/backend/trn/tile/compose_op/__init__.py
index 58c9a2c40..0b49900d5 100644
--- a/python/tvm_tirx/backend/trn/tile/compose_op/__init__.py
+++ b/python/tvm_tirx/backend/trn/tile/compose_op/__init__.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# Re-export the public namespace.
+# ruff: noqa: F403
+
from .binary_chain import *
from .binary_reduce import *
from .unary_reduce import *
diff --git a/python/tvm_tirx/backend/trn/tile/copy/__init__.py
b/python/tvm_tirx/backend/trn/tile/copy/__init__.py
index 358e44931..024d1f673 100644
--- a/python/tvm_tirx/backend/trn/tile/copy/__init__.py
+++ b/python/tvm_tirx/backend/trn/tile/copy/__init__.py
@@ -15,4 +15,4 @@
# specific language governing permissions and limitations
# under the License.
-from .default import *
+from .default import * # noqa: F403 - Re-export the public namespace.
diff --git a/python/tvm_tirx/backend/trn/tile/gemm/__init__.py
b/python/tvm_tirx/backend/trn/tile/gemm/__init__.py
index 358e44931..024d1f673 100644
--- a/python/tvm_tirx/backend/trn/tile/gemm/__init__.py
+++ b/python/tvm_tirx/backend/trn/tile/gemm/__init__.py
@@ -15,4 +15,4 @@
# specific language governing permissions and limitations
# under the License.
-from .default import *
+from .default import * # noqa: F403 - Re-export the public namespace.
diff --git a/python/tvm_tirx/backend/trn/tile/private_alloc.py
b/python/tvm_tirx/backend/trn/tile/private_alloc.py
index 43b63cad6..abc480614 100644
--- a/python/tvm_tirx/backend/trn/tile/private_alloc.py
+++ b/python/tvm_tirx/backend/trn/tile/private_alloc.py
@@ -83,7 +83,7 @@ def alloc_const_bias_trn(
if isinstance(max_inst_size, int | IntImm) and int(max_inst_size) == -1:
raise ValueError("Constant bias workspace allocation requires a finite
max_inst_size")
if bias_key in tensor_dict:
- bias_alloc, bias_init_stmt = tensor_dict[bias_key]
+ bias_alloc, _bias_init_stmt = tensor_dict[bias_key]
old_shape = bias_alloc.var.ty.shape
new_shape = [max(par_size, old_shape[0]), max(max_inst_size,
old_shape[1])]
if new_shape[0] == old_shape[0] and new_shape[1] == old_shape[1]:
diff --git a/python/tvm_tirx/backend/trn/tile/select/__init__.py
b/python/tvm_tirx/backend/trn/tile/select/__init__.py
index 358e44931..024d1f673 100644
--- a/python/tvm_tirx/backend/trn/tile/select/__init__.py
+++ b/python/tvm_tirx/backend/trn/tile/select/__init__.py
@@ -15,4 +15,4 @@
# specific language governing permissions and limitations
# under the License.
-from .default import *
+from .default import * # noqa: F403 - Re-export the public namespace.
diff --git a/python/tvm_tirx/backend/trn/tile/unary/__init__.py
b/python/tvm_tirx/backend/trn/tile/unary/__init__.py
index fa2b223d2..c074cf24d 100644
--- a/python/tvm_tirx/backend/trn/tile/unary/__init__.py
+++ b/python/tvm_tirx/backend/trn/tile/unary/__init__.py
@@ -15,6 +15,6 @@
# specific language governing permissions and limitations
# under the License.
-from .default import *
-from .utils import *
-from .with_bias_scale import *
+from .default import * # noqa: F403 - Re-export the public namespace.
+from .utils import * # noqa: F403 - Re-export the public namespace.
+from .with_bias_scale import * # noqa: F403 - Re-export the public namespace.
diff --git a/python/tvm_tirx/runtime/_tensor.py
b/python/tvm_tirx/runtime/_tensor.py
index c7cd312f1..0b98eb701 100644
--- a/python/tvm_tirx/runtime/_tensor.py
+++ b/python/tvm_tirx/runtime/_tensor.py
@@ -15,12 +15,10 @@
# specific language governing permissions and limitations
# under the License.
# pylint: disable=invalid-name, unused-import, redefined-outer-name
-# ruff: noqa: F401, RUF005
+
"""Runtime Tensor API"""
import ctypes
-import warnings
-from typing import Optional
import numpy as np
@@ -32,8 +30,6 @@ except ImportError:
import tvm_ffi
from tvm_ffi import DLDeviceType
-from tvm_tirx.runtime import Device
-
from . import _ffi_api
@@ -100,7 +96,7 @@ class Tensor(tvm_ffi.core.Tensor):
t = tvm_ffi.dtype(self.dtype)
shape, dtype = self.shape, self.dtype
if t.lanes > 1:
- shape = shape + (t.lanes,)
+ shape = (*shape, t.lanes)
t = t.with_lanes(1)
dtype = str(t)
@@ -159,7 +155,7 @@ class Tensor(tvm_ffi.core.Tensor):
shape, dtype = self.shape, self.dtype
old_dtype = dtype
if t.lanes > 1:
- shape = shape + (t.lanes,)
+ shape = (*shape, t.lanes)
t = t.with_lanes(1)
dtype = str(t)
if dtype == "int4":
diff --git a/python/tvm_tirx/script/__init__.py
b/python/tvm_tirx/script/__init__.py
index bcd71c957..3ed6db644 100644
--- a/python/tvm_tirx/script/__init__.py
+++ b/python/tvm_tirx/script/__init__.py
@@ -210,10 +210,10 @@ def __getattr__(name: str) -> Any:
# Shared source entry points participate in the same opaque namespace boundary.
-from .parser import register_namespace as _register_namespace
+from .parser import register_namespace as _register_namespace # noqa: E402
_register_namespace("script", sys.modules[__name__])
# Install the Script-owned display methods after the shared IR types exist.
-from . import printer
+from . import printer # noqa: E402 - Import after namespace initialization.
diff --git a/python/tvm_tirx/script/printer/highlight.py
b/python/tvm_tirx/script/printer/highlight.py
index bf1dd5d74..f8f923ef1 100644
--- a/python/tvm_tirx/script/printer/highlight.py
+++ b/python/tvm_tirx/script/printer/highlight.py
@@ -14,13 +14,15 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F821, RUF012
"""Highlight printed TVM script."""
import os
import sys
import warnings
-from typing import Any, Union
+from typing import TYPE_CHECKING, Any, ClassVar, Union
+
+if TYPE_CHECKING:
+ import pygments.style
def cprint(
@@ -142,7 +144,7 @@ def _get_pygments_style(
"""A Jupyter-Notebook-like Pygments style configuration (aka.
"light")"""
background_color = ""
- styles = {
+ styles: ClassVar[dict] = {
Keyword: "bold #008000",
Keyword.Type: "nobold #008000",
Name.Function: "#0000FF",
@@ -159,7 +161,7 @@ def _get_pygments_style(
"""A VSCode-Dark-like Pygments style configuration (aka. "dark")"""
background_color = ""
- styles = {
+ styles: ClassVar[dict] = {
Keyword: "bold #c586c0",
Keyword.Type: "#82aaff",
Keyword.Namespace: "#4ec9b0",
@@ -177,7 +179,7 @@ def _get_pygments_style(
"""The default style for terminal display with ANSI colors (aka.
"ansi")"""
background_color = ""
- styles = {
+ styles: ClassVar[dict] = {
Keyword: "bold ansigreen",
Keyword.Type: "nobold ansigreen",
Name.Class: "bold ansiblue",
diff --git a/python/tvm_tirx/support/popen_pool.py
b/python/tvm_tirx/support/popen_pool.py
index 41e925e55..731e17198 100644
--- a/python/tvm_tirx/support/popen_pool.py
+++ b/python/tvm_tirx/support/popen_pool.py
@@ -15,7 +15,7 @@
# specific language governing permissions and limitations
# under the License.
# pylint: disable=invalid-name
-# ruff: noqa: E731
+
"""Multiprocessing via Popen.
This module provides a multi-processing pool backed by Popen.
@@ -459,8 +459,10 @@ class PopenPoolExecutor:
future : concurrent.futures.Future
A future that can be used to access the result.
"""
- # pylint: disable=unnecessary-lambda
- worker = lambda *args: self._worker_run(*args)
+
+ def worker(*args):
+ return self._worker_run(*args)
+
return self._threadpool.submit(worker, fn, args, kwargs)
def map_with_error_catching(self, fn, iterator):
@@ -479,5 +481,8 @@ class PopenPoolExecutor:
out_iter : Iterator[MapResult]
The result iterator.
"""
- worker = lambda x: self._worker_run_with_error_catching(fn, (x,), None)
+
+ def worker(x):
+ return self._worker_run_with_error_catching(fn, (x,), None)
+
return self._threadpool.map(worker, iterator)
diff --git a/python/tvm_tirx/support/popen_worker.py
b/python/tvm_tirx/support/popen_worker.py
index 8ec7b6361..4ae7f9d89 100644
--- a/python/tvm_tirx/support/popen_worker.py
+++ b/python/tvm_tirx/support/popen_worker.py
@@ -31,7 +31,7 @@
# specific language governing permissions and limitations
# under the License.
# pylint: disable=invalid-name
-# ruff: noqa: RUF012
+
"""Internal PopenWorker for PopenPool."""
import logging
@@ -48,7 +48,7 @@ from tvm_tirx.support.popen_pool import StatusKind
class TimeoutStatus:
- __slot__ = ["status"]
+ __slots__ = ("status",)
def __init__(self):
self.status = StatusKind.RUNNING
diff --git a/python/tvm_tirx/support/utils.py b/python/tvm_tirx/support/utils.py
index ca0e70d81..c2f06b615 100644
--- a/python/tvm_tirx/support/utils.py
+++ b/python/tvm_tirx/support/utils.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: RUF005, RUF012
+
"""Common system utilities"""
import atexit
@@ -25,6 +25,7 @@ import pathlib
import shutil
import tempfile
import threading
+from typing import ClassVar
try:
import fcntl
@@ -71,7 +72,7 @@ class TempDirectory:
)
return cls._DEBUG_PARENT_DIR
- TEMPDIRS = set()
+ TEMPDIRS: ClassVar[set[str]] = set()
@classmethod
def remove_tempdirs(cls):
@@ -271,7 +272,7 @@ def which(exec_name):
path : str
The full path of executable if found, otherwise returns None
"""
- base_list = ["", "/bin"] + os.environ.get("PATH", "").split(os.pathsep)
+ base_list = ["", "/bin", *os.environ.get("PATH", "").split(os.pathsep)]
for path in base_list:
full_path = os.path.join(path, exec_name)
if os.path.isfile(full_path) and os.access(full_path, os.X_OK):
diff --git a/python/tvm_tirx/support/xcode.py b/python/tvm_tirx/support/xcode.py
index ddfaa0e66..741df4da1 100644
--- a/python/tvm_tirx/support/xcode.py
+++ b/python/tvm_tirx/support/xcode.py
@@ -15,7 +15,6 @@
# specific language governing permissions and limitations
# under the License.
# pylint: disable=invalid-name
-# ruff: noqa: E501, RUF005
"""Utility to invoke Xcode compiler toolchain"""
import json
@@ -39,7 +38,7 @@ def xcrun(cmd):
out : str
The output string.
"""
- cmd = ["xcrun"] + cmd
+ cmd = ["xcrun", *cmd]
proc = subprocess.Popen(cmd, stdout=subprocess.PIPE,
stderr=subprocess.STDOUT)
(out, _) = proc.communicate()
return out.strip()
@@ -136,7 +135,7 @@ def compile_metal(code, path_target=None, sdk="macosx",
min_os_version=None):
file_target = path_target if path_target else temp_target
# See:
- # -
https://developer.apple.com/documentation/metal/gpu_functions_libraries/building_a_library_with_metal_s_command-line_tools#overview
# pylint: disable=line-too-long
+ # -
https://developer.apple.com/documentation/metal/gpu_functions_libraries/building_a_library_with_metal_s_command-line_tools#overview
#
# xcrun -sdk macosx metal -c MyLibrary.metal -o MyLibrary.air
# xcrun -sdk macosx metallib MyLibrary.air -o MyLibrary.metallib
diff --git a/python/tvm_tirx/testing/__init__.py
b/python/tvm_tirx/testing/__init__.py
index 448d5a2a0..63be3b4cc 100644
--- a/python/tvm_tirx/testing/__init__.py
+++ b/python/tvm_tirx/testing/__init__.py
@@ -25,5 +25,5 @@ from ._ffi_api import (
test_raise_error,
)
from .locking import run_with_gpu_lock
-from .utils import *
+from .utils import * # noqa: F403 - Re-export the public namespace.
from . import env
diff --git a/python/tvm_tirx/testing/bench.py b/python/tvm_tirx/testing/bench.py
index e8519830e..bf5636d5f 100644
--- a/python/tvm_tirx/testing/bench.py
+++ b/python/tvm_tirx/testing/bench.py
@@ -927,7 +927,7 @@ class CudaProfiler:
profiler_buffer: T.Tensor,
write_stride: int,
num_groups: int,
- default_leader: None | tvm_tirx.tirx.Expr | bool = None,
+ default_leader: tvm_tirx.tirx.Expr | bool | None = None,
profiler_enabled: bool | tvm_tirx.tirx.Expr = True,
):
self.buffer = profiler_buffer
@@ -951,7 +951,7 @@ class CudaProfiler:
name="profiler_write_offset",
)
- def _leader(self, leader: None | tvm_tirx.tirx.Expr | bool):
+ def _leader(self, leader: tvm_tirx.tirx.Expr | bool | None):
if leader is not None:
if isinstance(leader, bool | np.bool_):
return T.bool(bool(leader))
@@ -972,7 +972,7 @@ class CudaProfiler:
)
@T.inline
- def start(self, event_type: Enum, leader: None | tvm_tirx.tirx.Expr | bool
= None):
+ def start(self, event_type: Enum, leader: tvm_tirx.tirx.Expr | bool | None
= None):
if self.profiler_enabled:
T.cuda.timer_start(
event_type,
@@ -984,7 +984,7 @@ class CudaProfiler:
)
@T.inline
- def end(self, event_type: Enum, leader: None | tvm_tirx.tirx.Expr | bool =
None):
+ def end(self, event_type: Enum, leader: tvm_tirx.tirx.Expr | bool | None =
None):
if self.profiler_enabled:
T.cuda.timer_end(
event_type,
@@ -996,7 +996,7 @@ class CudaProfiler:
)
@T.inline
- def finalize(self, leader: None | tvm_tirx.tirx.Expr | bool = None):
+ def finalize(self, leader: tvm_tirx.tirx.Expr | bool | None = None):
if self.profiler_enabled:
T.cuda.timer_finalize(
self.buffer.data,
diff --git a/python/tvm_tirx/tirx/__init__.py b/python/tvm_tirx/tirx/__init__.py
index 673846977..3975b1122 100644
--- a/python/tvm_tirx/tirx/__init__.py
+++ b/python/tvm_tirx/tirx/__init__.py
@@ -15,6 +15,8 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+# ruff: noqa: E402
+# Namespace registration must precede dependent imports to break
initialization cycles.
# pylint: disable=unused-import, redefined-builtin
"""Tensor-level IR nodes, layouts, operators, analysis, and compilation."""
diff --git a/python/tvm_tirx/tirx/analysis/__init__.py
b/python/tvm_tirx/tirx/analysis/__init__.py
index c750c3c75..e77f03c4b 100644
--- a/python/tvm_tirx/tirx/analysis/__init__.py
+++ b/python/tvm_tirx/tirx/analysis/__init__.py
@@ -17,4 +17,4 @@
"""Analyze tensor-level IR, including memory access, control flow, and
well-formedness."""
# pylint: disable=wildcard-import, invalid-name
-from .analysis import *
+from .analysis import * # noqa: F403 - Re-export the public namespace.
diff --git a/python/tvm_tirx/tirx/expr.py b/python/tvm_tirx/tirx/expr.py
index 2106659e8..2e3c85ccc 100644
--- a/python/tvm_tirx/tirx/expr.py
+++ b/python/tvm_tirx/tirx/expr.py
@@ -28,10 +28,13 @@ For example, you can use addexp.a to get the left operand
of an Add node.
assert(y.a == x)
"""
+# Historical compatibility imports are public exports.
+# ruff: noqa: F401
+
from tvm_tirx import ir
from tvm_tirx.ir import Expr as Expr
from tvm_tirx.ir import const as const
-from tvm_tirx.ir._prim_expr_overload import ( # noqa: F401
+from tvm_tirx.ir._prim_expr_overload import (
EqualOp,
ExprOp,
NotEqualOp,
@@ -40,7 +43,7 @@ from tvm_tirx.ir._prim_expr_overload import ( # noqa: F401
from tvm_tirx.ir.prim import convert as convert
# Retain historical imports as aliases of the canonical shared definitions.
-from tvm_tirx.ir.prim.expr import ( # noqa: F401
+from tvm_tirx.ir.prim.expr import (
EQ,
GE,
GT,
diff --git a/python/tvm_tirx/tirx/script/__init__.py
b/python/tvm_tirx/tirx/script/__init__.py
index 88d0ef8ae..2a71ac8a0 100644
--- a/python/tvm_tirx/tirx/script/__init__.py
+++ b/python/tvm_tirx/tirx/script/__init__.py
@@ -72,6 +72,7 @@ def __getattr__(name: str) -> _Any:
return
getattr(_importlib.import_module("tvm_tirx.tirx.script.ir_builder"), name)
-from .jit import OptionalAnnotation as Optional
+# Resolve JIT annotations after the lazy dialect namespace is installed.
+from .jit import OptionalAnnotation as Optional # noqa: E402
globals().pop("jit", None)
diff --git a/python/tvm_tirx/tirx/script/ir_builder/__init__.py
b/python/tvm_tirx/tirx/script/ir_builder/__init__.py
index 4100c142e..1f41384dc 100644
--- a/python/tvm_tirx/tirx/script/ir_builder/__init__.py
+++ b/python/tvm_tirx/tirx/script/ir_builder/__init__.py
@@ -17,17 +17,18 @@
"""Concrete TIRx construction operations over the shared native IRBuilder
stack."""
from tvm_tirx.ir.type import _init_type_api
-from tvm_tirx.script.ir_builder import *
+from tvm_tirx.script.ir_builder import * # noqa: F403 - Re-export the public
namespace.
from tvm_tirx.script.ir_builder import __all__ as _shared_exports
+from tvm_tirx.script.ir_builder import constexpr, dtype, dynamic
from . import ir as _native
-from .ir import *
+from .ir import * # noqa: F403 - Re-export the public namespace.
from .ir import __all__ as _ir_exports
from .ir import boolean as bool
-from .op import *
+from .op import * # noqa: F403 - Re-export the public namespace.
from .op import __all__ as _op_exports
from .op import _get_script_namespace as __getattr__
-from .parser_protocol import *
+from .parser_protocol import * # noqa: F403 - Re-export the public namespace.
from .parser_protocol import __all__ as _protocol_exports
from .utils import (
frame_scope as frame_scope,
diff --git a/python/tvm_tirx/tirx/script/ir_builder/external_kernel.py
b/python/tvm_tirx/tirx/script/ir_builder/external_kernel.py
index 0a4e14900..49cb4488a 100644
--- a/python/tvm_tirx/tirx/script/ir_builder/external_kernel.py
+++ b/python/tvm_tirx/tirx/script/ir_builder/external_kernel.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: E722
+
"""External kernel integration fro TIR"""
import json
@@ -161,7 +161,7 @@ class SourceKernel(BaseKernel): # pylint:
disable=too-few-public-methods
if source_path.is_file():
with open(source_path) as f:
source_code = f.read()
- except: # pylint: disable=bare-except
+ except Exception:
pass
from tvm_tirx.backend.cuda.compiler import compile_source
diff --git a/python/tvm_tirx/tirx/script/ir_builder/triton.py
b/python/tvm_tirx/tirx/script/ir_builder/triton.py
index 53cd6447d..6a749291d 100644
--- a/python/tvm_tirx/tirx/script/ir_builder/triton.py
+++ b/python/tvm_tirx/tirx/script/ir_builder/triton.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: RUF005
+
"""Triton kernel integration with TIR"""
from numbers import Integral
@@ -78,10 +78,11 @@ class TritonKernel(BaseKernel):
assert kernel_metadata.num_ctas == 1, "Cluster is not supported"
num_warps = kernel_metadata.num_warps
grid = launch_args
- launch_param_tags = ["threadIdx.x"] + ["blockIdx.x", "blockIdx.y",
"blockIdx.z"][
- : len(grid)
+ launch_param_tags = [
+ "threadIdx.x",
+ *["blockIdx.x", "blockIdx.y", "blockIdx.z"][: len(grid)],
]
- launch_args = [num_warps * 32] + list(grid)
+ launch_args = [num_warps * 32, *grid]
kernel_arg_types = []
for arg in kernel_args:
if isinstance(arg, int):
diff --git a/python/tvm_tirx/tirx/transform/__init__.py
b/python/tvm_tirx/tirx/transform/__init__.py
index 0850184e6..eef2adc5f 100644
--- a/python/tvm_tirx/tirx/transform/__init__.py
+++ b/python/tvm_tirx/tirx/transform/__init__.py
@@ -24,4 +24,4 @@ expected output and invocation.
# pylint: disable=wildcard-import, invalid-name
from .function_pass import function_pass, FunctionPass
-from .transform import *
+from .transform import * # noqa: F403 - Re-export the public namespace.
diff --git a/tests/lint/check_asf_header.py b/tests/lint/check_asf_header.py
index 86301e250..06a4a2109 100644
--- a/tests/lint/check_asf_header.py
+++ b/tests/lint/check_asf_header.py
@@ -14,7 +14,6 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: E741
"""Helper tool to check and fix ASF headers in source files tracked by git.
Ported from tvm-ffi. Replaces the previous RAT-based check_asf_header.sh
@@ -268,10 +267,10 @@ def check_header(fname: str, header: str) -> bool:
has_asf_header = False
has_copyright = False
- for i, l in enumerate(lines):
- if l.find("Licensed to the Apache Software Foundation") != -1:
+ for i, line in enumerate(lines):
+ if line.find("Licensed to the Apache Software Foundation") != -1:
has_asf_header = True
- elif copyright_line(l):
+ elif copyright_line(line):
has_copyright = True
if has_asf_header and not has_copyright:
@@ -331,10 +330,10 @@ def add_header(fname: str, header: str) -> None:
has_asf_header = False
has_copyright = False
- for i, l in enumerate(lines):
- if l.find("Licensed to the Apache Software Foundation") != -1:
+ for i, line in enumerate(lines):
+ if line.find("Licensed to the Apache Software Foundation") != -1:
has_asf_header = True
- elif copyright_line(l):
+ elif copyright_line(line):
has_copyright = True
lines[i] = ""
diff --git a/tests/python/backend/cuda/codegen/test_basic.py
b/tests/python/backend/cuda/codegen/test_basic.py
index f2596f0ba..cc6517ec0 100644
--- a/tests/python/backend/cuda/codegen/test_basic.py
+++ b/tests/python/backend/cuda/codegen/test_basic.py
@@ -500,7 +500,7 @@ def test_cuda_reduction_binding(*, backend_config):
B[m_0 * 32 + m_1] = T.float32(0.0)
B[m_0 * 32 + m_1] = B[m_0 * 32 + m_1] + A[m_0 * 32 +
m_1, k]
- func = tvm_tirx.tirx.compile(Module, target={"kind": "cuda",
"backend_config": backend_config})
+ tvm_tirx.tirx.compile(Module, target={"kind": "cuda", "backend_config":
backend_config})
@pytest.mark.gpu
diff --git a/tests/python/backend/cuda/codegen/test_fp8.py
b/tests/python/backend/cuda/codegen/test_fp8.py
index f94abb2f1..59b52b60c 100644
--- a/tests/python/backend/cuda/codegen/test_fp8.py
+++ b/tests/python/backend/cuda/codegen/test_fp8.py
@@ -14,9 +14,10 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F401, F841
-from itertools import product
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
import numpy as np
import pytest
@@ -175,7 +176,6 @@ def test_fp8_vector_conversions(native_dtype,
promoted_dtype, numpytype):
mod = _create_mod(native_dtype, promoted_dtype)
target = "cuda"
fadd = tvm_tirx.tirx.compile(mod, target=target).mod
- cuda_src = fadd.imports[0].inspect_source()
dev = tvm_tirx.cuda(0)
if "x" in native_dtype:
@@ -442,7 +442,6 @@ class
TestFP8e4x4QuantDequantScale(BaseFP8E4M3QuantScaleOnly):
weight_np = np.random.uniform(-100, 100,
weight_shape).astype(model_dtype)
weight = tvm_tirx.runtime.tensor(weight_np, device=dev)
quant_weight, scales = quant(weight)
- quant_weight_np, scales_np = quant_weight.numpy(), scales.numpy()
dequant_weight = dequant(quant_weight, scales)
dequant_weight_np = dequant_weight.numpy()
@@ -492,7 +491,7 @@ def test_copy(dtype, vec_len):
B[tx, i] = A[tx, i]
mod = tvm_tirx.IRModule({"main": func})
- rtmod = tvm_tirx.tirx.compile(mod, target="cuda")
+ tvm_tirx.tirx.compile(mod, target="cuda")
num_experts = 8
diff --git a/tests/python/backend/rocm/codegen/test_basic.py
b/tests/python/backend/rocm/codegen/test_basic.py
index c05666f9e..d3e85ceb7 100644
--- a/tests/python/backend/rocm/codegen/test_basic.py
+++ b/tests/python/backend/rocm/codegen/test_basic.py
@@ -14,7 +14,10 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
# ruff: noqa: F841
+
import numpy as np
import pytest
diff --git a/tests/python/codegen/llvm/test_basic.py
b/tests/python/codegen/llvm/test_basic.py
index 533a3c555..967362f7f 100644
--- a/tests/python/codegen/llvm/test_basic.py
+++ b/tests/python/codegen/llvm/test_basic.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: E501, E731, E741, F841
+# ruff: noqa: E501, E741
from __future__ import annotations
import math
@@ -104,7 +104,7 @@ def test_llvm_intrin():
)
)
- fcode = tvm_tirx.tirx.compile(Module)
+ tvm_tirx.tirx.compile(Module)
@pytest.mark.skipif(not env.has_llvm(), reason="need llvm")
@@ -116,7 +116,7 @@ def test_llvm_void_intrin():
# Create an intrinsic that returns void.
T.call_llvm_intrin("llvm.assume", T.bool(True), ty="void")
- fcode = tvm_tirx.tirx.compile(Module)
+ tvm_tirx.tirx.compile(Module)
@pytest.mark.skipif(not env.has_llvm(), reason="need llvm")
@@ -143,7 +143,7 @@ def test_llvm_overloaded_intrin():
def main(A: T.Tensor((1, 1), "int32"), C: T.Tensor((1, 1), "int32")):
C[0, 0] = T.call_llvm_pure_intrin("llvm.ctlz", A[0, 0], int1_zero,
ty="int32")
- f = tvm_tirx.tirx.compile(Module, target="llvm")
+ tvm_tirx.tirx.compile(Module, target="llvm")
@pytest.mark.skipif(not env.has_llvm(), reason="need llvm")
@@ -155,7 +155,7 @@ def test_llvm_lookup_intrin():
A_buf = T.decl_tensor(A, ty_args=[T.Tensor((1,), "uint8x8",
data=A)])
T.evaluate(T.call_llvm_pure_intrin("llvm.ctpop.v8i8", A_buf[0],
ty="uint8x8"))
- fcode = tvm_tirx.tirx.compile(Module, None)
+ tvm_tirx.tirx.compile(Module, None)
@pytest.mark.skipif(not env.has_llvm(), reason="need llvm")
@@ -623,14 +623,22 @@ def test_llvm_div(start, end, dstart, dend, dtype,
floor_div):
_dend = tvm_tirx.tirx.const(dend, dtype)
if start == end:
- clipa = lambda x: _start
+
+ def clipa(x):
+ return _start
else:
- clipa = lambda x: T.min(_end, T.max(_start, x))
+
+ def clipa(x):
+ return T.min(_end, T.max(_start, x))
if dstart == dend:
- clipb = lambda x: _dstart
+
+ def clipb(x):
+ return _dstart
else:
- clipb = lambda x: T.min(_dend, T.max(_dstart, x))
+
+ def clipb(x):
+ return T.min(_dend, T.max(_dstart, x))
@I.ir_module
class Module:
@@ -921,7 +929,7 @@ def test_llvm_scalar_concat():
# This will crash in LLVM codegen if CodeGenLLVM::CreateVecConcat doesn't
convert
# scalars to single-lane LLVM vectors.
with tvm_tirx.ir.transform.PassContext(config={"tirx.disable_assert":
True}):
- m = tvm_tirx.tirx.compile(Module, target="llvm")
+ tvm_tirx.tirx.compile(Module, target="llvm")
@pytest.mark.skipif(not env.has_llvm(), reason="need llvm")
@@ -1024,7 +1032,7 @@ def test_llvm_assume():
B_1 = T.decl_tensor(B.data, ty_args=[T.Tensor((14,), "int32",
data=B.data)])
B_1[i] = A_1[i] * 2
- m = tvm_tirx.tirx.compile(Module, target="llvm")
+ tvm_tirx.tirx.compile(Module, target="llvm")
@pytest.mark.skipif(not env.has_llvm(), reason="need llvm")
@@ -1081,7 +1089,6 @@ def test_subroutine_call():
A = T.decl_tensor(A_data, ty_args=[T.Tensor(1, "float32",
data=A_data)])
A[0] = 42.0
- target = "llvm"
dev = tvm_tirx.cpu()
built = tvm_tirx.tirx.compile(Module)
@@ -1119,7 +1126,7 @@ def test_call_packed_returning_void():
# Error occurred during build, as part of
# CodeGenCPU::MakeCallPackedLowered.
- built = tvm_tirx.tirx.compile(Module, target="llvm")
+ tvm_tirx.tirx.compile(Module, target="llvm")
@pytest.mark.skipif(not env.has_llvm(), reason="need llvm")
@@ -1143,7 +1150,7 @@ def test_call_packed_without_string_arg():
)
with pytest.raises(RuntimeError):
- built = tvm_tirx.tirx.compile(Module, target="llvm")
+ tvm_tirx.tirx.compile(Module, target="llvm")
@pytest.mark.skipif(not env.has_llvm(), reason="need llvm")
@@ -1160,7 +1167,7 @@ def test_call_extern_returning_void():
ty="void",
)
- built = tvm_tirx.tirx.compile(Module, target="llvm")
+ tvm_tirx.tirx.compile(Module, target="llvm")
def test_invalid_volatile_masked_buffer_load():
diff --git a/tests/python/codegen/llvm/test_riscv.py
b/tests/python/codegen/llvm/test_riscv.py
index 16450fb00..0c0ed986e 100644
--- a/tests/python/codegen/llvm/test_riscv.py
+++ b/tests/python/codegen/llvm/test_riscv.py
@@ -124,7 +124,7 @@ def test_rvv_vscale_llvm_dbginfo(target):
# tvm.RuntimeError: Can't fetch the lanes of a scalable vector at a
compile time.
with tvm_tirx.target.Target(target):
- f = tvm_tirx.tirx.compile(rvv_with_vscale, target).mod
+ tvm_tirx.tirx.compile(rvv_with_vscale, target).mod
@pytest.mark.skipif(not env.has_llvm_min_version(14), reason="need llvm >= 14")
diff --git a/tests/python/codegen/test_extern.py
b/tests/python/codegen/test_extern.py
index 44d68508e..13ec10902 100644
--- a/tests/python/codegen/test_extern.py
+++ b/tests/python/codegen/test_extern.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F841
+
import numpy as np
import pytest
import tvm_ffi
@@ -29,7 +29,6 @@ from tvm_tirx.script import tirx as T
def test_add_pipeline():
"""Test extern-style add pipeline with vectorized operations."""
nn = 64
- max_threads = 4
# CPU version: serial loop with vectorized operations
@I.ir_module
diff --git a/tests/python/ir/test_attrs.py b/tests/python/ir/test_attrs.py
index 4d27e46a7..9945018d3 100644
--- a/tests/python/ir/test_attrs.py
+++ b/tests/python/ir/test_attrs.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F841
+
import pytest
import tvm_ffi
@@ -25,7 +25,7 @@ def test_dict_attrs():
dattr = tvm_tirx.ir.make_node("ir.DictAttrs", x=1, y=10, name="xyz",
padding=(0, 0))
assert dattr.x == 1
assert dattr.dict["x"] == 1
- datrr = tvm_tirx.ir.load_json(tvm_tirx.ir.save_json(dattr))
+ dattr = tvm_tirx.ir.load_json(tvm_tirx.ir.save_json(dattr))
assert dattr.name == "xyz"
assert isinstance(dattr, tvm_tirx.ir.DictAttrs)
assert "name" in dattr
@@ -159,7 +159,7 @@ def test_op_attr_map_lookup_without_value():
assert unset not in attr_map
assert attr_map.get(unset) is None
assert attr_map.get(unset, 5) == 5
- with pytest.raises(KeyError, match="test.attr_map_lookup.unset"):
+ with pytest.raises(KeyError, match=r"test\.attr_map_lookup\.unset"):
attr_map[unset]
# A non-Op expression has no attributes, as in C++ OpAttrMap::get(Expr,
default).
diff --git a/tests/python/ir/test_roundtrip_runtime_module.py
b/tests/python/ir/test_roundtrip_runtime_module.py
index 96d77438d..6f792b4d4 100644
--- a/tests/python/ir/test_roundtrip_runtime_module.py
+++ b/tests/python/ir/test_roundtrip_runtime_module.py
@@ -14,12 +14,10 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F401
+
"""Test roundtrip of runtime modules"""
# pylint: disable=missing-docstring
-import pytest
-
import tvm_tirx
import tvm_tirx.testing
diff --git a/tests/python/script/test_basic_usage.py
b/tests/python/script/test_basic_usage.py
index 515e589be..09fafc6ca 100644
--- a/tests/python/script/test_basic_usage.py
+++ b/tests/python/script/test_basic_usage.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
"""Shared parser basic usage."""
from __future__ import annotations
@@ -25,7 +28,6 @@ import pytest
import tvm_ffi
# Script-local bindings are observed through the constructed IR.
-# ruff: noqa: F841
from minilang import Value
import tvm_tirx
diff --git a/tests/python/script/test_dynamic_shape.py
b/tests/python/script/test_dynamic_shape.py
index 56aa284ac..46860cb73 100644
--- a/tests/python/script/test_dynamic_shape.py
+++ b/tests/python/script/test_dynamic_shape.py
@@ -23,7 +23,6 @@ import sys
import pytest
# Script-local bindings are observed through the constructed IR.
-# ruff: noqa: F821
from tvm_tirx.script.parser import entry
@@ -55,7 +54,7 @@ def
test_signature_read_before_introduction_remains_unbound(language):
with pytest.raises(NameError, match="name 'n' is not defined"):
@M.function
- def main(x: M.Tensor((n,), "float32")):
+ def main(x: M.Tensor((n,), "float32")): # noqa: F821 - The parser
supplies this symbolic dimension.
return x
@@ -105,7 +104,7 @@ def
test_external_symbol_does_not_introduce_same_named_python_binding(language):
@M.function
def main(x: M.Tensor((symbol, symbol))):
- M.record(n)
+ M.record(n) # noqa: F821 - The parser supplies this symbolic
dimension.
@pytest.mark.skipif(sys.version_info < (3, 12), reason="PEP 695 requires
Python 3.12")
diff --git a/tests/python/script/test_error_handling.py
b/tests/python/script/test_error_handling.py
index 5b3093845..4a110a159 100644
--- a/tests/python/script/test_error_handling.py
+++ b/tests/python/script/test_error_handling.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
"""Shared parser error handling."""
from __future__ import annotations
@@ -23,9 +26,6 @@ import inspect
import traceback
import pytest
-
-# Script-local bindings are observed through the constructed IR.
-# ruff: noqa: F841
from minilang import Value
import tvm_tirx
@@ -346,7 +346,7 @@ def
test_store_rhs_failure_precedes_target_evaluation(language):
def test_namespace_rebinding_reports_source_location(language):
# A fixed source namespace cannot become a local binding; the error must
identify that target.
M = language.M
- with pytest.raises(SyntaxError, match="namespace|shadow|rebind") as caught:
+ with pytest.raises(SyntaxError, match=r"namespace|shadow|rebind") as
caught:
@M.function
def main():
@@ -425,7 +425,7 @@ def
test_reserved_captured_builtin_reports_its_use(language, name):
@pytest.mark.parametrize("binding", ["alias, *other = M, 1, 2", "value =
(alias := M)"])
def test_local_namespace_alias_is_rejected(language, binding):
source = "@M.function\ndef main():\n " + binding
- with pytest.raises(SyntaxError, match="namespace 'M'.*alias") as caught:
+ with pytest.raises(SyntaxError, match=r"namespace 'M'.*alias") as caught:
entry.parse(source, {"M": language.M}, filename="alias.py")
assert (caught.value.filename, caught.value.lineno) == ("alias.py", 3)
assert source.splitlines()[2][caught.value.offset - 1] == "M"
@@ -444,7 +444,7 @@ def
test_host_import_cannot_replace_reserved_names(language, prefix):
if prefix.startswith("class")
else "@M.function\ndef main():\n pass"
)
- with pytest.raises(SyntaxError, match="'range'.*reserved") as caught:
+ with pytest.raises(SyntaxError, match=r"'range'.*reserved") as caught:
entry.parse(
source, {"M": language.M, "effect": lambda: effects.append(1)},
filename="import.py"
)
diff --git a/tests/python/script/test_meta_programming.py
b/tests/python/script/test_meta_programming.py
index 9ade53d8e..f0c6a68fc 100644
--- a/tests/python/script/test_meta_programming.py
+++ b/tests/python/script/test_meta_programming.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
"""Shared parser meta programming."""
from __future__ import annotations
@@ -22,7 +25,6 @@ import gc
import weakref
# Script-local bindings are observed through the constructed IR.
-# ruff: noqa: F841
from contextlib import contextmanager
from types import SimpleNamespace
diff --git a/tests/python/support/test_popen_pool.py
b/tests/python/support/test_popen_pool.py
index fba5093b0..90f614136 100644
--- a/tests/python/support/test_popen_pool.py
+++ b/tests/python/support/test_popen_pool.py
@@ -14,11 +14,10 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F401, F841
+
"""Test PopenPoolExecutor."""
import os
-import time
import psutil
import pytest
@@ -88,8 +87,6 @@ def test_popen_worker_recycles():
def test_popen_pool_executor():
import tvm_ffi
- import tvm_tirx
-
pool = PopenPoolExecutor(max_workers=2, timeout=0.015)
value1 = pool.submit(identity_after, 1, 100)
value2 = pool.submit(terminate_self)
@@ -173,7 +170,7 @@ def test_popen_pool_executor_timeout():
while not f1.done():
pass
try:
- res = f1.result()
+ f1.result()
except Exception as ex:
assert isinstance(ex, TimeoutError)
diff --git a/tests/python/sym/test_canonical_simplify.py
b/tests/python/sym/test_canonical_simplify.py
index 5a9ac36f5..ef8e600ac 100644
--- a/tests/python/sym/test_canonical_simplify.py
+++ b/tests/python/sym/test_canonical_simplify.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F841
+
import tvm_ffi
import tvm_tirx
@@ -310,7 +310,6 @@ def test_complex_cases():
def test_simplify_cast():
ck = CanonicalChecker()
tcast = tvm_tirx.tirx.Cast
- fld = tvm_tirx.tirx.floordiv
flm = tvm_tirx.tirx.floormod
# cast(i + j + 1, i64) - cast(i, i64)
i = tvm_tirx.tirx.Var("i", "int32")
diff --git a/tests/python/sym/test_const_int_bound.py
b/tests/python/sym/test_const_int_bound.py
index 6699ff824..95e375268 100644
--- a/tests/python/sym/test_const_int_bound.py
+++ b/tests/python/sym/test_const_int_bound.py
@@ -14,9 +14,10 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: RUF012
+
import contextlib
+from typing import ClassVar
import pytest
@@ -143,10 +144,10 @@ class TestBoundsUsingReciprocals(BaseCompare):
achieve its minimum while `A*B` simultaneously achieves its maximum.
"""
- A, B, C = [tvm_tirx.tirx.Var(letter, "int64") for letter in "ABC"]
+ A, B, C = (tvm_tirx.tirx.Var(letter, "int64") for letter in "ABC")
- symmetric_bounds = {A: (1, 4095), B: (1, 4095), C: (2048, 2048)}
- asymmetric_bounds = {A: (1, 1024), B: (1, POS_INF), C: (2048, 2048)}
+ symmetric_bounds: ClassVar[dict] = {A: (1, 4095), B: (1, 4095), C: (2048,
2048)}
+ asymmetric_bounds: ClassVar[dict] = {A: (1, 1024), B: (1, POS_INF), C:
(2048, 2048)}
_test_case_params = (
TestCase((A + B) * C - A * B, (2048, None),
known_bounds=symmetric_bounds),
diff --git a/tests/python/sym/test_deduce_bound.py
b/tests/python/sym/test_deduce_bound.py
index 574c21a46..f6c3bc1de 100644
--- a/tests/python/sym/test_deduce_bound.py
+++ b/tests/python/sym/test_deduce_bound.py
@@ -14,8 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F401, F841
-import pytest
+
import tvm_tirx
import tvm_tirx.testing
@@ -125,11 +124,9 @@ def test_check():
a = tvm_tirx.tirx.Var("a", "int32")
b = tvm_tirx.tirx.Var("b", "int32")
c = tvm_tirx.tirx.Var("c", "int32")
- d = tvm_tirx.tirx.Var("d", "int32")
b_s = tvm_tirx.sym.IntervalSet(2, 3)
c_s = tvm_tirx.sym.IntervalSet(5, 7)
- d_s = tvm_tirx.sym.IntervalSet(-3, -1)
# no compare operator
res1 = tvm_tirx.sym.deduce_bound(a, a + b, {b: b_s}, {})
diff --git a/tests/python/sym/test_intset.py b/tests/python/sym/test_intset.py
index 732188e59..36afd1cee 100644
--- a/tests/python/sym/test_intset.py
+++ b/tests/python/sym/test_intset.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F841
+
import tvm_tirx
import tvm_tirx.testing
from tvm_tirx import tirx
@@ -141,7 +141,7 @@ def test_max_min():
def test_select():
ck = IntSetChecker()
- x, y = tvm_tirx.tirx.Var("x", "int32"), tvm_tirx.tirx.Var("y", "int32")
+ x = tvm_tirx.tirx.Var("x", "int32")
ck.verify(
tvm_tirx.tirx.Select(x > 0, x - 1, x + 1), {x:
tvm_tirx.sym.IntervalSet(0, 10)}, (-1, 11)
)
diff --git a/tests/python/sym/test_iter_affine_map.py
b/tests/python/sym/test_iter_affine_map.py
index 8c2aecc47..e7170c58f 100644
--- a/tests/python/sym/test_iter_affine_map.py
+++ b/tests/python/sym/test_iter_affine_map.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F841
+
import pytest
import tvm_tirx
@@ -91,7 +91,7 @@ def assert_iter_map_simplify(
expect_dict, dom_map, predicate=True, check_level="surjective",
simplify_trivial_iterators=True
):
keys = list(expect_dict.keys())
- imap = tvm_tirx.sym.detect_iter_map(
+ tvm_tirx.sym.detect_iter_map(
keys,
dom_map,
predicate=predicate,
@@ -828,7 +828,6 @@ def test_subspace_divide_accepts_external_analyzer():
def test_subspace_divide_trivial_iters():
x = tvm_tirx.tirx.Var("x", "int32")
y = tvm_tirx.tirx.Var("y", "int32")
- z = tvm_tirx.tirx.Var("z", "int32")
# trivial 1.1
res = tvm_tirx.sym.subspace_divide(
@@ -1422,7 +1421,6 @@ def test_normalize_to_iter_sum():
z = tvm_tirx.tirx.Var("z", "int64")
a = tvm_tirx.tirx.Var("a", "int64")
n = tvm_tirx.tirx.Var("n", "int64")
- flm = tvm_tirx.tirx.floormod
assert_normalize_to_iter_sum(
z + ((y + x * 4 + 2) * n) + 3,
diff --git a/tests/python/sym/test_modular_set.py
b/tests/python/sym/test_modular_set.py
index fc3e71926..7e34f654a 100644
--- a/tests/python/sym/test_modular_set.py
+++ b/tests/python/sym/test_modular_set.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F841
+
import pytest
import tvm_tirx
@@ -55,7 +55,7 @@ def test_mul():
def test_shift_left():
analyzer = tvm_tirx.sym.Analyzer()
- x, y = tvm_tirx.tirx.Var("x", "int32"), tvm_tirx.tirx.Var("y", "int32")
+ x = tvm_tirx.tirx.Var("x", "int32")
m = analyzer.modular_set((x * 4 + 2) << 2)
assert m.coeff == 16
assert m.base == 8
@@ -71,7 +71,7 @@ def test_floormod():
def test_div_shift():
analyzer = tvm_tirx.sym.Analyzer()
- x, y = tvm_tirx.tirx.Var("x", "int32"), tvm_tirx.tirx.Var("y", "int32")
+ x = tvm_tirx.tirx.Var("x", "int32")
# not sure if x is non-negative
tdiv = tvm_tirx.tirx.truncdiv
m = analyzer.modular_set(tdiv(x * 4 + 2, 2))
@@ -94,7 +94,7 @@ def test_div_shift():
def test_mod():
analyzer = tvm_tirx.sym.Analyzer()
- x, y = tvm_tirx.tirx.Var("x", "int32"), tvm_tirx.tirx.Var("y", "int32")
+ x = tvm_tirx.tirx.Var("x", "int32")
tmod = tvm_tirx.tirx.truncmod
fmod = tvm_tirx.tirx.floormod
# not sure if x is non-negative
diff --git a/tests/python/sym/test_rewrite_simplify.py
b/tests/python/sym/test_rewrite_simplify.py
index 05c015c46..2d3df7309 100644
--- a/tests/python/sym/test_rewrite_simplify.py
+++ b/tests/python/sym/test_rewrite_simplify.py
@@ -14,9 +14,10 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: RUF012
+
import inspect
+from typing import ClassVar
import pytest
import tvm_ffi
@@ -776,7 +777,7 @@ class TestDivIndex(BaseCompare):
tvm_tirx.tirx.Var("y", "int32"),
tvm_tirx.tirx.Var("z", "int32"),
)
- non_negative = [x >= 0, y >= 0, z >= 0]
+ non_negative: ClassVar[list] = [x >= 0, y >= 0, z >= 0]
_test_case_params = (
TestCase(tdiv(x, x), 1),
diff --git a/tests/python/target/test_parser_mprofile.py
b/tests/python/target/test_parser_mprofile.py
index 6c8c652c4..3ee6ea65d 100644
--- a/tests/python/target/test_parser_mprofile.py
+++ b/tests/python/target/test_parser_mprofile.py
@@ -14,13 +14,13 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F821
"""
Tests to verify Python interactions with Target Parsing
"""
import pytest
+import tvm_tirx.testing
from tvm_tirx.target import Target
diff --git a/tests/python/target/test_riscv_features.py
b/tests/python/target/test_riscv_features.py
index c81ca2e4c..261379bb5 100644
--- a/tests/python/target/test_riscv_features.py
+++ b/tests/python/target/test_riscv_features.py
@@ -14,12 +14,11 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: E501, F401
+# ruff: noqa: E501
import pytest
-import tvm_tirx
from tvm_tirx import codegen
-from tvm_tirx.codegen import llvm, target_has_features
+from tvm_tirx.codegen import llvm
from tvm_tirx.target import Target
LLVM_VERSION = codegen.llvm.version_major()
diff --git a/tests/python/target/test_x86_features.py
b/tests/python/target/test_x86_features.py
index c56d96a48..42b4c21c0 100644
--- a/tests/python/target/test_x86_features.py
+++ b/tests/python/target/test_x86_features.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F841
+
import pytest
import tvm_ffi
@@ -238,7 +238,6 @@ def test_x86_target_features(min_llvm_version, tvm_target,
x86_feature, is_suppo
##
with Target(tvm_target):
- mcpu = str(Target.current(False).attrs.get("mcpu", ""))
# check for feature via the python api (current context target)
assert target_has_features(x86_feature) == is_supported
# check for feature via the python api (with explicit target)
diff --git a/tests/python/tirx/analysis/test_check_well_defined.py
b/tests/python/tirx/analysis/test_check_well_defined.py
index cc075894b..afde1066e 100644
--- a/tests/python/tirx/analysis/test_check_well_defined.py
+++ b/tests/python/tirx/analysis/test_check_well_defined.py
@@ -68,7 +68,7 @@ def test_error_for_out_of_scope_usage():
with pytest.raises(
(ValueError, RuntimeError),
- match="Invalid use of undefined variable i at .* no longer in-scope.",
+ match=r"Invalid use of undefined variable i at .* no longer
in-scope\.",
):
tvm_tirx.tirx.analysis.check_well_formed(func)
diff --git a/tests/python/tirx/base/test_index_map.py
b/tests/python/tirx/base/test_index_map.py
index bd97dc387..0aa84e81e 100644
--- a/tests/python/tirx/base/test_index_map.py
+++ b/tests/python/tirx/base/test_index_map.py
@@ -14,14 +14,14 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: E741, F401
+# ruff: noqa: E741
import numpy as np
import pytest
import tvm_ffi
import tvm_tirx
import tvm_tirx.testing
-from tvm_tirx.ir import IntImm, assert_structural_equal, const
+from tvm_tirx.ir import IntImm, assert_structural_equal
from tvm_tirx.script import tirx as T
from tvm_tirx.tirx import IndexMap, Var, floordiv, floormod
diff --git a/tests/python/tirx/base/test_intrin.py
b/tests/python/tirx/base/test_intrin.py
index b2fdf2d95..e9a4811f5 100644
--- a/tests/python/tirx/base/test_intrin.py
+++ b/tests/python/tirx/base/test_intrin.py
@@ -14,8 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: E712, F401
-import ctypes
+# ruff: noqa: E712
import math
import numpy as np
@@ -23,10 +22,8 @@ import pytest
import tvm_tirx
import tvm_tirx.testing
-from tvm_tirx import tirx
from tvm_tirx.script import ir as I
from tvm_tirx.script import tirx as T
-from tvm_tirx.support import clang, utils
def _unary_kernel(op, dtype="float32", out_dtype=None, gpu=False):
diff --git a/tests/python/tirx/base/test_nodes.py
b/tests/python/tirx/base/test_nodes.py
index 9a6e04dae..281484b7f 100644
--- a/tests/python/tirx/base/test_nodes.py
+++ b/tests/python/tirx/base/test_nodes.py
@@ -63,7 +63,7 @@ def test_tir_const_dtype_inference():
def test_make():
x = tvm_tirx.tirx.const(1, "int32")
y = tvm_tirx.tirx.Var("x", "int32")
- z = x + y
+ assert isinstance(x + y, tvm_tirx.tirx.Add)
assert isinstance(tvm_tirx.tirx.max(x, y), tvm_tirx.tirx.Max)
assert isinstance(tvm_tirx.tirx.min(x, y), tvm_tirx.tirx.Min)
@@ -94,8 +94,9 @@ def test_ir2():
def test_let():
x = tvm_tirx.tirx.Var("x", "int32")
- y = tvm_tirx.tirx.Var("y", "int32")
stmt = tvm_tirx.ir.Bind(x, 10)
+ assert stmt.var.same_as(x)
+ assert stmt.value == 10
def test_cast():
@@ -153,7 +154,7 @@ def test_any():
y = tvm_tirx.tirx.Var("y", "int32")
z = tvm_tirx.tirx.Var("z", "int32")
try:
- t = x or x
+ x or x
assert False
except ValueError:
pass
@@ -176,7 +177,7 @@ def test_all():
y = tvm_tirx.tirx.Var("y", "int32")
z = tvm_tirx.tirx.Var("z", "int32")
try:
- t = x and x
+ x and x
assert False
except ValueError:
pass
diff --git a/tests/python/tirx/codegen/test_ampere.py
b/tests/python/tirx/codegen/test_ampere.py
index 07c05d1a8..4e8418384 100644
--- a/tests/python/tirx/codegen/test_ampere.py
+++ b/tests/python/tirx/codegen/test_ampere.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
# pylint: disable=missing-function-docstring
"""Codegen tests for Ampere (sm_80) warp-level ``mma.sync`` tensor cores.
diff --git a/tests/python/tirx/codegen/test_blackwell.py
b/tests/python/tirx/codegen/test_blackwell.py
index 9f4c302c8..d9b3a0395 100644
--- a/tests/python/tirx/codegen/test_blackwell.py
+++ b/tests/python/tirx/codegen/test_blackwell.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
# pylint: disable=missing-function-docstring
import math
diff --git a/tests/python/tirx/codegen/test_cuda.py
b/tests/python/tirx/codegen/test_cuda.py
index 652afbe8e..751a8ff6a 100644
--- a/tests/python/tirx/codegen/test_cuda.py
+++ b/tests/python/tirx/codegen/test_cuda.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
# pylint: disable=missing-function-docstring
import gc
import re
@@ -932,7 +935,7 @@ def test_cuda_thread_fence():
if tx == 0:
T.cuda.thread_fence()
- src, mod = _get_source(main)
+ src, _mod = _get_source(main)
assert "tvm_builtin_cuda_thread_fence" in src
@@ -945,7 +948,7 @@ def test_cuda_nano_sleep():
if tx == 0:
T.cuda.nano_sleep(1)
- src, mod = _get_source(main)
+ src, _mod = _get_source(main)
assert "tvm_builtin_cuda_nano_sleep" in src
@@ -958,7 +961,7 @@ def test_cuda_atomic_cas():
if tx == 0:
T.cuda.atomic_cas(A.data, T.int32(1), T.int32(2))
- src, mod = _get_source(main)
+ src, _mod = _get_source(main)
assert "tvm_builtin_cuda_atomic_cas" in src
@@ -1190,7 +1193,7 @@ def test_ptx_ldmatrix(trans, num):
B[row + tx // 4, col + tx % 4 * 2 + i % 2] = A_local[i]
# fmt: on
- src, mod = _get_source(main, target="cuda")
+ _src, mod = _get_source(main, target="cuda")
A_np = np.arange(16 * 16, dtype="float16").reshape((16, 16))
B_np = np.zeros((16, 16), dtype="float16")
B_ref = np.zeros((16, 16), dtype="float16")
diff --git a/tests/python/tirx/codegen/test_cuda_cta_reduce.py
b/tests/python/tirx/codegen/test_cuda_cta_reduce.py
index 9b0f71a7d..a018c4f0d 100644
--- a/tests/python/tirx/codegen/test_cuda_cta_reduce.py
+++ b/tests/python/tirx/codegen/test_cuda_cta_reduce.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
"""Tests for T.cuda.cta_reduce / cta_sum / cta_max / cta_min intrinsics."""
import numpy as np
diff --git a/tests/python/tirx/codegen/test_cuda_warp_reduce.py
b/tests/python/tirx/codegen/test_cuda_warp_reduce.py
index 6402bcd98..da27a11b0 100644
--- a/tests/python/tirx/codegen/test_cuda_warp_reduce.py
+++ b/tests/python/tirx/codegen/test_cuda_warp_reduce.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
"""Tests for T.cuda.warp_reduce / warp_sum / warp_max / warp_min intrinsics."""
import numpy as np
diff --git a/tests/python/tirx/codegen/test_dsmem.py
b/tests/python/tirx/codegen/test_dsmem.py
index 79119c653..f226682ba 100644
--- a/tests/python/tirx/codegen/test_dsmem.py
+++ b/tests/python/tirx/codegen/test_dsmem.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
# pylint: disable=missing-function-docstring
"""Tests for cp.async.bulk.shared::cluster.shared::cta PTX instruction
codegen."""
diff --git a/tests/python/tirx/codegen/test_hopper.py
b/tests/python/tirx/codegen/test_hopper.py
index 34426a662..b4685e6ac 100644
--- a/tests/python/tirx/codegen/test_hopper.py
+++ b/tests/python/tirx/codegen/test_hopper.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
# pylint: disable=missing-function-docstring
import math
@@ -72,7 +75,7 @@ def test_ptx_setmaxnreg(inc):
T.ptx[f"setmaxnreg.{'inc' if T.constexpr(inc) else
'dec'}.sync.aligned.u32"](32)
# fmt: on
- src, mod = _get_source(func)
+ src, _mod = _get_source(func)
assert "setmaxnreg" in src
if inc:
assert "inc" in src
@@ -302,7 +305,7 @@ def test_bar_arrive():
T.ptx.bar.arrive(0, 128)
# fmt: on
- src, mod = _get_source(func)
+ src, _mod = _get_source(func)
assert "tvm_builtin_ptx_bar_arrive_arrive((uint)0, (uint)128)" in src
assert 'bar.arrive %0, %1;" : : "r"(__a), "r"(__b) : "memory"' in src
@@ -319,7 +322,7 @@ def test_bar_sync():
T.ptx.bar.sync(0, 128)
# fmt: on
- src, mod = _get_source(func)
+ src, _mod = _get_source(func)
assert "tvm_builtin_ptx_bar_sync_count_sync((uint)0, (uint)128)" in src
assert 'bar.sync %0, %1;" : : "r"(__a), "r"(__b) : "memory"' in src
@@ -336,7 +339,7 @@ def test_barrier_sync_unaligned():
T.ptx.barrier.sync(0, 128)
# fmt: on
- src, mod = _get_source(func)
+ src, _mod = _get_source(func)
assert "tvm_builtin_ptx_barrier_sync_count_sync((uint)0, (uint)128)" in src
assert 'barrier.sync %0, %1;" : : "r"(__a), "r"(__b) : "memory"' in src
@@ -353,7 +356,7 @@ def test_fence_mbarrier_init_release_clsuter():
T.ptx.fence.mbarrier_init.release.cluster()
# fmt: on
- src, mod = _get_source(func)
+ src, _mod = _get_source(func)
assert "fence.mbarrier_init.release.cluster" in src
@@ -370,7 +373,7 @@ def test_ptx_elect_sync():
A[tx] = tx
# fmt: on
- src, mod = _get_source(func)
+ src, _mod = _get_source(func)
print(src)
assert "elect.sync _|%%px, %1;" in src
assert ".reg .b32 %%rx;" not in src
@@ -390,7 +393,7 @@ def test_ptx_fence(sem, scope):
T.ptx[f"fence.{sem}.{scope}"]()
# fmt: on
- src, mod = _get_source(func)
+ src, _mod = _get_source(func)
assert f"fence.{sem}.{scope};" in src
@@ -408,7 +411,7 @@ def test_fence_proxy_async():
# fmt: on
- src, mod = _get_source(func)
+ src, _mod = _get_source(func)
assert "fence.proxy.async.global" in src
assert "fence.proxy.async.shared::cta" in src
@@ -1410,7 +1413,7 @@ def test_mapa():
if cbx == 0 and tx == 0:
T.ptx.mapa.u64(mapped[0], A_smem.data, T.uint32(cbx))
- src, mod = _get_source(func)
+ src, _mod = _get_source(func)
print(src)
assert "tvm_builtin_ptx_mapa_u64(" in src
diff --git a/tests/python/tirx/codegen/test_ptx_addr.py
b/tests/python/tirx/codegen/test_ptx_addr.py
index e4bf18de4..bafa84088 100644
--- a/tests/python/tirx/codegen/test_ptx_addr.py
+++ b/tests/python/tirx/codegen/test_ptx_addr.py
@@ -81,7 +81,7 @@ def
test_ptx_addr_table_validation_rejects_wrong_operand_classes():
(OperandSlot("x", kind="addr", space="tmem",
allow_imm_offset=True),),
),
)
- with pytest.raises(ValueError, match="reg.x.*kind='reg'"):
+ with pytest.raises(ValueError, match=r"reg\.x.*kind='reg'"):
_validate_imm_offset_slots(bad_entries)
@@ -249,7 +249,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, match="does not support T.ptx.addr"):
+ with pytest.raises(ValueError, match=r"does not support T\.ptx\.addr"):
@T.function
def ptr_operand(src: T.Tensor((8,), "uint32")):
@@ -259,7 +259,7 @@ def test_ptx_addr_pointer_and_raw_address_validation():
def test_ptx_addr_tma_tmem_and_independent_immediate_rejections():
- with pytest.raises(ValueError, match="does not support T.ptx.addr"):
+ with pytest.raises(ValueError, match=r"does not support T\.ptx\.addr"):
@T.function
def tma(tmap: T.Tensor((8,), "uint64")):
@@ -270,7 +270,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, match="does not support T.ptx.addr"):
+ with pytest.raises(ValueError, match=r"does not support T\.ptx\.addr"):
@T.function
def tmem(raw: T.uint32):
diff --git a/tests/python/tirx/codegen/test_ptx_dialect.py
b/tests/python/tirx/codegen/test_ptx_dialect.py
index 0ab3dff69..64ea253aa 100644
--- a/tests/python/tirx/codegen/test_ptx_dialect.py
+++ b/tests/python/tirx/codegen/test_ptx_dialect.py
@@ -1223,7 +1223,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, match="hi.s32"):
+ with pytest.raises(ValueError, match=r"hi\.s32"):
@T.function
def sat_on_lo(out: T.Tensor((1,), "int32")):
@@ -1319,7 +1319,7 @@ def test_ptx_floating_point_dispatch():
tvm_tirx.ir.assert_structural_equal(kernel, reparsed)
# sqrt has no f64 approximation at any spelling, unlike rcp (9.7.3.14).
- with pytest.raises(ValueError, match="only on the .f32 line"):
+ with pytest.raises(ValueError, match=r"only on the \.f32 line"):
@T.function
def sqrt_approx_f64(out: T.Tensor((1,), "float64")):
@@ -1327,7 +1327,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, match="rcp.approx.ftz.f64"):
+ with pytest.raises(ValueError, match=r"rcp\.approx\.ftz\.f64"):
@T.function
def rcp_approx_f64(out: T.Tensor((1,), "float64")):
@@ -1419,7 +1419,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, match="takes no .ftz"):
+ with pytest.raises(ValueError, match=r"takes no \.ftz"):
@T.function
def abs_bf16_ftz(out: T.Tensor((1,), "uint16")):
@@ -1470,7 +1470,7 @@ def test_ptx_mixed_precision_dispatch():
tvm_tirx.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, match="takes no .ftz"):
+ with pytest.raises(ValueError, match=r"takes no \.ftz"):
@T.function
def mixed_with_ftz(out: T.Tensor((1,), "float32")):
@@ -1478,7 +1478,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, match="only exists on the .f32"):
+ with pytest.raises(ValueError, match=r"only exists on the \.f32"):
@T.function
def mixed_f64(out: T.Tensor((1,), "float64")):
@@ -1571,7 +1571,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, match="only to .f32 comparisons"):
+ with pytest.raises(ValueError, match=r"only to \.f32 comparisons"):
@T.function
def ftz_off_f32(out: T.Tensor((1,), "uint32")):
@@ -1710,7 +1710,7 @@ def test_ptx_half_comparison_dispatch():
T.device_entry()
T.ptx.set.eq.ftz.bf16.f32(out[0], T.float32(0), T.float32(0))
- with pytest.raises(ValueError, match="spells no .ftz"):
+ with pytest.raises(ValueError, match=r"spells no \.ftz"):
@T.function
def setp_ftz_bf16(out: T.Tensor((1,), "uint32")):
@@ -2011,7 +2011,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, match="applies to .add"):
+ with pytest.raises(ValueError, match=r"applies to \.add"):
@T.function
def acc_on_min(A: T.Tensor((1,), "uint32")):
@@ -2020,7 +2020,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, match="requires .sys"):
+ with pytest.raises(ValueError, match=r"requires \.sys"):
@T.function
def mmio_gpu(A: T.Tensor((1,), "uint32")):
@@ -2388,7 +2388,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, match="mxf4.*block32"):
+ with pytest.raises(ValueError, match=r"mxf4.*block32"):
@T.function
def invalid_mxf4_block16():
@@ -3691,7 +3691,7 @@ def test_ptx_helper_source_golden():
{"op": "add", "noftz": "noftz", "vec": "v2", "type": "bf16"},
),
):
- with pytest.raises(ValueError, match="bf16 atom bit-bucket
destination.*withholds"):
+ with pytest.raises(ValueError, match=r"bf16 atom bit-bucket
destination.*withholds"):
tokens_for(TABLE[entry_name], **by_name)
# red's half-word and packed vector syntax lines share this entry; the
# conditional policy operand follows the value group.
@@ -4070,7 +4070,7 @@ def test_ptx_single_instruction_invariant():
for entry in TABLE.values():
raw = entry.raw_render is not None
for tokens, dtypes, predicated, imms, sinks in renderings(entry):
- opcode, helper, source = render_variant(entry, tokens, predicated,
dtypes, imms, sinks)
+ opcode, _helper, source = render_variant(entry, tokens,
predicated, dtypes, imms, sinks)
asm_blocks = asm_re.findall(source)
assert len(asm_blocks) == 1, f"{opcode}: {len(asm_blocks)} asm
blocks, expected 1"
if raw:
@@ -4313,7 +4313,7 @@ def test_ptxas_gate_rejects_invalid():
"}\n"
)
source = "\n".join((_CERT_PRELUDE, _certification_kernel("f", bogus, 0)))
- with pytest.raises(Exception, match="bogus|error"):
+ with pytest.raises(Exception, match=r"bogus|error"):
_assert_ptxas_ok(source)
diff --git a/tests/python/tirx/codegen/test_ptx_ld_st_ops.py
b/tests/python/tirx/codegen/test_ptx_ld_st_ops.py
index 6b3e9568e..d77288e6a 100644
--- a/tests/python/tirx/codegen/test_ptx_ld_st_ops.py
+++ b/tests/python/tirx/codegen/test_ptx_ld_st_ops.py
@@ -75,7 +75,6 @@ def _shared_scratch_copy_kernel(num_bytes: int):
"""Build shared → local scratch → shared copy kernel for ``num_bytes``
width."""
spec = _SHARED_COPY_CASES[num_bytes]
smem_dtype = spec["smem_dtype"]
- tmp_dtype = spec["tmp_dtype"]
nelems = spec["nelems"]
fill_offset = spec.get("fill_offset")
fill_value = spec.get("fill_value")
diff --git a/tests/python/tirx/iket/test_orchestration.py
b/tests/python/tirx/iket/test_orchestration.py
index 29852d9b0..62b7e27df 100644
--- a/tests/python/tirx/iket/test_orchestration.py
+++ b/tests/python/tirx/iket/test_orchestration.py
@@ -35,7 +35,7 @@ def _write_artifacts(output_dir: Path, postprocess: str, *,
count: int = 1) -> N
for index in range(count):
if postprocess in ("json", "all"):
(output_dir / f"iket_pid_{index}.trace.json").write_text(
- '{"launches": [{"index": %d}]}' % index, encoding="utf-8"
+ f'{{"launches": [{{"index": {index}}}]}}', encoding="utf-8"
)
if postprocess in ("perfetto", "html", "all"):
(output_dir / f"iket_pid_{index}.pftrace").write_bytes(b"perfetto")
diff --git a/tests/python/tirx/iket/test_profiler.py
b/tests/python/tirx/iket/test_profiler.py
index dca8cbebd..f9f37fffb 100644
--- a/tests/python/tirx/iket/test_profiler.py
+++ b/tests/python/tirx/iket/test_profiler.py
@@ -666,13 +666,13 @@ def
test_rejects_token_payload_schema_mismatch(kernel_func, message):
def test_rejects_cross_kernel_payload_schema_conflicts():
- with pytest.raises(TypeError, match="changes payload type.*across
kernels"):
+ with pytest.raises(TypeError, match=r"changes payload type.*across
kernels"):
IketProfiler().compile(
tvm_tirx.IRModule({"i32": schema_i32, "u32": schema_u32}),
target=TARGET,
tir_pipeline="tirx",
)
- with pytest.raises(ValueError, match="changes payload presence.*across
kernels"):
+ with pytest.raises(ValueError, match=r"changes payload presence.*across
kernels"):
IketProfiler().compile(
tvm_tirx.IRModule({"i32": schema_i32, "none": schema_no_payload}),
target=TARGET,
@@ -713,7 +713,7 @@ def test_extended_declaration_limit(capfd):
capfd.readouterr()
assert "__iket_evt_decl_e0000_64_attrs" in source
assert "__iket_evt_decl_e4031_4095_attrs" in source
- with pytest.raises(ValueError, match="at most 4032.*got 4033"):
+ with pytest.raises(ValueError, match=r"at most 4032.*got 4033"):
_compile(_many_marks(4033))
diff --git a/tests/python/tirx/script/test_basic_usage.py
b/tests/python/tirx/script/test_basic_usage.py
index 9afb9ba4d..cde9a26d8 100644
--- a/tests/python/tirx/script/test_basic_usage.py
+++ b/tests/python/tirx/script/test_basic_usage.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
"""TIRx script basic usage."""
from __future__ import annotations
@@ -1748,9 +1751,11 @@ def test_loop_jump_statement():
jumps = []
tvm_ffi.structural_walk(
implicit.body,
- lambda node: jumps.append(node)
- if isinstance(node, tvm_tirx.ir.Break | tvm_tirx.ir.Continue)
- else None,
+ lambda node: (
+ jumps.append(node)
+ if isinstance(node, tvm_tirx.ir.Break | tvm_tirx.ir.Continue)
+ else None
+ ),
)
assert sum(isinstance(node, tvm_tirx.ir.Break) for node in jumps) == 1
assert sum(isinstance(node, tvm_tirx.ir.Continue) for node in jumps) == 1
diff --git a/tests/python/tirx/script/test_dynamic_shape.py
b/tests/python/tirx/script/test_dynamic_shape.py
index 13c59f524..c04bd9c7b 100644
--- a/tests/python/tirx/script/test_dynamic_shape.py
+++ b/tests/python/tirx/script/test_dynamic_shape.py
@@ -127,7 +127,7 @@ def
test_tir_return_annotation_does_not_define_symbolic_var():
@pytest.mark.skipif(sys.version_info < (3, 12), reason="PEP 695 requires
Python 3.12")
def test_type_vars_roundtrip():
M = T.dynamic("M")
- UNUSED = T.dynamic("UNUSED")
+ UNUSED = T.dynamic("UNUSED") # noqa: F841 - Keep this lexical-scope
fixture for parser capture.
@T.function(private=True)
def func(A: T.Tensor((M, M * 2), "float32")):
@@ -239,7 +239,7 @@ def test_captured_shape_requires_concrete_symbols():
n = T.dynamic("n")
function = build((n, 16))
assert function.params[0].ty.shape[0].same_as(n)
- with pytest.raises(TypeError, match="Expected `Array<ir.PrimExpr>`"):
+ with pytest.raises(TypeError, match=r"Expected `Array<ir\.PrimExpr>`"):
build(("n", 16))
diff --git a/tests/python/tirx/script/test_error_handling.py
b/tests/python/tirx/script/test_error_handling.py
index 69a822cec..1f79b7e73 100644
--- a/tests/python/tirx/script/test_error_handling.py
+++ b/tests/python/tirx/script/test_error_handling.py
@@ -33,7 +33,7 @@ from tvm_tirx.script import tirx as T
def test_optional_annotation_requires_jit_at_the_source_parameter():
# Ordinary argument validation must reject Optional at its original
parameter location.
- with pytest.raises(TypeError, match="^T.Optional is only supported by
@T.jit$") as caught:
+ with pytest.raises(TypeError, match=r"^T\.Optional is only supported by
@T\.jit$") as caught:
@T.function
def invalid(value: T.Optional(T.handle)):
diff --git a/tests/python/tirx/script/test_gpu.py
b/tests/python/tirx/script/test_gpu.py
index 796ad35b9..70b4114e2 100644
--- a/tests/python/tirx/script/test_gpu.py
+++ b/tests/python/tirx/script/test_gpu.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
"""TIRx script gpu."""
import pytest
diff --git a/tests/python/tirx/script/test_ir_builder.py
b/tests/python/tirx/script/test_ir_builder.py
index aa7da77c0..2b4315396 100644
--- a/tests/python/tirx/script/test_ir_builder.py
+++ b/tests/python/tirx/script/test_ir_builder.py
@@ -151,11 +151,11 @@ def test_alloc_tensor_annotated_binding(use_let):
def test_ir_builder_tir_for():
with IRBuilder() as ib:
- with T.serial(128) as a:
- with T.parallel(64) as b:
- with T.vectorized(32) as c:
- with T.unroll(16) as d:
- with T.parallel(8, annotations={"thread_binding":
"threadIdx.x"}) as e:
+ with T.serial(128):
+ with T.parallel(64):
+ with T.vectorized(32):
+ with T.unroll(16):
+ with T.parallel(8, annotations={"thread_binding":
"threadIdx.x"}):
T.evaluate(0)
# the for generated by IRBuilder
@@ -210,7 +210,7 @@ def test_ir_builder_tir_for():
def test_ir_builder_tir_for_uint():
with IRBuilder() as ib:
- with T.serial(tirx.const(128, "uint32")) as a:
+ with T.serial(tirx.const(128, "uint32")):
T.evaluate(0)
# the for generated by IRBuilder
@@ -567,7 +567,7 @@ def test_ir_builder_loc_applies_to_emitted_stmt():
def test_ir_builder_tir_thread():
with IRBuilder() as ib:
with T.function():
- with T.launch_thread("blockIdx.y", 1) as brow:
+ with T.launch_thread("blockIdx.y", 1):
T.evaluate(0)
# the function generated by IRBuilder
@@ -592,7 +592,7 @@ def test_ir_builder_tir_allocate():
with IRBuilder() as ib:
with T.function():
T.func_name_("test")
- buf = T.bind_(T.alloc_tensor(ty_args=[T.Tensor([10], "float32",
scope="local")]))
+ T.bind_(T.alloc_tensor(ty_args=[T.Tensor([10], "float32",
scope="local")]))
T.evaluate(1)
# the allocate generated by IRBuilder
@@ -611,7 +611,7 @@ def test_ir_builder_tir_decl_tensor():
with IRBuilder() as ib:
with T.function():
T.func_name_("test")
- buf = T.bind_(T.alloc_tensor(ty_args=[T.Tensor([128, 128],
"float32")]))
+ T.bind_(T.alloc_tensor(ty_args=[T.Tensor([128, 128], "float32")]))
T.evaluate(1)
# the decl_tensor generated by IRBuilder
diff --git a/tests/python/tirx/script/test_meta_programming.py
b/tests/python/tirx/script/test_meta_programming.py
index f095c44fe..c43bbf950 100644
--- a/tests/python/tirx/script/test_meta_programming.py
+++ b/tests/python/tirx/script/test_meta_programming.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
"""TIRx script meta programming."""
from __future__ import annotations
@@ -33,7 +36,7 @@ def test_meta_class_constructor_rejects_unowned_resource():
@T.meta_class
class Bad:
def __init__(self):
- tmp = T.bind_(T.alloc_tensor(ty_args=[T.Tensor((1,), "int32",
scope="local")]))
+ T.bind_(T.alloc_tensor(ty_args=[T.Tensor((1,), "int32",
scope="local")]))
with pytest.raises(ValueError):
diff --git a/tests/python/tirx/script/test_printer.py
b/tests/python/tirx/script/test_printer.py
index 5f68b5811..70ee79488 100644
--- a/tests/python/tirx/script/test_printer.py
+++ b/tests/python/tirx/script/test_printer.py
@@ -774,11 +774,8 @@ def test_printer_nvshmem_namespace():
def test_printer_ptx_more():
- r = tir.Var("r", "handle")
- s = tir.Var("s", "handle")
d = tir.Var("d", "handle")
a = tir.Var("a", "handle")
- b = tir.Var("b", "handle")
_assert_namespace_print(
cuda_op.tcgen05_encode_matrix_descriptor(d, a, 1, 2, 0),
"""
@@ -1092,7 +1089,6 @@ T.nki.affine_select(A[0], A[0], A[0], T.float32(1.0))""",
def test_printer_ptx_mma_and_wgmma():
- r = tir.Var("r", "handle")
d = tir.Var("d", "handle")
a = tir.Var("a", "handle")
tir.Var("b", "handle")
diff --git a/tests/python/tirx/script/test_printer_function.py
b/tests/python/tirx/script/test_printer_function.py
index cae67a165..5460c00e8 100644
--- a/tests/python/tirx/script/test_printer_function.py
+++ b/tests/python/tirx/script/test_printer_function.py
@@ -14,7 +14,7 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: E501
+# ruff: noqa: E501, F841
"""TIRx function printing and roundtrip coverage."""
import sys
diff --git a/tests/python/tirx/script/test_printer_structural_equal.py
b/tests/python/tirx/script/test_printer_structural_equal.py
index f5e2ed0dc..82ff00631 100644
--- a/tests/python/tirx/script/test_printer_structural_equal.py
+++ b/tests/python/tirx/script/test_printer_structural_equal.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
"""TIRx script printer structural equal."""
import pytest
diff --git a/tests/python/tirx/script/test_source_locations.py
b/tests/python/tirx/script/test_source_locations.py
index 0f818fdfc..f8a1bd88d 100644
--- a/tests/python/tirx/script/test_source_locations.py
+++ b/tests/python/tirx/script/test_source_locations.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
"""TIRx script source locations."""
from __future__ import annotations
@@ -162,9 +165,11 @@ def test_parser_attaches_loc_to_tile_op_call():
func = tile_call
call = _find_ir_node(
func,
- lambda node: isinstance(node, tvm_tirx.ir.Call)
- and isinstance(node.op, tvm_tirx.ir.Op)
- and node.op.name == "tirx.cuda.tile.apply",
+ lambda node: (
+ isinstance(node, tvm_tirx.ir.Call)
+ and isinstance(node.op, tvm_tirx.ir.Op)
+ and node.op.name == "tirx.cuda.tile.apply"
+ ),
)
assert _loc_range(call.loc) == _loc_range(source.to_loc(call_ast))
diff --git a/tests/python/tirx/script/test_tensors.py
b/tests/python/tirx/script/test_tensors.py
index 08753fce8..083ffb0ee 100644
--- a/tests/python/tirx/script/test_tensors.py
+++ b/tests/python/tirx/script/test_tensors.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
"""TIRx script tensors."""
import math
@@ -593,7 +596,7 @@ def test_tensor_physical_view_layout_overrides_roundtrip():
B_custom[0, 0] = T.float32(2)
# fmt: on
- _, b_buf, b_storage, b_custom = _collect_tensors(func)
+ _, _b_buf, b_storage, b_custom = _collect_tensors(func)
assert_structural_equal(b_storage.ty.layout, T.TileLayout(T.S[(2, 8, 2) :
(2, 4, 1)]))
assert not b_storage.ty.layout.is_trivial()
assert not b_custom.ty.layout.is_trivial()
@@ -1165,7 +1168,6 @@ def test_tensor_sub_swizzle_commutation():
a2, b2, c2, d2, e2, _ = _collect_tensors(func2)
base2 = a2.layout.apply(0).get("m", 0)
- shape2 = [2, 16, 8]
for name, child, to_parent in [
("B", b2, lambda c: (c[0], 1, c[1])),
("C", c2, lambda c: (c[0], c[1], 1)),
diff --git a/tests/python/tirx/test_base.py b/tests/python/tirx/test_base.py
index c6135e4cc..394914672 100644
--- a/tests/python/tirx/test_base.py
+++ b/tests/python/tirx/test_base.py
@@ -219,7 +219,7 @@ def test_continue_statement():
def test_exception():
with pytest.raises(TypeError):
- x = tirx.Var(name=1, ty="int")
+ tirx.Var(name=1, ty="int")
def test_eq_ops():
diff --git a/tests/python/tirx/test_control_flow.py
b/tests/python/tirx/test_control_flow.py
index 4e1a9877c..53f4a3f34 100644
--- a/tests/python/tirx/test_control_flow.py
+++ b/tests/python/tirx/test_control_flow.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
import numpy as np
import pytest
diff --git a/tests/python/tirx/test_jit.py b/tests/python/tirx/test_jit.py
index 10275a54e..7a2336b93 100644
--- a/tests/python/tirx/test_jit.py
+++ b/tests/python/tirx/test_jit.py
@@ -14,7 +14,6 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F821
"""Tests for ``@T.jit`` compile-time specialization."""
from __future__ import annotations
@@ -35,9 +34,9 @@ from tvm_tirx.tirx.script.jit import JITFunction
def test_int_constexpr_specializes_loop_bound():
@T.jit(private=True)
def add(
- A: T.Tensor((N,), "int32"),
- B: T.Tensor((N,), "int32"),
- C: T.Tensor((N,), "int32"),
+ A: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
+ B: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
+ C: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
*,
N: T.constexpr,
):
@@ -59,9 +58,9 @@ def test_int_constexpr_specializes_loop_bound():
def test_constexpr_in_2d_tensor_shape():
@T.jit(private=True)
def matadd(
- A: T.Tensor((M, K), "int32"),
- B: T.Tensor((M, K), "int32"),
- C: T.Tensor((M, K), "int32"),
+ A: T.Tensor((M, K), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
+ B: T.Tensor((M, K), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
+ C: T.Tensor((M, K), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
*,
M: T.constexpr,
K: T.constexpr,
@@ -86,8 +85,8 @@ def test_constexpr_in_2d_tensor_shape():
def test_constexpr_in_body_expression():
@T.jit(private=True)
def scaled_copy(
- A: T.Tensor((N,), "int32"),
- B: T.Tensor((N,), "int32"),
+ A: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
+ B: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
*,
N: T.constexpr,
SCALE: T.constexpr,
@@ -109,7 +108,7 @@ def test_constexpr_in_body_expression():
def test_specialize_cache_returns_same_instance():
@T.jit(private=True)
def k(
- A: T.Tensor((N,), "int32"),
+ A: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
*,
N: T.constexpr,
):
@@ -124,7 +123,7 @@ def test_specialize_cache_returns_same_instance():
def test_specialize_different_args_produce_different_funcs():
@T.jit(private=True)
def k(
- A: T.Tensor((N,), "int32"),
+ A: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
*,
N: T.constexpr,
):
@@ -137,7 +136,7 @@ def
test_specialize_different_args_produce_different_funcs():
def test_specialize_missing_constexpr_raises():
@T.jit(private=True)
def k(
- A: T.Tensor((N,), "int32"),
+ A: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
*,
N: T.constexpr,
SCALE: T.constexpr,
@@ -152,7 +151,7 @@ def test_specialize_missing_constexpr_raises():
def test_specialize_extra_kwarg_raises():
@T.jit(private=True)
def k(
- A: T.Tensor((N,), "int32"),
+ A: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
*,
N: T.constexpr,
):
@@ -166,7 +165,7 @@ def test_specialize_extra_kwarg_raises():
def test_jit_kernel_with_nested_inline_helper():
@T.jit(private=True)
def k(
- A: T.Tensor((N,), "int32"),
+ A: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
*,
N: T.constexpr,
):
@@ -190,7 +189,7 @@ def test_jit_kernel_with_nested_inline_helper():
def test_constexpr_default_value():
@T.jit(private=True)
def k(
- A: T.Tensor((N,), "int32"),
+ A: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
*,
N: T.constexpr,
SCALE: T.constexpr = 7,
@@ -214,7 +213,7 @@ def test_constexpr_default_value():
def test_specialize_returns_function():
@T.jit(private=True)
def k(
- A: T.Tensor((N,), "int32"),
+ A: T.Tensor((N,), "int32"), # noqa: F821 - The parser supplies this
symbolic dimension.
*,
N: T.constexpr,
):
@@ -369,7 +368,7 @@ def test_standard_optional_annotations_are_specializable():
def test_t_optional_is_restricted_to_jit():
- with pytest.raises(TypeError, match="only supported by @T.jit"):
+ with pytest.raises(TypeError, match=r"only supported by @T\.jit"):
@T.function(private=True)
def invalid(a: T.Optional(T.handle)):
diff --git a/tests/python/tirx/test_ops.py b/tests/python/tirx/test_ops.py
index 88a0e7ae5..03717478a 100644
--- a/tests/python/tirx/test_ops.py
+++ b/tests/python/tirx/test_ops.py
@@ -217,9 +217,9 @@ def test_if_then_else():
def test_llvm_intrin():
with pytest.raises(ValueError, match=r"Unknown llvm intrinsic function
llvm.dummy"):
- a = tvm_tirx.tirx.call_llvm_intrin("llvm.dummy", ty="int32x4")
+ tvm_tirx.tirx.call_llvm_intrin("llvm.dummy", ty="int32x4")
with pytest.raises(ValueError, match=r"Unknown llvm intrinsic function
llvm.dummy"):
- a = tvm_tirx.tirx.call_llvm_pure_intrin("llvm.dummy", ty="int32x4")
+ tvm_tirx.tirx.call_llvm_pure_intrin("llvm.dummy", ty="int32x4")
if __name__ == "__main__":
diff --git a/tests/python/tirx/test_verifier.py
b/tests/python/tirx/test_verifier.py
index 91fb02e9c..4929c929a 100644
--- a/tests/python/tirx/test_verifier.py
+++ b/tests/python/tirx/test_verifier.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
import pytest
from tvm_tirx.script import tirx as T
@@ -175,7 +178,7 @@ def test_preferred_cluster_validation():
T.cuda.LaunchConfig(grid=8, block=128, preferred_cluster=2)
with pytest.raises(ValueError, match="divisible"):
T.cuda.LaunchConfig(grid=(8, 3), block=128, cluster=(2, 1),
preferred_cluster=(2, 2))
- with pytest.raises(ValueError, match="cluster.z"):
+ with pytest.raises(ValueError, match=r"cluster\.z"):
T.cuda.LaunchConfig(grid=(8, 4, 2), block=128, cluster=(2, 1),
preferred_cluster=(2, 2, 2))
# fmt: on
diff --git a/tests/python/tirx/tile/cuda/test_add_sub_mul_div.py
b/tests/python/tirx/tile/cuda/test_add_sub_mul_div.py
index cad2a0562..3b32c74f5 100644
--- a/tests/python/tirx/tile/cuda/test_add_sub_mul_div.py
+++ b/tests/python/tirx/tile/cuda/test_add_sub_mul_div.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
"""Native arithmetic instructions, broadcasts, and caller-composed half
division."""
import math
@@ -108,7 +111,7 @@ def _reference(operation, a, b):
@pytest.mark.parametrize("operands_type", ["region_region", "region_const",
"const_region"])
@pytest.mark.parametrize("dtype", ["float16"])
def test_binary_op_shared(input, op_type, operands_type, dtype):
- g_shape, st_a, st_b, st_res, ext_a, ext_b, ext_res, thread_cnt, device_id
= input
+ g_shape, st_a, st_b, st_res, ext_a, ext_b, ext_res, thread_cnt, _device_id
= input
size, count = math.prod(g_shape), math.prod(ext_a)
layout = TileLayout(S[g_shape])
@@ -434,7 +437,7 @@ def test_binary_op_local_subcta_trivial(thread_scope,
rhs_kind, op_type):
@pytest.mark.parametrize("op_type", ["add", "sub", "mul", "fdiv"])
@pytest.mark.parametrize("dtype", ["float16"])
def test_binary_op_vectorized(input, storage_scope, thread_scope, op_type,
dtype):
- a_shape, b_shape, res_shape, thread_cnt, device_id = input
+ a_shape, b_shape, _res_shape, thread_cnt, _device_id = input
row_a, row_b = a_shape[1:], b_shape[1:]
n, m = math.prod(row_a), math.prod(row_b)
# Half pairs are adjacent, including in the m-axis broadcast.
diff --git a/tests/python/tirx/tile/cuda/test_cp_async.py
b/tests/python/tirx/tile/cuda/test_cp_async.py
index b1e0f4ba4..9e61bc9ea 100644
--- a/tests/python/tirx/tile/cuda/test_cp_async.py
+++ b/tests/python/tirx/tile/cuda/test_cp_async.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
# pylint: disable=invalid-name, missing-function-docstring
"""Tests for the non-bulk CTA-level copy_async dispatch (vectorized load)."""
diff --git a/tests/python/tirx/tile/cuda/test_cp_async_bulk_tensor.py
b/tests/python/tirx/tile/cuda/test_cp_async_bulk_tensor.py
index 4ff6c2464..265c6acbb 100644
--- a/tests/python/tirx/tile/cuda/test_cp_async_bulk_tensor.py
+++ b/tests/python/tirx/tile/cuda/test_cp_async_bulk_tensor.py
@@ -2059,7 +2059,7 @@ def test_descriptor_fill_rejects_invalid_format(case_id):
box = (1, *kwargs["s_shape"]) if rank == 3 else kwargs["s_shape"]
opcode = _store(rank) if kwargs.get("direction") == "s2g" else _load(rank)
fill = "bogus" if case_id == "reject-unknown-oob" else 1
- with pytest.raises(RuntimeError, match="invalid literal|floating-point"):
+ with pytest.raises(RuntimeError, match=r"invalid literal|floating-point"):
_direct(
opcode,
g_shape=kwargs["g_shape"],
@@ -2570,5 +2570,7 @@ def
test_native_compiler_rejects_unsupported_tma_target(arch, gather):
# ISA target requirements belong to the PTX compiler, not another tile
# registry's architecture policy. A wrong target must still fail to build.
- with pytest.raises(Exception, match="(?i)(requires|require|unsupported|not
supported|feature)"):
+ with pytest.raises(
+ Exception, match=r"(?i)(requires|require|unsupported|not
supported|feature)"
+ ):
_compile_module(kernel, arch=arch)
diff --git a/tests/python/tirx/tile/cuda/test_cvt.py
b/tests/python/tirx/tile/cuda/test_cvt.py
index 209a9f56c..1d55b5395 100644
--- a/tests/python/tirx/tile/cuda/test_cvt.py
+++ b/tests/python/tirx/tile/cuda/test_cvt.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
"""Native cvt instructions over explicit register views."""
import math
@@ -437,7 +440,7 @@ def test_cast_local_view_sliced(A_dtype, B_dtype,
slice_start, slice_end):
def test_cast_mixed_axes_and_subregion(slice_start, slice_end):
"""Test cast with mixed axes and subregion."""
- N_WARPS, LANES = 2, 32
+ N_WARPS = 2
LOCAL_LEN = 4
full_shape = (8, N_WARPS, 4, LOCAL_LEN)
g_layout = TileLayout(S[full_shape])
@@ -555,7 +558,7 @@ def test_cast_validate_extent_mismatch_rejected():
# span the full CTA), which fires first — either rejection is fine.
with pytest.raises(
Exception,
- match="incorrect .* owner",
+ match=r"incorrect .* owner",
):
tvm_tirx.tirx.compile(mod, target=target, pipeline="tirx")
diff --git a/tests/python/tirx/tile/cuda/test_ewise.py
b/tests/python/tirx/tile/cuda/test_ewise.py
index fc5ad16f0..6981a687d 100644
--- a/tests/python/tirx/tile/cuda/test_ewise.py
+++ b/tests/python/tirx/tile/cuda/test_ewise.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
"""Caller-defined scalar compositions and explicit memory staging."""
import math
@@ -85,7 +88,7 @@ def _run(kernel, arrays, expected, output=0, atol=1e-3,
rtol=1e-7, target="cuda"
def _shared_case(input, operation, src_dtype, dst_dtype, bias_type=None,
target="cuda"):
- shape, start, end, extent, result_extent, threads, device = input
+ shape, start, end, extent, result_extent, threads, _device = input
size, count = math.prod(shape), math.prod(extent)
in_place = src_dtype == dst_dtype
bias_constant = 0.88 if bias_type == "const" else 0
@@ -295,7 +298,7 @@ def test_unary_op_shared_with_bias_scale(input, op_type,
bias_type, src_dtype, d
"src_dtype,dst_dtype", [("float16", "float16"), ("float32", "float16"),
("float32", "bfloat16")]
)
def test_unary_op_local(input, op_type, src_dtype, dst_dtype):
- layout, N_GROUPS, N_WARPS, thread_cnt, device_id = input
+ layout, N_GROUPS, N_WARPS, _thread_cnt, device_id = input
assert layout == "wgmma", "logical tensor which is not WGMMA layout is not
supported"
# get shape info
@@ -436,7 +439,7 @@ def test_unary_op_local(input, op_type, src_dtype,
dst_dtype):
"src_dtype,dst_dtype", [("float32", "float32"), ("float32", "float16"),
("float32", "bfloat16")]
)
def test_unary_op_local_with_bias_scale(input, op_type, bias_type, src_dtype,
dst_dtype):
- layout, N_GROUPS, N_WARPS, thread_cnt, device_id = input
+ layout, N_GROUPS, N_WARPS, _thread_cnt, device_id = input
assert layout == "wgmma", "logical tensor which is not WGMMA layout is not
supported"
# get shape info
diff --git a/tests/python/tirx/tile/cuda/test_explicit_patterns.py
b/tests/python/tirx/tile/cuda/test_explicit_patterns.py
index a24368239..3a8b72d36 100644
--- a/tests/python/tirx/tile/cuda/test_explicit_patterns.py
+++ b/tests/python/tirx/tile/cuda/test_explicit_patterns.py
@@ -332,7 +332,7 @@ def test_cluster_pattern_requires_launch_cluster():
T.cuda.tile.apply(T.ptx[LD], r, A, scope="cluster")
with tvm_tirx.target.Target({"kind": "cuda", "arch": "sm_100a"}):
- with pytest.raises(RuntimeError, match="explicit
LaunchConfig.cluster"):
+ with pytest.raises(RuntimeError, match=r"explicit
LaunchConfig\.cluster"):
tvm_tirx.tirx.transform.TileDispatch()(tvm_tirx.IRModule({"main":
kernel}))
@@ -390,5 +390,5 @@ def test_cuda_index_region_bounds(index, block, size,
offset):
if offset == 0:
tvm_tirx.tirx.transform.TileDispatch()(tvm_tirx.IRModule({"main":
kernel}))
else:
- with pytest.raises(RuntimeError, match="in bounds|exceeds tensor
shape"):
+ with pytest.raises(RuntimeError, match=r"in bounds|exceeds tensor
shape"):
tvm_tirx.tirx.transform.TileDispatch()(tvm_tirx.IRModule({"main": kernel}))
diff --git a/tests/python/tirx/tile/cuda/test_fma.py
b/tests/python/tirx/tile/cuda/test_fma.py
index 98b2ca00d..45a6bfe81 100644
--- a/tests/python/tirx/tile/cuda/test_fma.py
+++ b/tests/python/tirx/tile/cuda/test_fma.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
"""Native fma instructions with scalar broadcasts and explicit register
views."""
import re
diff --git a/tests/python/tirx/tile/cuda/test_ld_st_cache_policy.py
b/tests/python/tirx/tile/cuda/test_ld_st_cache_policy.py
index d800fc49e..7626aa6b1 100644
--- a/tests/python/tirx/tile/cuda/test_ld_st_cache_policy.py
+++ b/tests/python/tirx/tile/cuda/test_ld_st_cache_policy.py
@@ -138,7 +138,7 @@ def test_copy_vec_nc_rejects_non_global_src():
target = tvm_tirx.target.Target("cuda")
with target:
mod = tvm_tirx.IRModule({"main": kernel})
- with pytest.raises(RuntimeError, match="expected storage
scope.*global"):
+ with pytest.raises(RuntimeError, match=r"expected storage
scope.*global"):
tvm_tirx.tirx.compile(mod, target=target, pipeline="tirx")
diff --git a/tests/python/tirx/tile/cuda/test_ld_st_gmem_smem.py
b/tests/python/tirx/tile/cuda/test_ld_st_gmem_smem.py
index 695fe6c1f..d129dcdd7 100644
--- a/tests/python/tirx/tile/cuda/test_ld_st_gmem_smem.py
+++ b/tests/python/tirx/tile/cuda/test_ld_st_gmem_smem.py
@@ -298,7 +298,7 @@ def _reject_explicit_staging(g_shape, s_shape, g_layout,
s_layout, start, word_d
g_shape, s_shape, g_layout, s_layout, start, word_dtype, words
)
target = tvm_tirx.target.Target("cuda")
- with target, pytest.raises(RuntimeError, match="consecutive
elements|byte.aligned"):
+ with target, pytest.raises(RuntimeError, match=r"consecutive
elements|byte[- ]aligned"):
tvm_tirx.tirx.transform.TileDispatch()(tvm_tirx.IRModule({"main":
kernel}))
diff --git a/tests/python/tirx/tile/cuda/test_ldmatrix_stmatrix.py
b/tests/python/tirx/tile/cuda/test_ldmatrix_stmatrix.py
index 81de1fce9..fd25ca8ab 100644
--- a/tests/python/tirx/tile/cuda/test_ldmatrix_stmatrix.py
+++ b/tests/python/tirx/tile/cuda/test_ldmatrix_stmatrix.py
@@ -102,7 +102,6 @@ def _build_warp_kernel(num, direction, trans,
swizzle=False):
r_layout = _r_layout_warp(num)
s_layout = _maybe_wrap_swizzle(_s_layout_warp(num, trans), swizzle)
s_shape = (8, 4, num, 2)
- full = (slice(0, 8), slice(0, 4), slice(0, num), slice(0, 2))
M, N = 8, num * 8
def _coord(row, cp, t, w):
@@ -192,7 +191,6 @@ def _build_warpgroup_kernel(num, direction, trans,
swizzle=False):
r_layout = _r_layout_warpgroup(num)
s_layout = _maybe_wrap_swizzle(_s_layout_warpgroup_or_cta(num, trans),
swizzle)
s_shape = (4, 8, 4, num, 2)
- full = (slice(0, 4), slice(0, 8), slice(0, 4), slice(0, num), slice(0, 2))
M, N = 32, num * 8
def _coord(wid, row, cp, t, w):
@@ -285,7 +283,6 @@ def _build_cta_kernel(num, direction, trans, swizzle=False):
r_layout = _r_layout_cta(num)
s_layout = _maybe_wrap_swizzle(_s_layout_warpgroup_or_cta(num, trans),
swizzle)
s_shape = (4, 8, 4, num, 2)
- full = (slice(0, 4), slice(0, 8), slice(0, 4), slice(0, num), slice(0, 2))
M, N = 32, num * 8
def _coord(wid, row, cp, t, w):
diff --git a/tests/python/tirx/tile/cuda/test_mma_sync.py
b/tests/python/tirx/tile/cuda/test_mma_sync.py
index 7b85c5084..160cdb68d 100644
--- a/tests/python/tirx/tile/cuda/test_mma_sync.py
+++ b/tests/python/tirx/tile/cuda/test_mma_sync.py
@@ -810,7 +810,7 @@ def test_cuda_gemm_mma_rejects_incompatible_instruction(a,
b, c, d):
"""
with pytest.raises(
(RuntimeError, ValueError, KeyError, AssertionError),
- match="mma|ptx instruction|packed into",
+ match=r"mma|ptx instruction|packed into",
):
_lower(_build_dtypes(a, b, c, d))
diff --git a/tests/python/tirx/tile/cuda/test_tcgen05_cp.py
b/tests/python/tirx/tile/cuda/test_tcgen05_cp.py
index 2513fb5cc..18909bd17 100644
--- a/tests/python/tirx/tile/cuda/test_tcgen05_cp.py
+++ b/tests/python/tirx/tile/cuda/test_tcgen05_cp.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
# pylint: disable=invalid-name, missing-function-docstring
"""Native tcgen05.cp layouts, regions, multicast and CTA-pair readbacks.
@@ -152,8 +155,6 @@ def _make_cp_kernel(
extra_cfg=None,
pre_zero=False,
):
- s_sl = tuple(slice(a, b) for a, b in s_region)
- t_sl = tuple(slice(a, b) for a, b in t_region)
s_full_sl = tuple(slice(0, e) for e in s_full_shape)
cfg = {"shape": shape, "multicast": multicast}
if extra_cfg:
@@ -488,7 +489,6 @@ def _make_cp_kernel_cta2(s_full, s_shape, t_full, t_shape,
dtype, cfg, W32, n_co
only the even CTA issues the cp (cta_group=2, commit cta_mask=3); both CTAs
dump their own tmem into B[cbx*128 + lane]."""
s_sl = tuple(slice(0, e) for e in s_shape)
- t_sl = tuple(slice(0, e) for e in t_shape)
shape = cfg["shape"]
opcode = _cp_opcode(shape, cfg.get("multicast"), cfg["cta_group"],
cfg.get("decompress"))
@@ -762,7 +762,7 @@ def test_cp_rejects_wrong_replica_for_multicast():
def test_cp_rejects_illegal_shape_multicast_combo():
- with pytest.raises((ValueError, RuntimeError), match="takes no
multicast|Invalid PTX"):
+ with pytest.raises((ValueError, RuntimeError), match=r"takes no
multicast|Invalid PTX"):
_build_case("128x256b", None, 3, "bfloat16", 1,
extra_cfg={"multicast": "warpx4"})
@@ -771,8 +771,8 @@ def test_cp_rejects_missing_multicast_for_64x128b():
C = 128 // bits
t_full, t_full_shape = _tmem_layout_for("64x128b", "warpx2::02_13", C)
s_full = mma_shared_layout("bfloat16", SwizzleMode(3), [64, 64])
- with pytest.raises((ValueError, RuntimeError), match="requires
multicast|Invalid PTX"):
- kernel = _make_cp_kernel(
+ with pytest.raises((ValueError, RuntimeError), match=r"requires
multicast|Invalid PTX"):
+ _make_cp_kernel(
s_full,
[64, 64],
[(0, 64), (0, C)],
@@ -812,7 +812,7 @@ def test_cp_rejects_non_128b_tcol_offset():
def test_cp_rejects_unknown_shape():
- with pytest.raises((KeyError, ValueError, RuntimeError),
match="96x128b|Invalid PTX"):
+ with pytest.raises((KeyError, ValueError, RuntimeError),
match=r"96x128b|Invalid PTX"):
_build_case("32x128b", "warpx4", 3, "bfloat16", 1, extra_cfg={"shape":
"96x128b"})
diff --git a/tests/python/tirx/tile/cuda/test_tcgen05_ld_st.py
b/tests/python/tirx/tile/cuda/test_tcgen05_ld_st.py
index e672b0dd2..e08b99017 100644
--- a/tests/python/tirx/tile/cuda/test_tcgen05_ld_st.py
+++ b/tests/python/tirx/tile/cuda/test_tcgen05_ld_st.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
# pylint: disable=invalid-name, missing-function-docstring
"""Tests for the tcgen05 TMEM load/store (ld/st) dispatch.
@@ -931,7 +934,6 @@ def _run_load_test(shape: str, rep: int, dtype: str):
# configs whose K_cols_fp32 exceeds 128 we split the stage into multiple
# chunks of CHUNK_FP32 fp32 cols each.
CHUNK_FP32 = 128
- chunk_elem = CHUNK_FP32 * elem_per_32b
num_chunks = tmem_col_width_32b // CHUNK_FP32 if tmem_col_width_32b >
CHUNK_FP32 else 1
chunk_width_32b = tmem_col_width_32b if num_chunks == 1 else CHUNK_FP32
chunk_width_elem = chunk_width_32b * elem_per_32b
@@ -942,9 +944,6 @@ def _run_load_test(shape: str, rep: int, dtype: str):
if stage_width_elem % VEC_LEN != 0:
pytest.skip(f"stage_width_elem {stage_width_elem} % VEC_LEN {VEC_LEN}
!= 0")
- g_layout = TileLayout(
- S[(128, stage_width_elem // VEC_LEN, VEC_LEN) : (stage_width_elem,
VEC_LEN, 1)]
- )
chunk_view = TileLayout(S[(128, chunk_width_elem) : (1 @ axis_tid_in_wg,
1)])
# Preserve the original register allocation; _atom_words only reorganizes
# its logical axes for the instruction explicitly selected below.
diff --git a/tests/python/tirx/tile/cuda/test_tcgen05_mma.py
b/tests/python/tirx/tile/cuda/test_tcgen05_mma.py
index 61648c0e8..d15680e1a 100644
--- a/tests/python/tirx/tile/cuda/test_tcgen05_mma.py
+++ b/tests/python/tirx/tile/cuda/test_tcgen05_mma.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
"""Native tcgen05.mma instructions with explicit operand views."""
import functools
@@ -478,8 +481,6 @@ def test_gemm_tcgen05_cta_group_2(task):
A_layout = mma_shared_layout(A_dtype, A_swizzle_mode, A_shape_per_cta)
B_layout = mma_shared_layout(B_dtype, B_swizzle_mode, B_shape_per_cta)
- r_smem_A_in = list(slice(0, A_shape_per_cta[i]) for i in
range(len(A_shape_per_cta)))
- r_smem_B_in = list(slice(0, B_shape_per_cta[i]) for i in
range(len(B_shape_per_cta)))
total_bytes = (
functools.reduce(operator.mul, A_shape, 1) * A_elem_bytes
+ functools.reduce(operator.mul, B_shape, 1) * B_elem_bytes
@@ -1114,11 +1115,11 @@ def test_gemm_block_scaled_fp8_cta_group_1(task):
(A_shape, A_dtype, A_region, A_swizzle_mode),
(B_shape, B_dtype, B_region, B_swizzle_mode),
SF_dtype,
- transA,
- transB,
+ _transA,
+ _transB,
) = task
- M, K = A_shape
+ M, _K = A_shape
N = B_shape[0]
width = C_region[1][1] - C_region[1][0]
assert C_shape[0] == 128
@@ -1457,8 +1458,6 @@ def test_gemm_block_scaled_fp8_cta_group_2(task):
A_layout = mma_shared_layout(A_dtype, A_swizzle_mode, A_shape_per_cta)
B_layout = mma_shared_layout(B_dtype, B_swizzle_mode, B_shape_per_cta)
- r_smem_A_in = list(slice(0, A_shape_per_cta[i]) for i in
range(len(A_shape_per_cta)))
- r_smem_B_in = list(slice(0, B_shape_per_cta[i]) for i in
range(len(B_shape_per_cta)))
total_bytes = (
functools.reduce(operator.mul, A_shape, 1) * A_elem_bytes
+ functools.reduce(operator.mul, B_shape, 1) * B_elem_bytes
@@ -3062,8 +3061,6 @@ def test_gemm_tcgen05_arbitrary_tiles(task, copy_perm,
a_major, b_major):
else mma_shared_layout(B_dtype, B_swizzle_mode, B_shape)
)
- r_gmem_A = list(slice(0, A_shape[i]) for i in range(len(A_shape)))
- r_gmem_B = list(slice(0, B_shape[i]) for i in range(len(B_shape)))
total_bytes = (
functools.reduce(operator.mul, A_shape, 1) * A_elem_bytes
+ functools.reduce(operator.mul, B_shape, 1) * B_elem_bytes
diff --git a/tests/python/tirx/tile/cuda/test_tcgen05_mma_codegen.py
b/tests/python/tirx/tile/cuda/test_tcgen05_mma_codegen.py
index 8a3b61687..ea53accd7 100644
--- a/tests/python/tirx/tile/cuda/test_tcgen05_mma_codegen.py
+++ b/tests/python/tirx/tile/cuda/test_tcgen05_mma_codegen.py
@@ -282,12 +282,12 @@ def
test_gemm_tcgen05_cta1_m64_packed_c_requires_explicit_weight_stationary():
def test_gemm_tcgen05_dense_descI_rejected():
- with pytest.raises(Exception, match="unknown operands.*descI"):
+ with pytest.raises(Exception, match=r"unknown operands.*descI"):
_compile(_stage_fixture(descI=T.uint32(0x04410490)))
def test_gemm_tcgen05_dense_descI_rejected_at_dispatch():
- with pytest.raises(Exception, match="unknown operands.*descI"):
+ with pytest.raises(Exception, match=r"unknown operands.*descI"):
_lower(_stage_fixture(descI=T.uint32(0x04410490)))
@@ -298,7 +298,7 @@ def test_gemm_smem_desc_modes_codegen(mode):
assert len(_calls(src)) == 4
assert "encode_matrix_descriptor" not in src
assert "shfl.sync" not in src
- with pytest.raises(Exception, match="unknown operands.*smem_desc"):
+ with pytest.raises(Exception, match=r"unknown operands.*smem_desc"):
_lower(_stage_fixture(smem_desc=mode))
@@ -482,8 +482,6 @@ def _scale_fixture(split_n=False, lane_offsets=False):
else
"tcgen05.mma.cta_group::1.kind::mxf4nvf4.block_scale.scale_vec::4X"
)
- b_shape = (K, N // cta)
-
@T.function
def kernel():
T.device_entry(launch=T.cuda.LaunchConfig(grid=cta, block=128,
cluster=cta))
diff --git a/tests/python/tirx/tile/test_dispatcher.py
b/tests/python/tirx/tile/test_dispatcher.py
index 0163d5df9..2215f390b 100644
--- a/tests/python/tirx/tile/test_dispatcher.py
+++ b/tests/python/tirx/tile/test_dispatcher.py
@@ -48,7 +48,7 @@ def
test_selected_lowerer_reports_context_and_cause(monkeypatch):
monkeypatch.setattr(SPECS[call.op.name], "lower", failing)
context = SimpleNamespace(target=tvm_tirx.target.Target("cuda"),
scope_kind="thread")
with pytest.raises(
- RuntimeError, match="op=tirx.cuda.tile.apply target=cuda scope=thread"
+ RuntimeError, match=r"op=tirx\.cuda\.tile\.apply target=cuda
scope=thread"
) as error:
run_dispatch(call, context)
assert "unsupported layout" in str(error.value)
@@ -180,7 +180,7 @@ def test_dispatch_requires_stmt_result():
register_op_attr(
call.op.name, "FTileOpDispatch", lambda call, context: Function([],
Evaluate(0))
)
- with pytest.raises(TypeError, match="ir.Stmt"):
+ with pytest.raises(TypeError, match=r"ir\.Stmt"):
_lower(Evaluate(call))
diff --git a/tests/python/tirx/tile/trn/test_copy.py
b/tests/python/tirx/tile/trn/test_copy.py
index 996f45f04..abe212fed 100644
--- a/tests/python/tirx/tile/trn/test_copy.py
+++ b/tests/python/tirx/tile/trn/test_copy.py
@@ -803,5 +803,5 @@ def test_tensor_copy_rejects_partition_transpose():
)
T.trn.tile.tensor_copy(b, a)
- with target, pytest.raises(RuntimeError, match="identity|transpose"):
+ with target, pytest.raises(RuntimeError, match=r"identity|transpose"):
tvm_tirx.backend.trn.transform.LowerTIRx()(tvm_tirx.IRModule({"main":
copy}))
diff --git a/tests/python/tirx/transform/test_convert_ssa.py
b/tests/python/tirx/transform/test_convert_ssa.py
index 0ef8c45ac..b102c6307 100644
--- a/tests/python/tirx/transform/test_convert_ssa.py
+++ b/tests/python/tirx/transform/test_convert_ssa.py
@@ -14,11 +14,13 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F401
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
import tvm_tirx
import tvm_tirx.testing
-from tvm_tirx import ir, tirx
from tvm_tirx.script import ir as I
from tvm_tirx.script import tirx as T
diff --git a/tests/python/tirx/transform/test_lower_tensor_layouts.py
b/tests/python/tirx/transform/test_lower_tensor_layouts.py
index 08ac5a11f..aa2f67074 100644
--- a/tests/python/tirx/transform/test_lower_tensor_layouts.py
+++ b/tests/python/tirx/transform/test_lower_tensor_layouts.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
from types import SimpleNamespace
import pytest
diff --git a/tests/python/tirx/transform/test_lower_tvm_builtin.py
b/tests/python/tirx/transform/test_lower_tvm_builtin.py
index 2a0b34ba8..14903ee39 100644
--- a/tests/python/tirx/transform/test_lower_tvm_builtin.py
+++ b/tests/python/tirx/transform/test_lower_tvm_builtin.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
import numpy as np
import pytest
import tvm_ffi
diff --git a/tests/python/tirx/transform/test_remove_no_op.py
b/tests/python/tirx/transform/test_remove_no_op.py
index 53705d2f0..22ec484c3 100644
--- a/tests/python/tirx/transform/test_remove_no_op.py
+++ b/tests/python/tirx/transform/test_remove_no_op.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
import pytest
import tvm_tirx
diff --git a/tests/python/tirx/transform/test_simplify.py
b/tests/python/tirx/transform/test_simplify.py
index 5bf482f39..9d9d7a670 100644
--- a/tests/python/tirx/transform/test_simplify.py
+++ b/tests/python/tirx/transform/test_simplify.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
import pytest
import tvm_tirx
@@ -685,7 +688,7 @@ def test_most_restrictive_conditional():
allow for later rewrites. For example, if it is known that `a <= b`,
then `a >= b` cannot be proven, but can be reduced to `a == b`.
"""
- i, j, k = [T.dynamic(name, "int32") for name in "ijk"]
+ i, j, _k = [T.dynamic(name, "int32") for name in "ijk"]
tir_int = T.int32(0)
test_cases = [
diff --git a/tests/python/tirx/transform/test_split_host_device.py
b/tests/python/tirx/transform/test_split_host_device.py
index 42c3b3a30..41cc3e025 100644
--- a/tests/python/tirx/transform/test_split_host_device.py
+++ b/tests/python/tirx/transform/test_split_host_device.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
from __future__ import annotations
import pytest
diff --git a/tests/python/tirx/transform/test_storage_rewrite.py
b/tests/python/tirx/transform/test_storage_rewrite.py
index 363c1f203..b7f54bcbf 100644
--- a/tests/python/tirx/transform/test_storage_rewrite.py
+++ b/tests/python/tirx/transform/test_storage_rewrite.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
import uuid
import pytest
diff --git a/tests/python/tirx/transform/test_tile_dispatch.py
b/tests/python/tirx/transform/test_tile_dispatch.py
index a9348279d..a5ebe2099 100644
--- a/tests/python/tirx/transform/test_tile_dispatch.py
+++ b/tests/python/tirx/transform/test_tile_dispatch.py
@@ -15,6 +15,9 @@
# specific language governing permissions and limitations
# under the License.
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
+
from contextlib import nullcontext
import pytest
diff --git a/tests/python/tirx/transform/test_unroll.py
b/tests/python/tirx/transform/test_unroll.py
index 127ba27d4..58c3ab473 100644
--- a/tests/python/tirx/transform/test_unroll.py
+++ b/tests/python/tirx/transform/test_unroll.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
from __future__ import annotations
import pytest
diff --git a/tests/python/tirx/transform/test_vectorize.py
b/tests/python/tirx/transform/test_vectorize.py
index 6ee3d996b..3ba3df580 100644
--- a/tests/python/tirx/transform/test_vectorize.py
+++ b/tests/python/tirx/transform/test_vectorize.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+
+# The parser consumes named IR bindings that Python sees as unused.
+# ruff: noqa: F841
import pytest
import tvm_tirx
