fix(hipblaslt): skip gfx1250 subtile codegen tests on other archs (#8893)

ISSUE ID : AIHPBLAS-3658

## Motivation

Address test failures on gfx950
```
FAILED Tensile/Tests/unit/test_subtile_gfx1250_codegen.py::TestGfx1250SubtileCodegen::test_gr_ptr_updates_tdm[A] - AssertionError: assert 's_add_u64' in '/* TDM addr update: A += 128 */\ns_add_u32 s[sgprAddressA], s[sgprAddressA], 128\n\ns_addc_u32 s[sgprAddressA+1], s[sgprAddressA+1], 128\ns_mov_b64 s[sgprtdmAGroup0+2:sgprtdmAGroup0+2+1], s[sgprAddressA:sgprAddressA+1] // sync ...
FAILED Tensile/Tests/unit/test_subtile_gfx1250_codegen.py::TestGfx1250SubtileCodegen::test_gr_ptr_updates_tdm[B] - AssertionError: assert 's_add_u64' in '/* TDM addr update: B += 128 */\ns_add_u32 s[sgprAddressB], s[sgprAddressB], 128\n\ns_addc_u32 s[sgprAddressB+1], s[sgprAddressB+1], 128\ns_mov_b64 s[sgprtdmBGroup0+2:sgprtdmBGroup0+2+1], s[sgprAddressB:sgprAddressB+1] // sync ...
```

## Technical Details

Rocisa emits two u32 instructions on gfx950 and a single u64 on gfx1250.

Use pytestmark skipif + a module-scoped fixture to pin rocisa to
gfx1250, matching the pattern in test_emitMfmaInstruction.py. On
machines with an older amdclang++ that does not know gfx1250, the entire
module is skipped instead of failing with wrong encodings.

## Test

Expect CI to pass.

Author: WorldofKerry

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

@@ -11,7 +11,6 @@
 
 import os
 import sys
-import shutil
 
 import pytest
 from types import SimpleNamespace
@@ -21,19 +20,33 @@
 TENSILE_ROOT = os.path.abspath(os.path.join(SCRIPT_DIR, "..", "..", ".."))
 sys.path.insert(0, TENSILE_ROOT)
 
+from gpu_test_helpers import init_rocisa
+
 
 GFX1250_ISA = (12, 5, 0)
 WAVESIZE_32 = 32
 
 
-def _init_rocisa_gfx1250():
-    from rocisa import rocIsa
-    from Tensile.Common.Architectures import gfxToIsa
-    ri = rocIsa.getInstance()
-    isa = gfxToIsa("gfx1250")
-    asmpath = shutil.which('amdclang++') or '/usr/bin/amdclang++'
-    ri.init(isa, asmpath)
-    ri.setKernel(isa, WAVESIZE_32)
+def _gfx1250_asm_supported():
+    """Return True if the host assembler supports gfx1250 instructions."""
+    try:
+        init_rocisa(target="gfx1250", wavesize=WAVESIZE_32)
+        from rocisa import rocIsa
+        caps = rocIsa.getInstance().getAsmCaps()
+        return bool(caps.get("s_add_u64", 0))
+    except Exception:
+        return False
+
+
+pytestmark = pytest.mark.skipif(
+    not _gfx1250_asm_supported(),
+    reason="assembler does not support gfx1250 instructions",
+)
+
+
+@pytest.fixture(scope="module", autouse=True)
+def _rocisa_once():
+    init_rocisa(target="gfx1250", wavesize=WAVESIZE_32)
 
 
 def _mock_dtype(num_bytes=2):
@@ -143,7 +156,6 @@ class TestGfx1250SubtileCodegen:
                              ids=[f"{a}x{b}_wg{w[0]}x{w[1]}" for a, b, w in CONFIGS_1x1 + CONFIGS_MULTI_WAVE])
     def test_tdm_skips_gr_offset_alloc(self, mt_a, mt_b, wg):
         """TDM-enabled kernel: GR offset registers should not be allocated."""
-        _init_rocisa_gfx1250()
         kernel = _create_gfx1250_kernel(mt_a, mt_b, mi_wave_group=wg)
         writer, tiA, tiB = _create_writer_gfx1250(kernel)
         tiA.allocOffsetRegisters(writer, kernel)
@@ -155,7 +167,6 @@ def test_tdm_skips_gr_offset_alloc(self, mt_a, mt_b, wg):
                              ids=[f"{a}x{b}_wg{w[0]}x{w[1]}" for a, b, w in CONFIGS_MULTI_WAVE])
     def test_lr_tile_assignment_multi_wave(self, mt_a, mt_b, wg):
         """LR tile assignment with multi-wave TDM produces valid assembly."""
-        _init_rocisa_gfx1250()
         from Tensile.Components.Subtile.SubtileLREmit import lraTileAssignment
         kernel = _create_gfx1250_kernel(mt_a, mt_b, mi_wave_group=wg)
         writer, tiA, tiB = _create_writer_gfx1250(kernel)
@@ -173,7 +184,6 @@ def test_lr_tile_assignment_multi_wave(self, mt_a, mt_b, wg):
                              ids=[f"{a}x{b}" for a, b, _ in CONFIGS_1x1])
     def test_ds_read_dual_load(self, mt_a, mt_b, wg):
         """Wave32 8-VGPR tiles emit two DSLoadB128 (lo + hi K-halves)."""
-        _init_rocisa_gfx1250()
         from Tensile.Components.Subtile.SubtileLREmit import emitSingleDsRead
         kernel = _create_gfx1250_kernel(mt_a, mt_b)
         writer, tiA, tiB = _create_writer_gfx1250(kernel)
@@ -201,7 +211,6 @@ def test_select_d_geometry_wave32(self):
 
     def test_zero_tiles_wmma(self):
         """gfx1250 tile zeroing uses v_wmma_f32_16x16x4_f32 with acc2_imm=0."""
-        _init_rocisa_gfx1250()
         from Tensile.Components.Subtile.Kernel import initVgprTilesToZero
         kernel = _create_gfx1250_kernel(32, 32)
         writer, tiA, tiB = _create_writer_gfx1250(kernel)
@@ -218,7 +227,6 @@ def test_zero_tiles_wmma(self):
     @pytest.mark.parametrize("tc", ['A', 'B'])
     def test_gr_lds_buffer_swap_tdm(self, tc):
         """TDM LDS buffer swap emits XOR on tracking SGPR."""
-        _init_rocisa_gfx1250()
         from Tensile.Components.Subtile.SubtileGREmit import globalReadLDSBufferSwap
         kernel = _create_gfx1250_kernel(64, 64, mi_wave_group=[2, 2])
         writer, tiA, tiB = _create_writer_gfx1250(kernel)
@@ -233,7 +241,6 @@ def test_gr_lds_buffer_swap_tdm(self, tc):
     @pytest.mark.parametrize("tc", ['A', 'B'])
     def test_gr_ptr_updates_tdm(self, tc):
         """TDM pointer update increments Address and syncs descriptor."""
-        _init_rocisa_gfx1250()
         from Tensile.Components.Subtile.SubtileGREmit import globalReadPtrUpdates
         kernel = _create_gfx1250_kernel(64, 64, mi_wave_group=[2, 2])
         writer, tiA, tiB = _create_writer_gfx1250(kernel)
@@ -250,7 +257,6 @@ def test_gr_ptr_updates_tdm(self, tc):
     @pytest.mark.parametrize("tc", ['A', 'B'])
     def test_buffer_load_tdm(self, tc):
         """TDM emitSingleBufferLoad emits tensor_load_to_lds."""
-        _init_rocisa_gfx1250()
         from Tensile.Components.Subtile.SubtileGREmit import emitSingleBufferLoad
         kernel = _create_gfx1250_kernel(64, 64)
         writer, tiA, tiB = _create_writer_gfx1250(kernel)