Repository navigation
hipblaslt: add support for gfx1250-strict build - #11777
jichangjichang merged 12 commits into
Conversation
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
b6cf90f to
e2a8f1f
Compare
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## develop #11777 +/- ##
===========================================
+ Coverage 70.66% 70.67% +0.01%
===========================================
Files 2939 2939
Lines 483878 483977 +99
Branches 69775 69781 +6
===========================================
+ Hits 341889 342014 +125
+ Misses 118127 118106 -21
+ Partials 23862 23857 -5
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
pmoutsias-amd
left a comment
There was a problem hiding this comment.
I checked the edit in the readme and it is a straight removal. LGTM.
bnemanich
left a comment
There was a problem hiding this comment.
mxdatagenerator change looks fine. Will need someone else to look over the rest of PR.
|
FFM gfx1250 failed is caused by another merged PR. |
hipblaslt: add support for gfx1250-strict build MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit # gfx1250-strict as an architecture of its own, plus a gfx1250v0 subtree for v0 parts that report gfx1250 Follow-up to [#10784](ROCm/rocm-libraries#10784), which introduced `gfx1250-strict` as an ordinary architecture name. This PR removes the special cases #10784 left in CMake. It also retires develop's `gfx1250v0` ASIC-revision workaround in favour of the stepping's own name, and keeps v0 silicon correct while ROCr still reports it as `gfx1250` by default. ## Motivation gfx1250 ships as two silicon steppings. They share the ISA tuple `(12, 5, 0)` but are separate compiler targets (ELF `EF_AMDGPU_MACH` `0x49` for gfx1250/v1, `0xEB` for gfx1250-strict/v0). A code object built for one is rejected on the other. On ROCm 10.2, what a v0 part reports depends on `HSA_DISABLE_GFX12_STRICT`: | `HSA_DISABLE_GFX12_STRICT` | `gcnArchName` on v0 | code objects that load | `asicRevision` | | --- | --- | --- | --- | | `0` | `gfx1250-strict` | `0xEB` only | 0 | | `1` or unset | `gfx1250` | `0x49` only | 0 | In the second row, v0 would otherwise load v1 kernels. Those use instructions v0 does not have, such as `v_wmma_scale_f32_32x16x128_f4`, and TDM multicast. Measured on v0: an FP4 v1 kernel returns `norm_error=nan`. That row is ROCr's default today, so it has to stay correct until the default flips. Tensile keys its capability map by ISA, so one `TensileCreateLibrary` run cannot describe both steppings. #10784 worked around that in CMake: a second `hipblaslt_create_device_library` call for `gfx1250-strict`, plus an extra install rule driven by a global property. This PR moves the split into `TensileCreateLibrary`, where the knowledge of which names collide already lives. CMake goes back to a single call with `${GPU_TARGETS}`. The ISA-keyed machinery itself is deliberately left alone. Re-keying it by name would touch ~88 `isaInfoMap[...]` subscripts across 24 files, plus two rocisa C++ sites and the nanobind ABI, on every architecture's code path. ## What gets built and loaded | Subtree | Contents | Compiler target / ELF | File-name token | Built by | | --- | --- | --- | --- | --- | | `library/gfx1250/` | v1 tuning | `gfx1250` / `0x49` | `gfx1250` | `-a gfx1250` | | `library/gfx1250v0/` | v0: `gfx1250-strict` logic and capability overrides | `gfx1250` / `0x49` | `gfx1250` | `-a gfx1250` (default, `HIPBLASLT_BUILD_GFX1250V0=ON`) | | `library/gfx1250-strict/` | v0 | `gfx1250-strict` / `0xEB` | `gfx1250-strict` | `-a gfx1250-strict` | The runtime picks the Tensile GEMM subtree from the reported name first, then the revision. This applies to both hipBLASLt and hipSPARSELt, when `*_TENSILE_LIBPATH` is unset: | Reported `gcnArchName` | `asicRevision` | Subtree | | --- | --- | --- | | `gfx1250-strict` | not read | `gfx1250-strict/` (unchanged) | | `gfx1250` | `0` (v0) | `gfx1250v0/`. If it is missing, **hard error**: no fallback to `gfx1250/`, whose kernels are wrong on v0 | | `gfx1250` | non-zero (v1) | `gfx1250/` (unchanged) | | anything else | not read | `<name>/` (unchanged) | ## Technical Details ### Fan-out `isaCollisionFreeGroups()` partitions the requested names so that no group holds two names spelling one ISA. A run handed more than one group re-executes itself once per group as a subprocess. They are separate processes rather than a loop, because a run settles arch-dependent state process-wide. A single group takes the old path untouched: no subprocess, and no change to parallelism, scratch naming or cleanup. That covers every build that names no gfx1250 variant. Jobs are split by architecture count rather than evenly. ### Scratch directories Kernel basenames come from the ISA, so both steppings name their `.s` and `.o` files identically. The fixed-name artifacts (`Kernels.cpp`, `Kernels.h`, the six static headers, `code_object_tmp/Kernels.o`) carry no architecture at all. Scratch is therefore `build_tmp/<STEM>-<stepping>` for a run asked for a stepping, and stays `build_tmp/<STEM>` otherwise. Fan-out children carry `TENSILE_GROUP_BUILD`, so each reclaims only its own subdirectory. ### Benchmark cache key (correctness fix) The tuning cache key gains the compiler targets; otherwise the second stepping would load the first stepping's code objects. `_loadLegacyCacheIfMatches` is removed, because it is keyed by directory alone and would bypass the new field. ### Architecture names - Tuning is selected by `ArchitectureName`, which a stepping spells as itself. `archMatch()` compares whole names, so `gfx1250` cannot claim a `gfx1250-strict` request, and it strips both qualifiers and bracketed predicates, so `gfx950[cu=64]` still matches its own logic. - `steppingArchOf()` derives a stepping rather than looking it up in a table: a name is a stepping exactly when it does not survive a round trip through its ISA. A future stepping needs no registration. - Detection keeps the stepping. The answer from `amdgpu-arch`, or from the enumerator, goes through `restore_steppings`, so a bare `Tensile` run on v0 builds gfx1250-strict rather than base gfx1250. Colon target features are stripped for `--gpu-targets`; the hyphenated stepping is kept. - A stepping inherits no marks from its base. Configs that must stay off gfx1250 silicon list both `skip-gfx1250` and `skip-gfx1250-strict`. ### gfx1250v0: v0 kernels for a runtime that reports gfx1250 `gfx1250v0` is a build-only alias. It is never in `GPU_TARGETS` (CMake's supported list rejects it), never a runtime name, and excluded from `all` and `archMacroNames`. Naming it directly is rejected, both by `TensileCreateLibrary --architecture` (before any code generation) and by `Tensile --gpu-targets` / a config's `Architecture`. A single declarative table, `ARCH_BUILD_ALIASES` in `Common/Architectures.py`, defines it: | Field | Value | Used by | | --- | --- | --- | | `tuningArch` | `gfx1250-strict` | Logic selection (`archMatch`), `ARCH_CAP_OVERRIDES` in `Capabilities.applyArchCapOverrides`, `StinkyTofuArchName` | | `target` | `gfx1250` | Assembler `-mcpu` and `.amdgcn_target`, bundle target id, helper `--offload-arch`, the declared `ArchitectureName` after parsing (`LibraryIO` `archRenames`), master/Mapping/`.co` names | | `deviceTargetFeatures` | `+needs-aligned-2addr-lds`, `-wmma-f4-insts`, `-block16-cvt-scale-insts` | The assembler (`-Xclangas`) and the HIP helper compile (`-Xclang`) | - **Codegen divergence.** gfx1250v0 goes through exactly gfx1250-strict's Tensile codegen: same logic, same capability overrides (no TDM multicast, no fp4 32x16 WMMA), same `StinkyTofuArchName`. Only the compiler target, file names and output directory differ. The added target features make the assembler reject the instructions v0 lacks, as it does for gfx1250-strict, and give the helper kernels strict's ISA. - **Expansion.** With `--gfx1250v0`, `withArchBuildAliases()` appends `gfx1250v0` when bare `gfx1250` is requested, including through `all`. It runs before `isaCollisionFreeGroups` and never in a fan-out child, so `-a "gfx1250;gfx1250-strict"` builds three groups. `gfx1250-strict` alone does not expand. If no gfx1250-strict logic matches (a filtered or custom logic build), the alias is skipped with a warning rather than producing an empty subtree. - **Knobs.** `TensileCreateLibrary --gfx1250v0` is opt-in, so direct callers (tox, the benchmark client, GenerateSummations) are unchanged. The CMake option `HIPBLASLT_BUILD_GFX1250V0` (default `ON`, in `hipblaslt_codegen.cmake`, so hipSPARSELt shares it) passes the flag, and adds `gfx1250-strict` to the `TensileLogic --check-all` architectures of a gfx1250 build. Turning v0 off for gfx1250 once TheRock builds gfx1250-strict is `-DHIPBLASLT_BUILD_GFX1250V0=OFF`. - **Output.** Files carry the `gfx1250` token and land in `library/gfx1250v0/`. It holds no ExtOp or Transform files; those load from `library/<gcnArchName>/` as before. `validate_library_layout.py` gains `REVISION_SUBTREES = {"gfx1250v0": "gfx1250"}`. ### Runtime selection - **hipBLASLt.** `tensile_host.cpp` reads `asicRevision` only for a device reporting `gfx1250`, and asks the pure helper `rocblaslt_revisioned_arch_name(processor, asicRevision)` for the subdirectory (restored from develop as `rocblaslt_arch_revision.hpp`). Master file names, `codeObjectTargets` and lazy loading keep using the reported name, so files named `gfx1250` inside `gfx1250v0/` resolve unchanged. - A missing v0 library (lazy or non-lazy master) prints an error naming the directory and the two ways out. Loading then fails through the existing error paths, and API calls return an error status. - `row_uniformity_gtest`'s library helper uses the same helper. - **hipSPARSELt.** The same rule and the same hard error in its `tensile_host.cpp` (`asicRevision` read only for `gfx1250`, guarded by `HIP_VERSION >= 307` as in its `handle.cpp`). ### Removed knobs `HIPBLASLT_ASIC_REVISION`, `--asic-revision`, `gpu_revision_probe`, `GpuRevisionTarget.py`, and the `gfx1250v0` overlay checks (`find_gfx1250v0_overlay_violations`, `--require-gfx1250v0-overlay`) are removed. The stepping is named in `GPU_TARGETS`, and the gfx1250 v0 subtree is controlled by `HIPBLASLT_BUILD_GFX1250V0`. ### Tuning logic `Logic/asm_full/gfx1250v0/` moves to `gfx1250-strict/` (226 files; directory, filename prefix, `ScheduleName` and `ArchitectureName` renamed, tuning values unchanged). hipSPARSELt gains its own `gfx1250-strict` spmm tree (64 files). ### Guards, packaging, reproducibility - **Source guards.** Compiling for gfx1250-strict predefines `__gfx1250_strict__` and not `__gfx1250__`. Eight guards (`memory_gfx.h`, the float4/bfloat6/ float6/e5m3 headers, `DataTypes_E5M3.hpp`, mxDataGenerator) accept both, and the extops/matrix-transform CMake arch regexes keep `-strict`. - **Assembler retarget.** `_retargetAssemblySource` matches the whole hyphenated processor name, so `gfx1250-strict` is not rewritten to `gfx1250-strict-strict`. - **Packaging checks.** `validate_library_layout.py` and `check_dat_integrity.py` accept hyphens in an architecture token. - **Reproducibility.** Logic files are merged in sorted order. - **Test isolation.** `test_MatrixInstructionConversion` updated the shared `defaultSolution` in place, which leaked into whichever test xdist ran next on that worker; it now works on a copy. ## Test Plan The full procedure is in `projects/hipblaslt/docs/gfx1250-strict-testing.md` (Suite 6 covers gfx1250v0). It runs on a real v0 machine with the newest ROCm nightly, through `pipeline-v0.sh`. 1. **TensileLite unit tests.** `pytest -m unit` over `Tests/unit` and `Tests/extras`. 2. **Architecture-name parsing** (`test_arch_steppings.py`). Every `architectureMap` name in every CMake spelling (bare, `:xnack±`, `[cu=64]`, `[id=74a0]`) must match its own logic header and no other. Also covers the gfx1250v0 alias table, expansion (with / without `--gfx1250v0`, fan-out child), rejection of the alias as a requested name, the three-group partition, caps equality with strict end to end through `KernelWriter`, compiler target and features, logic selection and keying, and the no-strict-logic skip. 3. **Parser diff against the base commit** for 121 spellings. 4. **Packaging script tests** (`tools/scripts/tests/`), including the v0 subtree. 5. **pre-commit.** Both the hipBLASLt config and the repo root config. 6. **stinkytofu `ctest`.** 7. **Build** `./install.sh -c -a "gfx1250;gfx1250-strict"`, then `verify_v0.py` on the install tree: - ELF per subtree (`0x49` / `0x49` / `0xEB`); - file-name tokens; - no fp4 32x16 WMMA, `v_cvt_scale_pk*`, multicast or helper `ds_*_2addr` in either v0 tree, while `gfx1250/` does have them, so the scan is not vacuous; - the kernel set of `gfx1250v0/` equals that of `gfx1250-strict/`. Plus `validate_library_layout.py`. 8. **hipblaslt-test, all and `*quick*`**, in both reporting modes. Each run checks that `ldd` resolves to the built `libhipblaslt.so`, and that the log shows the subtree loaded: `HSA_DISABLE_GFX12_STRICT=0` → `gfx1250-strict/`, `=1` → `gfx1250v0/`, unset → `gfx1250v0/`. 9. **FP4 regression.** `hipblaslt-bench` FP4 TN 256x256x512 with MX scales (`--requested_solution -1 -v`) and the `smoke_matmul_gemm_tn_mxf4` test, in both modes: every solution must pass. Negative control: forcing `HIPBLASLT_TENSILE_LIBPATH` at `library/gfx1250/` on v0 must bring the NaN back. 10. **No fallback.** In a copy of the install without `gfx1250v0/`, a v0 reporting gfx1250 must fail with the error, and no test may pass. 11. **tox config tests** (`Tests/common`), in strict mode. 12. **hipSPARSELt.** Build `-a "gfx1250;gfx1250-strict"`, run `verify_v0.py` on its library tree, check no-fallback, confirm with `strace` which subtree each mode opens, then run the full `hipsparselt-test` in both modes (one GPU each), with an `ldd` check. ## Test Result ### Committed series (without gfx1250v0) Measured on the series at `bdb85d2f61` (develop `670b375ba0`) with the fixes since committed as `8f11772cbb`. ROCm 10.2.0a20260926, `HSA_DISABLE_GFX12_STRICT=0`, 4× gfx1250 v0. | Check | Result | | --- | --- | | TensileLite unit (`-m unit`) | 9083 passed, 0 failed | | `test_arch_steppings.py` | 787 passed | | Packaging script tests | 15 passed | | pre-commit, hipBLASLt config (full unit) | 9084 passed, both hooks Passed | | pre-commit, repo root config | Passed | | stinkytofu `ctest` | 1573 / 1573 | | Build, both architectures | rc=0, 400 files each | | ELF machine flag | `0x49` (gfx1250), `0xEB` (gfx1250-strict) | | `validate_library_layout.py` | OK | | hipblaslt-test, all (gfx1250-strict) | **67515 passed, 0 failed**, 17 skipped | | hipblaslt-test `*quick*` (gfx1250-strict) | **18043 passed, 0 failed** | | tox `Tests/common` | **8783 passed, 0 failed** | | hipsparselt-test (gfx1250-strict) | **165405 passed, 0 failed** | ### With gfx1250v0 This change as staged on `8f11772cbb` (develop `d0e4bc5f98`), since committed as `1c366475cc`. The commit is identical file-for-file except `spmm_f16_sia4.yaml`, which is re-validated on its own below. ROCm 10.2.0a20260929, 4× gfx1250 v0, every GPU step on GPUs 2 and 3. **Tensile and build** | Check | Result | | --- | --- | | TensileLite unit (`-m unit -n 8`) | 9348 passed, 0 failed | | `test_arch_steppings.py` | 828 passed | | Packaging script tests | 18 passed | | pre-commit, hipBLASLt config (staged diff) | 9348 passed, both hooks Passed | | pre-commit, repo root config | Passed | | stinkytofu `ctest` | 1579 / 1579 | | Build `-a "gfx1250;gfx1250-strict"` | rc=0, 3 groups `[gfx1250]`, `[gfx1250-strict]`, `[gfx1250v0]` | | Files gfx1250 / v0 / strict | 400 / 397 / 400 (v0 has no ExtOp / Transform) | | ELF gfx1250 / v0 / strict | `0x49` / `0x49` / `0xEB` | | fp4 32x16 WMMA, gfx1250 / v0 / strict | 896 / 0 / 0 | | multicast setup, gfx1250 / v0 / strict | 45 / 0 / 0 | | helper `ds_*_2addr`, gfx1250 / v0 / strict | 78 / 0 / 0 | | Kernel set v0 vs strict | 2840 = 2840, none unique to either | | v0 vs strict code (smoke, 7 logic pairs) | identical apart from `.amdgcn_target` | | v0 assembler (`-mcpu=gfx1250` + v0 features) | rejects `v_wmma_f32_32x16x128_f4` and `v_cvt_scale_pk8_f32_fp4`, as `-mcpu=gfx1250-strict` does; plain gfx1250 accepts both | | `validate_library_layout.py` | OK | **hipBLASLt runtime** | Check | Result | | --- | --- | | `ldd hipblaslt-test` | the built `libhipblaslt.so.1` and `libtensilelite-host.so.1` | | Subtree loaded, `=0` / `=1` / unset (log) | `gfx1250-strict/` / `gfx1250v0/` / `gfx1250v0/` | | Subtree opened, `=0` / `=1` (`strace -e openat`) | only `gfx1250-strict/` / only `gfx1250v0/` | | hipblaslt-test all, `=0` (gfx1250-strict) | **67545 passed, 0 failed**, 17 skipped | | hipblaslt-test `*quick*`, `=0` | **18043 passed, 0 failed** | | hipblaslt-test all, `=1` (gfx1250v0) | **67540 passed, 1 failed**, 21 skipped (`nightly_matmul_deepbench_bf16_..._NN_1152_128_7000_..._3_t2`, intermittent; 20 / 20 reruns passed, see Known) | | hipblaslt-test `*quick*`, `=1` | **18043 passed, 0 failed** | | FP4 bench, `=0` / `=1` | 345 / 345 and 345 / 345, no nan; smoke passed in both | | FP4 negative control (v1 tree forced, `=1`) | 1 solution wrong (nan), as expected | | No fallback (v0 removed, `=1`) | error naming `library/gfx1250v0`; 0 passed, 128 failed; exit 1 | **TensileLite config tests and hipSPARSELt** | Check | Result | | --- | --- | | tox `Tests/common` (`=0`) | **9050 passed, 0 failed**, 1439 skipped, 2 xfailed | | `spmm_f16_sia4.yaml` as committed (`=0`) | 100 / 100 solutions passed, including the first group's DepthU 128 / PLR 1 / SIA 3 | | hipSPARSELt build `-a "gfx1250;gfx1250-strict"` | rc=0; 131 files in each of `gfx1250/`, `gfx1250v0/`, `gfx1250-strict/` | | hipSPARSELt ELF gfx1250 / v0 / strict | `0x49` / `0x49` / `0xEB`; fp4 32x16 WMMA and multicast 0 in all three | | hipSPARSELt kernel set v0 vs strict | 64 = 64, none unique to either | | hipSPARSELt subtree opened, `=0` / `=1` (`strace`) | only `gfx1250-strict/` / only `gfx1250v0/` | | hipSPARSELt no fallback (v0 removed, `=1`) | error naming `library/gfx1250v0`; every spmm test fails (status 10, no solutions); prune/compress, which do not use the GEMM library, pass (2448); run capped at 300 s | | hipsparselt-test, `=0` (GPU 2, gfx1250-strict) | **165405 passed, 0 failed** | | hipsparselt-test, `=1` (GPU 3, gfx1250 → gfx1250v0) | **165405 passed, 0 failed** | Before this change, the FP4 bench on v0 in `=1` mode loaded `library/gfx1250/` and gave 344 / 345 passes, with the MI32x16 kernel returning `norm_error=nan`. ## Deliberate behaviour changes 1. **gfx1250 builds produce a third subtree.** Any CMake build naming gfx1250 (including `all`) builds `library/gfx1250v0/` in an extra fan-out group, which roughly doubles gfx1250's Tensile build time and library size. `-DHIPBLASLT_BUILD_GFX1250V0=OFF` removes it. 2. **v0 reporting gfx1250 now loads `gfx1250v0/`, or fails.** Previously it silently ran v1 kernels. A package without the v0 subtree now fails on v0 with `HSA_DISABLE_GFX12_STRICT` unset or `1`; use `=0` with a gfx1250-strict build instead. 3. **Benchmark cache keys change for every architecture** (the architecture is folded into the hash), so existing tuning output directories miss once. 4. **Shipped `.dat` bytes change for every architecture** (deterministic logic merge order changes solution indices). 5. **The linker response file** moves to `<destPath>.linker_args`. 6. **`Kernels.cpp`, `Kernels.h` and the static headers** are written under the run's scratch directory rather than the output root. 7. **`_loadLegacyCacheIfMatches` is removed.** ## Known, not addressed - **Mixed v0 and v1 in one process.** The library is loaded once, from the first device to initialize. - **Other kernels on v0 in non-strict mode.** ExtOps, matrix-transform and the HIP kernels compiled into libhipblaslt / libhipsparselt load their gfx1250 (v1) builds, as on develop. hipblaslt-test and hipsparselt-test in `=1` mode exercise them. - **v1's `asicRevision`.** It is assumed non-zero; there was no v1 machine to confirm it. - **`*_TENSILE_LIBPATH`** is used verbatim. Pointing it at `library/gfx1250/` on v0 still loads v1 kernels, and a mismatched install is not validated. - **Predicated gfx1250 targets.** A `gfx1250[cu=…]` entry in `GPU_TARGETS` builds gfx1250v0 without its predicate, and does not add gfx1250-strict to the `TensileLogic --check-all` list. - **Helper-kernel cache.** The key is the arch name, not the device target features. Clear `~/.tensile/helper_cache` if those change. - **Stale v0 subtree.** Turning `HIPBLASLT_BUILD_GFX1250V0` off in an existing build directory leaves the old `library/gfx1250v0/` to be installed. - **Workflow.** `.github/workflows/gfx1250_hipblaslt_tensilelite_test.yml` still sets `TL_TARGET=gfx1250v0`, a name Tensile now rejects as a target; it needs porting to gfx1250-strict. - **Intermittent bf16 result on v0.** `nightly_matmul_deepbench_bf16_..._NN_1152_128_7000_..._3` returned a wrong value once in a full strict-mode run before this change (40 / 40 reruns passed). Its `_t2` variant did the same once in a full gfx1250v0-mode run after it (20 / 20 reruns passed). - **Sparse codegen bug.** On gfx1250-strict, `PrefetchLocalRead=1` with `DepthU=128`, `ScheduleIterAlg=4` and buffer_load sparse A returns wrong values. `spmm_f16_sia4.yaml`'s first group now uses `ScheduleIterAlg=3`, keeping DepthU 128 and PLR 1, which passes. The SIA 4 combination and the base `gfx1250` twin config are unverified. - **`isaInfoMap` annotation.** It is `Dict[str, IsaInfo]` while production keys it by `IsaVersion`. - **`Makefile:55`** gates `TRUE16_FEATURE` on an unset `TARGET_ARCH`, so `make co` never passes `+real-true16`. This predates the branch. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. JIRA ID : AIHPBLAS-4727 JIRA ID : AIHPBLAS-4728 JIRA ID : AIHPBLAS-4729 JIRA ID : AIHPBLAS-4730 JIRA ID : AIHPBLAS-4768 JIRA ID : AIHPSPAR-275 JIRA ID : ROCM-31226 JIRA ID : ROCM-30511
…20260901 Resolves a conflict in TensileLibLogicToYaml.py's import block: develop's #12148 rewrote TensileLite's internal imports as relative ones, and this branch had added imports to the same block. Keep this branch's import list in develop's relative form, and convert the branch's remaining absolute self-imports as well -- the three function-local "from Tensile import ClientWriter" and the import_module() of Tensile.SolutionStructs.Problem -- so the module no longer names its own package. Also adapts to develop's #11777, which merged without a textual conflict but broke the branch: it removed Common.Architectures.ARCH_COMPILER_TARGET, which the converter imported, and replaced the gfx1250v0 revision (identified by ScheduleName) with the gfx1250-strict stepping, whose logic records it as ArchitectureName, plus a gfx1250v0 build alias. The build target is now the logic's ArchitectureName, with an alias standing for its tuningArchOf(), and Architecture is emitted when steppingArchOf() says the target is a stepping -- the one thing GlobalParameters.ISA cannot express, which is how Tensile.py reads it. Update the build-target tests and ADR 0031 to match. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
TheRock Submodule Bump ActivityNewest first
|
test(tensilelite): normalize StinkyTofu label hashes before canonicalize_asm InsertClusterBarrierPass names its skip labels with a random 16-char [0-9a-zA-Z] hash. canonicalize_asm only renames suffixes that are all [A-Z0-9], so about one hash in 6000 took that path in one emit while its counterpart in the other emit was renamed by _stableLabels instead, making test_gfx1250v0_emits_what_gfx1250_strict_emits fail roughly once in 750 runs. Run _stableLabels on the raw source inside _emit, ahead of canonicalize_asm, so every hash is renamed by the same pass, and compare the emitted sources directly. ## Motivation `Tensile/Tests/unit/test_arch_steppings.py::test_gfx1250v0_emits_what_gfx1250_strict_emits`, added in #11777, fails intermittently in unrelated PRs' CI (for example `s_preliminary_hipblaslt_PR-12538`: 1 failed / 9822 passed). The emitted gfx1250v0 and gfx1250-strict kernels are actually identical; the failure comes from how the test normalizes random label names before comparing them. This PR removes the flake without weakening the assertion. ## Technical Details **Where the randomness comes from.** #9853 made rocisa's `LabelManager` deterministic (`prefix_<counter>`), but StinkyTofu's `InsertClusterBarrierPass::makeRandomHash()` (`shared/stinkytofu/src/transforms/asm/InsertClusterBarrierPass.cpp`) still suffixes `label_skipCBPreSignal_*` / `label_skipCBPreSignal_LCL_*` with a 16-char hash drawn from `[0-9a-zA-Z]`, seeded from `std::random_device`. This kernel (`MULTICAST_MI`, `ClusterDim=[2,1]`) carries 4 such labels per emit. **Why two normalization passes disagree.** The test normalized each emit twice, with different rules: | Pass | Matches | Renames to | |---|---|---| | `canonicalize_asm` (inside `_emit`) | `_[A-Z0-9]{16}\b` | `_LBL<n>` | | `_stableLabels` (in the test, afterwards) | `_[A-Za-z0-9]{16}(?![A-Za-z0-9])` | `_L<n>` | A hash is all `[A-Z0-9]` with probability (36/62)^16 ≈ 1/6000. When that happens in only one of the two emits, that label becomes `_LBL0` on one side and `_L3` on the other, and every later `_L<n>` shifts by one, so the comparison fails. With 8 hashes per test run, the failure rate is about 1 in 750. Fixing `_stableLabels` alone is not enough: one label is spelled `label_InitCIterWmma_label_skipCBPreSignal_<hash>_0`, where `\b` does not fire before `_`, so `canonicalize_asm` would rename an all-caps hash in some occurrences and not others. Normalization has to happen on the raw source. **Change.** - Move `_stableLabels` to module scope. - Apply it to the raw source in `_emit`, before `canonicalize_asm`. After that, `canonicalize_asm`'s label regex has nothing left to match, and it only strips `matrix_[ab]_reuse`. - `test_gfx1250v0_emits_what_gfx1250_strict_emits` now asserts `v0Src == strictSrc` directly. `test_gfx1250_emits_multicast_gfx1250_strict_does_not`, the only other `_emit` caller, is unaffected: its `!=` check holds because the multicast setup genuinely differs. The root fix is to make StinkyTofu's label naming deterministic (as #9853 did for rocisa); that is left to the StinkyTofu owners. ## Test Plan Run in a gfx1250 FFM container (ROCm 10.2 nightly), with rocisa built from this branch's tree: 1. Full file: `pytest Tensile/Tests/unit/test_arch_steppings.py`. 2. Fault injection on real emitted kernels: take raw strict and v0 sources, replace each of v0's 4 StinkyTofu hashes in turn with an all-`[A-Z0-9]` hash, and compare under the old and new normalization order. 3. Determinism check: emit 12 strict/v0 pairs (24 emits) and compare them after label normalization, against each other and across pairs, to confirm the hash is the only run-to-run variation. ## Test Result 1. `828 passed`. 2. Old order: all 4 injections mismatch (reproducing the CI failure signature `_L3` vs `_LBL0` with shifted numbering). New order: all 4 compare equal, including the hash that also appears with a `_0` suffix. 3. All 24 emits are identical after normalization; no other nondeterminism found. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. Co-authored-by: Cursor <cursoragent@cursor.com>
Follow-up to [#11777](#11777) ("hipblaslt: add support for gfx1250-strict build"). ## Motivation #11777 made eight source guards that tested only `__gfx1250__` accept `__gfx1250_strict__` too, because compiling for gfx1250-strict predefines only the latter. Two of those guards enable the hardware e5m3 conversion. That was valid with the compiler #11777 was tested on (10.2.0a20260930), but the compiler has since changed. The hipBLASLt client build for `gfx1250-strict` fails on ROCm nightlies from 10.2.0a20261004: ``` error: '__builtin_amdgcn_cvt_pk_fp8_f32_e5m3' needs target feature fp8e5m3-insts ``` ROCm/llvm-project [`7667f52dd687`](ROCm/llvm-project@7667f52dd687) cherry-picks [llvm/llvm-project#227861](llvm/llvm-project#227861) ("Disable e5m3 conversions on gfx1250-strict", LCOMPILER-2845). It moves `FeatureFP8E5M3Insts` out of the common gfx1250 feature set, so `gfx1250-strict` (v0) no longer has the e5m3 conversion instructions. `hipblaslt_e5m3.h` and `DataTypes_E5M3.hpp` still enable the hardware fast path for both `__gfx1250__` and `__gfx1250_strict__`, as #11777 left them. This PR makes the strict build work again on current nightlies and keeps it working on older ones. ## Technical Details `HIP_E5M3_CVT_FAST_PATH` is now enabled for `__gfx1250__` only, in both copies of the header: - `projects/hipblaslt/library/include/hipblaslt/hipblaslt_e5m3.h` - `projects/hipblaslt/tensilelite/include/Tensile/DataTypes_E5M3.hpp` ```diff -#if (defined(__gfx1250__) || defined(__gfx1250_strict__)) && __HIP_DEVICE_COMPILE__ +// gfx1250-strict (v0) has no e5m3 conversion instructions; it takes the software path. +#if defined(__gfx1250__) && __HIP_DEVICE_COMPILE__ #define HIP_E5M3_CVT_FAST_PATH 1 ``` - **gfx1250-strict:** takes the existing software conversion. This is also correct on nightlies up to 20261003. Their compiler still accepted the e5m3 builtins on strict, but v0 has no such instructions. - **gfx1250 (v1):** unchanged. It keeps the hardware `v_cvt_pk_fp8_f32`, and the condition evaluates the same for it as before. - **The remaining guards from #11777** (`memory_gfx.h`, the float4/bfloat6/float6 headers, mxDataGenerator) are unchanged. They still accept both names, and the client and libraries build with them on every nightly from 20260930 to 20261006. - **Tensile kernels:** not affected. TensileLite emits no e5m3 conversion, and the strict library tree assembles on 20261004 to 20261006. - **Not covered:** the other changes in the same cherry-pick ([#227426](llvm/llvm-project#227426), [#227475](llvm/llvm-project#227475)) re-enable `v_cvt_scale_pk8/pk16` on strict in their block32 forms. hipBLASLt and hipSPARSELt do not use these instructions or the HIP FP4/FP6 headers, so they need no change here. - **Still not addressed, as listed in #11777:** on a v0 part with `HSA_DISABLE_GFX12_STRICT` unset or `1`, the device reports `gfx1250`. The HIP kernels compiled into libhipblaslt and the clients then load their gfx1250 builds, which keep the hardware e5m3 path. This change does not alter that. It covers only the `gfx1250-strict` builds. ## Test Plan - Compile the real client TU (`hipblaslt_init_device.cpp`, flags from `compile_commands.json`) for `gfx1250` and `gfx1250-strict`, with the original and the new headers, on every nightly from 20260930 to 20261006. - Clean builds of hipBLASLt and hipSPARSELt, clients included (`install.sh -c -a "gfx1250;gfx1250-strict"`), on every nightly from 20260930 to 20261006, then verify the library layout. - Full test pass on gfx1250 (v1), ROCm 10.2.0a20261005, following `gfx1250-strict-testing.md`: - hipblaslt-test (all, and `*quick*`) - FP4 bench and mxf4 smoke - hipsparselt-test - TensileLite unit, arch-stepping and packaging tests - tox `Tests/common` - pre-commit (hipBLASLt and repo root configs) ## Test Result - **Client TU compile:** - Original headers fail on `gfx1250-strict` for 20261004, 20261005 and 20261006. - With this change, both targets build on every nightly from 20260930 to 20261006. - **Clean builds, 20260930 to 20261006:** rc=0 on every nightly. - hipBLASLt: 400 / 397 / 400 files for `gfx1250` / `gfx1250v0` / `gfx1250-strict`, layout OK. - hipSPARSELt: 131 files in each subtree. - **gfx1250 (v1), ROCm 10.2.0a20261005:** | Suite | Result | |---|---| | hipblaslt-test all | 67569 passed, 0 failed, 34 skipped | | hipblaslt-test `*quick*` | 18043 passed | | FP4 bench / mxf4 smoke | 345 solutions, no nan / passed | | hipsparselt-test | 165405 / 165405 passed | | TensileLite unit, arch-stepping, packaging | passed | | tox `Tests/common` | 9642 passed, 1 failed | | pre-commit, both configs | passed | The one tox failure is `gsu_gfx1250` (BF16 GSU, size 512x512x1x32767, Relu). It is flaky and unrelated to e5m3. Of three reruns on an idle GPU, two passed and one failed, and the failing kernel differs between runs. The tests ran on develop `d116965bed9` plus this change. The branch is rebased onto develop `9db04770489`. The 3 commits in between touch only rocALUTION and rocSPARSE. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
Handle the architecture spelling introduced by ROCm#11777, including feature suffixes, without accepting misleading substrings.
Handle the architecture spelling introduced by ROCm#11777, including feature suffixes, without accepting misleading substrings.
Handle the architecture spelling introduced by ROCm#11777, including feature suffixes, without accepting misleading substrings.
JIRA ID : N/A ## Motivation - Run hipBLASLt + TensileLite on the gfx1250-strict path (#11777) for every PR. - Per-PR TheRock builds don't produce gfx1250-strict yet (ROCm/TheRock#8350, draft), so the job builds from source instead of waiting for artifacts. - Add Nightly schedule ## Technical Details - Toolchain: pinned ROCm SDK nightly (`rocm_sdk_version`, default `10.2.0a20261007`). A warn-only step flags when the SDK's compiler differs from the one at ci-env's `therock-ref`. - Checks that the runtime reports `gfx1250-strict` (`HSA_DISABLE_GFX12_STRICT=0`), then builds with `install.sh -a gfx1250-strict --skip_rocroller`. - One runner per GPU: the job container gets the runner's `ROCR_VISIBLE_DEVICES` (`-e ROCR_VISIBLE_DEVICES`) and fails early unless it names exactly one GPU. Checked on the host during a CI run: all 4 hipBLASLt shards ran on the runner's own GPU. - hipBLASLt: loads `library/gfx1250-strict/`; fails if `hipblaslt-test` doesn't link the library built here or its code objects aren't labelled `gfx1250-strict`. Runs as 4 gtest shards. - TensileLite: Uses 4 workers sharing the GPU. Unit tests dropped (host-only, already run per PR in TheRock CI). Selection `-m "common and not skip-gfx1250-strict"`; skipped if setup failed, still runs after a hipBLASLt failure. - Log scrub on build and test output: hardware details and wall-clock times are removed before anything reaches the console or the uploaded logs. - Per PR: `pull_request` trigger behind a GitHub-hosted gate (non-draft, ROCm org label, changes under `projects/hipblaslt`, its shared deps, this workflow or `ci-env/action.yml`). One run per PR; a newer push cancels it. Manual dispatch and the nightly are ungated. A manual run and the nightly on the same branch don't cancel each other. - Nightly: develop at 03:00 UTC (7 pm PST, 8 pm PDT). - Runner state: a fresh venv per job from the image's Python 3.12 (no `actions/setup-python`, so nothing from the runner's tool cache); the runner's pip cache is used only by the SDK install and is read-only except for the nightly; the job hands the root-owned workspace back to the runner's user at the end. - Includes #13464 (merged into this branch): `skip-gfx1250-strict` on 4 TensileLite configs and the bf8 MX gtest rows limited to gfx950, + fixed comment saying gfx1250 has no MX kernels with bf8 inputs. - Skips `gsu_gfx1250-strict.yaml` on gfx1250-strict for now (`skip-gfx1250-strict`): it intermittently returns wrong values from GSU multi-buffer kernels (see Test Result). Followed up with the config owner; the skip comes out when that's fixed. - Removed: the `workflow_run` trigger, artifact resolution, commit-status posting, unused `GFX1250_HW_XFAIL`. ## Test Plan - This PR's own `gfx1250-strict hipBLASLt + TensileLite` check on the gfx1250 runners (the PR changes the workflow, so it runs here). ## Test Result - CI on the gfx1250 runners ([run 38319193893](https://github.com/ROCm/rocm-libraries/actions/runs/38319193893)): hipBLASLt 68,138 tests over 4 shards, 0 failed; TensileLite common 109 passed, 2 xfailed, 0 failed. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level - Low. Self-contained workflow. PR code runs on a self-hosted runner only for ROCm org-labelled PRs 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com>
…trict gfx1250-strict/ sparse yamls were copied from gfx1250/ (#11777) and were identical to it apart from Architecture and the mirrored skip marks. The SIA3->SIA4+TDMInst=3 migration and the TransposeLDS/1LDSBuffer/LDSTrInst/ ExpandPointerSwap fixes in this PR only touched gfx1250/, leaving strict on the old buffer_load+SIA3 configs. - Apply the same yaml changes to the shared files in gfx1250-strict/. - Bring spmm_tdm_all.yaml in line with the gfx1250 one (TDMSplit, extra Exact sizes and the TDM3 + GSU>1 / Batch>1 block it was missing). - Remove spmm_f16_sia4.yaml here as well, as in gfx1250/. - Add spmm_i8is.yaml / spmm_i8is_sb.yaml for gfx1250-strict (the shipped gfx1250-strict I8IS logic yamls had no tuning source either), derived from the gfx1250 ones with only Architecture and the skip mark swapped. After this, every sparse yaml in gfx1250-strict/ differs from its gfx1250/ counterpart only in Architecture and the mirrored skip mark. JIRA: <https://amd-hub.atlassian.net/browse/AIHPSPAR-279> Co-authored-by: Cursor <cursoragent@cursor.com>
…A4 sparse kernels (#12407) JIRA ID - AIHPSPAR-279 [AIHPSPAR-279](https://amd-hub.atlassian.net/browse/AIHPSPAR-279) ## Motivation [PR #11758](#11758) started migrating the gfx1250 sparse (2:4) yaml regression suite from `ScheduleIterAlg=3` to `ScheduleIterAlg=4` + `TDMInst=3`, but the migration was incomplete/buggy. **This PR completes the migration** for both `gfx1250/` and `gfx1250-strict/`. Problems found: several int8/fp8 yamls still explicitly forked `TransposeLDS: [0, 1]`, and `TransposeLDS=0` (dual TileMajor, both operands transpose-read) combined with a multi-wave `MIWaveTile`/`MIWaveGroup` split hits a real bug in `resetTDMDescriptorForTailWaveSeparated()` (`KernelWriterAssembly.py`) that miscomputes the per-wave `compId` offset for 3 of 4 transpose orientations (only NT, where `TLUA=TLUB=True`, happens to compute correctly). This is not sparse-specific - a minimal dense repro hits the same bug - but sparse's forced `TransposeLDS=[0, 1]` fork was the only place exercising the broken `TransposeLDS=0` path for those orientations. Separately, `Solution.py`'s `ExpandPointerSwap` derivation had a `Sparse and PrefetchGlobalRead > 0` carve-out that force-kept `EPS=1`, which is incompatible with `LDSTrInst=True` (required whenever TDM+TileMajor is picked) and caused solver rejects/incomplete solution sets for several gfx1250 sparse configs. `gfx1250-strict/` sparse yamls (added by [#11777](#11777)) were copied from `gfx1250/` and differed from it only in `Architecture` and the mirrored skip mark, so they must receive the same migration. Two gfx1250 yamls added later (`spmm_i8is*.yaml`) and two existing ones (`spmm_tdm_iteratemode.yaml`, `spmm_tdm_split_xfail.yaml`) were also missing the `skip-gfx1250-strict` mark. 0924 update: if `DirectToVgprSparseMetadata`, still turn on `ExpandPointerSwap` ## Technical Details - **`Solution.py` (~line 2774-2786):** removed the `or (state["ProblemType"]["Sparse"] and state["PrefetchGlobalRead"] > 0)` clause from the `ExpandPointerSwap` derivation, so sparse now follows the exact same EPS logic as dense. Verified this does not regress the only Sparse+PGR>0 yamls without TDM/LDSTrInst (`bf16_activation.yaml`, `f16_activation.yaml`). - **12 int8/fp8 sparse gfx1250 yamls** (`spmm_b8*.yaml`, `spmm_f8*.yaml`, `spmm_i8*.yaml`): changed `TransposeLDS: [0, 1]` -> `TransposeLDS: [-1]`. Auto-select correctly picks `TransposeLDS=1` for NN/TN/TT orientations (avoiding the buggy dual-TileMajor `TransposeLDS=0` path) and falls back to `TransposeLDS=0` only for NT, the one orientation where the multi-wave descriptor reset happens to be correct. - Completed the SIA3->SIA4+`TDMInst=3` conversion across the remaining gfx1250 sparse yamls that PR #11758 had left inconsistent (`spmm_bf16*.yaml`, `spmm_f16*.yaml`, `spmm_tdm_*.yaml`, `spmm_fp16_ml1.yaml`, `f16_activation.yaml`/`bf16_activation.yaml`), setting `LDSTrInst`/`ExpandPointerSwap`/`1LDSBuffer` consistently per-orientation to match the now-corrected `Solution.py` derivation. - Removed `spmm_f16_sia4.yaml` (both `gfx1250/` and `gfx1250-strict/`): it duplicated `spmm_f16.yaml`, which now covers sparse + SIA4 + `TDMInst=3` itself. `test_EnableESM2TrackValuVsrc_derived.py`, which used it as its sparse SIA4 config, now uses `spmm_f16.yaml` (test passes). - Added `spmm_i8is.yaml` / `spmm_i8is_sb.yaml` (gfx1250, `Sparse=1` / `Sparse=2`): the first tuning/test configs for the I8IS ProblemType (`DataType: i8`, `DestDataType: i`). Until now only the 8 shipped I8IS Logic yamls existed (added by hand in `1a6bbcf62f7`, with no matching tuning yaml). The new files are derived from `spmm_i8hs*.yaml` with only `DestDataType: h` -> `i` changed, so they use the same TDM+SIA4 setup as the rest of the suite. Both carry `skip-gfx1250-strict` (the `spmm_i8hs*.yaml` they were copied from predate that mark being added in #11777), and a `gfx1250-strict` copy of each is added with `Architecture`/skip mark swapped (the shipped `gfx1250-strict` I8IS logic yamls had no tuning source either). - Due to CI error on gfx942 when using `DirectToVgprSparseMetadata`, currently still turn on `ExpandPointerSwap` in this case. Note: `DirectToVgprSparseMetadata` is not used anymore. - **`Solution.py` (`fc39076ff1c`, re-enable sparse auto `VectorWidth`):** reverted `2d9ececd8a0` so `VectorWidthA/B: [-1]` auto-selects again for sparse on every arch, not only gfx1250 (the sparse operand stays capped by `findSparseVectorWidth`, the dense operand takes the full auto VW). This exposed an ordering bug: `enableLDSTrMetadata` (`MetadataLayout=1`) forced the sparse operand's VW to 1 only after `StoreVectorWidth` was already derived from the auto VW, so i8 kernels ended up with `VWA=1` but `SVW=16` and the epilogue ran out of accVgpr reads (`replaceHolder(None, ...)` in `GlobalWriteBatch._prolog`). `enableLDSTrMetadata` is now computed next to `enableLDSTrA/B` (reusing `isLDSTrEnabled`) and the `VectorWidth`/`VectorWidthMetadata = 1` forcing is applied there, so `StoreVectorWidth` etc. see the final VW; the late block is removed. - **`gfx1250-strict/` sparse yamls:** applied the same migration to the shared yamls (same `TransposeLDS`/`ScheduleIterAlg`/`TDMInst`/`1LDSBuffer`/`LDSTrInst`/`ExpandPointerSwap` changes) and brought `spmm_tdm_all.yaml` in line with the `gfx1250/` one (it was missing `TDMSplit`, extra `Exact` sizes and the TDM3 + GSU>1 / Batch>1 block). After this, every sparse yaml in `gfx1250-strict/` differs from its `gfx1250/` counterpart only in `Architecture` and the mirrored skip mark. `gfx1250-strict` yamls are skipped on gfx1250 hardware, so they are a mirror of the validated `gfx1250/` changes. - Added the missing `skip-gfx1250-strict` mark to `gfx1250/spmm_tdm_iteratemode.yaml` and `gfx1250/spmm_tdm_split_xfail.yaml`, so they are no longer picked up on a `gfx1250-strict` agent. - **gfx942/gfx950 sparse yamls (`165b48659cf`):** the revert above is not gfx1250-specific (`2d9ececd8a0` forced `VectorWidthA/B=1` for every sparse arch), so the gfx942/gfx950 sparse yamls also moved to auto VW and some of them lost all of their solutions (`VW>1 not compatible with StoreRemap`, `Your parameters resulted in 0 valid solutions`). - `gfx94x/f8n_gate_r.yaml` (the two Sparse A groups), `spmm_b8f8n`, `spmm_bf16`, `spmm_bf8n`, `spmm_f8b8n`, `spmm_f8n`, `spmm_i8bs`, `spmm_i8hs` and `spmm_i8is` get `VectorWidthA/B: [-1, 1]`. Each of the 8 single-group yamls goes from 0 back to 6 solutions (same as `develop`), and the `f8n_gate_r` Sparse A groups from 0 to 12 each. The Sparse B groups of `f8n_gate_r` are unchanged. - `gfx950/f8_gate_r.yaml` pins `VectorWidthA/B: [1]`, so its emitter statuses stay as before. - `test_setcover_sparse_char.py`: fingerprints, fork counts and emit statuses recomputed for the six configs whose yaml changed. `spmm_i8is` and `spmm_bf8n` now sample the whole 32-permutation search (per-config `_LIMITS`), because the new VW fork pushes every valid permutation past the default first 8. - Known side effect, not changed here: some gfx942/gfx950 yamls produce fewer solutions than on `develop` (gfx942 `spmm_gsuc` 40 -> 24 and `spmm_i8_mi16` 14 -> 8, 11 gfx950 yamls roughly halved). No test fails on them. Note: `spmm_gate.yaml`'s stale `xfail-gfx1250` marker (same root cause) was already removed upstream by [PR #12355](#12355), so this PR no longer needs to touch it or `test_config_helpers.py` after rebasing onto latest `develop`. ## Test Plan 1. gfx1250 tox 2. New `spmm_i8is.yaml` / `spmm_i8is_sb.yaml` run through Tensile on gfx1250 hardware. ## Test Result 1. gfx1250 tox passed ``` ===== 101 passed, 878 skipped, 1 xfailed, 66 warnings in 918.29s (0:15:18) ===== py3: OK (994.02=setup[12.66]+cmd[0.45,42.87,8.73,0.24,929.08] seconds) congratulations :) (994.05 seconds) ``` 2. `spmm_i8is.yaml` and `spmm_i8is_sb.yaml`: 16/16 PASSED each on gfx1250 (4 transpose kernels per file, all `clientExit=0`). 3. gfx1250 sparse tox after rebasing onto latest `develop` (`b614a6478bd`, includes `fc39076ff1c`, `VectorWidthA/B: [-1]` auto-select re-enabled): `tox -e py3 -- tensilelite/Tests -m common -k "gfx1250 and sparse"`, 0 failed (each `test_config` does build + GPU run with `NumElementsToValidate: -1`; includes `spmm_tdm_all.yaml`, `spmm_i8is.yaml`, `spmm_i8is_sb.yaml`) ``` ===== 40 passed, 361 skipped, 1 xfailed, 64 warnings in 371.88s (0:06:11) ====== py3: OK (539.76=setup[11.70]+cmd[0.39,43.50,9.31,0.23,474.63] seconds) congratulations :) (539.82 seconds) ``` The only `xfailed` is the expected `spmm_tdm_split_xfail.yaml`; `gfx1250-strict` yamls are skipped on gfx1250 hardware. Before the `enableLDSTrMetadata` fix, re-enabling auto VW failed `spmm_tdm_all.yaml` in the build step (2 i8 `MetadataLayout=1` kernels). 4. `spmm_tdm_all.yaml` run directly through the Tensile client with `PrintWinnersOnly: False`: 122/122 validation rows PASSED (42 kernels, `clientExit=0`), including the two i8 `MetadataLayout=1` kernels (`VWA=1`, `VWB=2`, `SVW=8`). 5. gfx942/gfx950 sparse yamls (CPU only: solution generation and kernel emit, no GPU run): the 8 single-group gfx94x yamls above produce 6 valid solutions each again, `f8n_gate_r` Sparse A groups 12 each, and `tox -e py3 -- tensilelite/Tests/unit/characterization/_codegen/test_setcover_sparse_char.py`: 15 passed. The gfx942 `test_config` build and run of these yamls is covered by CI only. Note: results 1, 2 and 4 above were taken before the rebase onto latest `develop`; result 3 was taken after it; result 5 is CPU only and was taken on the final tree. JIRA: https://amd-hub.atlassian.net/browse/AIHPSPAR-279 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. [AIHPSPAR-279]: https://amd-hub.atlassian.net/browse/AIHPSPAR-279?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ --------- Co-authored-by: Cursor <cursoragent@cursor.com>
generateKernelHelperObjects used to de-dup helpers with list(set(khos)). KernelWriterBase hashes its kernel name, and str hashes are salted per process, so the helper order in Kernels.cpp, and with it Kernels.so-000-<arch>.hsaco, changed between builds of the same source (AIHPBLAS-5171). #11777 already replaced the set() with dict.fromkeys, which fixes the order, but nothing tests it. Add test_kernel_helper_order.py. It uses fake helper writers and name callables, so it needs no compiler and no logic files. It pins the exact order and checks that order in subprocesses run with PYTHONHASHSEED set to 0..3. With set() put back, both tests fail. Also drop the khos.index(x) sort key. index() is O(n) and every comparison calls KernelWriterBase.__eq__, which builds two kernel names, so the sort was O(n^2) name builds. Stable-sorting reversed(khos) on "Enum" not in name gives the same order: Enum helpers first, then the rest, each group in reverse insertion order. Keeping that order, rather than switching to plain insertion order, means Kernels.cpp and the hsaco stay byte-identical to current develop output. JIRA ID : AIHPBLAS-5171 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…testing redesign (#12997) > **Reading this PR:** the diffs are hard to follow, so here are both files rendered as they stand on this branch. The links follow the branch, so they always show the latest commit. > > - [projects/hipblaslt/TESTING.md](https://github.com/ROCm/rocm-libraries/blob/users/tony-davis/AIHPBLAS-5028-testing-md-tracking/projects/hipblaslt/TESTING.md) > - [projects/hipblaslt/tensilelite/TESTING.md](https://github.com/ROCm/rocm-libraries/blob/users/tony-davis/AIHPBLAS-5028-testing-md-tracking/projects/hipblaslt/tensilelite/TESTING.md) JIRA ID : AIHPBLAS-5028 ## Motivation hipBLASLt is AMD's matrix-multiply (GEMM) library, and TensileLite is the Python code generator that writes its GPU kernels. Each has a TESTING.md that describes how it's tested today and ends with a "Known Risks and Gaps" table and an "Improvement Roadmap". While checking those against current CI, I found a lot of specific shortcomings, and most of them had no ticket. This PR records them and gives each one a ticket. It's the first step of a bigger effort, though, not the plan. Testing across these two products is pretty scattered, and I think it needs a top-down redesign rather than a longer list of fixes. That redesign comes in the PRs stacked on this one: a shared testing plan, an inventory of every suite, and a coverage model. So the roadmap sections here are framed as tactical findings. They aren't a commitment or a priority order. ## Technical Details Docs only. No code, test, or CI configuration changes. - **Every gap row and roadmap item names a ticket.** New work items are under the epic AIHPBLAS-5028, plus #12996 for the out-of-date required-check list in docs/gardening.md. Where an existing ticket already covered a gap, the row uses it. For example, the rows about Math CI's `preliminary` job point at AIHPBLAS-3354, which moves TensileLite testing to TheRock CI, since Math CI is being retired. - **The roadmap sections are reframed.** Each one now opens by saying it's not yet a plan, and its items are grouped by kind of work ("Process and policy fixes", "Structural changes", "Larger investments") instead of near, medium, and longer term. - **Six rows and three roadmap items are removed.** Four of the gaps are closed on develop: TheRock now runs hipBLASLt tiers through CTest, submodule-bump PRs run more tests rather than fewer, gfx950 runs per PR in Multi-Arch CI, and TheRock's installed-artifact lane now has the snapshot-test plugin (syrupy). `preliminary`'s target-branch condition no longer exists. The last one is a Math CI job name we don't plan to change, since Math CI is going away. - **Four rows are added** from existing tickets: large-size address arithmetic (AIHPBLAS-4988), tests of the public API itself (AIHPBLAS-3579), the library-logic check skipping derived-parameter assignment (AIHPBLAS-3575), and no sanitizer lane being a required check (AIHPBLAS-1079). - **Body text that no longer matched CI is corrected.** The ones most likely to change how someone reads a CI result: - The TensileLite coverage floor blocks a merge. Its lane rolls up to `Component CI Summary`, which the develop ruleset requires. The floor is 82%, not 75%. - The required checks are `Math CI Summary`, `pre-commit`, `Multi-Arch CI Summary` and `Component CI Summary`. Neither AddressSanitizer (ASAN) lane is one of them. - gfx950 runs per PR in TheRock Multi-Arch CI at the `standard` tier, as of the TheRock pin bump in #12826. - Math CI schedules a PR by its `project:` label, and the labeler doesn't give that label to `shared/rocroller` or `shared/mxdatagenerator`. So a PR that only touches those gets no hipBLASLt Math CI jobs at all. - Smaller fixes: Math CI runs multi-GPU tests on two gfx950 GPUs, the mutation-testing slice is ten files, one test file and one corpus check were removed in #11777, a dead cmake link, and the current rocPTS nightly details. - **Rebased onto develop.** develop changed both files after this PR opened: the TensileLite lowercase namespace move (#12150, #13399), and the ASAN lane change in #13126. I kept develop's paths and lane facts. One exception: #13126 marks the gfx90a HOST_ASAN lane as gating, but `hipBLASLt ASAN CI Summary` isn't a required check in the develop ruleset, so the table says it doesn't gate. I also updated six other passages that still described full ASAN on gfx90a per PR, and two TensileLite links that still used the old `Tensile/` path. ## Test Plan Docs only, so the checks are about accuracy and links: - `pre-commit run --files` on both files. - Resolve every relative link and in-document anchor in both files. - Resolve every Jira key and GitHub link through the Jira REST API and `gh api`. - Check each corrected statement against develop, TheRock at the pinned commit, rocJenkins at its live branch, and the develop ruleset (`gh api repos/ROCm/rocm-libraries/rules/branches/develop`). ## Test Result - `pre-commit` passes. The Markdown files fall outside every hook's file filter, so I also checked trailing whitespace, links, and anchors with a script. - All Jira keys and GitHub references resolved when this PR was first checked, and the rebase and reframing added none. - The original corrections were checked against develop at f4ec50c and TheRock c9bde5b. After the rebase I rechecked the ASAN lanes and the required checks against develop at 67b685e. I haven't re-audited every other statement against TheRock 76bd097, which develop pins now. - Automated GitHub checks are pending. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): documentation only. No product, test, CI, or device behavior changes. The main risk is that one of the corrections is itself wrong, which is what review is for. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
JIRA ID : AIHPBLAS-4764 ## Motivation `sparse/gfx1250/spmm_gate.yaml` doesn't run on the gfx1250 hardware lane. #12235 fixed its failure, and #12355 removed its xfail after it passed. Then #11777 marked the config `skip-gfx1250-strict` ("not supported yet"). ## Technical Details - Add `sparse/gfx1250-strict/spmm_gate.yaml`, a copy of the gfx1250 config following the pattern of the other 37 sparse strict copies: - `skip-gfx1250` instead of `skip-gfx1250-strict`. - The gfx1250 config is unchanged and keeps running where it runs today. ## Test Plan - This PR's `gfx1250 HW hipBLASLt + TensileLite Tests` run collects and passes `sparse/gfx1250-strict/spmm_gate.yaml`. ## Test Result - New strict test [passes](https://github.com/ROCm/rocm-libraries/actions/runs/39722965803/job/119640530321?pr=14175) on gfx1250 HW run. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. One new test config that runs only on gfx1250-strict. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude <noreply@anthropic.com>
…check The V0 stepping is now served only by gfx1250-strict: with HSA_DISABLE_GFX12_STRICT=0 it reports gfx1250-strict and loads library/gfx1250-strict/. hipBLASLt and hipSPARSELt no longer read hipDeviceProp_t::asicRevision, and the build no longer produces library/gfx1250v0/ (gfx1250-strict's logic built for the gfx1250 target). This undoes the gfx1250v0 parts of ROCm#11777 by hand; later commits touched the same files, so a git revert does not apply. The xnack handling built on top of it (compilerTargetOf, the xnack- capability overrides, outDirs) is kept. - tensile_host.cpp in hipBLASLt and hipSPARSELt, and validate_library_layout.py, are restored to their state before the gfx1250v0 change. - HIPBLASLT_BUILD_GFX1250V0 and TensileCreateLibrary --gfx1250v0 are removed, and with them ARCH_BUILD_ALIASES, tuningArchOf, deviceTargetFeaturesOf, withArchBuildAliases and archRenames. - The layout validator reports a library/gfx1250v0/ left in a reused install prefix. - Comments that named gfx1250v0 as the V0 build now name gfx1250-strict. Behavior changes: - A V0 part at HSA_DISABLE_GFX12_STRICT=1 or unset reports gfx1250 and loads library/gfx1250/ without an error. V0 needs =0 and a package that ships gfx1250-strict/. - A gfx1250-only build no longer runs --check-all on gfx1250-strict logic. Testing (FFM, ROCm 10.3.0a20261103): - gfx1250 and gfx1250-strict code objects match develop except for StinkyTofu's random label hashes; hipSPARSELt is byte-identical. - V0 topology at =0 loads gfx1250-strict/: hipblaslt-test *quick* 18043 passed, hipsparselt-test *quick* 36816 passed, tox Tests/common 198 passed and 0 failed. - TensileLite unit tests and both pre-commit configs pass. Co-authored-by: Cursor <cursoragent@cursor.com>
…check The V0 stepping is now served only by gfx1250-strict: with HSA_DISABLE_GFX12_STRICT=0 it reports gfx1250-strict and loads library/gfx1250-strict/. hipBLASLt and hipSPARSELt no longer read hipDeviceProp_t::asicRevision, and the build no longer produces library/gfx1250v0/ (gfx1250-strict's logic built for the gfx1250 target). This undoes the gfx1250v0 parts of #11777 by hand; later commits touched the same files, so a git revert does not apply. The xnack handling built on top of it (compilerTargetOf, the xnack- capability overrides, outDirs) is kept. - tensile_host.cpp in hipBLASLt and hipSPARSELt, and validate_library_layout.py, are restored to their state before the gfx1250v0 change. - HIPBLASLT_BUILD_GFX1250V0 and TensileCreateLibrary --gfx1250v0 are removed, and with them ARCH_BUILD_ALIASES, tuningArchOf, deviceTargetFeaturesOf, withArchBuildAliases and archRenames. - The layout validator reports a library/gfx1250v0/ left in a reused install prefix. - Comments that named gfx1250v0 as the V0 build now name gfx1250-strict. Behavior changes: - A V0 part at HSA_DISABLE_GFX12_STRICT=1 or unset reports gfx1250 and loads library/gfx1250/ without an error. V0 needs =0 and a package that ships gfx1250-strict/. - A gfx1250-only build no longer runs --check-all on gfx1250-strict logic. Testing (FFM, ROCm 10.3.0a20261103): - gfx1250 and gfx1250-strict code objects match develop except for StinkyTofu's random label hashes; hipSPARSELt is byte-identical. - V0 topology at =0 loads gfx1250-strict/: hipblaslt-test *quick* 18043 passed, hipsparselt-test *quick* 36816 passed, tox Tests/common 198 passed and 0 failed. - TensileLite unit tests and both pre-commit configs pass. Co-authored-by: Cursor <cursoragent@cursor.com>
AIHPBLAS-4768
gfx1250-strict as an architecture of its own, plus a gfx1250v0 subtree for v0 parts that report gfx1250
Follow-up to #10784, which
introduced
gfx1250-strictas an ordinary architecture name. This PR removes thespecial cases #10784 left in CMake. It also retires develop's
gfx1250v0ASIC-revision workaround in favour of the stepping's own name, and keeps v0
silicon correct while ROCr still reports it as
gfx1250by default.Motivation
gfx1250 ships as two silicon steppings. They share the ISA tuple
(12, 5, 0)butare separate compiler targets (ELF
EF_AMDGPU_MACH0x49for gfx1250/v1,0xEBfor gfx1250-strict/v0). A code object built for one is rejected on the other.
On ROCm 10.2, what a v0 part reports depends on
HSA_DISABLE_GFX12_STRICT:HSA_DISABLE_GFX12_STRICTgcnArchNameon v0asicRevision0gfx1250-strict0xEBonly1or unsetgfx12500x49onlyIn the second row, v0 would otherwise load v1 kernels. Those use instructions v0
does not have, such as
v_wmma_scale_f32_32x16x128_f4, and TDM multicast.Measured on v0: an FP4 v1 kernel returns
norm_error=nan. That row is ROCr'sdefault today, so it has to stay correct until the default flips.
Tensile keys its capability map by ISA, so one
TensileCreateLibraryruncannot describe both steppings. #10784 worked around that in CMake: a second
hipblaslt_create_device_librarycall forgfx1250-strict, plus an extra installrule driven by a global property. This PR moves the split into
TensileCreateLibrary, where the knowledge of which names collide alreadylives. CMake goes back to a single call with
${GPU_TARGETS}.The ISA-keyed machinery itself is deliberately left alone. Re-keying it by name
would touch ~88
isaInfoMap[...]subscripts across 24 files, plus two rocisaC++ sites and the nanobind ABI, on every architecture's code path.
What gets built and loaded
library/gfx1250/gfx1250/0x49gfx1250-a gfx1250library/gfx1250v0/gfx1250-strictlogic and capability overridesgfx1250/0x49gfx1250-a gfx1250(default,HIPBLASLT_BUILD_GFX1250V0=ON)library/gfx1250-strict/gfx1250-strict/0xEBgfx1250-strict-a gfx1250-strictThe runtime picks the Tensile GEMM subtree from the reported name first, then
the revision. This applies to both hipBLASLt and hipSPARSELt, when
*_TENSILE_LIBPATHis unset:gcnArchNameasicRevisiongfx1250-strictgfx1250-strict/(unchanged)gfx12500(v0)gfx1250v0/. If it is missing, hard error: no fallback togfx1250/, whose kernels are wrong on v0gfx1250gfx1250/(unchanged)<name>/(unchanged)Technical Details
Fan-out
isaCollisionFreeGroups()partitions the requested names so that no group holdstwo names spelling one ISA. A run handed more than one group re-executes itself
once per group as a subprocess. They are separate processes rather than a loop,
because a run settles arch-dependent state process-wide.
A single group takes the old path untouched: no subprocess, and no change to
parallelism, scratch naming or cleanup. That covers every build that names no
gfx1250 variant. Jobs are split by architecture count rather than evenly.
Scratch directories
Kernel basenames come from the ISA, so both steppings name their
.sand.ofiles identically. The fixed-name artifacts (
Kernels.cpp,Kernels.h, the sixstatic headers,
code_object_tmp/Kernels.o) carry no architecture at all.Scratch is therefore
build_tmp/<STEM>-<stepping>for a run asked for astepping, and stays
build_tmp/<STEM>otherwise. Fan-out children carryTENSILE_GROUP_BUILD, so each reclaims only its own subdirectory.Benchmark cache key (correctness fix)
The tuning cache key gains the compiler targets; otherwise the second stepping
would load the first stepping's code objects.
_loadLegacyCacheIfMatchesisremoved, because it is keyed by directory alone and would bypass the new field.
Architecture names
ArchitectureName, which a stepping spells as itself.archMatch()compares whole names, sogfx1250cannot claim agfx1250-strictrequest, and it strips both qualifiers and bracketedpredicates, so
gfx950[cu=64]still matches its own logic.steppingArchOf()derives a stepping rather than looking it up in a table: aname is a stepping exactly when it does not survive a round trip through its
ISA. A future stepping needs no registration.
amdgpu-arch, or from theenumerator, goes through
restore_steppings, so a bareTensilerun on v0builds gfx1250-strict rather than base gfx1250. Colon target features are
stripped for
--gpu-targets; the hyphenated stepping is kept.gfx1250 silicon list both
skip-gfx1250andskip-gfx1250-strict.gfx1250v0: v0 kernels for a runtime that reports gfx1250
gfx1250v0is a build-only alias. It is never inGPU_TARGETS(CMake'ssupported list rejects it), never a runtime name, and excluded from
allandarchMacroNames. Naming it directly is rejected, both byTensileCreateLibrary --architecture(before any code generation) and byTensile --gpu-targets/ a config'sArchitecture. A single declarative table,ARCH_BUILD_ALIASESinCommon/Architectures.py, defines it:tuningArchgfx1250-strictarchMatch),ARCH_CAP_OVERRIDESinCapabilities.applyArchCapOverrides,StinkyTofuArchNametargetgfx1250-mcpuand.amdgcn_target, bundle target id, helper--offload-arch, the declaredArchitectureNameafter parsing (LibraryIOarchRenames), master/Mapping/.conamesdeviceTargetFeatures+needs-aligned-2addr-lds,-wmma-f4-insts,-block16-cvt-scale-insts-Xclangas) and the HIP helper compile (-Xclang)Tensile codegen: same logic, same capability overrides (no TDM multicast, no
fp4 32x16 WMMA), same
StinkyTofuArchName. Only the compiler target, filenames and output directory differ. The added target features make the
assembler reject the instructions v0 lacks, as it does for gfx1250-strict, and
give the helper kernels strict's ISA.
--gfx1250v0,withArchBuildAliases()appendsgfx1250v0when baregfx1250is requested, including throughall. It runsbefore
isaCollisionFreeGroupsand never in a fan-out child, so-a "gfx1250;gfx1250-strict"builds three groups.gfx1250-strictalone doesnot expand. If no gfx1250-strict logic matches (a filtered or custom logic
build), the alias is skipped with a warning rather than producing an empty
subtree.
TensileCreateLibrary --gfx1250v0is opt-in, so direct callers(tox, the benchmark client, GenerateSummations) are unchanged. The CMake option
HIPBLASLT_BUILD_GFX1250V0(defaultON, inhipblaslt_codegen.cmake, sohipSPARSELt shares it) passes the flag, and adds
gfx1250-strictto theTensileLogic --check-allarchitectures of a gfx1250 build. Turning v0 off forgfx1250 once TheRock builds gfx1250-strict is
-DHIPBLASLT_BUILD_GFX1250V0=OFF.gfx1250token and land inlibrary/gfx1250v0/. Itholds no ExtOp or Transform files; those load from
library/<gcnArchName>/as before.
validate_library_layout.pygainsREVISION_SUBTREES = {"gfx1250v0": "gfx1250"}.Runtime selection
tensile_host.cppreadsasicRevisiononly for a devicereporting
gfx1250, and asks the pure helperrocblaslt_revisioned_arch_name(processor, asicRevision)for the subdirectory(restored from develop as
rocblaslt_arch_revision.hpp). Master file names,codeObjectTargetsand lazy loading keep using the reported name, so filesnamed
gfx1250insidegfx1250v0/resolve unchanged.directory and the two ways out. Loading then fails through the existing
error paths, and API calls return an error status.
row_uniformity_gtest's library helper uses the same helper.tensile_host.cpp(
asicRevisionread only forgfx1250, guarded byHIP_VERSION >= 307as inits
handle.cpp).Removed knobs
HIPBLASLT_ASIC_REVISION,--asic-revision,gpu_revision_probe,GpuRevisionTarget.py, and thegfx1250v0overlay checks(
find_gfx1250v0_overlay_violations,--require-gfx1250v0-overlay) areremoved. The stepping is named in
GPU_TARGETS, and the gfx1250 v0 subtree iscontrolled by
HIPBLASLT_BUILD_GFX1250V0.Tuning logic
Logic/asm_full/gfx1250v0/moves togfx1250-strict/(226 files; directory,filename prefix,
ScheduleNameandArchitectureNamerenamed, tuning valuesunchanged). hipSPARSELt gains its own
gfx1250-strictspmm tree (64 files).Guards, packaging, reproducibility
__gfx1250_strict__and not
__gfx1250__. Eight guards (memory_gfx.h, the float4/bfloat6/float6/e5m3 headers,
DataTypes_E5M3.hpp, mxDataGenerator) accept both, andthe extops/matrix-transform CMake arch regexes keep
-strict._retargetAssemblySourcematches the whole hyphenatedprocessor name, so
gfx1250-strictis not rewritten togfx1250-strict-strict.validate_library_layout.pyandcheck_dat_integrity.pyaccept hyphens in an architecture token.
test_MatrixInstructionConversionupdated the shareddefaultSolutionin place, which leaked into whichever test xdist ran nexton that worker; it now works on a copy.
Test Plan
The full procedure is in
projects/hipblaslt/docs/gfx1250-strict-testing.md(Suite 6 covers gfx1250v0). It runs on a real v0 machine with the newest ROCm
nightly, through
pipeline-v0.sh.TensileLite unit tests.
pytest -m unitoverTests/unitandTests/extras.Architecture-name parsing (
test_arch_steppings.py). EveryarchitectureMapname in every CMake spelling (bare,:xnack±,[cu=64],[id=74a0]) must match its own logic header and no other. Also covers thegfx1250v0 alias table, expansion (with / without
--gfx1250v0, fan-outchild), rejection of the alias as a requested name, the three-group
partition, caps equality with strict end to end through
KernelWriter,compiler target and features, logic selection and keying, and the
no-strict-logic skip.
Parser diff against the base commit for 121 spellings.
Packaging script tests (
tools/scripts/tests/), including the v0subtree.
pre-commit. Both the hipBLASLt config and the repo root config.
stinkytofu
ctest.Build
./install.sh -c -a "gfx1250;gfx1250-strict", thenverify_v0.pyonthe install tree:
0x49/0x49/0xEB);v_cvt_scale_pk*, multicast or helperds_*_2addrineither v0 tree, while
gfx1250/does have them, so the scan is notvacuous;
gfx1250v0/equals that ofgfx1250-strict/.Plus
validate_library_layout.py.hipblaslt-test, all and
*quick*, in both reporting modes. Each runchecks that
lddresolves to the builtlibhipblaslt.so, and that the logshows the subtree loaded:
HSA_DISABLE_GFX12_STRICT=0→gfx1250-strict/,=1→gfx1250v0/, unset →gfx1250v0/.FP4 regression.
hipblaslt-benchFP4 TN 256x256x512 with MX scales(
--requested_solution -1 -v) and thesmoke_matmul_gemm_tn_mxf4test, inboth modes: every solution must pass. Negative control: forcing
HIPBLASLT_TENSILE_LIBPATHatlibrary/gfx1250/on v0 must bring the NaNback.
No fallback. In a copy of the install without
gfx1250v0/, a v0reporting gfx1250 must fail with the error, and no test may pass.
tox config tests (
Tests/common), in strict mode.hipSPARSELt. Build
-a "gfx1250;gfx1250-strict", runverify_v0.pyon itslibrary tree, check no-fallback, confirm with
stracewhich subtree eachmode opens, then run the full
hipsparselt-testin both modes (one GPUeach), with an
lddcheck.Test Result
Committed series (without gfx1250v0)
Measured on the series at
bdb85d2f61(develop670b375ba0) with the fixessince committed as
8f11772cbb. ROCm 10.2.0a20260926,HSA_DISABLE_GFX12_STRICT=0,4× gfx1250 v0.
-m unit)test_arch_steppings.pyctest0x49(gfx1250),0xEB(gfx1250-strict)validate_library_layout.py*quick*(gfx1250-strict)Tests/commonWith gfx1250v0
This change as staged on
8f11772cbb(developd0e4bc5f98), since committed as1c366475cc. The commit is identical file-for-file exceptspmm_f16_sia4.yaml, which is re-validated on its own below.ROCm 10.2.0a20260929, 4× gfx1250 v0, every GPU step on GPUs 2 and 3.
Tensile and build
-m unit -n 8)test_arch_steppings.pyctest-a "gfx1250;gfx1250-strict"[gfx1250],[gfx1250-strict],[gfx1250v0]0x49/0x49/0xEBds_*_2addr, gfx1250 / v0 / strict.amdgcn_target-mcpu=gfx1250+ v0 features)v_wmma_f32_32x16x128_f4andv_cvt_scale_pk8_f32_fp4, as-mcpu=gfx1250-strictdoes; plain gfx1250 accepts bothvalidate_library_layout.pyhipBLASLt runtime
ldd hipblaslt-testlibhipblaslt.so.1andlibtensilelite-host.so.1=0/=1/ unset (log)gfx1250-strict//gfx1250v0//gfx1250v0/=0/=1(strace -e openat)gfx1250-strict// onlygfx1250v0/=0(gfx1250-strict)*quick*,=0=1(gfx1250v0)nightly_matmul_deepbench_bf16_..._NN_1152_128_7000_..._3_t2, intermittent; 20 / 20 reruns passed, see Known)*quick*,=1=0/=1=1)=1)library/gfx1250v0; 0 passed, 128 failed; exit 1TensileLite config tests and hipSPARSELt
Tests/common(=0)spmm_f16_sia4.yamlas committed (=0)-a "gfx1250;gfx1250-strict"gfx1250/,gfx1250v0/,gfx1250-strict/0x49/0x49/0xEB; fp4 32x16 WMMA and multicast 0 in all three=0/=1(strace)gfx1250-strict// onlygfx1250v0/=1)library/gfx1250v0; every spmm test fails (status 10, no solutions); prune/compress, which do not use the GEMM library, pass (2448); run capped at 300 s=0(GPU 2, gfx1250-strict)=1(GPU 3, gfx1250 → gfx1250v0)Before this change, the FP4 bench on v0 in
=1mode loadedlibrary/gfx1250/and gave 344 / 345 passes, with the MI32x16 kernel returning
norm_error=nan.Deliberate behaviour changes
(including
all) buildslibrary/gfx1250v0/in an extra fan-out group,which roughly doubles gfx1250's Tensile build time and library size.
-DHIPBLASLT_BUILD_GFX1250V0=OFFremoves it.gfx1250v0/, or fails. Previously itsilently ran v1 kernels. A package without the v0 subtree now fails on v0
with
HSA_DISABLE_GFX12_STRICTunset or1; use=0with a gfx1250-strictbuild instead.
folded into the hash), so existing tuning output directories miss once.
.datbytes change for every architecture (deterministic logicmerge order changes solution indices).
<destPath>.linker_args.Kernels.cpp,Kernels.hand the static headers are written under therun's scratch directory rather than the output root.
_loadLegacyCacheIfMatchesis removed.Known, not addressed
device to initialize.
HIP kernels compiled into libhipblaslt / libhipsparselt load their gfx1250
(v1) builds, as on develop. hipblaslt-test and hipsparselt-test in
=1modeexercise them.
asicRevision. It is assumed non-zero; there was no v1 machine toconfirm it.
*_TENSILE_LIBPATHis used verbatim. Pointing it atlibrary/gfx1250/onv0 still loads v1 kernels, and a mismatched install is not validated.
gfx1250[cu=…]entry inGPU_TARGETSbuildsgfx1250v0 without its predicate, and does not add gfx1250-strict to the
TensileLogic --check-alllist.features. Clear
~/.tensile/helper_cacheif those change.HIPBLASLT_BUILD_GFX1250V0off in an existingbuild directory leaves the old
library/gfx1250v0/to be installed..github/workflows/gfx1250_hipblaslt_tensilelite_test.ymlstillsets
TL_TARGET=gfx1250v0, a name Tensile now rejects as a target; it needsporting to gfx1250-strict.
nightly_matmul_deepbench_bf16_..._NN_1152_128_7000_..._3returned a wrongvalue once in a full strict-mode run before this change (40 / 40 reruns
passed). Its
_t2variant did the same once in a full gfx1250v0-mode runafter it (20 / 20 reruns passed).
PrefetchLocalRead=1withDepthU=128,ScheduleIterAlg=4and buffer_load sparse A returns wrongvalues.
spmm_f16_sia4.yaml's first group now usesScheduleIterAlg=3,keeping DepthU 128 and PLR 1, which passes. The SIA 4 combination and the base
gfx1250twin config are unverified.isaInfoMapannotation. It isDict[str, IsaInfo]while production keysit by
IsaVersion.Makefile:55gatesTRUE16_FEATUREon an unsetTARGET_ARCH, somake conever passes+real-true16. This predates the branch.Submission Checklist
JIRA ID : AIHPBLAS-4727
JIRA ID : AIHPBLAS-4728
JIRA ID : AIHPBLAS-4729
JIRA ID : AIHPBLAS-4730
JIRA ID : AIHPBLAS-4768
JIRA ID : AIHPSPAR-275
JIRA ID : ROCM-31226
JIRA ID : ROCM-30511