hipblaslt: replace the gfx1250v0 revision workaround with gfx1250-strict - #11777
Open
geotseng-amd wants to merge 10 commits into
Open
geotseng-amd wants to merge 10 commits into
geotseng-amd wants to merge 10 commits into
Conversation
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
geotseng-amd
requested review from
cycheng,
imcarsonliao and
jichangjichang
September 13, 2026 14:48
geotseng-amd
force-pushed
the
users/geotseng/develop-gfx1250-strict
branch
from
September 14, 2026 08:42
b6cf90f to
e2a8f1f
Compare
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## develop #11777 +/- ##
===========================================
+ Coverage 70.34% 70.35% +0.01%
===========================================
Files 2812 2812
Lines 463054 463104 +50
Branches 68187 68187
===========================================
+ Hits 325706 325806 +100
+ Misses 113761 113715 -46
+ Partials 23587 23583 -4
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
jichangjichang
approved these changes
Sep 14, 2026
cycheng
reviewed
Sep 14, 2026
cycheng
approved these changes
Sep 14, 2026
pmoutsias-amd
approved these changes
Sep 14, 2026
pmoutsias-amd
left a comment
pmoutsias-amd
left a comment
Contributor
There was a problem hiding this comment.
I checked the edit in the readme and it is a straight removal. LGTM.
bnemanich
approved these changes
Sep 14, 2026
bnemanich
left a comment
bnemanich
left a comment
Contributor
There was a problem hiding this comment.
mxdatagenerator change looks fine. Will need someone else to look over the rest of PR.
randyh62
approved these changes
Sep 14, 2026
geotseng-amd
force-pushed
the
users/geotseng/develop-gfx1250-strict
branch
4 times, most recently
from
September 17, 2026 04:13
35a9809 to
3dd8f22
Compare
gfx1250 ships as two silicon steppings. They share ISA (12,5,0) but are
separate compiler targets whose code objects are rejected on each other,
and ROCr now reports the stepping directly in the agent name. Treat the
A0 stepping as an ordinary architecture named gfx1250-strict -- the name
the compiler and the runtime both use -- instead of deriving it from an
ASIC-revision probe.
Detection no longer probes asicRevision. Tensile/GpuArch.py (was
GpuRevisionTarget.py) asks amdgpu-arch, falling back to rocminfo, and
uses the name it gets back. It stays free of Tensile.Common so the
invoke build path does not pull in rocisa just to name a target, and it
avoids rocm_agent_enumerator, whose gfx\d+ capture group truncates the
suffix and would silently build B0 kernels for A0 silicon. This removes
gpu_revision_probe.cpp, rocblaslt_arch_revision.hpp, the --asic-revision
option, and the arch-revision mapping in the runtime.
A stepping is now derived rather than declared: a name is a stepping
exactly when it does not survive a round trip through its ISA, so
steppingArchOf() answers gfx1250 for gfx1250-strict and None for every
ordinary arch. That retires ARCH_COMPILER_TARGET, gfxToCompilerTarget()
and REVISION_SUBTREES; a future stepping needs no registration. What
remains declared is ARCH_CAP_OVERRIDES, because the assembler accepts
the same instructions under both steppings and cannot tell them apart
by probing.
Compiling for gfx1250-strict predefines __gfx1250_strict__ and not
__gfx1250__, so guards testing only the latter fell through to their
else branch and produced non-functional kernels with no build error.
Eight such guards across memory_gfx.h, the hipblaslt float4/bfloat6/
float6/e5m3 headers, DataTypes_E5M3.hpp, and mxDataGenerator now accept
both. The CMake arch regexes in extops and matrix-transform were
likewise truncating "-strict" while stripping ":xnack+"; they now keep
it.
Because Tensile keys its capability map by ISA version, a single
TensileCreateLibrary invocation cannot hold both steppings. Rather than
name an architecture in the build system, the split lives in
TensileCreateLibrary: isaCollisionFreeGroups() partitions the requested
names so no group holds two that spell one ISA, and a run handed more
than one group re-executes itself once per group as a subprocess --
separate processes rather than a loop because a run settles
arch-dependent state process-wide. A single group, which is every build
that names no stepping, takes the old path untouched: no subprocess, no
change to parallelism, scratch naming or cleanup. CMake keeps one call
with ${GPU_TARGETS}.
Two runs sharing an output directory must not share intermediates:
kernel basenames come from the ISA, so both would name their .s and .o
identically while the machine code inside differs. Scratch therefore
becomes build_tmp/<STEM>-<stepping> for a run that was asked for one,
and stays build_tmp/<STEM> for every run that was not. Fan-out children
carry TENSILE_GROUP_BUILD so they reclaim only their own subdirectory
and leave a concurrent sibling's alone, while a standalone run still
removes build_tmp whole. Jobs are split by architecture count, so a
fourteen-architecture group no longer waits on a one-architecture
sibling holding half the cores.
The tuning cache key gains the compiler targets for the same reason.
Two steppings agree on every other field the key covers, so without
them the second run would load the first's code objects and hand the
wrong machine code to silicon that cannot run it. The legacy cache
fallback goes with it: that path is keyed by directory alone, so it
would bypass the new field and reintroduce exactly the mismatch.
Tuning is selected by ArchitectureName, which a stepping spells as
itself, so gfx1250-strict logic is no longer absorbed into gfx1250.
archMatch() compares whole names for that reason, and strips a
bracketed predicate as well as a qualifier -- TensileLogic filters with
GPU_TARGETS exactly as CMake passed it, having never split predicates
off, so a prefix-free comparison alone would fail gfx950[cu=64] against
its own logic and stop the build.
The generator scripts detect a name rather than an ISA, since deriving
the name back from the ISA builds gfx1250 code for either stepping.
They also pass the capability map to assignGlobalParameters, which has
required it since before this change.
With the stepping named for itself throughout, the
HIPBLASLT_ASIC_REVISION environment variable and the --asic-revision
flag have nothing left to select and are removed. Picking a revision is
now spelling it in GPU_TARGETS, which needs no separate knob and
generalises to the next stepping.
The packaging checks learn that an architecture token may contain a
hyphen. Their regexes accepted only alphanumerics, so a stepping's
files matched nothing, and an unmatched file leaves its subtree without
masters -- which the scanner reads as a subtree to skip, passing the
package silently rather than reporting it.
Logic files are now merged in sorted order. Set iteration and readdir
order were leaking salted string hashing into the solution indices
written into the master library, so a build was not reproducible.
Fixing it changes the shipped .dat bytes for every architecture, which
is worth knowing when diffing against a reference build.
stinkytofu keeps Gfx1250v0 as its C++ identity, since "gfx1250-strict"
is not spellable as an identifier; ArchHelper aliases the toolchain
spelling onto it and a FileCheck test pins that both spellings reach the
same cost table.
The stepping tests land as test_arch_steppings.py: the cases are about
steppings in general rather than about one name.
Note: the tuning logic move from Logic/asm_full/gfx1250v0/ to
gfx1250-strict/ is deliberately not part of this commit.
Co-authored-by: Cursor <cursoragent@cursor.com>
…rlay SUPPORTED_ISA cannot name a stepping, since a stepping shares the ISA it steps from, so "all" expanded to every architecture except gfx1250-strict. A default build therefore left A0 silicon with no library to load, and the one spelling that was supposed to mean "everything supported" was the one spelling that missed it. supportedSteppings() reads the names off architectureMap instead, restricted to steppings whose base architecture is already covered -- being a stepping does not make an unsupported architecture supported. The expansion now collides with itself by design, a stepping and its base being unbuildable in one run; isaCollisionFreeGroups partitions it as it already does for names given explicitly. The gfx1250v0 overlay machinery goes with the revision it described. That workaround kept v0 logic in a subtree of gfx1250's corpus tagged by ScheduleName, which needed find_gfx1250v0_overlay_violations to check the tagging, --require-gfx1250v0-overlay to opt a corpus in, and a matching CMake keyword to reach it -- all of it there because the two revisions could not be told apart by name. gfx1250-strict has its own architecture directory, so there is no overlay left to validate and nothing for the caller to opt into. The .amdgcn_target directive has to be brought back in step with the assembler. rocisa derives it from the ISA alone, which spells the base name for a stepping, while the build assembles that stepping with -mcpu=<stepping>. Assemblers through ROCm 10.1 accepted the mismatch; 10.2 rejects it -- "target id '...--gfx1250' specifies a processor that is not valid for subarch 'amdgpu12.50s'" -- so the build broke on a toolchain move rather than on a code change. Rewriting the directive before assembling is a no-op for every ordinary architecture, where the target already equals the name derived from the ISA. The packaging check learns that a hyphen is not enough to make a stepping. _stepping_base accepted any hyphenated subtree, which would have exempted library/gfx942-xnack+/ from the very check that reports it; the trailing sign separates a target feature from a stepping there, as it already does for filenames. The test corpus marks a stepping as itself. configMarks reads the architecture out of a filename, and without the suffix bf16_CLS_gfx1250-strict.yaml would mark itself gfx1250 and claim to be a config for the architecture it exists to stay off of; pytest.ini registers the stepping's marker alongside its xfail- and skip- forms. Co-authored-by: Cursor <cursoragent@cursor.com>
…cripts amdgpu-arch reports a configuration (gfx950:sramecc+:xnack-), not a CMake target. tox -e py3 forwards that string into --gpu-targets, so configure rejects the whole name. Strip colon-delimited features on both the get-gpu-arch and explicit --gpu-targets paths, and keep hyphenated steppings such as gfx1250-strict. Also install SoftmaxGenerator.py and LayerNormGenerator.py into the tensilelite test artifact tree. Installed unit tests spawn those scripts next to AMaxGenerator.py, which was the only generator previously shipped. Co-authored-by: Cursor <cursoragent@cursor.com>
gfx1250-strict is a separate architecture name, so a skip list naming only gfx1250 does not cover it: every config an architecture cannot run has to name the stepping too, or tox runs it on v0 silicon and fails there. Add skip-gfx1250-strict to the 410 configs already skipping the architectures they do not support, and add the gemm, gradient and sparse configs the stepping is meant to run, tuned for what v0 supports -- no TDM-multicast and no fp4 32x16 WMMA. Co-authored-by: Cursor <cursoragent@cursor.com>
gfx1250-strict is a separate compiler target, so the architecture allowlist has to name it or configure rejects a build asked for it, and the logic tree has to carry its own tuning or such a build ships an empty spmm library. The logic is declared under the stepping's own name, not gfx1250's: the two share ISA 12.5.0 but reject each other's code objects, so the selection must not let one claim the other's files. Co-authored-by: Cursor <cursoragent@cursor.com>
The stepping is an architecture of its own now, so its logic has to live under that name in all three places the build reads: the directory it is globbed from, the filename's codename prefix, and the ArchitectureName its own header declares. The declared name is what keys the master library, and the per-arch writes are keyed by the requested name, so a header still saying gfx1250 would key a library no write ever addresses -- a build that reports success having shipped an empty subtree. Declaring the stepping also keeps a plain gfx1250 build from globbing this tuning, which was derived under capabilities v1 does not share. Tuning values are unchanged; only the names are. Co-authored-by: Cursor <cursoragent@cursor.com>
…est data and under the coverage gate Origami matches gcnArchName exactly against a table holding base architectures only, so a part reporting gfx1250-strict matched nothing, left HipAMDGPU::analyticalHardware null, and threw on the assert guarding every query that consults the analytical model. That path is not opt-in: skDynamicGrid defaults to k_split_aware, so any Stream-K solution reaches the assert unless TENSILE_STREAMK_DYNAMIC_GRID is set. Hand Origami a copy of the properties naming the base architecture instead. The name is rebuilt through toProcessor/toString rather than trimmed as a string, so the revisions folded here stay the ones Tensile recognises and an unrecognised name still resolves to something Origami rejects. Both entry points read gcnArchName out of the properties they are handed, so the same copy reaches each; renaming only the support check would trade the null for a throw out of get_default_num_xcds. Only the copy is renamed -- archName() and the code-object subtree still spell the revision, and Origami consumes machine constants rather than code objects, so the base entry's instruction timings approximate the revision's performance and affect nothing else. match_test_category compares known_bug_platforms token-for-token with strcasecmp against the name hipblasLtGetArchName reports, which keeps the revision suffix. gfx1250-strict therefore missed the gfx1250 token and ran the TF32 Inf cases ROCM-1545 already quarantines on gfx1250 and gfx950; the defect is in silicon the revision shares, so list the token too. Three subtile bf16 configs were missed when the rest of the corpus got its skip-gfx1250-strict mark. gfx1250-strict is its own target, not a spelling of gfx1250, so a config written for gfx1250 has to opt out of it the same way it opts out of every other arch it was not written for; left unmarked they run on the strict stepping and build code objects the silicon rejects. The sibling-DeviceNames check reads ScheduleName, gfx arch and CU count out of each logic file's own header, and treats a file it cannot read as an unknown of each. Nothing exercised that: the corpus the unit tests run against is readable, so the three except arms in _arch_variant_key and _chip_id_dir_suffix's empty-arch guard were never entered, and the coverage ratchet caught the module dropping to 93.33%. Point the key at a path that does not exist and at a header with no ArchitectureName, and assert what the fallback is for: files it cannot read group together rather than each becoming its own singleton, and a corpus containing one still completes instead of raising. test_a_detected_name_yields_a_target_the_build_accepts read cmake/tensilelite_supported_architectures.cmake through parents[4] to check the probe's vocabulary against the list GPU_TARGETS is validated against. That path resolves in a checkout but reaches outside the tensilelite tree, which is installed on its own as a test artifact without that file, so all seven parametrizations failed on FileNotFoundError while a checkout stayed green. Remove the test rather than add an install rule whose only consumer is a test; cmake_gpu_target keeps the two tests beside it, and what is lost is the guarantee that the two lists cannot drift apart. Verified on a part reporting gfx1250-strict under ROCm 10.2.0a20260914: analyticalHardware goes from NULL to present, while the base stepping behaves exactly as before. Co-authored-by: Cursor <cursoragent@cursor.com>
…logic The 64 Gridbased logic files shipped InternalSupportParams with KernArgsVersion: 2. That key is bound to the generator version rather than being a tuning result, so library logic is meant to follow defaultInternalSupportParams instead of pinning a layout that goes stale; reorderSolutionDictForDictMerge drops it on write for exactly that reason. Every other tree already reflects this -- gfx950's 129 files and gfx1250's 64 carry no occurrence, and gfx1250-strict was the only one that did. Drop the key here too. The .s metadata and benchmark solution files still carry it, so nothing that needs the generator's layout loses it. The remaining line now matches gfx1250's byte for byte, so a later regeneration will not churn these 64 lines again. Co-authored-by: Cursor <cursoragent@cursor.com>
_detectArchNames asks a device enumerator first so that a target.lst or an HSA_OVERRIDE_GFX_VERSION pin still wins, and took that answer verbatim. As of ROCm 10.2 amdgpu-arch -- the default enumerator -- is a trampoline that execs offload-arch, which names an agent from the KFD node's gfx_target_version alone and loads neither ROCr nor HIP. Both gfx1250 steppings publish 120500 there, since the revision lives in the node's capability bits 25:22 that offload-arch never reads, so an A0 part came back as a bare "gfx1250" and the detect_gpu_archs fallback never ran. Configs that name no ISA take their target from here, so on strict silicon they built and tuned as base. Nothing failed to say so: a base code object still loads and still runs on A0, the strict hazard workarounds were simply absent, and the only sign was the architecture in the artifact path. Route the enumerator's answer through restore_steppings. ROCr does apply the revision rule and rocminfo reports through ROCr, so it still answers gfx1250-strict. Letting rocminfo only ever lengthen a name keeps each tool doing what it is good for: rocminfo needs read-write /dev/kfd, so a caller outside the render group gets nothing from it and has to keep enumerating with amdgpu-arch. A base rocminfo reports under more than one spelling is left alone rather than guessed at, since a box holding both steppings has no single right answer to substitute, and a base the enumeration already spells with a stepping is left alone for the same reason from the other direction -- that answer came from a tool that can tell the two apart on this box, so the device rocminfo happens not to be reporting would otherwise be renamed to its neighbour's stepping. Three testenvs had been adding --gpu-targets to the Tensile options whenever the detected name ended in -strict. That covered for the detection defect above while causing one of its own, forcing the strict stepping onto configs pinned to other ISAs; with detection fixed it is both unnecessary and wrong, so all three are gone. Adding --gpu-targets to make a strict test pass is a symptom of this bug rather than a fix for it. The same three testenvs exported TENSILE_ARCHITECTURE from the detected name, which would put an unnormalized string back into Tensile downstream of the normalization below. Nothing else reads the variable, so those are gone too. HSA_DISABLE_GFX12_STRICT picks which stepping the runtime reports on a revision-0 part, and tox does not inherit a variable that is not listed, so it read as unset inside the testenv -- which on ROCm 10.2 means the base one. That name drives both test selection through skip-<arch> and the client build, so dropping it did not narrow the run, it ran the base half of the corpus on strict silicon and reported green. An installed ROCm ships libtensilelite-host.so.1 too, and $ROCM_PATH/lib ahead of the build directory let that copy win. The two are not interchangeable: CMakeLists.txt defines TENSILE_YAML or TENSILE_MSGPACK and never both, so a msgpack-only host library has no reader for the YAML libraries Tensile writes by default. Every hardware config then failed at once, the client dying on "Failed to load solution library" before it ran a single GEMM, so put the build directory first. build-client only overrides the compilers when it is given a ROCm path and has no detection fallback of its own, so CMake's search won and found /opt/rocm's amdclang++ -- which rejects gfx1250-strict outright on a host whose system ROCm predates the target. Pin $ROCM_PATH at each call site. The quotes are escaped because tox parses the command with shlex and would otherwise strip them before the inner shell ever sees them. The test-side detection called rocm_agent_enumerator, which parses rocminfo with a capture group ending at gfx\d+ and so truncates a suffix in exactly the way described above. That list is what config_helpers. configMarks keys its skip-<arch> marks off, and the two spellings of a config carry mirrored marks -- a strict config says skip-gfx1250, a base config says skip-gfx1250-strict. Handing that comparison a truncated name therefore reverses the selection rather than narrowing it. Read through Tensile.GpuArch instead, the same detection the build uses, with TENSILE_ROCM_PATH applied around the call and taken back out again so the test-specific override does not leak into anything the tests launch. That detection answers with a configuration rather than an architecture, so the answer needs normalizing before a mark name is built from it: amdgpu-arch names a gfx90a agent gfx90a:sramecc+:xnack-, listing the target features that agent happens to have, while a mark is written for the architecture. Since configMarks builds its names by concatenation -- "skip-%s" % arch -- an unnormalized answer looks for skip-gfx90a:sramecc+:xnack- and finds nothing any config file spells, unskipping every config pinned off that architecture: 446 tests on gfx90a, 380 on gfx942, 324 on gfx950, each announced beforehand by an Unknown pytest.mark warning carrying the feature string that nobody reads as an error. Route the answer through cmake_gpu_target, the same normalization the build spells GPU_TARGETS with, so only the colon-delimited features come off. The hyphenated stepping stays, for the mirrored-mark reason above: truncating gfx1250-strict to gfx1250 does not narrow that selection either, it reverses it. The enumerator this replaces had been taking the features off as a side effect of taking the stepping off, which is why both axes have to be spelled out here. _rocmShimReporting shims the detection tools through a fake ROCM_PATH so a test cannot fall through to PATH and answer from the machine running it. rocminfo is now consulted to restore a stepping, so it needs shimming for the same reason; left out, it reported this machine's stepping onto whatever architecture the caller asked for. It is the one tool here asked in its own format, read for indented Name: lines rather than bare ones. Co-authored-by: Cursor <cursoragent@cursor.com>
The rebase onto develop picked up the true16 NoSDWA change, whose arch test "^gfx1[12][0-9][0-9]$" does not match a stepping suffix. Merging the two cleanly left gfx1250-strict extops assembled without +real-true16 while gfx1250 got it, with no error to show for it. Widen it the way the wavefront test above it is already widened. A stepping shares its ISA with the architecture it steps, so true16 comes with it; only what A0 genuinely lacks is gated per-stepping.
mengzcai
force-pushed
the
users/geotseng/develop-gfx1250-strict
branch
from
September 18, 2026 08:51
3dd8f22 to
43cdd0e
Compare
Alex-Vasile
reviewed
Sep 18, 2026
Comment on lines
+2
to
+3
| - gfx1250-strict | ||
| - gfx1250-strict |
Contributor
There was a problem hiding this comment.
Why are both of these gfx1250-strict?
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
8 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Fan out the build when two architectures share an ISA
Follow-up to #10784, which
introduced
gfx1250-strictas an ordinary architecture name. That PR left onespecial case behind in CMake; this one removes it.
Motivation
gfx1250-strictis a silicon stepping ofgfx1250. The two share the ISA tuple(12, 5, 0)but are separate compiler targets producing different machine code(ELF
EF_AMDGPU_MACH0x49vs0xEB), and a code object built for one isrejected on the other.
Tensile keys its capability map by ISA, so a single
TensileCreateLibraryinvocation cannot describe both. #10784 worked around that in the build system:
device-library/CMakeLists.txtpeeledgfx1250-strictinto a secondhipblaslt_create_device_librarycall with its ownOUTPUT_DIR, and thetop-level lists carried an extra install rule driven by a global property. That
put an architecture name in the build system and would need extending by hand
for the next stepping.
This PR moves the split into
TensileCreateLibrary, where the knowledge ofwhich 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
architecture name would touch ~88
isaInfoMap[...]subscripts across 24 filesplus two rocisa C++ sites and the nanobind ABI, on every architecture's code
path. This PR separates the two architectures only where their outputs would
otherwise be written to the same place.
Technical Details
Fan-out
isaCollisionFreeGroups()partitions the requested names so no group holds twothat spell one ISA. A run handed more than one group re-executes itself once per
group as a subprocess -- 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 naming no
stepping, which is every build that exists today apart from gfx1250's.
Jobs are split by architecture count rather than evenly. What forces a second
group is one stepping colliding, which leaves that group holding a single
architecture and the other holding the rest; an even split would run the large
group at half speed and idle half the machine once the small one finished.
Scratch directories
The groups run concurrently into one shared
--output-path, and everything theywrite outside
library/would collide there. Kernel basenames are derived fromthe ISA --
getParameterValueAbbreviationrenders(12, 5, 0)as"1250"forboth architectures -- so both runs name their
.sand.oidentically while themachine code inside differs. The fixed-name artifacts (
Kernels.cpp,Kernels.h, the six static headers,code_object_tmp/Kernels.o) carry noarchitecture at all, and the static headers
#includeeach other by literalname inside shipped C headers, so no naming convention can separate them.
Scratch therefore becomes
build_tmp/<STEM>-<stepping>for a run that was askedfor a stepping, and stays
build_tmp/<STEM>for every run that was not -- soarchitecture sets that predate steppings keep the directory they had. Fan-out
children carry
TENSILE_GROUP_BUILDso they reclaim only their ownsubdirectory and leave a concurrent sibling's alone; a standalone run still
removes
build_tmpwhole.Benchmark cache key (correctness fix)
The tuning cache key gains the compiler targets. Two steppings agree on every
other field the key covers, so without them the second run would load the
first's code objects and hand the wrong machine code to silicon that cannot run
it.
_loadLegacyCacheIfMatchesis removed rather than kept as a fallback: that pathis keyed by directory alone, so it would bypass the new field and reintroduce
exactly the mismatch.
Architecture name matching
Tuning is selected by
ArchitectureName, which a stepping spells as itself, sogfx1250-strictlogic is no longer absorbed intogfx1250.archMatch()compares whole names for that reason -- a prefix comparison wouldlet
gfx1250claim a request forgfx1250-strict. It also strips a bracketedpredicate as well as a qualifier:
TensileLogicfilters withGPU_TARGETSexactly as CMake passed it, having never split predicates off, so a whole-name
comparison alone would fail
gfx950[cu=64]against its own logic and stop thebuild. Both halves are load-bearing, and the test suite pins each direction
independently.
steppingArchOf()derives a stepping rather than taking it from a table: a nameis a stepping exactly when it does not survive a round trip through its ISA. A
future stepping needs no registration.
Removed knobs
With the stepping named for itself throughout,
HIPBLASLT_ASIC_REVISIONand--asic-revisionhave nothing left to select and are removed. Picking arevision is now spelling it in
GPU_TARGETS.Packaging checks
validate_library_layout.pyandcheck_dat_integrity.pyaccepted onlyalphanumerics in an architecture token, so a stepping's files matched nothing.
An unmatched file leaves its subtree without masters, which the scanner reads as
a subtree to skip -- so the failure mode was a package passing silently rather
than an error. The token now admits hyphens.
Reproducibility
Logic files are merged in sorted order. Set iteration and
readdirorder wereleaking salted string hashing into the solution indices written into the master
library, so a build was not reproducible run-to-run. See "Deliberate behaviour
changes" below for the consequence.
Test Plan
Full procedure, including how to reproduce the environment, is in
docs/developer/gfx1250-strict-testing.md(added in this PR).pytest -m unitoverTests/unitandTests/extras.architectures. Every name in
architectureMapis expanded into everyspelling CMake may emit (bare,
:xnack+,:xnack-,[cu=64],[id=74a0])and checked in a matched pair: each spelling must match its own logic header,
and must be claimed by no other architecture.
gfxToIsa,baseArchName,steppingArchOf,isaToGfx,archNamesByIsa,splitArchsFromPredicates,isaCollisionFreeGroups,expandAllArchitecturesandarchMacroNamesaredumped for every architecture x every spelling from both this branch and a
worktree at the base commit, and the JSON diffed.
tools/scripts/tests/.ctest, run serially../install.sh -c -a "gfx1250;gfx1250-strict", then file countsper architecture, the ELF machine flag of both sets of code objects, and
validate_library_layout.pyon the install tree.*quick*against the locally built install tree, runtwice on one host: once with the runtime reporting
gfx1250-strict, oncewith
HSA_DISABLE_GFX12_STRICT=1so it reportsgfx1250. This is the A/Bthat shows the two steppings select different libraries.
The guard in step 2 was validated by fault injection rather than assumed: making
archMatchcompare prefixes fails 7 tests, allgfx1250-strictspellings;dropping the predicate strip fails 48, across every predicated legacy target.
Neither direction alone would have caught the original regression.
Test Result
-m unit)test_arch_steppings.py)ctest0x49(gfx1250),0xEB(gfx1250-strict)validate_library_layout.pyhipblaslt-test *quick*, gfx1250-stricthipblaslt-test *quick*, gfx1250The 121-spelling parser diff shows the only behaviour changes are the intended
ones:
archMatchmatches 120/120 spellings against their own logic header, nolegacy architecture is claimed by a
gfx1250*header, and thesplitArchsFromPredicatesdifference is ordering-only over the same 34-elementset.
The 144-test gap between the two
*quick*runs is accounted for entirely by theknown_bug_matmul_small_bf16quarantine: 144 such tests ongfx1250, 0 ongfx1250-strict, and 18043 - 17899 = 144. Verified directly rather thanassumed, since a gap of any other size would mean something else changed.
Deliberate behaviour changes
These are intended and should be called out in review.
folded into the hash unconditionally, matching
HelperKernelCache, soexisting tuning output directories miss once and recompile. A conditional key
would have kept the old value for non-stepping builds, but would also have
left the same collision in place for any two architectures tuned into one
output directory. Stale directories are inert, not invalid.
.datbytes change for every architecture. Making logic-filemerge order deterministic changes the solution indices written into the
master library. Worth knowing when diffing against a reference build.
<destPath>.linker_args, so two concurrent runs cannot overwrite eachother's.
Kernels.cpp,Kernels.hand the six static headers no longer survive inthe output root. They are written under the run's scratch directory
instead, because their names carry no architecture and a concurrent sibling
would otherwise overwrite them.
_loadLegacyCacheIfMatchesis removed (see above).Known, not addressed
isaInfoMapis annotatedDict[str, IsaInfo]across many signatures whileproduction keys it by
IsaVersion. The fakes inDuctile/test_benchmark_problems.pyfollowed the annotation and used stringkeys; they are corrected here, but the annotation itself is left alone since
fixing it is a wide, unrelated sweep.
Makefile:55gatesTRUE16_FEATUREonTARGET_ARCH, which is assignednowhere in the repository, so
make conever passes+real-true16for anyarchitecture, while
Component.py:173passes it for gfx1100, gfx1200,gfx1201, gfx1250 and gfx1250-strict. A Makefile rebuild therefore does not
reproduce what Tensile assembled. On the 10.1 nightly the feature is already
the default for these targets and the objects are byte-identical either way;
on ROCm 7.13 three of four measured b8b8s kernels fail to assemble without it,
loudly rather than as a silent miscompile. This predates the branch.
one.
Validators.py:121picksrocm_agent_enumeratorwhenisRhel8() or inFFMEnv, and that tool'sgfx\d+capture truncatesgfx1250-stricttogfx1250._detectArchNamesfalls through toamdgpu-archonly when theenumerator answers with nothing, and a truncated name is not nothing, so the
fallback never runs. Measured on A0 silicon: the enumerator says
gfx1250,amdgpu-archand the fallback path both saygfx1250-strict. A build thatnames its target --
--gpu-targets,-a, or the CMake default, which listsboth -- is unaffected; only a bare
Tensile <yaml> <out>is, and it thenbuilds the shipping stepping and reports success. Ordinary Linux hosts select
amdgpu-archand are correct.Not in this PR
The tuning logic move from
Logic/asm_full/gfx1250v0/togfx1250-strict/isdeliberately left out and will follow separately.
Submission Checklist
JIRA ID : AIHPBLAS-4727
JIRA ID : AIHPBLAS-4728
JIRA ID : AIHPBLAS-4729
JIRA ID : AIHPBLAS-4730
JIRA ID : AIHPSPAR-275