Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
70 changes: 54 additions & 16 deletions tests/prims_ts/conftest.py
Original file line number Diff line number Diff line change
Expand Up @@ -16,35 +16,73 @@
Shared test setup for vendored Prims-TS kernels.
"""

import importlib.util
from pathlib import Path

import pytest

from flashinfer.prims_ts.cutlass_dsl import (
ensure_cutlass_dsl_experimental,
get_cutlass_dsl_bootstrap_error,
require_cutlass_dsl_experimental,
)

require_cutlass_dsl_experimental()
# Prims-TS needs ``cutlass.experimental.primitives`` and
# ``cutlass.experimental.task_scheduling``, which exist only in CUTLASS DSL 4.7+.
# FlashInfer's declared floor is deliberately lower so it stays co-installable
# with vLLM and SGLang, so an environment that cannot provide them is supported
# and must degrade rather than fail -- the same contract the other 4.7-gated
# features honor (``flashinfer/attention/cute_dsl/sm120_fmha.py``,
# ``flashinfer/kda.py``). ``ensure_...`` records the import failure for
# ``pytest_report_header`` instead of raising.
HAS_PRIMS_TS_RUNTIME = ensure_cutlass_dsl_experimental()

# Modules that reach ``cutlass.experimental`` at IMPORT time cannot be marked
# during collection -- they raise before an item exists -- so they have to be
# dropped before collection instead.
#
# The reach is transitive, not direct: ``batched_gemm_config`` does an unguarded
# module-level ``from cutlass.experimental import primitives``, so every module
# importing it inherits the requirement. That is all ten ``test_batched_gemm_*``
# modules plus the two MoE support modules below.
#
# The glob is deliberate: new ``test_batched_gemm_*`` modules are the likely
# growth, and over-ignoring is harmless here (on a pre-4.7 lane the whole
# directory is skipped anyway). ``test_prims_ts_directory_still_collects_without_
# cutlass_experimental`` in tests/test_prims_ts_import_isolation.py enforces that
# this stays complete.
if not HAS_PRIMS_TS_RUNTIME:
collect_ignore_glob = ["test_batched_gemm_*.py"]
collect_ignore = [
"test_moe_bf16_support.py",
"test_moe_nvfp4_support.py",
]


def pytest_collection_modifyitems(config, items):

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The mechanism itself is right and the exit-code-5 rationale in this docstring is the part to keep exactly as-is β€” module-level
skips trading "red for red" on directory-scoped runs is a trap most people would have walked into. The gap is only that the
collect_ignore list is a load-bearing claim with no enforcement: the next test module added with an eager DSL import silently
resurrects #5213. Suggest adding the simulation as a CI-runnable meta-check β€” stub plugin (~15 lines) + pytest tests/prims_ts
--collect-only, assert exit 0:

# dsl46_stub.py β€” load with `pytest -p dsl46_stub` to emulate CUTLASS DSL 4.6
import importlib.abc, importlib.machinery, sys

class _EmptyLoader(importlib.abc.Loader): 
    def create_module(self, spec): return None
    def exec_module(self, module): module.__path__ = []

class _Stub(importlib.abc.MetaPathFinder):
    def find_spec(self, fullname, path=None, target=None):
        if fullname == "cutlass.experimental":
            return importlib.machinery.ModuleSpec(fullname, _EmptyLoader(), is_package=True)
        return None

sys.meta_path.insert(0, _Stub())

(A naive blocker that fails cutlass.experimental outright doesn't work as a simulator β€” on the 4.7 wheel, import cutlass itself
eagerly chains into experimental.primitives, so the stub has to satisfy the parent while leaving the submodules missing.)

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 9b5a3c81 β€” took your simulator approach essentially as written.

test_prims_ts_directory_still_collects_without_cutlass_experimental writes the stub to a tmp dir, runs pytest tests/prims_ts --collect-only -p dsl46_stub, and asserts exit 0. It lives in tests/test_prims_ts_import_isolation.py rather than under tests/prims_ts/ so it still runs on lanes where that directory is skipped.

Your parenthetical was the load-bearing part and I verified it rather than trusting it: with a package whose __init__ does from .experimental import primitives (emulating 4.7's eager chain), the stub yields ImportError: cannot import name 'primitives' from 'cutlass.experimental' while cutlass.experimental itself still resolves with __path__=[]. ensure_cutlass_dsl_experimental catches that via its except BaseException, so it lands in the unavailable path β€” which is the state we want to simulate. A finder that failed cutlass.experimental outright would have broken import cutlass itself and simulated the wrong thing, exactly as you said.

Also worth recording why exit 0 is the right assertion: a missed module raises during collection (exit 2), and a module-level skip collects nothing and exits 5 β€” so the check catches both the regression and the tempting wrong fix.

"""Skip this directory when the installed DSL cannot provide Prims-TS.

def _has_module(name: str) -> bool:
try:
return importlib.util.find_spec(name) is not None
except ModuleNotFoundError:
return False
Marking collected items rather than skipping at module level is deliberate:
a module-level skip collects nothing, and a run scoped to this directory
would then exit with pytest's "no tests collected" code 5 -- trading a red
for a red. Marked items are collected first, so the run reports skips and
exits 0.
"""

del config
if HAS_PRIMS_TS_RUNTIME:
return

HAS_PRIMS_TS_RUNTIME = all(
_has_module(name)
for name in (
"cutlass",
"cutlass.cute",
"cutlass.experimental.primitives",
"cutlass.experimental.task_scheduling",
here = Path(__file__).parent
skip_prims_ts = pytest.mark.skip(
reason=(
"Prims-TS needs cutlass.experimental from CUTLASS DSL 4.7+; the "
f"installed wheel cannot provide it: {get_cutlass_dsl_bootstrap_error()!r}"
)
)
)
for item in items:
# `in ... parents` rather than `== parent`: a future subdirectory of
# tests/prims_ts/ would otherwise escape the skip marker.
if here in Path(str(item.path)).parents:
item.add_marker(skip_prims_ts)
Comment thread
coderabbitai[bot] marked this conversation as resolved.


@pytest.fixture(autouse=True)
Expand Down
82 changes: 82 additions & 0 deletions tests/test_prims_ts_import_isolation.py
Original file line number Diff line number Diff line change
Expand Up @@ -114,3 +114,85 @@ def test_prims_ts_bootstrap_scopes_work_tile_info_customization():
assert WorkTileInfo.__init__ is original_init
"""
)


# Emulates a CUTLASS DSL 4.6 wheel: `cutlass.experimental` resolves as an empty
# package so the eager chain inside `import cutlass` is satisfied, while every
# submodule Prims-TS needs (`.primitives`, `.task_scheduling`) stays missing. A
# finder that fails `cutlass.experimental` outright would not simulate 4.6 -- on
# the 4.7 wheel it breaks `import cutlass` itself, which is a different failure.
_DSL46_STUB = """
import importlib.abc
import importlib.machinery
import sys


class _EmptyLoader(importlib.abc.Loader):
def create_module(self, spec):
return None

def exec_module(self, module):
module.__path__ = []


class _Stub(importlib.abc.MetaPathFinder):
def find_spec(self, fullname, path=None, target=None):
if fullname == "cutlass.experimental":
return importlib.machinery.ModuleSpec(
fullname, _EmptyLoader(), is_package=True
)
return None


sys.meta_path.insert(0, _Stub())
"""


def test_prims_ts_directory_still_collects_without_cutlass_experimental(tmp_path):
"""``tests/prims_ts`` must COLLECT cleanly on a DSL that lacks 4.7.

``collect_ignore`` in ``tests/prims_ts/conftest.py`` is hand-maintained: a new
test module that imports the vendored kernels at module scope cannot be marked
during collection, only dropped before it, and forgetting to list it silently
reintroduces #5213 (``ERROR tests/prims_ts``) on every pre-4.7 lane.

Exit code 0 is the assertion that matters, and it is stricter than it looks:
a missed module raises during collection (exit 2), and a module-level skip --
the obvious alternative fix -- collects nothing and exits 5.
"""

stub = tmp_path / "dsl46_stub.py"
stub.write_text(_DSL46_STUB)

env = os.environ.copy()
env.pop("PYTEST_ADDOPTS", None)
env["PYTHONPATH"] = os.pathsep.join(
filter(None, (str(tmp_path), str(_REPO_ROOT), env.get("PYTHONPATH")))
)
env["FLASHINFER_WORKSPACE_BASE"] = str(tmp_path / "workspace")

result = subprocess.run(
[
sys.executable,
"-m",
"pytest",
"tests/prims_ts",
"--collect-only",
"-q",
"--color=no",
"-p",
"dsl46_stub",
],
cwd=_REPO_ROOT,
env=env,
check=False,
capture_output=True,
text=True,
)

assert result.returncode == 0, (
"tests/prims_ts failed to collect under a simulated CUTLASS DSL 4.6.\n"
"If this is a newly added module that imports the Prims-TS kernels at "
"module scope, add it to collect_ignore in tests/prims_ts/conftest.py.\n"
f"exit={result.returncode}\n{result.stdout}\n{result.stderr}"
)
Loading