forked from ccf-ai-infra/TileOPs-Metax
267 lines
9.9 KiB
Python
267 lines
9.9 KiB
Python
from collections import defaultdict
|
|
|
|
import pytest
|
|
import torch
|
|
|
|
from tests.test_base import _check_result
|
|
|
|
|
|
def _under_repo_tests(item: pytest.Item) -> bool:
|
|
path = str(item.path)
|
|
return "tests/" in path and "benchmarks/tests/" not in path
|
|
|
|
|
|
def pytest_addoption(parser: pytest.Parser) -> None:
|
|
"""Register the opt-in in-kernel timeline-trace flag.
|
|
|
|
Off by default: when ``--trace-kernel`` is absent the process-local trace
|
|
switch stays off, so trace-dump tests no-op and normal runs are zero cost.
|
|
(``--trace`` itself is reserved by pytest for its pdb tracer.)
|
|
"""
|
|
parser.addoption(
|
|
"--trace-kernel",
|
|
action="store_true",
|
|
default=False,
|
|
help="Build instrumented kernels with in-kernel tracing and dump their "
|
|
"timeline (HTML + Chrome JSON) for the trace-dump tests.",
|
|
)
|
|
|
|
|
|
def pytest_configure(config: pytest.Config) -> None:
|
|
"""Flip the in-process trace switch on when ``--trace-kernel`` is passed.
|
|
|
|
Runs once at startup, before any kernel is built, so the traced build is the
|
|
one that gets cached. No environment variable is involved — the switch lives
|
|
in this pytest process only.
|
|
"""
|
|
if config.getoption("--trace-kernel"):
|
|
from tileops.trace import trace
|
|
|
|
trace.enable() # dumps to debug/ (gitignored)
|
|
|
|
|
|
@pytest.fixture(autouse=True)
|
|
def setup() -> None:
|
|
torch.manual_seed(1235)
|
|
if torch.cuda.is_available():
|
|
torch.cuda.manual_seed_all(1235)
|
|
|
|
|
|
@pytest.fixture
|
|
def isolated_dynamo():
|
|
"""Reset torch._dynamo state around a test that calls ``torch.compile``.
|
|
|
|
Dynamo's recompile cache is keyed per code object, and every
|
|
``torch.compile``-d plain callable (any non-``nn.Module``, e.g. a TileOps
|
|
``Op`` instance) shares torch's single wrapper frame. Each compiled op
|
|
instance therefore consumes one slot of that frame's shared
|
|
``cache_size_limit`` (default 8) for the whole pytest process, so compile
|
|
tests pollute each other's cache and later ``fullgraph=True`` tests fail
|
|
with ``FailOnRecompileLimitHit``. Request this fixture from every test
|
|
that calls ``torch.compile``.
|
|
"""
|
|
torch._dynamo.reset()
|
|
yield
|
|
torch._dynamo.reset()
|
|
|
|
|
|
NON_RUNTIME_OPS_TIER_FILES = {
|
|
"tests/ops/test_elementwise_caching_autotune.py",
|
|
"tests/ops/test_elementwise_compile.py",
|
|
"tests/ops/test_elementwise_config_dtype.py",
|
|
}
|
|
|
|
TILELANG_019_SKIP_REASON = (
|
|
"Skipped under TileLang 0.1.9: known regressions in autodiff/codegen "
|
|
"lowering produce incorrect numerics or compile failures; re-enable "
|
|
"when these tests pass against the current tilelang."
|
|
)
|
|
|
|
TILELANG_019_KNOWN_FAILING_PATH_SUFFIXES = ()
|
|
|
|
TILELANG_019_KNOWN_FAILING_NODEIDS = set()
|
|
|
|
TILELANG_019_KNOWN_FAILING_PREFIXES = ()
|
|
|
|
def _get_callspec_params(item: pytest.Item) -> dict | None:
|
|
callspec = getattr(item, "callspec", None)
|
|
if callspec is None:
|
|
return None
|
|
return getattr(callspec, "params", None)
|
|
|
|
|
|
def _freeze_value(value: object) -> object:
|
|
if isinstance(value, dict):
|
|
return tuple(sorted((key, _freeze_value(val)) for key, val in value.items()))
|
|
if isinstance(value, (list, tuple)):
|
|
return tuple(_freeze_value(item) for item in value)
|
|
if isinstance(value, set):
|
|
return tuple(sorted((_freeze_value(item) for item in value), key=str))
|
|
return value
|
|
|
|
|
|
def _without_dtype(params: dict) -> tuple[tuple[str, object], ...]:
|
|
return tuple(
|
|
sorted((key, _freeze_value(value)) for key, value in params.items() if key != "dtype")
|
|
)
|
|
|
|
|
|
def pytest_collection_modifyitems(items: list[pytest.Item]) -> None:
|
|
"""Validate explicit test tier assignments."""
|
|
tier_errors: list[str] = []
|
|
tier_names = ("smoke", "full", "nightly")
|
|
tilelang_019_skip = pytest.mark.skip(reason=TILELANG_019_SKIP_REASON)
|
|
|
|
for item in items:
|
|
path = str(item.path)
|
|
if not _under_repo_tests(item):
|
|
continue
|
|
if (
|
|
item.nodeid in TILELANG_019_KNOWN_FAILING_NODEIDS
|
|
or any(path.endswith(suffix) for suffix in TILELANG_019_KNOWN_FAILING_PATH_SUFFIXES)
|
|
or any(
|
|
item.nodeid.startswith(prefix) for prefix in TILELANG_019_KNOWN_FAILING_PREFIXES
|
|
)
|
|
):
|
|
item.add_marker(tilelang_019_skip)
|
|
|
|
tiers = [name for name in tier_names if item.get_closest_marker(name) is not None]
|
|
if len(tiers) != 1:
|
|
tier_errors.append(
|
|
f"{item.nodeid}: expected exactly one tier marker, found {tiers or 'none'}"
|
|
)
|
|
|
|
ops_groups: dict[tuple[str, str], list[pytest.Item]] = defaultdict(list)
|
|
for item in items:
|
|
path = str(item.path)
|
|
if "tests/ops/" not in path or "benchmarks/tests/" in path:
|
|
continue
|
|
test_name = getattr(item, "originalname", item.name)
|
|
ops_groups[(path, test_name)].append(item)
|
|
|
|
for (_path, _test_name), group in ops_groups.items():
|
|
if any(_path.endswith(path) for path in NON_RUNTIME_OPS_TIER_FILES):
|
|
continue
|
|
|
|
non_xfail_items = [
|
|
item for item in group if item.get_closest_marker("xfail") is None
|
|
]
|
|
smoke_items = [
|
|
item for item in group if item.get_closest_marker("smoke") is not None
|
|
]
|
|
|
|
# Smoke cases must never be xfail (checked before tune gate)
|
|
for item in smoke_items:
|
|
if item.get_closest_marker("xfail") is not None:
|
|
tier_errors.append(f"{item.nodeid}: smoke cases must not be xfail")
|
|
|
|
# For count and ordering checks, only consider non-xfail smoke cases
|
|
valid_smoke_items = [
|
|
item for item in smoke_items if item.get_closest_marker("xfail") is None
|
|
]
|
|
|
|
if non_xfail_items:
|
|
if len(valid_smoke_items) < 1:
|
|
tier_errors.append(
|
|
f"{non_xfail_items[0].nodeid}: each test must have at least one smoke case"
|
|
)
|
|
else:
|
|
# All smoke cases must appear as the first N non-xfail items
|
|
expected_smoke = non_xfail_items[: len(valid_smoke_items)]
|
|
if valid_smoke_items != expected_smoke:
|
|
tier_errors.append(
|
|
f"{non_xfail_items[0].nodeid}: all smoke cases must appear "
|
|
f"as the first {len(valid_smoke_items)} non-xfail cases of each test"
|
|
)
|
|
|
|
dtype_supported: set[object] = set()
|
|
dtype_smoke: set[object] = set()
|
|
smoke_signatures: set[tuple[tuple[str, object], ...]] = set()
|
|
dtype_cases_present = False
|
|
|
|
for item in non_xfail_items:
|
|
params = _get_callspec_params(item)
|
|
if not params or "dtype" not in params:
|
|
continue
|
|
|
|
dtype_cases_present = True
|
|
dtype_supported.add(params["dtype"])
|
|
|
|
if item.get_closest_marker("smoke") is not None:
|
|
dtype_smoke.add(params["dtype"])
|
|
smoke_signatures.add(_without_dtype(params))
|
|
|
|
if dtype_cases_present:
|
|
missing_smoke_dtypes = dtype_supported - dtype_smoke
|
|
if missing_smoke_dtypes:
|
|
tier_errors.append(
|
|
f"{non_xfail_items[0].nodeid}: each dtype must have at least one smoke case; "
|
|
f"missing smoke for {sorted(str(dtype) for dtype in missing_smoke_dtypes)}"
|
|
)
|
|
|
|
for item in non_xfail_items:
|
|
if item.get_closest_marker("full") is None:
|
|
continue
|
|
|
|
params = _get_callspec_params(item)
|
|
if not params or "dtype" not in params:
|
|
continue
|
|
|
|
if _without_dtype(params) in smoke_signatures:
|
|
tier_errors.append(
|
|
f"{item.nodeid}: full cases must not differ from a smoke case only by dtype"
|
|
)
|
|
|
|
first_tuned_item: pytest.Item | None = None
|
|
full_tuned_items: list[pytest.Item] = []
|
|
for item in group:
|
|
params = _get_callspec_params(item)
|
|
if params is None or "tune" not in params:
|
|
continue
|
|
|
|
tune = params["tune"]
|
|
is_smoke = item.get_closest_marker("smoke") is not None
|
|
if is_smoke and tune is True:
|
|
tier_errors.append(f"{item.nodeid}: smoke cases must use tune=False")
|
|
if tune is True:
|
|
if first_tuned_item is None:
|
|
first_tuned_item = item
|
|
if item.get_closest_marker("full") is not None:
|
|
full_tuned_items.append(item)
|
|
if first_tuned_item is not None:
|
|
if not full_tuned_items:
|
|
tier_errors.append(
|
|
f"{first_tuned_item.nodeid}: the first tune=True case must be marked full"
|
|
)
|
|
elif len(full_tuned_items) > 1:
|
|
tier_errors.append(
|
|
f"{group[0].path}::{group[0].originalname}: at most one tune=True case may be full"
|
|
)
|
|
elif full_tuned_items[0] is not first_tuned_item:
|
|
tier_errors.append(
|
|
f"{first_tuned_item.nodeid}: the first tune=True case must be the only full tuned case"
|
|
)
|
|
|
|
if tier_errors:
|
|
raise pytest.UsageError(
|
|
"Invalid explicit test tier assignments detected:\n" + "\n".join(tier_errors)
|
|
)
|
|
|
|
|
|
@pytest.hookimpl(hookwrapper=True)
|
|
def pytest_runtest_call(item):
|
|
"""After test execution, attach Op metadata from TestBase.check() to the item."""
|
|
yield
|
|
op_name = getattr(_check_result, "op_name", None)
|
|
if op_name:
|
|
item.user_properties.append(("op", op_name))
|
|
op_module = getattr(_check_result, "op_module", None)
|
|
if op_module:
|
|
item.user_properties.append(("op_module", op_module))
|
|
max_err = getattr(_check_result, "max_abs_err", None)
|
|
if max_err is not None:
|
|
item.user_properties.append(("max_abs_err", f"{max_err:.2e}"))
|
|
_check_result.op_name = None
|
|
_check_result.op_module = None
|
|
_check_result.max_abs_err = None
|