Repository navigation
fix(rocke): make static ISA analysis work on RDNA and gfx12 - #13965
Conversation
parse_isa only recognized CDNA instruction names, so on gfx11/gfx12 it reported zero matrix ops, zero loads and zero LDS reads, double-counted gfx12 split barriers and missed gfx12 split memory waits. parse_isa: - count v_wmma/v_swmmac (new `wmma`), global_load/global_store and global_load_lds, and scratch_load/scratch_store (spills) - count gfx11+ ds_load/ds_store as LDS reads/writes - count a gfx12 split barrier (s_barrier_signal + s_barrier_wait) once - count gfx12 s_wait_*cnt as memory waits; s_wait_alu gets its own field - keep WMMA out of VALU and waits/barriers out of SALU global_load/global_store were not counted on any arch before; existing CDNA fields are unchanged. ResourceInfo: replace workgroup_size_hint, which never held a size, with max_flat_workgroup_size from the code-object metadata. compare_variant_reports adds wg_launch and wg_compiled; a compiled bound above the launch size means LLVM ignored the flat-work-group-size attribute and budgeted registers for a larger workgroup. probe_isa_inspect: take headline counts from parse_isa instead of a separate CDNA-only classifier, and compile for `mcpu` (it always built for gfx950, so gfx11/gfx12 kernels failed to build). CDNA output is unchanged; gfx12 wait patterns keep the counter name.
github-actions
Bot
added
documentation
project: hip-kernel-provider
rocke: platform
✅ 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. |
probe_occupancy, probe_lowering_compare, probe_config_sweep and probe_intrinsic_counts built or lowered every kernel for gfx950 whatever target was requested, so gfx942/gfx90a results described a gfx950 kernel and gfx11/gfx12 kernels failed to build or lower. - probe_occupancy: compile for arch.name (caps exist for CDNA only) - probe_lowering_compare: LLVM-direct side compiles for `arch`, like the HIP side already did - probe_config_sweep, probe_intrinsic_counts: new `arch` parameter, default gfx950 - probe_intrinsic_counts: count WMMA and SWMMAC intrinsics gfx950 output is unchanged.
…ntrinsics Now that the dsl_probes compile for the requested arch, a kernel that needs an intrinsic the target cannot lower (e.g. a gfx950-only MFMA built for gfx942) emits ArchDomainWarning and then aborts the whole process inside LLVM, so the per-entry BUILD-FAIL handling never runs. probe_isa_inspect, probe_occupancy, probe_lowering_compare and probe_config_sweep now treat ArchDomainWarning as an error around compile_kernel. The warning only fires on a measured, definite negative, so the entry fails with a message naming the intrinsic and the run continues. Kernels that build are unaffected.
…counts LLVM names the older MFMA intrinsics without a dot before the type (mfma.f32.16x16x16f16, mfma.f32.32x32x8f16, mfma.f32.16x16x16bf16.1k), while the gfx950 ones have it (mfma.f32.16x16x32.f16). The 16x16x16 f16, 16x16x16 bf16, 32x32x8 f16 and 4x4x4 f16 patterns required the dot and never matched; make it optional. gfx950 counts are unchanged.
…ested arch compile_hip_backend passed arch to hipcc but not to lower_kernel_to_hip, so the HIP source was generated for gfx950 and gfx11/gfx12 kernels failed before hipcc ran. gfx950 output is unchanged.
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #13965 — make static ISA analysis work on RDNA and gfx12 — 🟢 First-round: no closer look needed — counts verified against real gfx1151 / gfx1201 code objects
Net-new: none. Utility: 🟢 Positive — the static-inspection step now reports real matrix, memory, LDS, barrier and wait counts on RDNA / gfx12 instead of zeros and double counts. Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted First-round review: a local qwen3.6 model made the independent first pass; the orchestrator (Claude Opus 5.5) verified it — merge simulation with develop, |
|
I reviewed this PR again. The following review threads are still unresolved, or comments have no author reply: If you have already fixed these, please reply on each comment to say so and resolve the threads. Please request my review again once all comments and threads are resolved. |
|
No changes needed from the first-round review: the counts match my own checks (independent mnemonic tally on gfx1151/gfx1201/gfx1250, and the TESTING.md Tier 1 gate on ROCm 10 shows the same pre-existing failures on this branch and on develop). The branch merges cleanly with develop |
CongMa13
left a comment
CongMa13
left a comment
There was a problem hiding this comment.
I found two actionable issues in the requested-architecture probe path; details are in the inline comments.
…-rdna-isa-analysis
…-rdna-isa-analysis
…orted histograms Review follow-ups: - dsl_probe_attention_demos.py forwards --arch to the config_sweep and intrinsic_counts demos; both previously used the probes' gfx950 default. --arch now accepts every known_arches() target. Only occupancy needs resource caps, so it alone rejects targets missing from ARCH_BY_NAME, with a CLI error instead of a KeyError. - probe_intrinsic_counts promotes ArchDomainWarning to an error per entry, like the other probes, so a kernel the target can't lower reports LOWER-FAIL instead of a histogram for an unsupported target. Later entries still run, and the caller's warning filter is restored. - README: document both behaviors and the HIP side's arch argument. - Tests: CLI target forwarding/choices (library) and per-entry LOWER-FAIL with real lowering (platform). Both skip when the dsl_docs tools are not installed.
github-actions
Bot
added
the
rocke-lib: attention
CongMa13
left a comment
CongMa13
left a comment
There was a problem hiding this comment.
I found one actionable issue in the new workgroup comparison.
…-rdna-isa-analysis
compare_variant_reports labelled KernelDef.max_workgroup_size as wg_launch. That value is the declared upper bound the flat-work-group-size attribute is lowered from (256 by default), not the launch size, so a kernel launched with fewer threads showed launch == compiled and hid the oversized register budget the column was meant to flag. - The row now has wg_bound (the declared bound), wg_launch (the actual threads per workgroup, None unless the caller supplies it) and wg_compiled (the size the compiler budgeted for). - VariantReport gains an optional launch_block_size, also accepted by from_artifact and included in as_dict. - Tests cover a bound-only report (wg_launch stays None) and a kernel left at the 256 default but launched with 64 threads.
…_probe The probe changes were only checked by hand before. Add host-only regressions in platform/tests/tools/test_dsl_probes.py: - probe_isa_inspect, probe_occupancy, probe_lowering_compare and probe_config_sweep compile for the requested target and fail one entry (not the whole run) when the target can't lower the kernel. Driven with a gfx942 kernel that needs ds.read.tr16.b64, so each entry fails during lowering and no compiler runs. - probe_intrinsic_counts: the per-entry failure test moves here from test_probe_intrinsic_counts.py. - count_intrinsics matches the pre-gfx950 MFMA names (no dot before the type) and WMMA/SWMMAC, using names rocKE's lowering emits. - categorize_isa counts verbatim gfx12 disassembly (WMMA, global, scratch, ds_load/store, split waits and barriers) and keeps gfx12 wait names. All seven fail against develop's probes and pass with this branch. They skip when the dsl_docs tools are not installed.
Like the WaveScope sidecar and LDS analyzer tests, it loads tools from the dsl_docs source tree, which is not installed.
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #13965 — make static ISA analysis work on RDNA and gfx12 — 🟢 Follow-up: no closer look needed — all three of CongMa13's items are fixed with tests, the counts we checked last time are untouched, and CI is green at head
Utility: 🟢 Positive — the static-inspection step now reports real matrix, memory, LDS, barrier and wait counts on RDNA / gfx12. The probes honour the requested target, and the report no longer passes a bound off as the launch size. Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted Follow-up review: a local qwen3.6 model made the independent first pass. The orchestrator (Claude Opus 5.5) verified it. The checks were a range-diff against the head we reviewed, the new and changed probe and report tests on a merge with today's develop and on develop, a check of their CTest enrollment, and the CI logs at head. The orchestrator owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection. |
|
I reviewed this PR again. The following review threads are still unresolved, or comments have no author reply: If you have already fixed these, please reply on each comment to say so and resolve the threads. Please request my review again once all comments and threads are resolved. |
|
Checked with a real cmake --install: the dsl_docs probes aren't installed (no dsl_probes dir, no probe_*.py), and the demo driver imports them at load time, so the driver itself can't run in the installed tree. Run there, this test skips all 17 cases, so enrolling it in a CTest lane would add a permanently skipped entry rather than coverage. It's source-only for the same reason as test_dsl_probes.py (TESTING.md gap G6). It runs from a source checkout and fails against the pre-fix driver. |
…-rdna-isa-analysis # Conflicts: # dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #13965 — make static ISA analysis work on RDNA and gfx12 — 🟢 Follow-up: no closer look needed — only a develop merge since our last review, and the author's answer on our N1 checks out, so we retract it
No net-new findings. The worker flagged Utility: 🟢 Positive — the static-inspection step reports real matrix, memory, LDS, barrier and wait counts on RDNA / gfx12, and the probes honour the requested target. Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted Follow-up review: a local qwen3.6 model made the independent first pass. The orchestrator (Claude Opus 5.5) verified it. The checks were a range-diff against the head we reviewed and the author's install claim, tested by running the demo tests with the |
cgmillette
left a comment
cgmillette
left a comment
There was a problem hiding this comment.
LGTM, pending @CongMa13 's approval
|
I reviewed this PR again. The following review threads are still unresolved, or comments have no author reply: If you have already fixed these, please reply on each comment to say so and resolve the threads. Please request my review again once all comments and threads are resolved. |
|
I reviewed this PR again. The following review threads are still unresolved, or comments have no author reply: If you have already fixed these, please reply on each comment to say so and resolve the threads. |
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #13965 — make static ISA analysis work on RDNA and gfx12 (JohnUzoka, AICK-2653) — 🟢 Follow-up: CI red is develop's
Net-new: none. Utility: 🟢 Positive — real RDNA / gfx12 ISA counts; probes honour the requested target. Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted Follow-up review: a local qwen3.6 model made the independent first pass. The orchestrator (Claude Opus 5.5) verified it. The checks were the PR's own diff at the old and new heads, including the hand-resolved CMake conflict, and a merge with today's develop: native build and ctest, the differential golden checks that CI runs, and the PR's three test files. They also covered the CI logs of all five attempts at head, traced to the develop commit that CI built on. The orchestrator owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection. |
TheRock Submodule Bump ActivityNewest first
|
Motivation
rocKE's static ISA analysis (
rocke.analysis.parse_isa) only recognizes CDNA instruction names.On gfx11/gfx12 it reports zero matrix ops, zero global loads and zero LDS reads, double-counts
gfx12 split barriers, and misses gfx12 split memory waits. The runbook's dsl_probes also build or
lower every kernel for gfx950 regardless of the requested arch, so gfx942/gfx90a results describe
a gfx950 kernel and gfx11/gfx12 kernels fail to build. The static-inspection step of the runbook
therefore gives no usable signal on RDNA.
Technical Details
parse_isa:v_wmma/v_swmmac(newwmma),global_load/global_store/global_load_lds, andspills (
scratch_load/scratch_store)ds_load/ds_storeas LDS reads/writess_barrier_signal+s_barrier_wait) onces_wait_*cntas memory waits;s_wait_alugets its own fieldglobal_*was not counted on any arch beforeResourceInfo.max_flat_workgroup_sizereplacesworkgroup_size_hint, which never held a size(
Nonefor code objects, the WGP-mode flag in assembler text).compare_variant_reports: addwg_bound(the declaredmax_workgroup_size),wg_launch(the real threads per workgroup, from the new optionalVariantReport.launch_block_size;Nonewhen not given) andwg_compiled(the size the compiler budgeted for). A compiled size above the bound means
LLVM ignored the
amdgpu-flat-work-group-sizeattribute; above the launchsize, the register budget is smaller than the launch needs.
probe_isa_inspect.py: headline counts come fromparse_isa; compile formcpu; gfx12 waitpatterns keep the counter name.
probe_occupancy.py: compile forarch.name(caps exist for CDNA only).probe_lowering_compare.py: both the LLVM-direct and the HIP-debug side lower forarch.probe_config_sweep.py,probe_intrinsic_counts.py: newarchparameter (defaultgfx950).probe_intrinsic_counts.py: count WMMA/SWMMAC intrinsics; the 16x16x16 f16/bf16, 32x32x8 f16and 4x4x4 f16 MFMA patterns never matched (pre-gfx950 intrinsic names have no dot before the
type) and now do.
aborting the process inside LLVM or reporting results for an unsupported target. All five
probes, including
probe_intrinsic_counts, raise onArchDomainWarningper entry; laterentries still run and the caller's warning filter is restored.
library/builders/common/dsl_probe_attention_demos.py) forwards--archto all five demos.--archaccepts anyknown_arches()target; onlyoccupancyneeds resource caps, so it alone rejects targets without them, with a CLI error.
Not in scope:
tensor_load_to_lds(gfx1250) is not counted; transpose loads have no counter oftheir own;
probe_occupancyhas no RDNA/wave32 resource caps.Test Plan
llvm-objdumplines andllvm-readelf --noteslines.platform/tests/tools/test_dsl_probes.py,library/tests/test_dsl_probe_attention_demos.py): every changed probebuilds for the requested target and fails per entry when it can't; the
intrinsic and ISA counters match MFMA/WMMA names and gfx12 disassembly;
the attention demo forwards
--arch. They fail against develop's probes,and skip when the dsl_docs tools are not installed.
platform/tests/run_all.py(Tier 1) on ROCm 10.2 (clang 24), on this branch and ondevelopat the same commit, each in a fresh
TMPDIR.kernels, GEMMs) for identical output, and every other known target through the demo CLI.
probe_lowering_comparewas run with real hipcc.parse_isacounts compared with an independent tally ofllvm-objdumpmnemonics.Test Result
fails the same pre-existing set on this branch and on
develop. The IR-validity gate isUNVALIDATED on both (no llvm24 flavor for ROCm 10's clang).
isa_inspectonly reorders rows withequal counts). For any other target each probe compiles or lowers for that target; kernels that
can't build for it fail per entry, and valid CDNA/RDNA/gfx12 kernels still produce results.
JIRA ID: AICK-2653