[tensilelite] Enable tensilelite tests in TheRock CI matrix

.github/scripts/therock_matrix.py

@@ -48,7 +48,13 @@
     },
     "blas": {
         "cmake_options": ["-DTHEROCK_ENABLE_BLAS=ON"],
-        "projects_to_test": ["hipblaslt", "rocblas", "hipblas", "rocroller"],
+        "projects_to_test": [
+            "hipblaslt",
+            "rocblas",
+            "hipblas",
+            "rocroller",
+            "tensilelite",
+        ],
     },
     "miopen": {
         "cmake_options": [

projects/hipblaslt/tensilelite/Tensile/Tests/unit/gpu_test_helpers.py

@@ -22,13 +22,21 @@
 import subprocess
 import tempfile
 
+import pytest
+
+hip = None
+try:
+    from hip import hip
+except ImportError:
+    pass
+
+requires_hip = pytest.mark.skipif(hip is None, reason="hip module not installed")
+
 # Add tensilelite to path
 SCRIPT_DIR = os.path.dirname(os.path.abspath(__file__))
 TENSILE_ROOT = os.path.abspath(os.path.join(SCRIPT_DIR, "..", "..", ".."))
 sys.path.insert(0, TENSILE_ROOT)
 
-from hip import hip, hiprtc  # type: ignore
-
 from unittest.mock import MagicMock
 from types import SimpleNamespace
 from dataclasses import dataclass
@@ -79,6 +87,10 @@ def _detect_gfx_target():
 # ---- Constants ----
 GFX_TARGET = _detect_gfx_target()
 HAS_GFX950 = GFX_TARGET == "gfx950"
+requires_gpu = pytest.mark.skipif(
+    hip is None or not HAS_GFX950,
+    reason=f"requires hip module and gfx950 (found hip={'yes' if hip else 'no'}, arch={GFX_TARGET})",
+)
 WAVESIZE   = 64
 NUM_WAVES  = 4
 NUM_THREADS = WAVESIZE * NUM_WAVES  # 256

projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_gr_lr_roundtrip.py

@@ -26,6 +26,7 @@
     setup_roundtrip_writer,
     build_roundtrip_inner_asm,
     alloc_export_vgprs,
+    requires_gpu,
 )
 
 
@@ -320,7 +321,7 @@ def print_grid_diff(label, actual, expected):
 # Pytest tests
 # ---------------------------------------------------------------------------
 
-@pytest.mark.skipif(not HAS_GFX950, reason=f"GPU tests require gfx950, found {GFX_TARGET}")
+@requires_gpu
 class TestGrLrRoundtrip:
 
     @pytest.fixture(params=CONFIGS, ids=lambda c: c.label)

projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_gr_lr_roundtrip_fp8.py

@@ -36,6 +36,7 @@
     setup_roundtrip_writer,
     build_roundtrip_inner_asm,
     alloc_export_vgprs,
+    requires_gpu,
 )
 
 BPE = 1  # fp8: 1 byte per element
@@ -271,7 +272,7 @@ def compare_tiles_fp8(actual_bytes, expected_tiles, tileInfoA, tileInfoB, wave_i
 # Pytest tests
 # ---------------------------------------------------------------------------
 
-@pytest.mark.skipif(not HAS_GFX950, reason=f"GPU tests require gfx950, found {GFX_TARGET}")
+@requires_gpu
 class TestGrLrRoundtripFP8:
 
     @pytest.fixture(params=TILE_CONFIGS_FP8, ids=lambda c: c.label)

projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_gr_offset.py

@@ -27,6 +27,7 @@
     assemble_and_run,
     generate_kernel_asm,
     generate_load_params,
+    requires_gpu,
 )
 
 from Tensile.Components.Subtile.SubtileGREmit import graTileAssignment
@@ -319,7 +320,7 @@ def verify_gr_output(output_bytes, ti, kernel, dest_vgprs, mt, stride, debug=Fal
 # Pytest tests
 # ---------------------------------------------------------------------------
 
-@pytest.mark.skipif(not HAS_GFX950, reason=f"GPU tests require gfx950, found {GFX_TARGET}")
+@requires_gpu
 class TestGrOffset:
 
     @pytest.fixture(params=CONFIGS, ids=lambda c: c.label)

projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_mfma_fp8.py

@@ -42,6 +42,7 @@
     setup_roundtrip_writer,
     build_roundtrip_inner_asm,
     collect_tile_vgprs,
+    requires_gpu,
 )
 
 # ---------------------------------------------------------------------------
@@ -380,7 +381,7 @@ def compare_mfma_output(actual_bytes, expected_pairs, debug=False):
 # Pytest tests
 # ---------------------------------------------------------------------------
 
-@pytest.mark.skipif(not HAS_GFX950, reason=f"GPU tests require gfx950, found {GFX_TARGET}")
+@requires_gpu
 class TestMfmaFP8:
     """GPU tests for FP8 GR->LDS->LR->MFMA roundtrip."""
 

projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_storeD_roundtrip.py

@@ -22,8 +22,6 @@
 TENSILE_ROOT = os.path.abspath(os.path.join(SCRIPT_DIR, "..", "..", ".."))
 sys.path.insert(0, TENSILE_ROOT)
 
-from hip import hip  # type: ignore
-
 from unittest.mock import MagicMock
 
 from rocisa.code import Module, TextBlock
@@ -39,19 +37,20 @@
 
 from gpu_test_helpers import (
     TileConfig,
-    HAS_GFX950,
     GFX_TARGET,
     WAVESIZE,
     NUM_THREADS,
     create_writer,
     generate_kernel_asm,
     assemble_and_run,
     assemble_kernel,
+    hip,
     hip_check,
     init_rocisa,
+    requires_gpu,
 )
 
-pytestmark = pytest.mark.skipif(not HAS_GFX950, reason=f"GPU tests require gfx950, found {GFX_TARGET}")
+pytestmark = requires_gpu
 
 # ---------------------------------------------------------------------------
 # Test configurations: (mt_a, mt_b, depth_u)

projects/hipblaslt/tensilelite/rocisa/rocisa/__init__.py

@@ -3,6 +3,17 @@
 
 import sys
 import types
+
+from pathlib import Path as _Path
+
+if any(_Path(__file__).parent.glob("_rocisa.abi3.*")) and sys.version_info < (3, 12):
+    raise ImportError(
+        f"rocisa stable-ABI extension requires Python >= 3.12 "
+        f"(running {sys.version_info.major}.{sys.version_info.minor}). "
+        f"Install a non-stable-ABI build or upgrade Python."
+    )
+del _Path
+
 from ._rocisa import *
 from . import _rocisa