mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-07-17 17:19:12 +00:00
[rocm-libraries] ROCm/rocm-libraries#9520 (commit 8676d27)
fix(ck-tile): close GEMM_GROUPED entry in arch_filter OPERATOR_TILE_CONSTRAINTS (#9520) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Related to the grouped-GEMM bridge merge that introduced the malformed entry. ISSUE ID : #9000 ## Summary I mistakenly deleted } when I was fixing merge conflicts of PR9000. This PR simply fixes that. Fixes a `SyntaxError` in `dispatcher/codegen/arch_filter.py` that currently breaks all GEMM codegen on `develop`. The `OperatorType.GEMM_GROUPED` entry in `OPERATOR_TILE_CONSTRAINTS` was opened with `{` but never given a body or a closing `}`, so the outer dict literal is never terminated: ``` arch_filter.py, line 65: OPERATOR_TILE_CONSTRAINTS = { SyntaxError: '{' was never closed ``` ## Design note `OPERATOR_TILE_CONSTRAINTS` maps each `OperatorType` to a tile-shape gate. The `GEMM_GROUPED` entry (from the grouped-GEMM bridge, #9000) and the `GEMM_STREAMK` entry (from the stream-K bridge, #9028) got interleaved: the `GEMM_GROUPED: {` value was left empty and the following `NOTE` comment block plus `OperatorType.GEMM_STREAMK` ran on *inside* it. The fix restores `GEMM_GROUPED` as its own dict (the same tile-shape gates as plain GEMM — grouped is multi-problem but each sub-problem obeys plain-GEMM tile validity) and lets the `NOTE` comment precede `GEMM_STREAMK` as originally intended. Because `arch_filter` is imported transitively by `unified_gemm_codegen`, the import-time crash cascades: `make dispatcher_gemm_lib` fails (`generate_gemm_fallback_kernel` Error 1) → no `libdispatcher_gemm_lib.so` → `tests/test_library_caching.py` fails 4/5 (only the pure source-string check survives). ## Changes **Modified** - `codegen/arch_filter.py` — give `GEMM_GROUPED` its own constraints dict + closing `},`; move the `NOTE` comment to precede `GEMM_STREAMK`. - `tests/CMakeLists.txt` — register the new test as `dispatcher_test_arch_filter_constraints`. **New** - `tests/test_arch_filter_constraints.py` — regression test that imports `arch_filter` and asserts every `OperatorType` has a well-formed constraints entry (required integer keys, positive values). Importing at all would have failed on the malformed table. ## Validation Reproduced on MI355X / gfx950 / ROCm 7.2 against `develop`: - Before: `python3 -m pytest tests/test_library_caching.py -v` → **4 failed / 1 passed**. - After: `arch_filter.py` parses (`PARSE OK`), `python3 tests/test_arch_filter_constraints.py` → **3 passed**, `make dispatcher_gemm_lib` builds `libdispatcher_gemm_lib.so`, and `test_library_caching.py` → **5 passed**. ## Test plan - [x] `arch_filter.py` parses cleanly - [x] `tests/test_arch_filter_constraints.py` passes (3/3) - [x] `make dispatcher_gemm_lib` builds the default `.so` - [x] `tests/test_library_caching.py` passes (5/5) - [ ] CI: `dispatcher_test_arch_filter_constraints` + `test_library_caching.py` pass Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com>
This commit is contained in:
committed by
assistant-librarian[bot]
parent
e44760f220
commit
c401998b69
@@ -88,6 +88,13 @@ OPERATOR_TILE_CONSTRAINTS = {
|
||||
"tile_k_alignment": 8,
|
||||
},
|
||||
OperatorType.GEMM_GROUPED: {
|
||||
"min_tile_m": 16,
|
||||
"min_tile_n": 16,
|
||||
"min_tile_k": 8,
|
||||
"tile_m_alignment": 16,
|
||||
"tile_n_alignment": 16,
|
||||
"tile_k_alignment": 8,
|
||||
},
|
||||
# NOTE: these are copied from plain GEMM and only gate tile *shape* validity.
|
||||
# They do NOT express Stream-K's real feasibility requirement -- that a problem
|
||||
# has enough output tiles to partition K-work across the CUs. That gate is
|
||||
|
||||
@@ -27,6 +27,19 @@ set_tests_properties(dispatcher_test_autocorrect PROPERTIES
|
||||
ENVIRONMENT "PYTHONPATH=${CMAKE_CURRENT_SOURCE_DIR}/../python:${CMAKE_CURRENT_SOURCE_DIR}/../codegen:${CMAKE_CURRENT_SOURCE_DIR}/../scripts"
|
||||
)
|
||||
|
||||
# Operator tile-constraints table regression test
|
||||
add_test(
|
||||
NAME dispatcher_test_arch_filter_constraints
|
||||
COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/test_arch_filter_constraints.py
|
||||
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/..
|
||||
)
|
||||
|
||||
set_tests_properties(dispatcher_test_arch_filter_constraints PROPERTIES
|
||||
LABELS "dispatcher;python;validation"
|
||||
TIMEOUT 60
|
||||
ENVIRONMENT "PYTHONPATH=${CMAKE_CURRENT_SOURCE_DIR}/../python:${CMAKE_CURRENT_SOURCE_DIR}/../codegen:${CMAKE_CURRENT_SOURCE_DIR}/../scripts"
|
||||
)
|
||||
|
||||
# Verbose version of the test
|
||||
add_test(
|
||||
NAME dispatcher_test_autocorrect_verbose
|
||||
|
||||
80
dispatcher/tests/test_arch_filter_constraints.py
Normal file
80
dispatcher/tests/test_arch_filter_constraints.py
Normal file
@@ -0,0 +1,80 @@
|
||||
#!/usr/bin/env python3
|
||||
|
||||
# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
||||
# SPDX-License-Identifier: MIT
|
||||
|
||||
"""
|
||||
Regression tests for OPERATOR_TILE_CONSTRAINTS in arch_filter.py.
|
||||
|
||||
Guards against a malformed constraints table (e.g. an operator entry left
|
||||
without a body/closing brace), which raises a SyntaxError on import and
|
||||
breaks all GEMM codegen -- and therefore every downstream build/test that
|
||||
imports unified_gemm_codegen.
|
||||
|
||||
Can be run as:
|
||||
python3 tests/test_arch_filter_constraints.py
|
||||
ctest -R test_arch_filter_constraints
|
||||
"""
|
||||
|
||||
import sys
|
||||
import unittest
|
||||
from pathlib import Path
|
||||
|
||||
# Setup paths
|
||||
SCRIPT_DIR = Path(__file__).parent.resolve()
|
||||
DISPATCHER_DIR = SCRIPT_DIR.parent
|
||||
sys.path.insert(0, str(DISPATCHER_DIR / "codegen"))
|
||||
|
||||
# Importing at all fails if arch_filter.py has a syntax error.
|
||||
from arch_filter import OPERATOR_TILE_CONSTRAINTS, OperatorType # noqa: E402
|
||||
|
||||
REQUIRED_KEYS = {
|
||||
"min_tile_m",
|
||||
"min_tile_n",
|
||||
"min_tile_k",
|
||||
"tile_m_alignment",
|
||||
"tile_n_alignment",
|
||||
"tile_k_alignment",
|
||||
}
|
||||
|
||||
|
||||
class TestOperatorTileConstraints(unittest.TestCase):
|
||||
"""Validate the OPERATOR_TILE_CONSTRAINTS table is well-formed."""
|
||||
|
||||
def test_every_operator_has_an_entry(self):
|
||||
"""Every OperatorType must have a constraints entry."""
|
||||
for op in OperatorType:
|
||||
self.assertIn(
|
||||
op,
|
||||
OPERATOR_TILE_CONSTRAINTS,
|
||||
f"{op} is missing from OPERATOR_TILE_CONSTRAINTS",
|
||||
)
|
||||
|
||||
def test_every_entry_is_a_well_formed_dict(self):
|
||||
"""Each entry must be a dict with the required integer keys."""
|
||||
for op, constraints in OPERATOR_TILE_CONSTRAINTS.items():
|
||||
self.assertIsInstance(
|
||||
constraints, dict, f"{op} constraints should be a dict"
|
||||
)
|
||||
self.assertEqual(
|
||||
REQUIRED_KEYS,
|
||||
set(constraints.keys()),
|
||||
f"{op} constraints keys mismatch",
|
||||
)
|
||||
for key, value in constraints.items():
|
||||
self.assertIsInstance(
|
||||
value, int, f"{op}[{key}] should be an int, got {type(value)}"
|
||||
)
|
||||
self.assertGreater(value, 0, f"{op}[{key}] should be positive")
|
||||
|
||||
def test_gemm_grouped_entry_present(self):
|
||||
"""Regression: GEMM_GROUPED must have its own distinct constraints dict."""
|
||||
self.assertIn(OperatorType.GEMM_GROUPED, OPERATOR_TILE_CONSTRAINTS)
|
||||
grouped = OPERATOR_TILE_CONSTRAINTS[OperatorType.GEMM_GROUPED]
|
||||
self.assertEqual(REQUIRED_KEYS, set(grouped.keys()))
|
||||
# GEMM_GROUPED and GEMM_STREAMK are separate entries, not a merged one.
|
||||
self.assertIn(OperatorType.GEMM_STREAMK, OPERATOR_TILE_CONSTRAINTS)
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
unittest.main(verbosity=2)
|
||||
Reference in New Issue
Block a user