Skip to content

feat(hip-kernel-provider): add gfx942 dense attention AOT ingestor engine - #14908

Merged
AnaghaRaoAMD merged 86 commits into
developfrom
users/anarao/gfx942-dense-aot
Nov 7, 2026
Merged

AnaghaRaoAMD merged 86 commits into
developfrom
users/anarao/gfx942-dense-aot

Conversation

@AnaghaRaoAMD

@AnaghaRaoAMD AnaghaRaoAMD commented Nov 2, 2026 •

Copy link
Copy Markdown
Contributor

ISSUE ID : AICK-2535

Problem

gfx942 has no ahead-of-time packaged dense SDPA forward engine. The rocKE dense prefill kernel (rocke/library/kernels/gfx942/attention_dense.py) is reachable only through rocKE's own launch path, not as a hipDNN engine.

Fix

  • New packaged ingestor engine hipkernel:Gfx942AttentionDense, generated through projects/hipdnn/tools/IngestorGenerator/configs/gfx942_attention_dense.engine.yaml. Descriptors live under descriptors/rocKE/gfx942_attention_dense/; the KDP is DVC-tracked (artifacts.dvc, pushed to the ingestor remote).
  • One shared dense pack for gfx942 and gfx950: packs/AttentionDenseCommon.hpp and packs/AttentionDenseNative.cpp hold graph and kernel matching, signatures and the dispatch handler, reading the graph through feat(hip-kernel-provider): add gfx950 unified attention 2D/3D ingestor engines #14242's SdpaGraphMatch.hpp; an AttentionDenseArch descriptor carries what differs. gfx942 plans save and restore under hipkernel.gfx942_attention_dense.dispatch.v1, frozen by golden samples of each launch form. gfx950 saved-plan contracts are unchanged. Both packs now also decline SDPA compute types other than FLOAT/HALF/BFLOAT16 (the kernels accumulate in FLOAT) and scales outside [2^-64, 2^4], and both restore paths apply the same saved mask and shape rules.
  • Scope: bf16/fp16, head size 64/128, BSHD Q/K/V/O; top-left causal, unmasked, and causal sliding windows of 128/512/1024/2048 keys (graph left_bound W keeps W keys, SEM-001); the SdpaFwd stats (LSE) output, bound like the FlyDSL CDNA pack does. Batch, both sequence lengths and both head counts are runtime kernel arguments. Everything else is declined; see Follow-ups. The ragged decline also lands in gfx950, which had the same gap.
  • Grids: the non-persistent qb_minor, bt_hkv_minor and hq_minor_swz orders, and the persistent D128 grid (304 CTAs, a pinned decode order, host-computed fast-division constants). Persistent is the cold default for D128 when 304 is a multiple of the device's CU count; the cold score then ranks tile, lds_row_pad, pv_priority and the body levers with their policy arm first.
  • rocKE: opt-in registry candidates gfx942_dense_qbminor and gfx942_dense_{persist,btminor,hqswz}_aot pin the grids the native dispatch launches (the tuned default resolves its grid from the build-time shape). supports_attention_dense refuses windows that leave query rows no key unless sinks or the LSE path handle them. The gfx942 and gfx950 dense runners check the same scale range as the packs.
  • Catalog: every lever that varies within a cell is a KMD field and an engine knob, so candidate enumeration (hipdnn_bench enumerate) works. variant_reduce keeps 224 of the manifest's 1,500 kernels from an autotune sweep on MI300X (two timed rounds over the 814 non-windowed served graphs and 153 stats graphs, eps 0.01, run-to-run drift 0.5%): 144 floor kernels (tile divisibility and declared defaults; the 32 window cells are floor-only) and 80 kept for measured wins. The sweep ran on a build of cde4c65; all 1,500 kernels lower to byte-identical LLVM IR with that tree's rocKE, with develop's at 3ab1173 and with this head's, so the ranked kernels are the shipped ones. kv_segments is a KMD field, and kernel_match serves 1 only. The manifest names the models behind each window size.
  • Tests and tooling: three golden-backed BSHD stats bundles with support claims; the knob suite runs every shipped kernel, one generated row each, against the CPU reference (those 247 rows in the comprehensive tier; its coverage, filter, swizzle-legality and save/load cases in every tier); the gfx942/gfx950 unit suites share TEST_P files for geometry, signature, dispatch and the operand rules; the IngestorGenerator gates check knob addressability over the swept space (g8); generated C++ goldens for the fast-division, grid, legality, score and knob-row tables, with a drift check. factorise_config.py now carries a per-arm policy resolution, which the reduced selection needed. The five engine-pinned attention engines (both gfx942 engines and the three gfx950 ones) register their external checks, GPU-reference targets and census from one table, HKP_PINNED_ATTENTION_ENGINES in src/CMakeLists.txt, which test_attention_bundle_filters.py checks against the categories YAMLs; the knob suites are gated from one table, and feat(hipdnn): gfx942 unified 3D split-KV decode as a hipDNN ingestor engine #14259's Unified3d knob suite runs on the shared harness. The census template wraps long lines as clang-format does, so all five attention census tests regenerate byte-for-byte. The two SdpaFwd Default support.json files carry this engine's claims alongside gfx942 Unified3d's.

Follow-ups

Not covered by this AOT engine. A graph-level gap is declined with an INFO line naming the cause; a window size the catalog lacks, or a sequence length no shipped tile divides, passes the graph match and then matches no kernel.

Item Status
Attention sinks (gpt-oss: D64, full causal and W=128) Not yet supported by this pack; the gfx942 kernel builds them. When this work started, hipDNN's SDPA references could not validate sinks; #14358 added that on 2026-11-02. gpt-oss's W=128 layers need sink variants of the shipped W=128 cells.
Sliding window W=4096 (Mistral-7B-v0.1, StarCoder2, Command R7B / Command A, OLMo 3, EXAONE 4.0; all D128, no sinks) and other sizes Not in the catalog. W is a compile-time kernel parameter; the shipped sizes are 128, 512, 1024 and 2048.
Bottom-right causal with seqlen_q != seqlen_kv (causal_bottom_right) Not yet supported by this pack; the kernel builds it, no variant ships it.
Split KV (kv_segments > 1) Not yet supported by this pack: it needs a workspace and a combine launch; kernel_match serves 1 only.
Windows that leave query rows no key (seqlen_q - seqlen_kv >= W) Declined with or without a stats output; the kernel zeroes key-less rows only on its LSE path.
Variable sequence lengths, ragged (packed) Q/K/V/O, paged K/V Not implemented in the gfx942 kernel.
Sequence lengths no shipped tile divides; bottom-right windows; non-causal bands Not implemented in the gfx942 kernel (no partial-tile path; top-left windows with right_bound 0 only).
Head sizes other than 64 and 128; FP32 and FP8 data types; dropout, additive / ALiBi / block masks, scale tensors; Max and Sum_exp outputs Not implemented.
Q/O of 2^31 elements or more, K/V of 2^31 bytes or more Addressing limit: the kernel uses 32-bit offsets.
gfx950 dense: stats (LSE) output, persistent and non-qb_minor grids Not shipped for gfx950: the shared pack declines a stats output there, and the gfx950 catalog has no such kernels.

Other follow-ups:

  • UHD catalog ranker for this engine, trained from this PR's sweep timings (98,097 candidate rows from sweep run 20261104T174818Z_68270754).
  • Register spills: 21 of the 224 kernels spill: 17 by 1-3 VGPRs (8-16 B private segment), 4 by 10-20 SGPRs (two with a 36 B private segment). 20 are in window cells, which the sweep never times; one is the fp16 D64 causal LSE 64/64 kernel. A spill check in the selection, and window-cell tile tuning (10 of the 20 are their cell's default 256/64 tile).
  • Port feat(hip-kernel-provider): add gfx950 unified attention 2D/3D ingestor engines #14242's gfx950 unified knob suite (IntegrationGpuGfx950AttentionUnifiedKnobs.cpp) onto ServedEngineHarness and the shared SDPA test helpers; it needs a gfx950 run.
  • rocKE: the dense spec factory (gfx942_dense._make_spec) refuses unmasked graphs with seqlen_q != seqlen_kv, which the pack serves. A Python/C++ parity gap already on develop; unpinned rocKE routing sends these graphs to the unified path, and this PR does not change that.

Verification

Rows ran on gfx942 (MI300X, 304 CUs) at 13dd161, after merging develop at 886472e, unless the row says otherwise; develop's gfx942 Unified3d (KDP and trained ranker) and gfx950 conv dgrad DVC artifacts were pulled so their bundles pack. The 224 shipped kernels lower to the same LLVM IR at this head as at the sweep.

Check Result
Kernel-ingestor, dense and unified attention, launch-value, FlyDSL, code-object-load, plugin-serialization and conv unit suites 743 passed, 2 skipped (develop's gfx950 dgrad catalog tests: no gfx950 tree in a gfx942 build)
Census, build tree (2 suites + 8 controls) and installed 10/10; 9/9 + 1/1
hipdnn_validate_descriptors; verify_variant_sets.py --mode full; shipped catalog regenerated byte-for-byte (g2) success; 224/224; identical
Device knob suite: every shipped kernel forced to its tile, grid and levers, plus shape, persistent, cold, scale, save/load and coverage cases, CPU reference (the per-kernel rows run in the comprehensive tier) 257/257, 0 skipped; 10 of them in the standard tier
Engine-pinned SdpaFwd bundles, both the GPU-reference registration and the external-integration-check target, claims enforced 70 ran: 48 confirmed (incl. the 3 stats bundles), 22 skipped (graphs the engine declines, none covered by a claim); no claim failures
gfx942 Unified3d (#14259): knob suite on the shared harness; engine-pinned bundles and external check 6/6, the same kernels as before the port; 67 ran: 18 confirmed, no claim failures, as before
Registrations from HKP_PINNED_ATTENTION_ENGINES: CTest listings (gfx942 and gfx942;gfx950 builds, install), build targets, compiled knob suites and install rules, against the tree before the refactor, at 2a1516b identical
Served corpus, engine-pinned and autotuned (814 non-windowed served graphs + 153 stats graphs), PyTorch reference; on a 9dee6cb install carrying the shipped 224-kernel KDP. Nothing committed since changes which graphs are served or by which kernel (same kernel IR; the bundles and knob suite above pass at this head) 967/967, worst max abs 3.9e-3 (bf16)
gfx942 and gfx950 census, and the dense, launch-value, ingestor-plan, unified and conv dgrad unit suites, on a gfx942;gfx950 build without a GPU; gfx950 knob suite compiles 45/45; 636 passed, 5 skipped (GPU-only)
LSE against an fp64 reference: rocKE probe of the gfx942_dense_qbminor default spec with emit_lse (bf16/fp16, D64/D128, causal/full, 3 sizes, MHA/GQA), 2026-11-03 48/48 within the 1e-5 stats tolerance, worst 1.8e-6
Persistent NP=304 parity, attention_dense_prefill.py --persistent --np 304 (bf16/fp16, causal/full, bt_hkv_minor, hq_minor_swz and qb_major decodes), 2026-11-03 20/20, worst max abs 2.2e-3 (bf16), 2.7e-4 (fp16)
rocKE dispatch/attention, test_attention_dense_gfx942.py and the library layering test (CPU); the rocke_library_attention_pytest CTest entry, which now collects the scale and key-less-window tests (at 2a1516b) 1239 passed, 253 skipped; 858 passed
IngestorGenerator suite with hipdnn_strict_skips; hip-kernel-provider hkp-pack-quick (descriptor-packaging -m quick) 1924 passed (develop's dgrad tests with their DVC instance lists pulled; #15102's three Unified3d ranker gates also pass with pandas and flatbuffers installed), 1 strict xfail (#14242's gfx950 unified 2D launch surface, with its reason); the 36 strict errors are prerequisites CI provides (validator path, flatbuffers headers, gfx950 unified KDPs, a repo-root build), and with them provided those tests pass (50/50, 14/14) or skip on the absent KDPs; 811 passed, 90 skipped
clang-tidy-20, provider .clang-tidy, FlyDSL off (as the Linux superbuild): the changed sources and every TU that depends on a changed header, each round through 2a1516b, including the knob suites; FlyDSL natives and test TUs that include a changed header syntax-checked with FlyDSL on 0 findings

AnaghaRaoAMD and others added 7 commits October 22, 2026 13:42
Add the registry-backed engine manifest for hipkernel:Gfx942AttentionDense and
the selection factorised from it. The gfx942 dense body reads batch, both
sequence lengths and both head counts at runtime, so cells are head size x
dtype x mask/window only; head counts sit in enumerate.fixed. Scope is the
non-persistent grid with causal, full and causal sliding windows of
128/512/1024/2048 keys. Sinks are out of scope: neither SDPA reference plan
accepts sink_token_tensor_uid.

Knobs are block_m, block_n, lds_row_pad and pv_priority: 288 kernels over 24
cells, within the 1000-instance AOT budget.

Co-Authored-By: Claude <noreply@anthropic.com>
…r engine

Ship the gfx942 dense attention forward as a packaged rocKE engine:
descriptors under descriptors/rocKE/gfx942_attention_dense (KDP tracked by
DVC), the native pack (graph/kernel match, score, 10-argument dispatch) and
its restated geometry, plus the registration, census, external-integration
and test-category wiring gfx950 has.

The gfx942 body reads batch, both sequence lengths and both head counts at
runtime, so candidates are matched on dtype, head size, mask and window only.
A causal band of left_bound L is served by the kernel built for
sliding_window L + 1; every non-causal band, sinks and stats are declined.
Tile divisibility and the sliding-window zero-trip guard are checked per
candidate. The cold score orders every block_m/block_n/lds_row_pad/
pv_priority combination, policy arms first.

Co-Authored-By: Claude <noreply@anthropic.com>
…l tile set

Add IntegrationGpuGfx942AttentionDenseKnobs: every (head size, dtype, mask,
window, tile) class the gfx942 catalog compiles runs once under a forced
tile, checked for the exact policy-arm kernel id and against the CPU
reference, plus cold selection, B > 1, cross-length, both alignments, the
deprecated causal flag, window spellings and the scale. A coverage case
keeps the table equal to the installed catalog.

Correct the tile-legality expectation: supports_attention_dense also admits
the non-power-of-two tiles (96/96, 192/96, ...) when built at seqlens they
divide, which is what isSupportedGfx942AttentionDenseTile accepts.

Co-Authored-By: Claude <noreply@anthropic.com>
…tentionDense

Record the 8 quick and 37 standard SdpaFwd/Default cases the gfx942 dense
attention engine served and verified against the GPU reference on gfx942
(written with --write-support-claims), and regenerate the support matrix
overview.

The knob device test now expects the lds_row_pad 0 arm for D128 256/128, the
only pad at which that tile fits LDS and so the kernel the score ranks first.

Co-Authored-By: Claude <noreply@anthropic.com>
variant_reduce could not read a sweep produced by the current hipDNN:

- WinnerCacheFile.hpp writes line format v=2 (#13028 added the DeviceKey
  memory fields), while variant_reduce accepted only v=1. Accept both, each
  with its exact device field set.
- hipDNN now mints Graph.id at finalization, a cache_ignore field that
  GraphContentKey leaves out, but variant_reduce keyed rankings on the raw
  serialized graph. The same graph got a new key every round, so no graph
  matched across rounds and the reduction refused the sweep. Zero Graph.id
  before keying.

Co-Authored-By: Claude <noreply@anthropic.com>
…00 kernels

Apply the stage-6 reduction: a 2-round autotune sweep of the 288-kernel
selection over the 814 served graphs of the dnn-benchmarking attention
corpus, all correctness phases passing, keeps 100 kernels at eps 0.01 (no
ranked graph loses more than 1% against its best packed kernel; measured
noise floor 0.6%). Causal and unmasked cells keep 52 (24 floor, 28 measured
wins); the 16 sliding-window cells, which no corpus graph exercises, keep
their 48 floor kernels.

Regenerate the bundle (KDP in DVC), its census suite, and the device knob
test's class table for the reduced catalog.

Co-Authored-By: Claude <noreply@anthropic.com>
… changes

Format the new gfx942 dense-attention sources and the variant_reduce change
with the versions pre-commit pins (clang-format 18.1.4, black 25.12.0). No
functional change.

Co-Authored-By: Claude <noreply@anthropic.com>
@therock-pr-bot

therock-pr-bot Bot commented Nov 2, 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

therock-pr-bot Bot commented Nov 2, 2026 •

Copy link
Copy Markdown

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

AnaghaRaoAMD and others added 3 commits November 2, 2026 16:59
Keeps the 100 gfx942 dense instances; the KDP specs pin persistent=false, so
develop's persistent-D128 default does not change the shipped ABI. Conflicts
resolved by keeping both sides; SUPPORT_MATRIX_OVERVIEW.md regenerated.

Co-Authored-By: Claude <noreply@anthropic.com>
…_bound and shard rules

develop moved left_bound to cuDNN semantics (SEM-001): left_bound = L keeps L keys
including the diagonal, 0 is rejected, and resolveDiagonalBand returns L - 1. The
engine's sliding_window = band.leftBound + 1 is therefore still the key count; the
overflow guard is unreachable and goes, and the matcher and knob tests now spell a
W-key window as left_bound W.

The census shard now has to reserve every TOML it names, so the gfx942 external
integration check passes hipkernel_Gfx942AttentionDense.toml to
hkp_reserve_census_shard(... FILES), as the gfx950 entry does.

Co-Authored-By: Claude <noreply@anthropic.com>
…b_minor grid

On develop the gfx942 dense spec's nonpersist_decode "auto" resolves from the
build-time shape, and for causal MHA it picks bt_hkv_minor. The ingestor's native
dispatch launches the (nqb, Hq, B) qb_minor grid at run time, so every packed causal
kernel decoded the wrong tile and faulted (illegal memory access in the knob suite).

Add gfx942_dense_qbminor, the twin of gfx950's qbminor variant: a non-persistent,
unsplit qb_minor default on the shared 256-row tile, and pin the manifest to it. The
manifest and selection now carry develop's new spec fields (nonpersist_decode,
kv_segments, window_wave_skip, v_reg_prefetch, chiplet_num_xcds), held at that
default. The 100 shipped kernels keep their names and ids; only those five spec
fields change, so the KDP moves to md5 83e609d4ef67960b6a7e71d4e62602c7.dir.

Co-Authored-By: Claude <noreply@anthropic.com>
@AnaghaRaoAMD
AnaghaRaoAMD requested a review from a team as a code owner November 2, 2026 23:52
@github-actions github-actions Bot added the rocke-lib: attention Touches rocKE library attention/FMHA kernels label Nov 2, 2026
… tests

The HipDNN superbuild lane compiles with clang-tidy-20 and warnings as errors:
modernize-avoid-c-arrays (Geometry witnesses, knob head table),
misc-const-correctness (Matchers tensors/nodes) and modernize-pass-by-value
(SelectedKernelReport). Test-only; unit 55/55 and the device knob suite 95/95
still pass.

Co-Authored-By: Claude <noreply@anthropic.com>
@cgmillette

Copy link
Copy Markdown
Contributor

Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max

PR #14908 — Gfx942AttentionDense AOT ingestor engine (AnaghaRaoAMD, AICK-2535) — 🔴 First-round: superbuild red, ragged gap

Prior item Status
Reviewers 🟢 None yet
• no reviews or threads; policy bot only
• needs hipdnn-reviewers and rocke-lib owner approvals
Bots / CI 🔴 Failing
• fails: hipDNN Linux + Windows superbuilds (+ summary) — N1
• fails: Linux math-libs gfx94X — build pod killed (infra); gfx94X tests skipped
• running: Windows gfx110X tests; 25 ✓ · 1 running · 4 ✗
• merge: conflicts with develop
Our checks 🟡 Conflict
• merge vs d1dd5a81: conflict in generated SUPPORT_MATRIX_OVERVIEW.md (ran)
• vs f44ee058: build + ctest 35 ✓, goldens 0 changed (ran)
• rocKE dispatch tests 135 ✓; rules 1–3 clean (ran)
• DVC fetch ✓ in CI; window L → W matches the kernel (read)
Code quality 🟠 Should-fix
• N3: pack, tests and wiring copied from gfx950 (read)
• N4: qbminor knobs leave the ingestor grid (ran)
• N5: no CI gate for the gfx942 manifest (read)
# Sev file:line Finding
N1 🔴 TestKernelIngestorEngine.cpp:627, IngestorPacks.cpp:47 New pack missing from packDispatchSymbols
• impact: EveryHandlerReportsSupportsSaving… fails both superbuilds, 5 vs 6 (read: CI logs)
• fix: add "hipkernel.gfx942_attention_dense.dispatch" to the list
N2 🔴 Gfx942AttentionDenseNative.cpp:184 Matcher accepts ragged (packed) Q/K/V/O once develop is merged
• impact: 14202 lets ragged graphs reach this provider; its hd128_causal_ragged bundle matches; kernel reads and writes past packed buffers (read)
• fix: decline ragged_offset_tensor_uid() as Gfx1151WmmaAttentionNative.cpp:280 does; gfx950 twin too
N3 🟠 Gfx942AttentionDenseNative.cpp:151, :310 Copies the gfx950 pack instead of sharing it
• impact: over 80% of the native code, 90%+ of three tests repeat gfx950's; N2's fix must land twice (read)
• fix: one shared AttentionDenseCommon.hpp, as FlydslSdpaCdnaCommon.hpp; checkedProduct from PackFeatures.hpp:99; TEST_P bodies
N4 🟠 gfx942_dense.py:300, dense_rules.py:130 gfx942_dense_qbminor knobs leave the grid it promises
• impact: 580 of 2432 production specs persistent, 64 reordered, 48 split KV (ran)
• impact: only the manifest's ignored list holds the pin
• fix: scope persistent, nonpersist_decode, kv_segments out for it, as gfx950's entry scopes nonpersist_decode (dense_rules.py:105)
N5 🟠 test_registry_regeneration.py:40, test_launch_surface.py:767 No CI gate covers the gfx942 manifest, selection or descriptors
• impact: both gates name gfx950 only; the body's generator gates are author-run (read)
• fix: parametrise both over the two manifests; g1, g4 and --check pass on gfx942 now (ran)
N6 🟡 gfx942_attention_dense/README.md:14, SdpaSemanticsRegister.md:47 README and SEM-001 disagree with the engine
• impact: left_bound 127–2047 match no kernel; README lists four knobs, the UED two; SEM-001 says rocKE kernels get L − 1
• fix: README to 128/512/1024/2048 and two knobs; this pack's L → W in SEM-001; candidate in dispatch/AGENTS.md

Utility: 🟡 Mixed — first packaged gfx942 dense SDPA engine; N1 and N2 block it.


Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted

First-round review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: merge with develop, build + ctest, differential goldens, rocKE dispatch tests, CI logs, DVC fetch, window semantics, generator gates, code quality. Flags whether a closer look is needed; not an approval or rejection.

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

Fully masked rows produce incorrect outputs, packed ragged tensors can cause out-of-bounds accesses, and the new pack registration breaks an existing contract test.

8 open issues
What changed in this PR

Adds ahead-of-time dense SDPA forward support for gfx942 through hipDNN’s kernel-ingestor path, reusing rocKE kernels with runtime shape arguments.

Changes:

  • Defines a 100-kernel catalog covering bf16/fp16, head sizes 64/128, causal and unmasked attention, and causal sliding windows.
  • Adds the opt-in gfx942_dense_qbminor candidate, native integration wiring, and host, census, device, and support-claim coverage.
  • Updates variant reduction to accept winner-cache format v2 and ignore minted graph IDs when matching timing rounds.
File Description
projects/​hipdnn/​tools/​IngestorGenerator/​tools/​variant_reduce.py Updated as part of this pull request.
projects/​hipdnn/​tools/​IngestorGenerator/​tests/​test_variant_reduce.py Updated as part of this pull request.
projects/​hipdnn/​tools/​IngestorGenerator/​configs/​gfx942_attention_dense.yaml Updated as part of this pull request.
projects/​hipdnn/​tools/​IngestorGenerator/​configs/​gfx942_attention_dense.engine.yaml Updated as part of this pull request.
dnn-providers/​integration-tests/​SUPPORT_MATRIX_OVERVIEW.md Updated as part of this pull request.
dnn-providers/​integration-tests/​integration-test-bundles/​standard/​SdpaFwd/​Default/​support.json Updated as part of this pull request.
dnn-providers/​integration-tests/​integration-test-bundles/​quick/​SdpaFwd/​Default/​support.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​packs/​TestGfx942AttentionDenseSignature.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​packs/​TestGfx942AttentionDensePacks.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​packs/​TestGfx942AttentionDenseMatchers.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​packs/​TestGfx942AttentionDenseGeometry.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​packs/​TestGfx942AttentionDenseDispatch.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​packs/​TestGfx942AttentionDenseCatalogTiles.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​integration_tests/​kernel_ingestor_engine/​IntegrationGpuGfx942AttentionDenseKnobs.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​integration_tests/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​packs/​Gfx942AttentionDenseGeometry.hpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​IngestorPacks.hpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​IngestorPacks.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​IngestorKernelCode.hpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​gfx942_attention_dense/​README.md Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​gfx942_attention_dense/​kernel_dtype_matches_graph.umd.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​gfx942_attention_dense/​gfx942_attention_dense.uhd.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​gfx942_attention_dense/​gfx942_attention_dense.ued.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​gfx942_attention_dense/​gfx942_attention_dense.udd.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​gfx942_attention_dense/​gfx942_attention_dense.kmd.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​gfx942_attention_dense/​artifacts.dvc Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​gfx942_attention_dense/​.gitignore Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​README.md Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​dispatch/​attention/​test_tuning_space.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​dispatch/​attention/​test_gfx942_dense_wiring.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​dispatch/​attention/​test_attention.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​dispatch/​attention/​gfx942_dense.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​hipkernel_Gfx942AttentionDense_test_categories_integration.yaml Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​descriptor-packaging/​cmake/​HkpPackaging.cmake Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​config/​hipkernel_Gfx942AttentionDense.toml Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​CMakeLists.txt Updated as part of this pull request.

Review effort: Max

SUPPORT_MATRIX_OVERVIEW.md regenerated with feature_matrix.py.

Co-Authored-By: Claude <noreply@anthropic.com>
AnaghaRaoAMD and others added 2 commits November 3, 2026 01:35
… the dense packs

Review follow-ups for hipkernel:Gfx942AttentionDense:
- Q/K/V/O with a ragged_offset_tensor_uid (packed sequences) are declined; the
  kernel indexes them as dense and would read and write past each sequence. The
  gfx950 twin had the same gap.
- A top-left window with seqlen_q - seqlen_kv >= W leaves the last query rows no
  key; the kernel has no empty-row path and averaged V there. Declined at graph
  match. Every per-block zero-trip failure on a tile-divisible shape is such a
  case, so the kernel-match test of it moves to graph match.
- The workgroup-limit check divides instead of multiplying, so a 2^62 block_m
  in descriptor metadata no longer overflows.
- packDispatchSymbols lists the gfx942 dispatch symbol (pack count was 5 vs 6).
- The knob suite's unmasked rows use hq, hkv, so GQA/MQA reach them.
- checkedProduct comes from PackFeatures.hpp; README window bounds, knob list and
  SEM-001/SEM-005 match the engine.

Co-Authored-By: Claude <noreply@anthropic.com>
The qbminor candidate promised the hipDNN ingestor's grid, but its knob space
still swept persistent, nonpersist_decode and kv_segments; only the manifest's
ignored list kept them out of the AOT catalog. They are now out of scope
whenever the base spec pins the block order, as gfx950 scopes
nonpersist_decode. The tuned gfx942_dense space is unchanged, and the manifest
still enumerates 288 kernels containing all 100 shipped specs. The dispatch
docs now list both gfx942 dense candidates and gfx950's qbminor.

Co-Authored-By: Claude <noreply@anthropic.com>
@GarbyGilbertAMD

Copy link
Copy Markdown
Contributor

Found by AI:

The hipkernelprovider gfx94X standard-tier Test step has a 30-minute (1800s) timeout_minutes (fetch_test_configurations.py). Reading recent presubmit job logs for this PR's shards, the step currently runs about 27m38s and 27m57s -- roughly 2 to 2.3 minutes of headroom left, down from about 3 minutes measured earlier in this PR's history. The new host-side pytest entries this PR adds (rocke_library_*_pytest for the KV-split, scale-range and bundle-filter tests) each add a small amount (about 1 to 17 seconds per entry, by comparing recent job logs); the device knob suite's own entry time is flat. None of this is a projected breach today, but the margin is thin for a shared CTest entry that other PRs will also add to.

Would it be worth splitting the new gfx942 device knob suite (or the new pytest entries) into their own CTest entry with its own timeout budget, rather than growing the shared standard-tier entry further?

AnaghaRaoAMD and others added 4 commits November 6, 2026 09:09
…est from the engine table

The census registration listed the five engine-pinned attention engines
by hand (a gfx950 loop plus a block per gfx942 engine) and repeated the
nine Packs cases four times; the four knob suites were gated by four
blocks that disagreed (one used an arch-aware predicate meant for
per-arch registrations, one had no SDPA-off warning); and
test_attention_bundle_filters.py kept its own copy of each engine's
filter and categories YAML, so a wrong table column passed it.

The census now loops over HKP_PINNED_ATTENTION_ENGINES with one shared
Packs case list, the knob suites come from one table gated on
hkp_product_offers_engine, and the filter test reads the table rows. The
two gfx942 availability wrappers lose their callers and go. CTest
listings, build targets and install rules are unchanged on the gfx942
and gfx942;gfx950 builds. The RUNBOOK and docs describe the table row.

Co-Authored-By: Claude <noreply@anthropic.com>
…d strides

The dense, unified and FlyDSL RDNA test graphs each had their own layout
enum and stride functions with the same formulas. SdpaLayoutTestSupport.hpp
holds them; every test graph keeps the same dims and strides (checked by
a host probe over every layout and test graph) and the same 396 tests.

Co-Authored-By: Claude <noreply@anthropic.com>
…ration suites

IngestorIntegrationSupport.hpp had its own SdpaLayout and sdpaStrides
beside SdpaLayoutTestSupport.hpp, and the gfx942 Unified3d test graphs a
decodeQoStrides equal to the BSHD strides at one query. The integration
header now uses SdpaLayoutTestSupport.hpp, which gains the head-row
padded BHSD layout the Unified3d knob suite needs; an unknown layout
throws. Strides and test graphs are unchanged (host probe), and the
Unified3d knob suite selects the same kernels, 6/6.

Co-Authored-By: Claude <noreply@anthropic.com>
…entry

A typo in an HKP_PINNED_ATTENTION_ENGINES row's arch or engine name, or
in a CatalogTiles case list's key, silently dropped that engine's census,
external check and GPU-reference target. test_attention_bundle_filters.py
now checks each row's arch, TOML and Packs census (which must pin that
arch's inventory), and configuration fails on a CatalogTiles key that
matches no row. hkp_gfx950_attention_dense_available goes like the gfx942
wrappers; its tests call hkp_product_ships_engine_for_arch. The docs
scope the table to the single-arch engines pinned against the SdpaFwd
bundles.

Co-Authored-By: Claude <noreply@anthropic.com>
@CongMa13

CongMa13 commented Nov 6, 2026

Copy link
Copy Markdown
Contributor

I reviewed this PR again. The following review threads are still unresolved, or comments have no author reply:

For reassessment, the PR author should reply to top-level feedback or state that thread feedback is addressed. Before approval, top-level comments need explicit resolution from their author or that author's approval of the PR.

@AnaghaRaoAMD

Copy link
Copy Markdown
Contributor Author

On the 11-06 follow-ups (#14908 (comment), #issuecomment-6679807408, #issuecomment-6681173299):

  • N17: bfdd6ea, 6c69c25. The [2^-64, 2^4] range has one home per layer: isSupportedScale and UNSUPPORTED_SCALE_REASON in AttentionDenseCommon.hpp, used by both arches' graph match and restore (the per-arch scale hook is gone), and check_dense_scale in kernels/common/attention_dense_spec.py, called by both rocKE runners before compiling. fec2a79 updates the comments that still named gfx950's runner as the only check or listed a scale hook.
  • N18: 6b1d376. The KV split runs in the gfx942 tuned default (now tuned_base_spec), and the split_kv flag is gone. f6343a9 gives the split's exclusions one home, kv_split_admits in attention_dense_kv_split.py, called by check_kv_segments and both arches' tuning scopes; the tuned default keeps a split only where supports() admits it. Same specs as before for 184,320 requests × 5 factories, and the same tuning space for all 9 dense candidates.
  • N19: 02f1ed6. test_attention_dense_scale.py and the two gfx942 key-less-window tests move under tests/dispatch/attention/, so rocke_library_attention_pytest runs them (27 scale tests and both window tests). No other test this PR adds sits outside a CI entry.
  • N20: a70e6d0. The census template wraps the knob and score-symbol lines when they do not fit, as it already did for the source-kind line. All five attention census tests regenerate byte-for-byte, the three known failures are gone, and a boundary test pins the wrap point.
  • Filters in two CMake files: 90b1b02 and 7c32457, then ccdda14, 9076359 and 2a1516b: HKP_PINNED_ATTENTION_ENGINES in src/CMakeLists.txt lists the five engine-pinned attention engines with their filter, TOML and categories YAML, and one loop each registers the external checks, the GPU-reference targets and the census. test_attention_bundle_filters.py reads the table, so a wrong filter, YAML, arch or engine name in a row fails it. CTest listings, targets and install rules are unchanged on the gfx942 and gfx942;gfx950 builds.
  • Knob-suite port: 16d6004. feat(hipdnn): gfx942 unified 3D split-KV decode as a hipDNN ingestor engine #14259's Unified3d knob suite derives from ServedEngineHarness and uses the shared graph builder, tolerance and served-kernel walk; the builder gains the decode KV layouts. Same 6 cases and assertions, 6/6 on gfx942. 9076359 gates the four knob suites from one table with hkp_product_offers_engine. feat(hip-kernel-provider): add gfx950 unified attention 2D/3D ingestor engines #14242's gfx950 unified suite keeps its own copies; porting it needs a gfx950 run, so it is a follow-up in the body.
  • Private _candidate_space call: adeba44, with e5db957.
  • feat(hipdnn): gfx942 unified 3D split-KV decode as a hipDNN ingestor engine #14259: merged in 639e197; SdpaShape is the one struct and UnifiedAttentionProblem an alias.
  • feat(hip-kernel-provider): gfx950 unified SDPA rankers and LSE-on kernels #15113: whichever lands second reconciles the filters: its quick_SdpaFwd_BshdStats.* patterns need the gfx950 rows of HKP_PINNED_ATTENTION_ENGINES to name a filter that includes them, which test_attention_bundle_filters.py will ask for.

Also from a self-review with the same lenses: 050ccfd (saved-plan mask and shape refusals on both arches), 0f773b7 and f6343a9 (public names for what the gfx942 manifest names), 4d9064f (sinks/LSE no-split rule on both arches), 1eaf329 and 01de338 (saved-launch rules documented once; integration-tests docs), 23ee321 and b178943 (one home for the SDPA test layouts and strides, unit and integration).

For the unified engines' owners (#14242, #14259), not changed here: gfx942 Unified3d has its own SDPA matcher and accepts any finite scale; the unified packs do not apply computesInFloat; the unified manifests name private rocKE functions (_get_3d_pipeline, _attention_3d_workspace_specs, _num_segments, _3d_signature, _attn_signature, _num_q_blocks).

N15 (spills) stays a follow-up in the body. The Verification table is refreshed at 2a1516b. The two rocke/platform docs still need a rocke-platform owner; your approval covers that.

@AnaghaRaoAMD
AnaghaRaoAMD enabled auto-merge (squash) November 6, 2026 16:44
@cgmillette

Copy link
Copy Markdown
Contributor

Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: high

PR #14908 — Gfx942AttentionDense AOT ingestor engine (AnaghaRaoAMD, AICK-2535) — 🔴 Follow-up: conflicts with develop since 14822

Prior item Status
amd-khushbu (11-05): approval 🟢 APPROVED 0b2202f358
• older head; still counts on develop
• does not cover the 14259 merge, engine table, 14822 merge
• left: rocke-platform owner for the two rocke/platform dispatch docs
amd-khushbu (11-03) 18 items; Copilot (11-03) ×10, (11-05) C1–C4 🟢 FIXED (holds)
• delta reopens none (read)
Our 11-03 – 11-06 N1–N14, N16–N18 🟢 FIXED (holds)
• scale range and KV-split rule: one home each (read)
• C++ scale home is AttentionDenseNative.cpp:129-135, not AttentionDenseCommon.hpp as replied (read)
Our 11-05 N15 🟡: shipped kernels spill 🟡 DEFERRED (holds)
• body keeps the spill check as a follow-up (read)
• differential golden 0 changed; parity ×4 (ran)
Our 11-06 N19 🟠: scale test in no CI lane 🟢 FIXED 02f1ed6828
• both saved probes FIXED at head (ran)
• 27 scale and both key-less-window tests ✓ in gfx94X shards (read)
• no other added test outside a CI entry (read)
Our 11-06 N20 🟡: census known failures 🟢 FIXED a70e6d0d56 at head
• all five census tests regenerate; list gone (ran)
• breaks again on the develop merge: N22
Our 11-04 overlap: knob-suite builder (with 14259) 🟢 FIXED 16d60043f0
• Unified3d suite on ServedEngineHarness; ✓ on gfx94X shards (read)
• gfx950 unified port: body follow-up, needs a gfx950 run (read)
This push (11-06): 9 commits, 02f1ed6828 – 2a1516bdc0 🟢 FIXED what they say
• table typos fail the filter test: 10 of 10 mutants (ran)
• SDPA test layouts and strides: one header (read)
cgmillette (11-03, 11-04): overlap notices 🟢 FIXED (holds)
• launchWithParams: 14822 landed; the merge keeps one (ran)
• its conflict is the doc comment only
cgmillette (11-03): queued for the next merge wave 🟡 OPEN
• left: merge develop (N21, N22); rocke-platform approval
• 15113: whichever lands second reconciles; the filter test catches either order (ran)
GarbyGilbertAMD (11-06): gfx94X standard step near its limit 🟡 OPEN
• at head both shards end nearer the limit than measured (read)
• one CTest run covers every standard entry: a new entry still counts (read)
• no author reply yet; its gfx942 knob suites add the most (read)
CongMa13 (11-06): author replies to five top-level comments 🟡 OPEN
• author answered our posts at 15:30Z (read)
• GarbyGilbertAMD's comment has no reply yet
Bots / CI 🟢 Green at head
• 54 ✓, 0 ✗, none running (read)
• merge: conflicts; no new run until develop is merged
Our checks 🔴 Conflicts
• merge vs 97abc0b1b3: 9 conflicting files (ran)
• vs e3c0247c74: clean; build + ctest 56 ✓ / 1 skip; parity ×4; goldens; differential 0 changed (ran)
• rocKE tests 814 ✓; installed 297 ✓ / 1 skip (ran)
• trial resolution: generator suite, filter test, table mutants (ran)
• rules 1–3 clean in scope (ran)
Code quality 🟡 Minor
• engine table, test layouts, scale range: one home each (ran)
• knob suites keep a second, unchecked engine list: N23 (ran)
• still open: ragged and varlen split clauses untested (ran); docs restate values and the table (read)
# Sev file:line Finding
N21 🔴 src/integration_tests/CMakeLists.txt, IngestorIntegrationSupport.hpp, descriptors/README.md, config_loader.py + 5 Conflicts with develop in 9 files since 14822 landed
• impact: cannot merge; no new CI run
• keep both: knob-suite loop plus 14822's dgrad AOT block; both helpers in IngestorIntegrationSupport.hpp; develop's README table plus a gfx942 dense row
• take develop's: launchWithParams comment, _check_knobs_addressable / _check_knobs_int_typed split and tests, census inventory arrays
• fix: merge develop, then N22
N22 🟠 projects/hipdnn/tools/IngestorGenerator/templates/test_packs.cpp.j2:85-104 Develop's census template fails four census tests on the merge
• impact: gfx942 dense and unified 3d, gfx950 unified 2D/3D census files and the unslotted golden fail (ran)
• fix: regenerate those four Test*Packs.cpp; re-record unslotted_render_golden.json; then no failure develop lacks (ran)
N23 🟡 dnn-providers/hip-kernel-provider/src/integration_tests/CMakeLists.txt:86-113 Knob suites name the pinned engines again, in _knob_suites
• impact: a misspelt engine compiles no suite, silently; the filter test passes it (ran)
• fix: check _knob_suites against HKP_PINNED_ATTENTION_ENGINES in test_attention_bundle_filters.py

Utility: 🟢 Positive — N19, N20 and the knob-suite port done; develop merge left


Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted

Follow-up review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: delta since 01de338724 (9 commits), merge with develop (build + ctest, parity, goldens, differential golden, rocKE tests, installed), a trial merge with today's develop and its resolution, the generator suite, the saved probes, engine-table mutants, CI shard logs, code quality. Flags whether a closer look is needed; not an approval or rejection.

@cgmillette
cgmillette disabled auto-merge November 6, 2026 19:30
@cgmillette

Copy link
Copy Markdown
Contributor

Automated merge-queue status (rocKE merge monitor)

Hi @AnaghaRaoAMD, I turned auto-merge off on this PR for now. Before a rocKE PR lands, the rocKE merge monitor checks it on top of today's develop branch, and this version would not land cleanly:

  • our latest review of this exact version found blocking problems, and they are still open

The open items from our review, each with its fix:

  • N21 (src/integration_tests/CMakeLists.txt, IngestorIntegrationSupport.hpp, descriptors/README.md, config_loader.py + 5): Conflicts with develop in 9 files since 14822 landed
    • impact: cannot merge; no new CI run
    • keep both: knob-suite loop plus 14822's dgrad AOT block; both helpers in IngestorIntegrationSupport.hpp; develop's README table plus a
  • N22 (projects/hipdnn/tools/IngestorGenerator/templates/test_packs.cpp.j2:85-104): Develop's census template fails four census tests on the merge
    • impact: gfx942 dense and unified 3d, gfx950 unified 2D/3D census files and the unslotted golden fail (ran)
    • fix: regenerate those four Test*Packs.cpp; re-record `

What to do: address the open items in our latest review (#14908 (comment)). After that, turn auto-merge back on; the monitor re-checks every push.

Nothing else changed: the monitor does not approve, reject or merge PRs. It only turned auto-merge off to keep develop green. This comment is updated in place.

AnaghaRaoAMD and others added 2 commits November 6, 2026 14:50
…2-dense-aot

Conflicts with #14822 (gfx950 conv dgrad AOT engine), resolved:
- IRunnableKernel.hpp: both added launchWithParams with one signature; keep
  develop's.
- test_packs.cpp.j2: develop's constexpr-array inventory replaces this
  branch's long-name fence; this branch's line-wrap rules stay. The four
  attention census tests develop did not regenerate are regenerated (same
  pack and kernel names).
- config_loader.py, test_config_loader.py, test_slots.py: develop's split
  (unslotted knobs int or string, slot knobs int) replaces this branch's
  single check, and this branch's slot string-knob test goes.
- integration_tests/CMakeLists.txt: this branch's knob-suite table plus
  develop's dgrad test block.
- kernel_ingestor_engine/CMakeLists.txt, descriptors/README.md,
  IngestorIntegrationSupport.hpp: both sides kept; the README table gains
  Gfx942AttentionDense and spells Gfx942AttentionUnified3d as the engine
  and pack are named.
- unslotted_render_golden.json: recomputed from the merged generator.

Co-Authored-By: Claude <noreply@anthropic.com>
…ention engines

#14822's Gfx950ConvDgradAot census block repeated the nine generated Packs
cases that src/tests/CMakeLists.txt keeps in one list, and gated on its
own per-engine wrapper over hkp_product_ships_engine_for_arch. It now uses
the shared list and the predicate directly, as the attention engines do,
and the wrapper goes.

Co-Authored-By: Claude <noreply@anthropic.com>

@GarbyGilbertAMD GarbyGilbertAMD 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.

LGTM

AnaghaRaoAMD and others added 3 commits November 6, 2026 16:54
…rehensive tier

The gfx94X standard-tier test step runs every standard entry in one CTest
call with a 30-minute limit, and was at about 28 minutes; the gfx942 dense
knob suite's 247 per-kernel rows take about 6 minutes of it. They are now
instantiated as Comprehensive/, which HIP_MLOPS_ENGINE_test_categories.yaml
keeps out of the quick, standard and ffm-quick tiers (the comprehensive
and full tiers run them). The suite's coverage, filter, swizzle-legality
and save/load cases still run in every tier. No other test in either
provider binary carries the prefix.

Co-Authored-By: Claude <noreply@anthropic.com>
… engine table

_knob_suites in src/integration_tests/CMakeLists.txt names the pinned
engines again, and a misspelt one compiled no suite with no error.
test_attention_bundle_filters.py now checks each row's source exists and
every engine it pins is an HKP_PINNED_ATTENTION_ENGINES row.

Co-Authored-By: Claude <noreply@anthropic.com>
The shared rule's test covered the persistent, window, sinks and LSE
clauses of kv_split_admits, not ragged or varlen. gfx942's supports()
refuses those layouts on other grounds, so the new test checks
check_kv_segments and the knob's scope directly, on both arches. The
gfx950 algorithm notes name the scale constants instead of their values.

Co-Authored-By: Claude <noreply@anthropic.com>
@AnaghaRaoAMD
AnaghaRaoAMD enabled auto-merge (squash) November 6, 2026 21:54
@AnaghaRaoAMD

Copy link
Copy Markdown
Contributor Author

On the 11-06 follow-ups (#14908 (comment), #issuecomment-6693458311) and the merge-queue note (#issuecomment-6693575535):

  • N21: f21b18f merges develop at 97abc0b, resolved as the review suggests: the knob-suite loop plus feat(hipdnn): rocke conv gfx950 dgrad bf16/fp16 engine and ranker #14822's dgrad block; both helpers in IngestorIntegrationSupport.hpp; develop's README table plus a gfx942 dense row; develop's launchWithParams comment, its _check_knobs_addressable / _check_knobs_int_typed split and tests (this branch's slot string-knob test goes), and its census inventory arrays.
  • N22: the same merge regenerates the four census tests (same pack and kernel names) and re-records unslotted_render_golden.json. a411a13 then points feat(hipdnn): rocke conv gfx950 dgrad bf16/fp16 engine and ranker #14822's dgrad census at the shared Packs case list and hkp_product_ships_engine_for_arch, and drops its wrapper.
  • N23: 5ff1ff3. test_attention_bundle_filters.py checks each _knob_suites row: its source exists and every engine it pins is an HKP_PINNED_ATTENTION_ENGINES row; a misspelt engine fails it.
  • Ragged and varlen split clauses: e8a73bc. test_ragged_and_varlen_splits_are_barred checks check_kv_segments and the knob's scope on both arches (gfx942's supports() refuses those layouts on other grounds); dropping either clause from kv_split_admits fails it.
  • Docs restating values: the gfx950 ALGORITHM.md names DENSE_MIN_SCALE / DENSE_MAX_SCALE instead of the numbers. The table docs (descriptor-packaging README, RUNBOOK, adding-tests.md) say how to add a row and copy none.
  • Correction: the C++ scale range's one home is isSupportedScale and UNSUPPORTED_SCALE_REASON in AttentionDenseNative.cpp:129-135, not AttentionDenseCommon.hpp as I wrote.
  • GarbyGilbertAMD's CI time: 00846da moves the gfx942 dense knob suite's 247 per-kernel rows to the comprehensive tier; details in the reply there.

At e8a73bc the knob suite runs 257/257 on gfx942 (10 of them in the standard tier), and the rest of the Verification table ran after the merge. Auto-merge is back on.

@AnaghaRaoAMD

Copy link
Copy Markdown
Contributor Author

On the gfx94X standard-tier time (#14908 (comment)): a separate CTest entry would not buy headroom, because the step runs every standard entry in one CTest call, so the step time is their sum either way. Instead, 00846da moves the gfx942 dense knob suite's 247 per-kernel rows (about 350-400 s of that step on MI300X) to a Comprehensive/ instantiation, which HIP_MLOPS_ENGINE_test_categories.yaml now keeps out of the quick, standard and ffm-quick tiers; the comprehensive and full tiers still run them. The suite's coverage, filter, swizzle-legality and save/load cases (about 100 s) stay in every tier. No other test in either provider binary carries the prefix, so nothing else moves. The new host pytest entries take seconds each and stay where they are.

@AnaghaRaoAMD

Copy link
Copy Markdown
Contributor Author

Replies to the top-level comments listed in #14908 (comment):

  • #issuecomment-6674329204, #issuecomment-6679807408, #issuecomment-6681173299 (cgmillette): answered in #issuecomment-6688552696.
  • #issuecomment-6685343773 (cgmillette), and his later #issuecomment-6693458311 and #issuecomment-6693575535: answered in #issuecomment-6696817788.
  • #issuecomment-6686557039 (GarbyGilbertAMD): answered in #issuecomment-6696824524.

…2-dense-aot

SUPPORT_MATRIX_OVERVIEW.md conflicted; take develop's. Since #15293 a
claim PR no longer regenerates the overview: it is refreshed in a PR of
its own, and no check fails when it lags.

Co-Authored-By: Claude <noreply@anthropic.com>
@cgmillette

Copy link
Copy Markdown
Contributor

Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: high

PR #14908 — Gfx942AttentionDense AOT ingestor engine (AnaghaRaoAMD, AICK-2535) — 🟡 Follow-up: develop merged; CI running

Prior item Status
amd-khushbu (11-05): approval 🟢 APPROVED 0b2202f358
• older head; still counts on develop
• does not cover the 14259 and 14822 merges, engine table, tier move
• left: rocke-platform owner for the two rocke/platform dispatch docs
GarbyGilbertAMD (11-06): approval 🟢 APPROVED a411a134cc
• older head; still counts on develop
• does not cover the tier move, two new tests, last develop merge
amd-khushbu (11-03) 18 items; Copilot (11-03) ×10, (11-05) C1–C4 🟢 FIXED (holds)
• delta reopens none (read)
Our 11-03 – 11-06 N1–N14, N16–N20; knob-suite overlap (with 14259) 🟢 FIXED (holds)
• delta touches none of them (read)
Our 11-05 N15 🟡: shipped kernels spill 🟡 DEFERRED (holds)
• body keeps "a spill check in the selection" as a follow-up (read)
• differential golden 0 changed; parity ×4 (ran)
Our 11-06 N21 🔴: conflicts with develop 🟢 FIXED f21b18f3ff
• merge with d38d6b661e: 0 conflicting files (ran)
• resolved as the row listed (read)
Our 11-06 N22 🟠: census tests fail on the merge 🟢 FIXED f21b18f3ff
• saved probe: five census files and golden pass (ran)
• a411a134cc: dgrad census on the shared case list (read)
Our 11-06 N23 🟡: knob suites name the engines again 🟢 FIXED 5ff1ff3607
• misspelt engine and missing source fail the test (ran)
• runs in hkp-pack-quick; Linux superbuild ✓ at head (read)
Our 11-06 code quality: split clauses; docs restate values 🟢 FIXED e8a73bce2a, partly
• dropping ragged or varlen fails the new test (ran)
• Python doc names the constants; C++ docs still spell values (read)
cgmillette (11-03, 11-04): overlap notices 🟢 FIXED (holds)
• launchWithParams: one copy on the merge (read)
cgmillette (11-03): queued for the next merge wave 🟡 OPEN
• left: rocke-platform code-owner review, still requested (read)
• CI still running; auto-merge armed (read)
GarbyGilbertAMD (11-06): gfx94X standard step near its limit 🟢 FIXED 00846da06a
• 247 per-kernel rows move to Comprehensive/ (read)
• those rows now run nightly only, not on PRs (read)
• step time at head: Linux shards not run yet
CongMa13 (11-06): author replies to five top-level comments 🟢 FIXED
• author mapped all five to replies, 21:55Z (read)
Bots / CI 🟡 Running
• 30 ✓, 0 ✗ (read)
• queued: Linux math-libs stages, Windows test jobs; no rocKE shard yet
• both hipDNN superbuilds ✓
Our checks 🟢 Clean
• merge vs d38d6b661e: 0 conflicts, 3 behind (ran)
• build + ctest 57 ✓ / 1 skip; parity ×4; goldens; differential 0 changed (ran)
• rocKE tests 818 ✓; installed 301 ✓ / 1 skip; census probe FIXED (ran)
• knob-table and split-clause mutants; rules 1–3 clean in scope (ran)
Code quality 🟡 Minor
• N23 and split-clause tests fixed: mutants fail them (ran)
• 15066 repeats the tier YAML change and SDPA tensor helper: N24 (read)
• C++ docs still spell scale and tile values (read)
# Sev file:line Finding
N24 🟡 HIP_MLOPS_ENGINE_test_categories.yaml:24-26; IngestorIntegrationSupport.hpp:539 15066 repeats this PR's tier YAML change and SDPA tensor helper
• impact: whichever lands second conflicts in the tier comment (ran); its bshdTensor / bshdStrides repeat makeSdpaTensor / sdpaStrides for BSHD
• here: "matching miopen-provider's convention" is wrong (no Comprehensive* prefix there); the suite list goes out of date
• fix: owner this PR (lands first): keep both; take 15066's wording (integration-tests precedent, no suite list). 15066 then drops its copies

Utility: 🟢 Positive — N21–N23 fixed, develop merged, standard tier trimmed; platform review left


Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted

Follow-up review: first pass by qwen3.6 (local), verified by the orchestrator (Claude Opus 5.5). Checks: delta since 2a1516bdc0 (4 commits, 2 develop merges), merge with develop (build + ctest, parity, goldens, differential golden, rocKE tests, installed), the saved census probe, knob-table and split-clause mutants, the hipDNN superbuild log, the overlap with 15066, code quality. Flags whether a closer look is needed; not an approval or rejection.

@cgmillette cgmillette 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.

LGTM

…2-dense-aot

src/tests/CMakeLists.txt conflicted with #15102, which registered the
gfx942 Unified3d census and its new ModelBinding suite in a hand-written
block. The HKP_PINNED_ATTENTION_ENGINES loop already registers the Packs
census; its companion suites now come in two kinds, CatalogTiles and
ModelBinding, each listed per row under _hkp_<stem>_cases_<arch>_<slug>,
and the unmatched-entry guard covers both. The ModelBinding source keeps
its HKP_REQUIRE_DVC_DESCRIPTORS definition. adding-tests.md and the
RUNBOOK describe both kinds.

Co-Authored-By: Claude <noreply@anthropic.com>
@AnaghaRaoAMD
AnaghaRaoAMD merged commit a6b0022 into develop Nov 7, 2026
106 of 107 checks passed
@AnaghaRaoAMD
AnaghaRaoAMD deleted the users/anarao/gfx942-dense-aot branch November 7, 2026 03:00
BradPepersAMD added a commit that referenced this pull request Nov 7, 2026
…dsl-aotriton-parity

Conflicts:
- HIP_MLOPS_ENGINE_test_categories.yaml: took develop's (#14908 owns the
  Comprehensive* tier); quick/standard/ffm-quick still exclude Comprehensive*.
- src/integration_tests/CMakeLists.txt: kept develop's attention knob-suite
  table (it absorbs the Unified2D/3D and gfx942 unified 3D blocks) plus the
  FlyDSL frontend-lane sources and HIPKERNELPROVIDER_FLYDSL_SDPA_ARCHS.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD added a commit that referenced this pull request Nov 7, 2026
FlydslSdpaFrontendGraphs.hpp drops its bshdStrides/bshdTensor copies and
builds Q, K, V with detail::makeSdpaTensor(..., SdpaLayout::BSHD) and O's
strides with sdpaStrides(SdpaLayout::BSHD, ...) from
IngestorIntegrationSupport.hpp, which #14908 owns.

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

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants