Skip to content

feat(rocke): add LLVM24 compiler lane - #14415

Merged
shumway merged 4 commits into
developfrom
users/shumway/aick-2740-llvm24-profile
Oct 29, 2026
Merged

shumway merged 4 commits into
developfrom
users/shumway/aick-2740-llvm24-profile

Conversation

@shumway

@shumway shumway commented Oct 26, 2026

Copy link
Copy Markdown
Contributor

Summary

Add LLVM 24 as an explicitly qualified rocKE compiler lane across the Python and native lowering implementations.

This separates compiler qualification from structural LLVM-IR emission: LLVM23 and LLVM24 are distinct lanes that intentionally share one emission profile because toolchain characterization found no datalayout, declaration, or call-shape delta. Future compiler majors now fail closed instead of silently inheriting the newest known profile.

Risk Assessment

Risk level 3. This changes core compiler selection and LLVM lowering policy, including the public native flavor enumeration, but does not change LLVM20/22/23 emitted output or kernel numerics. The principal risks are cross-engine drift, incorrect target-operation admission, and incomplete compiler-version handling; focused profile tests, all-flavor byte identity, real LLVM24 compilation, native CTests, and unchanged existing goldens cover those risks.

Related

ISSUE ID : AICK-2740

Related work: AICK-2708, AICK-2221, AICK-2741, AICK-2742.

Device / Architecture Coverage

This is compiler-lane and host-side lowering work rather than an architecture-specific kernel change. The LLVM24 emitted-IR validity corpus compiled and linked all 284 representative modules across the wired targets.

Runtime and numerical GPU qualification are intentionally outside this change. LLVM24 arch-domain artifact generation and broader CI enrollment remain separate follow-up work.

Testing Summary

  • Full developer pre-merge runner against an LLVM24 toolchain.
  • Python/C++ LLVM byte identity at every supported compiler lane.
  • Real LLVM24 datalayout validation across all wired targets.
  • Focused Python and native tests for lane resolution, profile capabilities, scaled MFMA, target operations, unknown values, and future compiler majors.
  • Representative and library golden checks using profile-aware aliases.
  • Adversarial multi-agent review followed by two verification passes.

Testing Checklist

  • Full platform gate — ROCM_PATH=<LLVM24 ROCm root> python platform/tests/run_all.py — Status: Passed
    • 4,632 passed, 1,421 skipped
    • 31/31 native CTests passed
    • 284/284 representative LLVM24 modules compiled and linked
    • 83 byte-identity families passed, including the LLVM24 target-operation families
  • LLVM20 byte identity — ROCKE_LLVM_FLAVOR=llvm20 python platform/tools/check_byte_identity.py — Status: Passed
  • LLVM22 byte identity — ROCKE_LLVM_FLAVOR=llvm22 python platform/tools/check_byte_identity.py — Status: Passed
  • LLVM23 byte identity — ROCKE_LLVM_FLAVOR=llvm23 python platform/tools/check_byte_identity.py — Status: Passed
  • LLVM24 byte identity — ROCKE_LLVM_FLAVOR=llvm24 python platform/tools/check_byte_identity.py — Status: Passed
  • Changed library golden suites — Status: 68 passed, 94 capability/configuration skips
  • Repository pre-commit hooks over all changed files — Status: Passed
  • Pybind runtime matrix — local environment lacks Python development headers; PR CI is expected to exercise the built extension
  • LLVM24 arch-domain artifact generation — deferred to AICK-2742; the current runner reports a named unvalidated skip rather than probing with the wrong profile
  • PR CI — GitHub PR checks — Status: Pending

Flags / Guardrails

No feature flag is required. LLVM24 maps to an explicit qualified lane while sharing the characterized LLVM23 emission profile. Compiler majors newer than LLVM24 fail with a diagnostic instead of silently selecting that profile.

Adjacent Tests Considered

  • Native compiler discovery and retained COMGR provenance
  • Emitted-IR verification and linked-HSACO validity
  • Datalayout drift against real compiler output
  • Architecture-domain warning and artifact-reproduction behavior
  • Portable-IR replay flavor parsing
  • Representative and workload-specific LLVM-IR goldens
  • Python/native explicit-flavor rejection contracts

Runtime/numeric qualification and generation of the LLVM24 architecture-domain column are deliberately handled by separate qualification work.

Technical Changes

  • Add llvm24 to Python, native C, C++, pybind, and replay CLI flavor surfaces.
  • Introduce neutral Python and C++ compiler-lane/emission-profile registries shared by lowerers and ISA adapters.
  • Map LLVM major 24 to llvm24, preserve 23 to llvm23, and reject newer unqualified majors.
  • Replace exact LLVM23 operation checks with profile capabilities or minimum-lane requirements.
  • Keep LLVM23 and LLVM24 on the same characterized datalayout, declaration ABI, scaled-MFMA ABI, and target-operation capability set.
  • Make explicit unknown-flavor behavior consistent across Python, native, and pybind paths.
  • Validate golden aliases against the authoritative emission-profile mapping while still lowering and checking every compiler lane.
  • Avoid running an absent LLVM24 arch-domain column as though it were measured; report a fast named skip until the dedicated artifact work lands.

@therock-pr-bot

therock-pr-bot Bot commented Oct 26, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

@shumway
shumway force-pushed the users/shumway/aick-2740-llvm24-profile branch from 8ea3b8a to 6fe923f Compare October 26, 2026 22:05
@shumway
shumway marked this pull request as ready for review October 27, 2026 02:08
@shumway
shumway requested review from a team as code owners October 27, 2026 02:08
@shumway
shumway force-pushed the users/shumway/aick-2740-llvm24-profile branch 2 times, most recently from 6c465a5 to 5d8352e Compare October 28, 2026 14:34
@tenpercent
tenpercent requested a balanced review from Copilot October 28, 2026 18:55

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

A non-string minimum_llvm_flavor is treated as absent, allowing malformed serialized IR to bypass the LLVM20 compiler restriction while Python rejects it.

1 open issue
What changed in this PR

Adds LLVM24 as an explicit rocKE compiler lane across Python and native lowering, sharing LLVM23’s emission profile while rejecting newer, unqualified compiler majors.

Changes:

  • Centralizes compiler-lane and emission-profile definitions, including capability-based operation checks.
  • Extends compiler-selection, lowering, and parity tests to LLVM24.
  • Adds profile-aware golden aliases and reports missing architecture-domain baselines as unvalidated.

Paths below are relative to dnn-providers/hip-kernel-provider/rocke/.

File Description
platform/​tools/​gen_arch_domain.py Uses profile-aware declarations and skips absent default baselines.
platform/​tools/​check_byte_identity.py Selects supplemental checks by profile capability.
platform/​tests/​tools/​test_arch_domain_artifact.py Checks LLVM24 declaration-profile sharing.
platform/​tests/​test_rocke.py Moves unknown-flavor rejection cases to LLVM25.
platform/​tests/​test_compiler_detection.py Tests LLVM24 resolution and future-major rejection.
platform/​tests/​test_binding_compiler_selection.py Tests binding flavor validation and LLVM24 acceptance.
platform/​tests/​portable_ir/​replay_cli.cpp Adds LLVM24 to CLI help text.
platform/​tests/​instances/​test_rocke_multiarch.py Applies datalayout checks by shared profile.
platform/​tests/​instances/​rocke_ir_parity_harness.py Writes profile baselines with lane aliases.
platform/​tests/​instances/​differential/​test_gate_verdict.py Checks environment restoration after golden recording.
platform/​tests/​instances/​differential/​test_check_byte_identity_progress.py Tests capability-based gate selection.
platform/​tests/​instances/​differential/​run_diff.py Restores the caller’s flavor environment setting.
platform/​tests/​instances/​differential/​golden/​llvm_gfx_all.json Updates one differential golden hash.
platform/​tests/​golden/​rocke_representative_ir_sha256.json Adds the LLVM24-to-LLVM23 baseline alias.
platform/​tests/​core/​test_llvm24_profile.py Adds lane, profile, and offline-emission tests.
platform/​tests/​core/​test_ir_golden.py Tests alias validation and all-lane comparison.
platform/​tests/​core/​test_gfx1250_scaled_wmma.py Checks scaled-WMMA ABI under LLVM24.
platform/​tests/​core/​test_gfx1250_llvm_bridge.py Checks LLVM23/24 bridge output identity.
platform/​tests/​core/​future_intrinsic_lowering.cpp Adds native shared-profile emission checks.
platform/​tests/​core/​flavor_datalayout_ssot.cpp Tests native lane mapping and invalid-value handling.
platform/​python/​rocke/​runtime/​comgr.py Rejects future compiler majors before layout checks.
platform/​python/​rocke/​core/​lower_llvm.py Uses shared profiles and minimum-lane requirements.
platform/​python/​rocke/​core/​llvm_profile.py Defines Python compiler lanes and emission profiles.
platform/​python/​rocke/​core/​isa/​backend.py Gates scaled WMMA by profile capability.
platform/​python/​rocke/​core/​ir.py Records minimum compiler requirements for scalar controls.
platform/​python/​rocke/​core/​ir_golden.py Supports and validates profile-aware golden aliases.
platform/​dsl_docs/​reference/​env_flags.md Documents LLVM24 selection and fail-closed behavior.
platform/​dsl_docs/​ir_lowering/​backend_details.md Explains compiler lanes, profiles, and golden aliases.
platform/​dsl_docs/​development/​multi_toolchain_strategy.md Marks earlier strategy tables as historical.
platform/​cpp/​include/​rocke/​lower_llvm.h Adds the public LLVM24 flavor enumerator.
platform/​cpp/​include/​rocke/​lower_llvm_internal.h Delegates profile declarations to the new header.
platform/​cpp/​include/​rocke/​llvm_profile_internal.h Defines internal native profile interfaces.
platform/​cpp/​include/​rocke/​isa_backend.h Documents invalid-flavor accessor behavior.
platform/​cpp/​core/​lower_llvm/​mma.cpp Uses profile capabilities for scaled WMMA.
platform/​cpp/​core/​lower_llvm/​data.cpp Selects declaration overrides by profile ABI.
platform/​cpp/​core/​lower_llvm/​crosslane.cpp Adds minimum-lane checks for inline assembly.
platform/​cpp/​core/​lower_llvm/​core.cpp Uses centralized flavor resolution and capabilities.
platform/​cpp/​core/​llvm_profile.cpp Implements the native lane and profile registry.
platform/​cpp/​core/​isa/​backend.cpp Shares datalayout resolution and rejects unknown flavors.
platform/​cpp/​core/​ir/​ir_tile.cpp Mirrors minimum compiler requirements in native IR.
platform/​cpp/​bindings/​rocke_engine.cpp Reports invalid explicit flavors as ValueError.
platform/​AGENTS.md Adds LLVM24 parity commands and lane documentation.
library/​tests/​test_mla_prefill_gfx942_golden.py Makes golden generation and coverage alias-aware.
library/​tests/​test_kda_gfx950_golden.py Writes distinct profile baselines with aliases.
library/​tests/​test_kda_gfx942_golden.py Writes distinct profile baselines with aliases.
library/​tests/​test_gdn_decode_golden.py Makes baseline generation and coverage alias-aware.
library/​tests/​test_dsa_indexer_golden.py Checks every compiler lane using profile aliases.
library/​tests/​test_attention_dense_gfx942_golden.py Writes profile baselines with lane aliases.
library/​tests/​test_attention_dense_bwd_ir.py Adds alias-aware comparisons and coverage checks.
library/​tests/​golden/​mla_prefill_gfx942_ir_sha256.json Adds the LLVM24 baseline alias.
library/​tests/​golden/​kda_gfx950_ir_sha256.json Adds the LLVM24 baseline alias.
library/​tests/​golden/​kda_gfx942_ir_sha256.json Adds the LLVM24 baseline alias.
library/​tests/​golden/​gdn_decode_gfx950_ir_sha256.json Adds the LLVM24 baseline alias.
library/​tests/​golden/​dsa_indexer_ir_sha256.json Adds the LLVM24 baseline alias.
library/​tests/​golden/​attention_dense_gfx942_ir_sha256.json Adds the LLVM24 baseline alias.
library/​tests/​golden/​attention_dense_bwd_ir_sha256.json Adds the LLVM24 baseline alias.
AGENTS.md Adds the LLVM24 byte-identity command.

Review effort: Balanced

@afagaj afagaj left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Bug-focused review of the LLVM 24 lane. I found no output change for LLVM 20, 22 or 23, and no Python/C++ drift. Two LLVM 24-only issues are below, inline.

@CongMa13 CongMa13 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Jira evidence was not consulted, so this review is based on the PR description and repository code.

@brockhargreaves-amd

Copy link
Copy Markdown
Contributor

AI-assisted review: This review was performed by Codex against commit b169db8, including source inspection, local tests, and compatibility checks with ready-for-review PRs.

No standalone correctness defect was found in this PR, but there are integration conflicts that should be coordinated:

  1. Important: feat(rocke): add MLA decode kernel for gfx942 and gfx950 #13923’s MLA decode admission rejects the new LLVM24 lane. Its admission list

    accepts only llvm22 and llvm23. This PR changes automatic selection for compiler major 24 from llvm23 to llvm24.

    Reproduced on the combined Python source using mocked compiler-version evidence: the same valid request selects MLA decode with major 23 and raises ValueError: no candidate supports request with major 24, on both gfx942 and gfx950. The dispatch code merges without conflict. Please coordinate LLVM24 admission and its dispatch, golden, numerical, and resource checks with feat(rocke): add MLA decode kernel for gfx942 and gfx950 #13923. Its numeric tests also filter on the admission list, so an LLVM24 run can skip the affected
    coverage.

  2. Important: feat(rocke): backend completeness for every production target and dense MMA instruction #14328’s declaration resolver needs reconciliation with the emission-profile model. Its resolver

    for inherited in LLVM_FLAVORS[: LLVM_FLAVORS.index(flavor) + 1]:
    walks compiler lanes through an override map containing only LLVM20/22/23. A compatibility probe adding this PR’s four-lane enumeration makes even empty-kernel lowering raise KeyError('llvm24').

    The pair also has textual conflicts. Please preserve the qualified-lane/emission-profile distinction while incorporating feat(rocke): backend completeness for every production target and dense MMA instruction #14328’s declaration composition and generated matrix declarations in both engines, then validate all four lanes.

  3. Preserve the convolution-family additions from feat(rocke): wave32 WMMA stem dgrad for gfx1250 (patch embeddings, small-C stems) #14387 and feat(rocke): wave32 WMMA direct and small-channel forward convolution for gfx1250 #14392 when resolving the runner conflicts. Both extend LLVM23_ONLY_FAMILIES, renamed here to TARGET_OP_PROFILE_FAMILIES. The resolution needs the union of their family lists, retaining supplementary LLVM23 comparisons for LLVM20/22 runs and direct execution for LLVM23/24.

No semantic conflict was demonstrated with #14437: it merges cleanly, and 28 focused tests passed on the combined source. #14218 already records LLVM24 compiler-option support; its 11 tests passed with this PR.

Local validation of this PR passed 307 Python tests, 32 native tests, and the LLVM20/22/23 plus canonical-IR golden checks. Some configuration-dependent tests and architecture-domain reproduction skipped. Real LLVM24 compilation, the full four-lane byte-identity suite, and the built-pybind matrix were not independently rerun. Provider CI was still pending when checked.

These findings require coordinated changes before the respective PR pairs are integrated; they are not standalone blockers reproduced in this PR.

@shumway
shumway force-pushed the users/shumway/aick-2740-llvm24-profile branch from b169db8 to 1d6e872 Compare October 28, 2026 22:50
@shumway

shumway commented Oct 28, 2026

Copy link
Copy Markdown
Contributor Author

Thanks for checking the interactions with the other in-flight rocKE PRs. These are sequencing considerations rather than standalone defects in #14415, and the referenced PRs are still open rather than present on develop.

#14415 establishes LLVM24 as a qualified lane, keeps compiler lanes distinct from shared emission profiles, and validates all four lanes. Once this lands, the affected PRs should rebase and reconcile against that model:

The current PR has passed the LLVM20/22/23/24 byte-identity gates and focused native/Python validation. To keep this foundational change reviewable, I do not plan to pull unrelated in-flight feature work into it. Please flag if any of these interactions represents a standalone defect in #14415 itself; otherwise they should be handled when those PRs rebase.

@shumway
shumway enabled auto-merge (squash) October 29, 2026 00:12
@shumway
shumway merged commit d695709 into develop Oct 29, 2026
97 checks passed
@shumway
shumway deleted the users/shumway/aick-2740-llvm24-profile branch October 29, 2026 03:19
AviralGoelAMD added a commit that referenced this pull request Oct 29, 2026
Resolve the representative IR golden conflict with #14394 and #14415 by
regeneration: take develop's file, then re-record it with
rocke_ir_parity_harness.py --write on the merged tree. Every entry equals
either develop's or this branch's hash; develop's aliases are kept.
jakpiase added a commit that referenced this pull request Oct 29, 2026
Picks up the LLVM24 lane (#14415) and gfx1250 split-K wgrad (#13936).

- check_byte_identity.py: #14415 renamed LLVM23_ONLY_FAMILIES to
  TARGET_OP_PROFILE_FAMILIES. Take develop's version and add
  conv_direct_fwd_wmma and conv_fwd_small_wmma to the new tuple, so the
  two TDM families still run only in the target-op-profile lane.
- test_check_byte_identity_progress.py: check the split --exclude value
  instead of a whole command element, and check every family in the
  tuple, so the test holds for a multi-family tuple.
- golden: counters recomputed from the merged cases (gfx1250 29/29/47,
  ok_cases 307/307/325); every case hash equals develop's or this
  branch's.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
jakpiase added a commit that referenced this pull request Oct 29, 2026
Conflicts:
- platform/tools/check_byte_identity.py: develop (#14415) renamed
  LLVM23_ONLY_FAMILIES to TARGET_OP_PROFILE_FAMILIES and gates it on the
  selected LLVM profile. Took develop's code and added this PR's
  conv_dgrad_patch_wmma and conv_dgrad_smallc_wmma to the tuple.
- platform/tests/golden/rocke_representative_ir_sha256.json: summary
  counters only. Re-blessed with rocke_ir_parity_harness.py --write; every
  case hash equals develop's or this branch's.

Also fix test_check_byte_identity_progress.py to check the split --exclude
value (and every TARGET_OP_PROFILE_FAMILIES entry) instead of requiring
TARGET_OP_PROFILE_FAMILIES[0] to be a whole command element, which only
held for a one-family tuple (cgmillette N1).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
jakpiase added a commit that referenced this pull request Oct 29, 2026
… for gfx1250 (#14392)

## Summary

Forward convolution on gfx1250 had no direct kernel: the direct forward
families are wave64 MFMA, so every forward layer ran implicit GEMM.

This PR adds wave32 WMMA direct forward kernels for gfx1250, in both
engines, and serves them from the direct-conv AOT cache, like the WMMA
direct dgrad.

JIRA ID: AICK-2777

## Changes

**`conv_direct_fwd_wmma`** — a wave32 WMMA direct forward.
- The input halo and the weights are staged in LDS with TDM transfers,
and every filter tap is a shifted halo read.
- Dense and grouped convolutions; small groups are packed into
block-diagonal super-groups.
- Stride 2, with an optional phase-split halo, and per-axis dilation.
- Weight rings per filter row or per tap.
- Schedule options: a ping-pong wave schedule and B-fragment reuse
across filter rows.

**Tiny channel counts**
- `direct_fwd_smallc_wmma` serves stems with C <= 4, including the 4x4,
14x14 and 16x16 patch-embedding stems.
- `direct_fwd_smallk_wmma` serves K <= 8.

**Depthwise forward:** a channel-packed `cpl` knob, on the wave32 grid
only.

**Benchmark:** `benchmark_direct_conv --run-from-cache` runs problems
past the 2 GiB offset limit as batch chunks.

**AOT:** every family is registered for gfx1250 only, with tuned AOT
tables and per-problem tile-efficiency filtering. gfx942 and gfx950
caches are unchanged.

**Both engines:** every new family and the `cpl` knob are mirrored in
the C++ engine:
- `instance_conv_direct_fwd_wmma.h`
- `instance_conv_fwd_small_wmma.h`
- `conv_direct_grouped_build_depthwise_packed.cpp`
- pybind and backend routing

**Parity and golden:**
- Parity pairs cover every knob: `conv_direct_fwd_wmma_emit`,
`conv_fwd_small_wmma_emit`, and `conv_direct_grouped_emit` configs
171–178 (renumbered after develop's grouped/depthwise wgrad configs).
- The TDM families run only in the LLVM23-generation target-op-profile
lane (`TARGET_OP_PROFILE_FAMILIES` in `check_byte_identity.py`, llvm23
and llvm24).
- The representative IR golden gains 12 cases; no existing hash changes.

**Shared WMMA forward gate:** the arch / wave32 / TDM / dtype /
WMMA-atom checks live in one helper per engine (`conv_fwd_wmma_gate.py`,
`conv_fwd_wmma_gate.cpp`), used by the direct and the small-C/small-K
validators. It returns the resolved `ArchTarget`, so it is not resolved
twice.

**Docs:** support matrix, instance index and convolution docs.

## Review follow-ups

- Both C++ validators reject `stride < 1` before computing output
extents, and describe reports 0 LDS for a rejected stride-0 spec, so bad
input cannot crash Python.
- Explicitly null dtypes are rejected consistently in the Python and C++
bindings.
- Cached-benchmark batch chunking handles unequal chunk sizes;
cache-backed selection tests cover the small-C and small-K paths.
- The duplicated capability gate is factored out (see above); the
emitted IR is unchanged.
- Merged `develop` three times (merge commits, no rebase). The first
merge conflicted in the variant registry, the grouped-conv parity
emitters and the representative golden; the second, with the gfx1250
Winograd forward (#14394), in `core/backend.py`, the parity harness,
`CMakeLists.txt`, the docs and the golden, all insert-vs-insert (both
sides kept). Each time the golden was re-blessed with
`rocke_ir_parity_harness.py --write` on the merge result and equals the
3-way union of base, develop and this branch.
- The third merge picks up the LLVM24 lane (#14415) and gfx1250 split-K
wgrad (#13936). #14415 renamed `LLVM23_ONLY_FAMILIES` to
`TARGET_OP_PROFILE_FAMILIES`; the two forward families now sit in the
new tuple, and `test_check_byte_identity_progress.py` checks the split
`--exclude` value and every family in the tuple, so it holds for more
than one family. Only the golden counters conflicted; they are
recomputed from the merged cases, and every case hash equals develop's
or this branch's.

## Testing

- `check_byte_identity.py --only conv,gfx1250` is GREEN at llvm20,
llvm22, llvm23 and llvm24 on the latest develop merge (at llvm20/22 the
two forward families run in the llvm23 supplemental lane: 22 and 15
configs).
- The representative golden checks clean at every flavor
(`rocke_ir_parity_harness.py --check`, llvm20/22/23/24).
- CPU tests, registered in CTest: `test_direct_fwd_wmma.py`,
`test_fwd_small_wmma.py` and `test_direct_depthwise_fwd_cpl.py`.
- They cover the validators, the AOT grid, and C++/Python agreement on
every AOT configuration and rejected spec.
- gfx1250 GPU numeric tests against a torch reference for all three
families.
- The stem caps were checked on gfx1250 over every cached configuration.

🤖 Generated with [Claude Code](https://claude.com/claude-code)

---------

Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

jakpiase added a commit that referenced this pull request Oct 30, 2026
…all-C stems) (#14387)

## Motivation

Stem convolutions have 3 or 4 dX channels. On gfx1250 their dgrad runs
on implicit GEMM, with the dX channels as the GEMM N tile, so almost
every MMA lane is idle.

This PR adds two wave32 WMMA dgrad families for gfx1250 that fold the
filter taps into the GEMM instead:
- `conv_dgrad_patch_wmma` for patch embeddings (stride == filter, no
padding);
- `conv_dgrad_smallc_wmma` for overlapping stems with `C <= 4`.

Both are in both engines with byte-identical IR. Like
`conv_direct_dgrad_wmma`, they are served by the direct-conv AOT cache;
there is no dispatcher wiring.

JIRA ID: AICK-2776

## Technical Details

**`conv_dgrad_patch_wmma`**
(`library/kernels/common/conv_dgrad_patch_wmma.py`; C++ twin
`instance_conv_dgrad_patch_wmma.h` / `conv_dgrad_patch_wmma.cpp`)
- A patch embedding has no tap reduction, so its dgrad is one GEMM,
`dY[px, K] x W[K, P*P*C]`. The KRSC weight tensor already is that
operand, so there is no weight pre-pass.
- Operand staging is that of `conv_direct_dgrad_wmma` with a 1x1 filter:
- per dY channel chunk, one 3-D TDM transfer of the block's dY pixels
and one of its weight columns, into an LDS ring of `stages` slots;
  - weights are transposed on the LDS read (`ds_load_tr16_b128`).
- A block owns `tile_h x tile_w` dY pixels and `r_blk` patch rows. Its
`r_blk*P*C` GEMM columns are padded to whole 16-column tiles, with at
most an eighth of padding.
- The epilogue pixel-shuffles the tile in LDS as `[i][r][j][s, c]` and
stores it with one rank-3 TDM transfer, clipped to the image.
- When `P` does not divide the image, `N` extra workgroups zero the dX
border no patch covers.

**`conv_dgrad_smallc_wmma`**
(`library/kernels/common/conv_dgrad_smallc_wmma.py`; C++ twin
`instance_conv_dgrad_smallc_wmma.h` / `conv_dgrad_smallc_wmma.cpp`)
- Computes `Z[q, (r, s, c)] = dY[q, :] x W[:, (r, s, c)]` for the dY
window that reaches a block's `out_h x out_w` dX tile, with the same
operand staging.
- `Z` is staged in LDS as fp32 and overlap-added phase by phase (output
row / column mod S), with a compile-time tap list per phase.
- The dX tile is stored by one TDM transfer.

**Common to both**
- Batch, extents and the dY channel count are kernargs. Filter, stride,
padding, the dX channel count, dtype and tile geometry are baked.
- Constraints: ungrouped, `kpg` a multiple of 8, every tensor below 2
GiB, and the LDS ring and staged tiles within the 320 KB LDS.
- The C++ twins share the TDM descriptor, the per-lane LDS read offsets
and the chunk loop through a private header
(`instance_conv_dgrad_stem_wmma_internal.hpp`).
- Plumbing:
- `rocke_engine` bindings
`conv_dgrad_{patch,smallc}_wmma_{lower_llvm,serialize_ir,verify,describe}`;
  - routes `rocke.core.backend.lower_conv_dgrad_{patch,smallc}_wmma`.

**AOT** (`kernels/conv/direct_variants.py`, `direct_aot_jobs.py`,
`benchmarks/common/direct_kernel_sweep.py`)
- New variants `direct_dgrad_patch_wmma` (patch sizes 4 / 14 / 16,
`DGRAD_PATCH_KNOBS`) and `direct_dgrad_smallc_wmma` (the 3x3 s1 / s2 and
7x7 s2 stems, `DGRAD_SMALLC_KNOBS`).
- Both are built on wave32 targets only, so the gfx950 and gfx942 caches
are unchanged.
- `direct_spec_for_problem` offers a binary only where at least half of
its padded tile grid is real work.

**Gate and docs**
- Both families join `check_byte_identity.TARGET_OP_PROFILE_FAMILIES`:
every config moves data with TDM, so under a profile without the
LLVM23-generation target ops they are compared in the llvm23 lane only.
- New parity pairs
`library/tests/parity/conv_dgrad_{patch,smallc}_wmma_emit.{py,c}` and
four golden cases.
- `test_dgrad_stem_wmma.py` is added to the library conv CTest.
- Docs: `SUPPORT_MATRIX.md`, `dsl_docs/instances/index.md`, and a new
section in `dsl_docs/instances/convolution.md`.

## Test Plan

- `check_byte_identity.py`, all families, at llvm20, llvm22 and llvm23
(fresh build roots), plus the `--ir` diagnostic for the new families.
After the latest develop merge: `--only conv,gfx1250` at llvm20 / 22 /
23 / 24.
- `rocke_ir_parity_harness.py --check` on all three flavors.
- `library/tests/test_dgrad_stem_wmma.py` on gfx1250, CPU and GPU:
  - spec validators and the AOT grid;
  - numerics against `torch.nn.grad.conv2d_input`:
- patch sizes 4 / 14 / 16, with and without an uncovered border, padded
GEMM columns, several patch-row blocks;
    - 3x3 stride 1 / 2 and 7x7 stride 2 small-C stems with image edges;
    - bf16 and fp16, prefetch with the `iterative-ilp` scheduler.
- Cross-engine agreement (CPU, when the `rocke_engine` binding is
importable):
- the C++ validator, reason, kernel name, grid and LDS size match Python
over the whole AOT grid and a set of rejected specs, on gfx1250 and
gfx950;
  - the binding rejects a malformed spec.
- Targeted library pytest set (direct-conv, dgrad, AOT cache,
builder-signature and layering tests).

## Test Result

- Byte identity: GREEN at llvm20, llvm22, llvm23 and (after the latest
develop merge) llvm24. `conv_dgrad_patch_wmma` and
`conv_dgrad_smallc_wmma`: 9 configs each, identical at llvm23; the 3
rejected configs per family are rejected by both engines.
- Golden IR: 4 new gfx1250 cases blessed at llvm23, with
expected-failure entries at llvm20 / llvm22. Every existing hash is
reproduced.
- `test_dgrad_stem_wmma.py`: green on gfx1250, GPU numeric cases
included. The cross-engine tests pass with the engine binding.
- Targeted pytest: no new failures. The remaining failures are identical
on `develop`:
  - `test_conv_abi::test_cache_applies_runtime_launch_limits`;
  - six `test_conv_batch_and_miopen_parse` cases;
-
`test_conv_dgrad_patch::TestPatchIdentity::test_off_keeps_the_historical_hash`.
- No performance results are included in this PR.

## Review Follow-ups

- Both validators (Python and C++) reject `stride < 1` before any stride
arithmetic, and the C++ bindings reject an explicit `dtype=None` like
Python does (8b2da4c).
- Cache-backed coverage: `direct_plans` cache-selection tests on CPU and
numerical checks of the cached plans on gfx1250 (2a051c3). Every
shipped 7x7/s2/pad3 C=3 stem knob and every shipped /14 patch-embedding
job (r_blk 2 / 7 / 14) now runs as a cached binary, on images with
partial edge tiles and an uncovered border; `test_every_shipped_knob`
pins that set (1c096a7).
- `conv_dgrad_patch_wmma` needs portable-IR `scf_if_else` support for
its main vs. border workgroups; that support is now on develop (#14394),
so it is no longer part of this diff.
- Merged `origin/develop` (037e89e). The only conflict was the
summary counters in `rocke_representative_ir_sha256.json`; it was
re-blessed on the merge result with `rocke_ir_parity_harness.py
--write`. Every case matches develop's golden except the four new
`conv_dgrad_{patch,smallc}_wmma` cases.
- Merged `origin/develop` again (15afe15, Winograd fwd) in
bc773ef. Same conflict (summary counters only): develop's new cases
and the four stem cases are disjoint, so the cases were unioned and the
summaries recomputed. On the merge: golden `--check` OK at llvm20/22/23,
byte identity (`--only conv,gfx1250`) GREEN at llvm20/22/23,
`test_dgrad_stem_wmma.py` 8 passed on CPU with the engine binding and
the 3 GPU tests passed on gfx1250.
- Merged `origin/develop` again (4423f01, develop 9d729d3).
Conflicts: `check_byte_identity.py`, where #14415 renamed
`LLVM23_ONLY_FAMILIES` to `TARGET_OP_PROFILE_FAMILIES` (the two families
now go in the renamed tuple), and the representative golden's summary
counters (re-blessed with `--write`; every case hash equals develop's or
this branch's). `test_check_byte_identity_progress.py` now checks the
split `--exclude` value, so it passes with more than one family in the
tuple. On the merge (CPU only): byte identity `--only conv,gfx1250`
GREEN at llvm20 / 22 / 23 / 24; golden `--check` OK; differential golden
(ll at llvm20 / 22 / 23 / 24, canonical ir) OK; library targeted pytest
153 passed, 6 skipped, 1 failed (the pre-existing
`test_off_keeps_the_historical_hash`, which also fails on develop);
platform portable-IR / CI-static / byte-identity-driver tests 52 passed.
- Merged `origin/develop` again (78ff090, develop 050acb2, which
brings in #14392's wave32 WMMA forward families). All 10 conflicts were
both-sides additions and take the union: `TARGET_OP_PROFILE_FAMILIES`
now lists `conv_dgrad_patch_wmma`, `conv_dgrad_smallc_wmma`,
`conv_direct_dgrad_wmma`, `conv_direct_fwd_wmma` and
`conv_fwd_small_wmma`; the engine bindings, `backend.py`,
`direct_variants.py`, `direct_aot_jobs.py`, the parity harness, CMake
and two tests keep both PRs' entries; the representative golden was
re-blessed with `--write` (every case hash equals develop's or this
branch's, develop's expected failures kept, llvm23 `ok_cases` 329). On
the merge (CPU only): byte identity `--only conv,gfx1250` GREEN at
llvm20 / 22 / 23 (all five target-op families GREEN in the llvm23 lane);
golden `--check` OK at llvm20 / 22 / 23 / 24; native build and ctest 33
/ 33; library targeted pytest (stem, direct-conv cache, kernel label,
builder signature, #14392's forward tests, dgrad patch,
`test_check_byte_identity_progress.py`) 201 passed, 16 skipped, 1 failed
(the pre-existing `test_off_keeps_the_historical_hash`), with
`test_check_byte_identity_progress.py` 3 / 3; `test_dgrad_stem_wmma.py`
with this tree's engine 8 passed, 3 skipped (gfx1250 GPU tests, not run
this round).
- Merged `origin/develop` again (87929d6, develop a0ff542, which
brings in #14385's strided / dilated / packed-group WMMA direct dgrad).
The shared dgrad files (`direct_variants.py`, `direct_aot_jobs.py`,
`backend.py`, `rocke_engine.cpp`, the parity harness,
`SUPPORT_MATRIX.md`) merged without conflicts and keep both PRs'
entries; the five-family `TARGET_OP_PROFILE_FAMILIES` tuple is
unchanged. The only conflict was the representative golden's summary
counters, re-blessed with `--write` (every case hash equals develop's or
this branch's; llvm23 `ok_cases` 335, llvm20 / 22 expected failures 28 =
develop's 24 + the 4 stem cases). On the merge (CPU only): byte identity
`--only conv,gfx1250` GREEN at llvm20 / 22 / 23 (llvm23 lane: patch 9,
smallc 9, direct dgrad 32, direct fwd 22, fwd small 15 configs); golden
`--check` OK at llvm20 / 22 / 23 / 24; differential golden (ll at llvm20
/ 22 / 23 / 24, canonical ir) OK; native build and ctest 33 / 33;
library targeted pytest (stem, direct-conv cache, kernel label, builder
signature, layering, dgrad patch, direct dgrad / fwd WMMA, fwd small,
byte-identity driver) 225 passed, 17 skipped, 1 failed
(`test_off_keeps_the_historical_hash`, which also fails on pristine
develop a0ff542); platform `tests/portable_ir` +
`test_rocke_ci_static.py` 49 passed; `test_dgrad_stem_wmma.py` with this
tree's engine 8 passed, 3 skipped (gfx1250 GPU tests, not run this
round); black 25.12.0 clean.
- Merged `origin/develop` again (fc2b655, develop 7e3d78a, which
brings in the gfx942 fp64 GEMM cases from #14357). The only conflict was
the representative golden's summary counters, re-blessed with `--write`
(every case hash equals develop's or this branch's, develop's expected
failures kept; llvm23 `ok_cases` 337 = develop's 333 + the 4 stem cases,
llvm20 / 22 expected failures 28 = develop's 24 + the 4 stem cases).
`TARGET_OP_PROFILE_FAMILIES` still lists all five families. On the merge
(CPU only): byte identity `--only conv,gfx1250` GREEN at llvm20 / 22 /
23 (llvm23 lane: patch 9, smallc 9, direct dgrad 32, direct fwd 22, fwd
small 15 configs, 0 bad); golden `--check` OK at llvm20 / 22 / 23 / 24;
differential golden (ll at llvm20 / 22 / 23 / 24, canonical ir) OK;
native build and ctest 33 / 33; library targeted pytest (same set as
above) 225 passed, 17 skipped, 1 failed
(`test_off_keeps_the_historical_hash`, which also fails on pristine
develop 7e3d78a); platform `tests/portable_ir` +
`test_rocke_ci_static.py` + byte-identity driver test 52 passed;
`test_dgrad_stem_wmma.py` with this tree's engine 8 passed, 3 skipped
(gfx1250 GPU tests, not run this round); black 25.12.0 clean.

## Submission Checklist

- [x] Both engines updated, byte identity GREEN at every flavor in
`LLVM_FLAVORS`
- [x] Parity pairs and golden cases added for both new families
- [x] On-GPU numeric tests (gfx1250) and CPU cross-engine tests
- [x] CTest entry, `SUPPORT_MATRIX.md`, `instances/index.md`,
`instances/convolution.md`
- [x] black 25.12.0 / clang-format 18.1.4 on changed files
- [ ] Dispatcher wiring: not in this PR. The families are served by the
direct-conv AOT cache, like `conv_direct_dgrad_wmma`.
- [ ] gfx950 / gfx942 GPU runs: not applicable. These arches build no
stem dgrad kernel; they are pinned by the parity and golden gates.

---------

Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
aosewski added a commit that referenced this pull request Nov 2, 2026
The LLVM 24 compiler lane (#14415) added an llvm24 flavor, and the
LLVM 24 compiler that the CI test lane installs resolves to it. Decode
served only llvm22 and llvm23, so on such a host dispatch refused every
decode request, while the decode resource gate, the numeric lane and
the IR golden check skipped instead of failing: the refusal went
unnoticed.

Decode now serves llvm24. It lowers to the same IR there as at llvm23:
llvm24 changes only how named barriers are declared and called, and
decode uses none. Equal IR says nothing about the backend; the resource
gate, which no longer skips on such a host, checks the compiled kernels.
The golden still gets its own llvm24 column, so a later divergence
shows up.

The tests now catch a flavor that decode does not serve, whatever
compiler the host has:
- the IR golden check compares every served flavor on any host, since
  lowering takes the flavor as an argument; it used to check only the
  host's flavor and skip when that one had no golden;
- the dispatch admission test runs at every platform flavor but llvm20,
  read from the platform instead of from decode's own list;
- a new test fails when the platform has a flavor that decode neither
  serves nor refuses on purpose.

The READMEs and test strings name the served flavors, and the READMEs'
resource tables name the llvm24 IR that the CI test lane compiles.

Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants