Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
WorldofKerry merged 1 commit into from
Jun 30, 2026
Merged

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

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
38 changes: 22 additions & 16 deletions projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_subtile_gfx1250_codegen.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,6 @@

import os
import sys
import shutil

import pytest
from types import SimpleNamespace
Expand All @@ -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):
Expand Down Expand Up @@ -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)
Expand All @@ -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)
Expand All @@ -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)
Expand Down Expand Up @@ -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)
Expand All @@ -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)
Expand All @@ -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)
Expand All @@ -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)
Expand Down
Loading