diff --git a/tests/py/kernels/conftest.py b/tests/py/kernels/conftest.py index 52be6fd4d7..71786626fe 100644 --- a/tests/py/kernels/conftest.py +++ b/tests/py/kernels/conftest.py @@ -37,16 +37,84 @@ def _has_cuda_core() -> bool: ) +def _decode_cuda_driver_version(version: int) -> tuple[int, int]: + """Decode CUDA's integer driver API version into ``(major, minor)``.""" + if version <= 0: + raise ValueError(f"invalid CUDA driver version: {version}") + return version // 1000, (version % 1000) // 10 + + +def _cuda_driver_nvrtc_versions() -> tuple[tuple[int, int], tuple[int, int]]: + """Return the maximum CUDA version supported by the driver and NVRTC's version.""" + try: + from cuda.bindings import driver, nvrtc + except ImportError: + # cuda-python < 12.8 exposed the same APIs from these modules. + from cuda import cuda as driver + from cuda import nvrtc + + driver_status, encoded_driver_version = driver.cuDriverGetVersion() + if int(driver_status) != 0: + raise RuntimeError( + f"cuDriverGetVersion failed with CUDA error {int(driver_status)}" + ) + + nvrtc_status, nvrtc_major, nvrtc_minor = nvrtc.nvrtcVersion() + if int(nvrtc_status) != 0: + raise RuntimeError(f"nvrtcVersion failed with CUDA error {int(nvrtc_status)}") + + return _decode_cuda_driver_version(encoded_driver_version), ( + nvrtc_major, + nvrtc_minor, + ) + + +def _ptx_compatibility_skip_reason( + driver_version: tuple[int, int], nvrtc_version: tuple[int, int] +) -> str | None: + """Explain why NVRTC PTX cannot be loaded, or return ``None`` if compatible.""" + if driver_version >= nvrtc_version: + return None + + driver = ".".join(map(str, driver_version)) + backend = ".".join(map(str, nvrtc_version)) + return ( + "CUDA kernel tests require a driver capable of loading the PTX emitted " + f"by NVRTC, but the active driver supports CUDA {driver} and NVRTC is " + f"CUDA {backend}. Upgrade the CI runner driver to >= {backend}, or use " + f"a CUDA toolkit/runtime <= {driver}." + ) + + def pytest_collection_modifyitems(config, items): - """Skip the whole kernels suite when cuda-core is missing. + """Skip the kernels suite when its CUDA PTX toolchain cannot be used. The QDP kernel tests compile CUDA via the ``cuda.core`` NVRTC API; without cuda-core they all error on import. Skip (not fail) so a plain test run is green — `just install-test-ext` pulls cuda-core so they actually run. + + NVRTC also emits PTX that must be JIT-loaded by the installed driver. CUDA's + minor-version compatibility does not cover PTX generated by a newer toolkit, + so do not run the suite when the driver API version predates NVRTC. """ - if _HAS_CUDA_CORE: + if not _HAS_CUDA_CORE: + reason = "cuda-core (cuda.core) not installed" + elif not torch.cuda.is_available(): return - skip = pytest.mark.skip(reason="cuda-core (cuda.core) not installed") + else: + try: + driver_version, nvrtc_version = _cuda_driver_nvrtc_versions() + reason = _ptx_compatibility_skip_reason(driver_version, nvrtc_version) + except Exception as exc: + # Only run kernel tests after positively verifying that their PTX can + # be loaded. This turns an opaque collection of kernel failures into + # one actionable CI skip reason if version discovery itself breaks. + reason = f"could not verify CUDA driver/NVRTC compatibility: {exc}" + + if reason is None: + return + + skip = pytest.mark.skip(reason=reason) for item in items: item.add_marker(skip)