feat(rocke): add MLA decode kernel for gfx942 and gfx950 by aosewski · Pull Request #13923 · ROCm/rocm-libraries

aosewski · 2026-10-19T10:34:37Z

ISSUE ID : AICK-1504

Problem

rocKE has no MLA decode path. Decode is one query per sequence over a long paged latent cache, so the query tile that prefill puts on the MFMA axis disappears; the per-head W_UK expansion that prefill absorbs into its query also dominates if done per key tile. Design: builders/mla/DESIGN.md (the README below records where the built kernel departs from it).

Fix

Kernel (gfx942, gfx950). Weight-absorbed bf16 decode: takes the absorbed query qa = [Q_nope @ W_UK_nopeᵀ | Q_rope] and returns the normalized latent (before W_UV) plus the LSE; W_UK_nope and W_UV run as batched GEMMs outside the kernel. Heads fill the MFMA n axis: 16 heads per wave, 64 per workgroup, sharing every staged key tile. One emitter serves both arches; the gfx950 spec subclass uses the 16x16x32 score atom and ds_read_tr16_b64 for the PV operand. Long key ranges split into segments merged by an arch-neutral combine kernel; an empty segment writes only (m, l) and its staging descriptors cover zero bytes, so it reads nothing.
Dispatch: one request, one op, regime by max_seqlen_q. Decode registers for max_seqlen_q == 1 on gfx942 and gfx950; prefill keeps max_seqlen_q >= 2. Decode serves only the absorbed-query contract and requires absorbed_query=True (prefill serves both contracts), so a one-query request without the flag is refused rather than handed a kernel that reads qa-shaped inputs. Decode admits 64 and 128 heads (the counts with GPU evidence), the IR flavors whose backend honors agpr_alloc (llvm22 and later; llvm20 ignores it and spills the kernel), one rule shared with prefill's agpr_acc gate, batches within the kernels' i32 offsets, and grids within the platform's per-axis limit (MAX_GRID_DIM; the host runner checks the device's own at launch); the split count fills the CUs, at least 32 key tiles per segment.
API change to feat(rocke): add gfx942 bf16 MLA prefill kernel and dispatch #12464's dispatch: MLARequest.max_seqlen_q is required; a decode request (max_seqlen_q == 1) must also set absorbed_query=True, so MLARequest(..., max_seqlen_q=1) without it is refused (all in-tree decode callers pass it); the family op is mla_fwd (was mla_prefill_fwd); prefill, including its opt-in fixed-layout candidates, refuses max_seqlen_q == 1. Prefill callers, tests, the shape loader and the score probe are updated. Happy to move this part into feat(rocke): add gfx942 bf16 MLA prefill kernel and dispatch #12464 if that is preferred.
Platform. rocke.runtime.hip_module.get_device_max_grid_size() (cached hipDeviceGetAttribute, like get_device_num_cus) and the shared grid limit rocke.helpers.spec.MAX_GRID_DIM (split_k.MAX_GRID_DIM_Z aliases it): decode selects against that constant, so selection does not depend on the device, and the host runner checks the device's own limit at launch.
Shared reference fix. gather_paged_kv no longer crashes on a sequence with S_k = 0.
Verification tooling. fp64 decode reference (single pass and split-KV merge); MlaDecodeRunner (host launcher that refuses mismatched, misaligned, aliased or off-device tensors; bind_latent for timed loops); verify_mla_decode (latent, whole op, and the prefill absorbed_query kernel as an independent GPU oracle), with NaN-poisoned unused cache rows, workspace and outputs; benchmark_mla_decode_live (resolve / build / verify / time; time mode gates correctness first and reports median and spread). Gate: elementwise 4e-2 + 2e-2·|ref|, per-sequence max_abs/peak <= 2e-2, LSE <= 1.7e-2.
Tests and CI. CPU tests (lowering and admission, routing, grid limits, empty segments, golden IR at llvm22, llvm23 and llvm24 (all three compared on every host), runner and harness checks, shape files) registered in the MLA CTest lane; the numeric test in the GPU lane.
Docs. library/examples/{gfx942,gfx950}/mla_decode/README.md: contract, preconditions, usage, gate, limits, and the correctness record.

No performance claims: this is functional bring-up; tuning is a separate piece of work.

Test plan

At the current head, with develop merged:

MLA CTest lane in CI's container, with the LLVM 24 compiler of CI's latest run and ROCKE_REQUIRE_COMPILER_TOOLS=1 -- 489 passed, 3 skipped (two need torch, one needs the C++ engine). The resource gate runs on both arches and reports zero scratch; the golden check and dispatch run at llvm24.
MLA, dispatch, runner, layering and platform-runtime tests with the C++ engine built -- 563 passed, 48 skipped (GPU numerics; the host llvm-objdump has no AMDGPU target).
Decode IR goldens at llvm22, llvm23 and llvm24, and C++/Python lowering parity (test_mla_decode_cpp_python_byte_identity) at each.
Resources, 128 and 64 heads, single pass and split: zero scratch on gfx942 and gfx950 under both LLVM 24 builds CI has used and under LLVM 22, within 256 VGPRs (two workgroups per CU).
GPU gfx942 and gfx950, after the split-kernel spill fix: numeric lane 20/20 (gfx942) and 19 passed, 1 skipped (gfx950; the oracle is gfx942-only); latent kernel 10/10 scenarios at num_seg auto/1/2/7/16/64; whole op 10/10. Repeated with the decode kernels compiled by CI's LLVM 24 and swapped into the harness (torch pins its own LLVM 22): same results, and a deliberately wrong binary fails the lane (7/7), so the swap is what runs.
GPU gfx942 and gfx950, after the review follow-ups (runner input checks): numeric and runner tests 81 passed (gfx942), 80 passed and 1 skipped (gfx950); latent 10/10; whole op 10/10. Later changes leave every IR golden unchanged.
Earlier heads, on GPU: prefill-kernel oracle 10/10 (gfx942) and the model-shape file 26/26 on both arches (the two 16-head tensor-parallel rows refused by design).
From rocke/platform/ (head of 4 November): tests/test_rocke.py -- 368 tests OK, 128 skipped; tests/run_all.py -- byte-identity gate green for all 90 families; its 150 pytest failures are in five platform files this PR does not touch and also fail on develop on the same host (no ml_dtypes; host llc / llvm-objdump without gfx950 support).
Deferred: the MLA kernels have no C++ builder twin (dsl_docs/development/known_gaps.md section 6); lowering parity is checked instead.

Known limits and follow-ups

Head counts below 64 (a tensor-parallel rank, e.g. 16 heads per device) are refused; they need a smaller per-workgroup head group. 192, 256, ... lower but are not admitted until validated on GPU.
qa, out and the split workspace use i32 element offsets, and a sequence's tile count is i32: oversized batches and max_seqlen_k above 2**31 - block_k are refused (the KV cache is 64-bit addressed per page).
Page ids in the block table are not range-checked on the device; check_contents=True checks them on the host (opt-in, it syncs), as the verify harness does.
Tolerances are fixed values; accumulation-aware thresholds (CK-Tile check_err) are a follow-up.
Decode now uses the shared lane setup, softmax constants and LSE helper, but still emits its own masked key tail, lazy-rescale softmax row and kernel-argument table, and declares two scalar arguments it never reads (num_seqs, max_k). Moving those onto _mla_body changes the decode IR (golden re-bless, device numerics on both arches), so it is a follow-up after refactor(rocke): MLA prefill on shared emitters; odd head counts on the packed layout #14502: AICK-3039, noted in known_gaps.md section 6.
The running max starts from a finite -1e30 (shared with prefill), so a score at or below it reads as masked; that needs |q . k| past 1e30.
Note for feat(rocke): add gfx942 bf16 MLA prefill kernel and dispatch #12464: the prefill absorbed_query kernel returns NaN when unused rows of a sequence's partial last page hold NaN (masked in the score, contracted in the PV). Repro: python -m builders.mla.verify_mla_decode --mode oracle --poison-unused --scenario ragged_h128 --num-seg 1 (gfx942). Decode zeroes those rows before its PV.

🤖 Generated with Claude Code

Copilot

🟡 Changes recommended

Unresolved validation gaps permit overflowing table offsets, unwritten outputs and silently reinterpreted requests.

7 open issues

Validate block-table stride and offset i32 bounds · New
Validate query offsets cover the full query range · New
Validate projection shapes before GEMM execution · New
Reject non-integer max_seqlen_q before request admission · New
Reject non-integer num_seg overrides · New
Guard ROCm property access on CUDA-only torch · New
Catch HipError when HIP library loading fails · New

What changed in this PR

Adds MLA decode support for gfx942 and gfx950 alongside the existing prefill path, with shared dispatch and verification tooling.

Changes:

Routes single-query requests to decode, including split-KV combine launches.
Adds device grid limits, host-side validation and numerical references.
Extends CPU/GPU tests, benchmarks and documentation.

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

File	Description
platform/tests/runtime/test_device_max_grid_size.py	Tests grid-limit queries and caching.
platform/python/rocke/runtime/hip_module.py	Adds cached device grid limits.
platform/python/rocke/instances/SUPPORT_MATRIX.md	Lists decode support.
platform/python/rocke/helpers/split_k.py	Reuses the shared grid limit.
platform/python/rocke/helpers/spec.py	Defines the fallback grid limit.
platform/python/rocke/helpers/__init__.py	Exports the grid limit.
platform/dsl_docs/instances/index.md	Indexes MLA kernels.
platform/dsl_docs/instances/attention.md	Documents MLA contracts and routing.
platform/dsl_docs/development/known_gaps.md	Records C++ parity limitations.
platform/CMakeLists.txt	Registers decode test suites.
library/tests/test_mla_prefill_gfx942.py	Updates prefill request fields.
library/tests/test_mla_prefill_gfx942_golden.py	Updates golden-test requests.
library/tests/test_mla_decode.py	Tests specs, lowering and resources.
library/tests/test_mla_decode_runner.py	Tests runner validation and binding.
library/tests/test_mla_decode_numeric.py	Adds GPU correctness coverage.
library/tests/test_mla_decode_golden.py	Checks IR stability and lowering parity.
library/tests/test_mla_decode_benchmark_shapes.py	Tests benchmark routing and reporting.
library/tests/test_builder_signature_contract.py	Registers decode spec contracts.
library/tests/golden/mla_decode_ir_sha256.json	Pins decode IR hashes.
library/tests/dispatch/mla/test_mla_prefill_gfx942_wiring.py	Updates prefill routing tests.
library/tests/dispatch/mla/test_mla_decode_wiring.py	Tests decode admission and launches.
library/kernels/mla/mla_decode_gfx950.py	Supplies gfx950 operand choices.
library/kernels/mla/mla_decode_combine.py	Merges split-KV partials.
library/examples/gfx950/mla_decode/README.md	Records gfx950 correctness results.
library/dispatch/mla/gfx950.py	Registers gfx950 decode.
library/dispatch/mla/gfx942.py	Registers decode alongside prefill.
library/dispatch/mla/decode.py	Implements decode selection and launch metadata.
library/dispatch/mla/common.py	Updates requests and device-property helpers.
library/dispatch/mla/__init__.py	Exposes unified MLA dispatch.
library/builders/mla/verify_score_probe_gfx942.py	Pins probe routing to prefill.
library/builders/mla/ref_mla_decode.py	Adds fp64 decode references.
library/builders/mla/ref_mla_attn.py	Handles empty key ranges.
library/builders/mla/mla_decode_runner.py	Launches the decode pipeline.
library/benchmarks/gfx950/attention/decode/mla_shapes.json	Updates shapes and unsupported-head cases.
library/benchmarks/gfx942/attention/decode/mla_shapes.json	Updates shapes and unsupported-head cases.
library/benchmarks/common/mla_shapes.py	Supplies the query-length routing field.
library/benchmarks/common/benchmark_mla_decode_live.py	Adds resolve, build, verify and timing modes.

_{Review effort: Balanced}

therock-pr-bot · 2026-10-20T19:56:58Z

✅ 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 · 2026-10-20T20:12:01Z

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

reshape(-1) cannot infer an axis from a zero-element tensor, so a sequence with S_k = 0 crashed the shared MLA reference. Pass the latent and rope widths explicitly. Co-Authored-By: Claude <noreply@anthropic.com>

Weight-absorbed bf16 decode: takes the absorbed query and returns the normalized latent plus LSE, with split-KV partials merged by an arch-neutral combine kernel. MLARequest gains a required max_seqlen_q field: it is what separates decode from prefill and cannot be derived from total_q and num_seqs. The family now serves one op, mla_fwd, and routes regimes by candidate: decode registers at priority 5 for max_seqlen_q <= 1 on gfx942 and gfx950, with prefill as the fallback. resolve_num_cus moves to dispatch.mla.common so both candidates share it. Prefill callers and tests state max_seqlen_q explicitly, and the prefill score probe pins its algorithm so it keeps grading the prefill kernel now that decode outranks it. Co-Authored-By: Claude <noreply@anthropic.com>

Adds the fp64 decode reference (single pass and split-KV merge), a runner that launches the dispatched latent + combine kernels, the verify_mla_decode correctness harness with its gpu-marked pytest lane, and a resolve/build/verify/time driver over the MLA decode benchmark shapes. Co-Authored-By: Claude <noreply@anthropic.com>

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

Prefill and decode take different inputs (projected query + W_UK vs the absorbed query), so prefill now admits max_seqlen_q >= 2 only: a one-query request decode cannot serve is refused instead of being handed a kernel that expects different tensors. Each kernel's input contract is stated in MLARequest and the dispatch modules. Empty split-KV segments no longer store their 512-wide partial row, and the combine skips rows with l == 0 (golden re-blessed: only the split and combine cases move). Decode is refused at llvm20, which ignores agpr_alloc and spills the kernel. mla_decode_launch / MlaDecodeLaunch describe both launches of a split selection and the workspace between them; with_num_seg forces a split count through the same gate dispatch applies. Decode owns its schedule knobs (DECODE_*), and the i32 extent check reuses INT32_LIMIT as check_mla_decode_extents. Tests: block-level IR guards for the empty-segment store and read, both flavor sides of the llvm20 gate, the launch description, the C++ parity check at every golden flavor, and the comgr resource probe in a subprocess so an LLVM option-parser exit cannot end the run. Co-Authored-By: Claude <noreply@anthropic.com>

The runner launches what mla_decode_launch describes and refuses, on the host, tensors whose shape, dtype, layout, device or block-table width does not match the dispatched request; check_contents adds an opt-in host check of the offsets and the page ids in use. A forced num_seg goes through MlaDecodeLaunch.with_num_seg, so it passes the same gate as dispatch's own choice and gets a matching kernel_id. The parity harness gains a spike scenario that lifts half of every wave's heads past the lazy-rescale threshold on a live accumulator, and poisons block-table slots past each sequence's last page with -1. The prefill oracle asks for prefill's layout at max_seqlen_q = 2 and states its page pool, both of which prefill now requires. The benchmark driver imports torch before compiling, with or without --arch, so the IR flavor resolves against torch's bundled comgr. Co-Authored-By: Claude <noreply@anthropic.com>

The decode README says which inputs each regime takes and that a one-query request decode cannot serve is refused, lists every point where the kernel departs from builders/mla/DESIGN.md, and links its code references. The results and static-resource tables are regenerated from a gfx942 and a gfx950 run of the current kernels, with the new spike scenarios; attention.md points at the README as the as-built record. Co-Authored-By: Claude <noreply@anthropic.com>

Dispatchers that size a grid axis from the problem need the limit of the device they launch on. get_device_max_grid_size reads hipDeviceAttributeMaxGridDim{X,Y,Z} through the same hipDeviceGetAttribute path as get_device_num_cus and caches it per device. MAX_GRID_DIM (65535) is the one fallback when no device can be asked; split_k's MAX_GRID_DIM_Z now aliases it instead of restating the number. The tests mock the HIP query and need no GPU. Co-Authored-By: Claude <noreply@anthropic.com>

A split-KV segment with no tiles (or a sequence with S_k = 0) still staged a page from slot 0, trusting that page 0 is always mapped; with an empty cache that read out of bounds. Its two staging descriptors now cover zero bytes, so those loads return 0 and touch no memory. The latent IR changes accordingly (goldens re-blessed; combine unchanged). Admission now checks every launch against the device's per-axis grid limit -- the live device's when it is the requested arch, else MAX_GRID_DIM -- so batches with more sequences than grid y allows are refused, and the split count is clamped to grid z. The LDS budget comes from the arch catalog instead of per-class constants. Decode admits 64 and 128 heads only, the counts with GPU evidence; larger multiples of 64 lower but wait for validation. Decode does not serve llvm20, so its goldens are recorded at llvm22 and llvm23 only, pinned to DECODE_LLVM_FLAVORS. Comments on the lazy rescale bound, the PV axis and the k_rope staging width are corrected. Co-Authored-By: Claude <noreply@anthropic.com>

The gate adds the MLA prefill harness's two bars: per sequence max_abs / peak <= 2e-2, so a lost key tile cannot hide under an elementwise bound that grows with the output, and LSE <= 1.7e-2, which one dropped page exceeds. Negative tests pin both. The runner refuses a k_rope whose page count differs from c_kv's, exposes bind_latent (checked and packed once, for timed loops), and runs every stage of forward on one stream; the torch absorb and W_UV steps are module functions with CPU tests against the fp64 reference. The benchmark builds the combine with the latent kernel, gates each shape's correctness before timing it, reports a median and spread over repeated windows, and treats only the known head-count gap as an expected refusal. The shape files gain a tensor-parallel rank of DeepSeek-V3 (16 heads), which decode does not serve yet. Co-Authored-By: Claude <noreply@anthropic.com>

Adds the decode CPU tests to the library MLA lane and the numeric test to the GPU lane, so installed-test runs cover decode like prefill. Co-Authored-By: Claude <noreply@anthropic.com>

The gfx942 README lists every launch argument and the workspace, the preconditions the kernel cannot check, a usage example, the IR flavor and comgr a run needs, and the head-count limits (a tensor-parallel rank is refused). The gate section describes the three bars, and both arches' result tables are regenerated from new gfx942 and gfx950 runs with the gated per-sequence column. Co-Authored-By: Claude <noreply@anthropic.com>

…ests The prefill tests added with the request-level absorbed_query field build MLARequest without max_seqlen_q, which the request requires since decode routes on it. They now state it. A mistyped total_q/num_seqs override is a malformed request, not a new shape, so _req no longer demands max_seqlen_q for one. max_seqlen_q joins the exact-int type tests; its mistyped-value stand-in is decode's regime, so the decode candidate names the error. Co-Authored-By: Claude <noreply@anthropic.com>

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>

cgmillette · 2026-11-02T17:44:49Z

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

PR #13923 — add MLA decode kernel for gfx942 and gfx950 (aosewski, AICK-1504) — 🟠 Follow-up: decode duplicates MLA family code

Prior item	Status
Copilot (10-19) ×7: block-table i32 bounds; `cu_seqlens_q` span; projection shapes; `max_seqlen_q` / `num_seg` truncation; non-ROCm torch; `HipError`	🟢 FIXED, unchanged • no fix site touched by `7c85919c` (ran: diff) • 7/7 threads resolved
Our 10-22 N1: no-spill gate skips in every CI shard	🟢 FIXED `7c85919c` • gate runs at llvm24: 4/4, 0 scratch (ran) • CI log at head pending: test shards not started
CongMa13 (10-22) APPROVED at `fa8e6110`	🟢 APPROVED `fa8e6110` • older head; still counts on develop
cgmillette (10-28) APPROVED at `87fb6ff2`	🟢 APPROVED `87fb6ff2` • older head; still counts on develop • since: gate rule `d493d329`; llvm24 serving `7c85919c`
Our 11-02 N1: decode refused at llvm24; CI skips golden + gate	🟢 FIXED `7c85919c` • served at llvm24 on both arches; lane 440 ✓ / 3 skip (ran) • tests fail on llvm24 drift, missing column, unserved lane (ran) • left: gfx942 on-device run at LLVM 24 (author: pending)
Our 11-02 N2: stale "Known limits" bullet in the body	🟢 FIXED, body edited • bullet now matches the code; commits 21–23 listed (read)
Bots / CI	🟡 Running • running: test stages; provider test results not in yet • 30 ✓ · 4 running · 0 ✗ · 1 cancelled (superseded bot run) • merge: blocked; Multi-Arch summary not posted
Our checks	🟢 Clean • merge vs `b088ac5a`: clean, 25 behind / 29 ahead (ran) • build + ctest 34 ✓ / 1 skip; parity, goldens, byte identity ✓ (ran) • PR tests 407 ✓ / 3 skip (torch); rules: 2 old encoding hits, ASCII (ran)
Code quality	🟠 Should-fix • DRY: N1 family emitters, N2 flavor rule, N3 gate rule (ran) • layers: N4 device probe in selection (read) • holds: actionable caller errors; launch path on platform primitives (ran)

#	Sev	file:line	Finding
N1	🟠	`kernels/mla/mla_decode.py:885`, `:958`, `:592`, `mla_decode_combine.py:274`	Decode re-inlines the MLA family's shared emitters • impact: second copy of softmax, rescale, k-loops, lanes; fixes miss one • author defers ("would move the decode IR"): swaps byte-identical 24/24 (ran) • fix: compose `_mla_body` emitters; re-bless only the masked tail
N2	🟠	`dispatch/mla/decode.py:96`, `tests/test_mla_decode_golden.py:39`, `tests/dispatch/mla/test_mla_decode_wiring.py:45`	Served flavors are a list; prefill gates the same fact by rule • impact: the list refused llvm24 (our 11-02 N1); fact in six places (read) • fix: one family predicate in `_mla_body` for both gates and the tests
N3	🟠	`tests/test_mla_decode.py:514`, `tests/test_mla_prefill_gfx942.py:43`	Resource gate's fail-or-skip rule copied from prefill's tests • impact: two copies of a CI rule; diverged until `d493d329` (read) • fix: one `tests/_mla_resource_gate.py` both gates call
N4	🟡	`dispatch/mla/common.py:374`	Decode selection reads the live device's grid limits • impact: selection differs with or without the device; dispatch §6.2 rule 3 (read) • fix: select against `MAX_GRID_DIM`; check the device at launch

Utility: 🟢 Positive — adds MLA decode on gfx942/gfx950; costs a second copy of family emitters.

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: the delta since our last review, merge with develop, build + ctest, goldens, PR tests, the MLA lane at LLVM 24 with negative controls, CI at head, code quality. Flags whether a closer look is needed; not an approval or rejection.

Review follow-up for the MLA decode PR. Three findings had the same fact in more than one place, or broke a dispatch rule. One flavor rule. flavor_honors_agpr_alloc in kernels/mla/_mla_body.py says which LLVM flavors the MLA kernels serve: those whose backend honors agpr_alloc, the modern flavors. Decode's admission, prefill's agpr_acc gate and the tests all ask it, so decode's own list of flavors is gone. The decode golden test derives its flavors from the rule and fails until every flavor the rule serves has a blessed column. Prefill's llvm20 refusal no longer names only llvm22 and llvm23; the two messages the prefill golden records for it are updated to match. Pure selection. Decode selection read the live device's grid limits, which dispatch ARCHITECTURE.md section 6.2, rule 3 forbids: selection must not depend on the host. It now checks the platform's MAX_GRID_DIM, and MlaDecodeRunner asks the device for its own limit at launch. A test fails if selection asks the device for anything. The tests of the launch check are torch-free, so CI runs them. One gate rule. The fail-or-skip rule of the resource gates was copied between prefill's tests and decode's. It now lives in tests/_mla_resource_gate.py, and a truth table covers every probe outcome with the compiler tools optional and required. The one device read selection keeps, the CU count behind an unset num_cus, has its own tests again. The runner tests no longer depend on the host's GPU, the READMEs say the rule instead of a list, and comments that described the old behavior are corrected. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>

aosewski · 2026-11-03T13:51:43Z

@cgmillette thanks for the follow-up. N2, N3 and N4 are fixed in c41875b; N1 is deferred. In detail:

N2, the flavor rule: fixed. One predicate, flavor_honors_agpr_alloc in kernels/mla/_mla_body.py, now decides which LLVM flavors the MLA kernels serve. Decode's admission, prefill's agpr_acc gate and the tests all ask it, so DECODE_LLVM_FLAVORS is gone. The golden test compares its columns with the served set, so a new flavor the rule serves fails until its column is blessed. Prefill's llvm20 refusal text no longer names only llvm22 and llvm23; its two recorded messages in the prefill golden are updated.
N3, the resource-gate rule: fixed. tests/_mla_resource_gate.py holds the fail-or-skip rule and both gates call it. A truth table covers every probe outcome with the compiler tools optional and required, including "a host mismatch skips even where the tools are required".
N4, selection reading the device: fixed. Selection now checks the platform's MAX_GRID_DIM, so it does not depend on the host (dispatch ARCHITECTURE.md section 6.2, rule 3). MlaDecodeRunner checks the device's own grid limit at launch. A new test fails if selection asks the device for anything, and the launch check's tests are torch-free so CI runs them.
N1, decode re-inlining the shared emitters: deferred to a follow-up PR. The pieces that move decode's IR (the masked tail) need a golden re-bless and device numerics on gfx942 and gfx950 (the gfx942 run at LLVM 24 is still pending), and the shared module is still moving. I would rather land it separately with that evidence than widen this PR.

yraparti

rocKE checklist review

ROCm/rocm-libraries#13923 · c41875b42c05 · model gemini-4-argon, fallback gpt-5.5

PASS · graded 238/238 · BLOCKER 0, HIGH 5, MEDIUM 2, LOW 2, NEEDS-INFO 35

The diff was truncated before grading. Items that needed a missing hunk are NEEDS-INFO.

Findings

HIGH §B · dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:194 · In MlaDecodeSpec.__post_init__ (mla_decode.py:201-214) and MlaDecodeCombineSpec.__post_init__ (mla_decode_combine.py:86-91), validate if type(value) is not int: (including num_warps) before range checks and raise ValueError.
HIGH §D · dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:28 · Put the knob on the KernelSpec and delete the environment read.
HIGH §H · dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:424 · Change check_contents: bool = False to check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward (leaving bind_latent with check_contents=False for the hot/timed path).
HIGH §H · dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:678 · Default check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward (or range-check page against num_pages on device before b.zext(b.readfirstlane(page), I64)), because folding page64 * stride into the 64-bit base pointer bypasses num_records and check_contents is currently off by default.
HIGH §O · dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:337 · Replace the string split comparison in _live_device_is with the existing target-id/base-arch resolution helper or ArchTarget-based comparison, and use the same resolved base arch for both the live device and request.
MEDIUM §E · dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:169 · Query the device through rocke.runtime.hip_module instead of importing torch.
MEDIUM §E · dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:591 · Remove the unused num_seqs and max_k scalar parameters from mla_decode_signature and build_mla_decode (and pass them only in run_prefill_oracle when invoking the prefill kernel), or use them in-kernel.
LOW §H · dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:617 · Enforce spec.block_k & (spec.block_k - 1) == 0 in supports_mla_decode (or MlaDecodeSpec.__post_init__), since b.div(..., b.const_i32(BK)) is used in the kernel and supports_mla_decode currently only checks spec.block_k % MFMA_N == 0.
LOW §S · platform/python/rocke/helpers/spec.py:264 · Run PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py from rocke/platform/ and record them in the PR test plan.

Needs info

§A.1 The diff for dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md (+12 -0 lines) is omitted from the prompt, so the exact entry documenting MLA decode as Python-only cannot be verified by line number.
§A.5 The diff for dnn-providers/hip-kernel-provider/rocke/library/tests/golden/mla_prefill_gfx942_ir_sha256.json (+2 -2 lines) is omitted, so it is not visible what changed in the existing prefill golden fixture.
§E.7 The diffs for SUPPORT_MATRIX.md, platform/dsl_docs/instances/attention.md, platform/dsl_docs/instances/index.md, and platform/dsl_docs/development/known_gaps.md are omitted from the prompt budget; additionally library/tests/_mla_resource_gate.py:23 introduces ROCKE_REQUIRE_COMPILER_TOOLS while dsl_docs/reference/env_flags.md is not in the file list.
§F.6 No diff-citable evidence shows that pre-commit was run locally and passed before push.
§F.7 GitHub bot/Copilot review-thread state, fixes, won't-fix justifications, and any deferred-ticket IDs are not present in the diff.
§F.11 No diff-citable evidence shows the PR title/body were read back from the GitHub API field after editing.
§G.5 Several test files in the PR (test_mla_decode.py, test_mla_decode_reference.py, test_mla_decode_runner.py, test_device_max_grid_size.py) are omitted from the diff budget, so reachability and falsifiability of all tests cannot be verified.
§G.6 dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md, dnn-providers/hip-kernel-provider/rocke/library/examples/gfx950/mla_decode/README.md, and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py are omitted from the diff, preventing verification of prose claims and resource tables against code/tests.

Section notes

§A: MLA decode is a Python-only instance family documented in known_gaps.md; lowering parity is tested via test_mla_decode_cpp_python_byte_identity and golden hashes are recorded in mla_decode_ir_sha256.json.
§D: Items 1-11 of section D
§E: Structure & hygiene checks on MLA decode PR.
§F: Several merge-policy items depend on PR metadata, local command results, GitHub review threads, or API state that is not present as diff-citable evidence.
§G: Items 1-17 of 34 graded; several test and doc files were truncated in the diff.
§G: Items 18-34 of Section G

Graded with the rocKE PR checklist, one model pass per section. FAIL rows that sit on a diff line are inline comments.

§A graded 14/14

#	Verdict	Evidence
1	NEEDS-INFO	The diff for `dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md` (+12 -0 lines) is omitted from the prompt, so the exact entry documenting MLA decode as Python-only cannot be verified by line number.
2	N/A	No core lowering-op or op-handler files (`core/lower_*`, `core/ir`) are modified in this PR.
3	PASS	Changes in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:1` and `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:1` are for the Python-only MLA family and dispatch layer, which have no C++ selector or builder twin…
4	N/A	No C++ kernel builder or helper twin exists or is modified in this PR.
5	NEEDS-INFO	The diff for `dnn-providers/hip-kernel-provider/rocke/library/tests/golden/mla_prefill_gfx942_ir_sha256.json` (+2 -2 lines) is omitted, so it is not visible what changed in the existing prefill golden fixture.
6	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:212` (`test_mla_decode_cpp_python_byte_identity`) explicitly tests cross-engine C++/Python lowering byte-identity in addition to the Python-drift golden hash check at line 193.
7	PASS	Existing prefill kernel emitter logic in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_prefill_gfx942.py:745` only refactors the host admission check to use `flavor_honors_agpr_alloc` without altering emitted IR ops.
8	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:212` provides `test_mla_decode_cpp_python_byte_identity` and the PR test plan verifies all 14 cases across llvm22, llvm23, and llvm24 with the C++ engine built.
9	N/A	This PR does not contain a shared op-handler fix or lowerer-only arch divergence.
10	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:119` defines golden hash cases across both arches (`gfx942`, `gfx950`), single-pass and `split_kv` (`latent_split_h*`), `combine_h128`, and dispatched configs across all served LLVM…
11	N/A	No new IR op is introduced in this PR.
12	N/A	`core/backend.py` and `ROCKE_BACKEND` fallback logic are not touched in this PR.
13	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:138` lowers kernels via `_lower_kernel_to_llvm_python` for deterministic SHA256 hashing without debug locations or build stamps.
14	PASS	Changes in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264` and `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/split_k.py:168` only define/alias the host-side constant `MAX_GRID_DIM`, and neither `spec.py` nor…

§B graded 6/6

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:81-87 documents the head-count scope (`DECODE_NUM_HEADS = (64, 128)`) with a TODO for validating 192/256 and smaller-BLOCK_H tensor-parallel shards vs hardware/layout limits in…
2	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:271-275 calls `spec_for(req)` inside `support(req)` and catches `ValueError` from `spec_cls` (`MlaDecodeSpec` / `MlaDecodeGfx950Spec`), then runs `_decode_spec_gate(spec, req)` before returning.
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:327 sets `ShapeRange("max_seqlen_q", min=2)` on prefill so `max_seqlen_q == 1` is never silently demoted to prefill, and…
4	FAIL	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:194-224 and `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:81-91` check `value <= 0` on integer fields (`num_heads`, `num_seg`, etc.) without checking `type(value) is int`; for direct-spec callers bypassing dispatch, `MlaDecodeSpec(num_heads=64, num_seg=True)` or `num_seg=2.0` is accepted (`True > 0` is `False` for `split_kv`, `2.0 > 1` enables the split-KV ABI).
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:87-216 scopes admission to `DECODE_NUM_HEADS = (64, 128)` and `flavor_honors_agpr_alloc`, which are verified in…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:303-373 (`supports_mla_decode`) and `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:108-128` (`supports_mla_decode_combine`) check only functional/hardware…

§C graded 10/10

#	Verdict	Evidence
1	PASS	Shared predicates and logic are extracted into single named helpers: `flavor_honors_agpr_alloc` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:154`, `resolve_num_cus` in…
2	PASS	One reusable builder per kernel family (`build_mla_decode(spec)` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:512` and `build_mla_decode_combine(spec)` in…
3	PASS	Builders and `kernel_name()` are named by operation and codegen knobs, never workload sizes (`num_seg` is excluded from `kernel_name()`, keeping only `split_kv`): `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:271` and…
4	PASS	The shared MLA decode body in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:165` uses spec class attributes (`ARCH`, `SCORE_K`, `PV_TRANSPOSE_READ`) with the gfx950 subclass in…
5	PASS	Every touched file belongs directly to the MLA decode kernel, dispatch, verification, tests, or the shared `MAX_GRID_DIM` / `get_device_max_grid_size` runtime helpers (`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264`…
6	PASS	Shared helpers (`resolve_num_cus` in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:346`, `flavor_honors_agpr_alloc` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:154`, `MAX_GRID_DIM` in…
7	N/A	This PR introduces a new MLA decode feature rather than performing a standalone no-op cleanup PR.
8	PASS	Reuses existing helpers from `develop` including `read_ahead` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:86`), `TransposeLdsReader` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:93`), `safe_inv_l`…
9	PASS	Bug fixes and gates are applied across sibling sites: `flavor_honors_agpr_alloc` unifies decode and prefill (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:266`…
10	PASS	Widening `MLARequest` with `max_seqlen_q` and `MLA_OP = 'mla_fwd'` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:54`, `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:137`) and `MLA_REGISTRY` with `gfx950`…

§D graded 21/21

#	Verdict	Evidence
1	FAIL	dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:28 reads the environment
2	PASS	`make_mla_decode_candidate` declares `Capability(arches=(arch,), ...)` where `arch = spec_cls.ARCH` (`gfx942` or `gfx950`), and `DECODE_NUM_HEADS = (64, 128)` is routing scope documented as GPU-validated counts…
3	PASS	`support` in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:248-276` runs cheap request/selector/flavor checks before `spec_for` and `_decode_spec_gate`, does not duplicate `Capability` checks, and `select` plus all tests call…
4	PASS	`decode_spec_for` and `_check_grids` take `arch` from `req` and only query device CU count when `req.num_cus == 0` via `resolve_num_cus(req)` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:127-194`, `common.py:346-359`), pinned by…
5	PASS	`MlaDecodeSpec.identity()` removes launch-only `num_seg` and includes `split_kv` plus class emission attributes (`arch`, `score_k`, `pv_transpose_read`), matching `kernel_name()` and the emitted IR…
6	PASS	Single-pass vs split-KV routing keeps `decode_spec_for`, `mla_decode_grid`, `mla_decode_signature`, `MlaDecodeLaunch`, and `MlaDecodeSpec.identity()` in agreement (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:127-166, 328-418`…
7	N/A	This PR does not introduce machine-regenerable heuristics or hand-pinned tuning overrides above them.
8	PASS	The LLVM flavor cohort rule `flavor_honors_agpr_alloc` is defined once in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:154-169` and imported by both `dispatch/mla/decode.py:32, 266` and `kernels/mla/mla_prefill_gfx942.py:105, 748`.
9	PASS	`make_mla_decode_candidate` registers `mla_decode_gfx942` and `mla_decode_gfx950` on `MLA_REGISTRY` with `Capability`, `build`, `grid`, `block`, and `signature` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:289-324`, `gfx942.py:385-393`…
10	PASS	The new `MLARequest.max_seqlen_q` field (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:137, 228`) and `MlaDecodeSpec.num_seg` (`decode.py:152-165`) are wired from request through dispatch to launch and tested at the dispatcher level in…
11	PASS	`mla_decode_gfx942` and `mla_decode_gfx950` are registered for the disjoint `max_seqlen_q == 1` decode regime where no prior MLA candidate existed (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:296-300`, `gfx942.py:327`).
12	PASS	Priority `DECODE_PRIORITY = 5` is justified against prefill (10) in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:73-75` and tested in `dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:85-87`.
13	PASS	`dispatch_mla` delegates directly to `MLA_REGISTRY.select` without catching or retrying on `PinRefused` in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/__init__.py:107`; the harness in…
14	PASS	`MlaDecodeSpec.identity()` omits the problem-dependent launch argument `num_seg` and keeps only `split_kv` and arch emission choices so `compile_key` does not fork per batch shape (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:253-270`)…
15	PASS	The PR intentionally partitions the MLA regime by `max_seqlen_q` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:325-327`), which is a deliberate routing change graded and tested in…
16	PASS	The existing `mla` family already exposes `registered_mla_combos`, `mla_sweep_space`, and `dispatch_mla_all` with `_request_errors` short-circuiting (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/__init__.py:89-94`), and `MLARequest` updates…
17	PASS	CPU routing tests in `dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:94-245` check cohort selection across `gfx942` and `gfx950`, priority over prefill, regime separation, and out-of-cohort rejection through `admits()`.
18	PASS	`MLA_REGISTRY` keeps `require_build=True` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/__init__.py:48-50`) and no ratchets or exempt sets are loosened.
19	N/A	This PR adds MLA decode kernels and dispatch under `dispatch/mla/` and does not touch `run_unified_attention_torch` or unified…
20	PASS	The PR explicitly documents in `dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:14-18` that MLA has no C++ instance mirror while lowering parity is verified via `test_mla_decode_cpp_python_byte_identity`…
21	N/A	This PR introduces a new MLA decode kernel and its dispatcher registration rather than working around a selection bug in an emitter.

§E graded 17/17

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/hip_module.py:343 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264 add platform-only code with zero library imports; library imports platform downward.
2	FAIL	dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:169 adds `import torch`
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_numeric.py:23 gates torch with pytest.importorskip('torch'), and dnn-providers/hip-kernel-provider/rocke/library/builders/mla/ref_mla_decode.py:24 uses numpy only.
4	PASS	Layering is respected downward: kernels (dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:67) -> dispatch (dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:32) -> builders…
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_prefill_gfx942.py:19 uses tempfile.TemporaryDirectory and pathlib.Path; platform changes in hip_module.py:343 and spec.py:264 contain no absolute or repo-escaping paths.
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:212 verifies C++/Python byte-identity and IR SHA256 goldens across all served LLVM flavors.
7	NEEDS-INFO	The diffs for SUPPORT_MATRIX.md, platform/dsl_docs/instances/attention.md, platform/dsl_docs/instances/index.md, and platform/dsl_docs/development/known_gaps.md are omitted from the prompt budget; additionally library/tests/_mla_resource_gate.py:23 introduces ROCKE_REQUIRE_COMPILER_TOOLS while dsl_docs/reference/env_flags.md is not in the file list.
8	PASS	Visible diffs (e.g. dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:1 and dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:1) contain no Jira IDs, PR numbers, or GitHub handles; tracking ID AICK-1504 is kept in the PR…
9	PASS	All touched files in the diff (e.g. dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:1, dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/hip_module.py:351…
10	PASS	The PR body notes explicit coordination with #12464 and merge simulation with develop, and dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_prefill_gfx942_wiring.py:42 updates existing prefill tests to match MLARequest.max_seqlen_q.
11	PASS	Only deliverable source, tests, shape JSONs, and required instance/example READMEs are committed (e.g. dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:1…
12	PASS	Both gfx942 and gfx950 are covered: dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:385 and dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx950.py:30 register decode candidates for both arches, and…
13	PASS	Renamed/hoisted helpers (_resolve_num_cus -> resolve_num_cus in dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:346, _require_tools_or_skip -> require_tools_or_skip in…
14	PASS	Removed lines in modified files (e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:95, dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/split_k.py:168) only replace local definitions with shared helpers without dropping…
15	FAIL	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:591-593 declares scalar kernel parameters `num_seqs` and `max_k` (`b.param("num_seqs", I32)` and `b.param("max_k", I32)`) whose return handles are discarded and never read in the kernel body, while dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:377-381 passes `req.num_seqs` and `req.max_seqlen_k` into them on every launch.
16	N/A	This item applies to how review feedback is written rather than the code changes in the pull request diff.
17	PASS	no absolute throughput, bandwidth, or latency figure in added lines

§F graded 11/11

#	Verdict	Evidence
1	PASS	branch `users/aosewski/aick-1504-mla-decode` uses the users// prefix
2	PASS	title `feat(rocke): add MLA decode kernel for gfx942 and gfx950` matches type(scope): summary and is 56 characters
3	PASS	description is at least 30 characters and contains an issue reference
4	PASS	test file `dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py` is in the diff
5	PASS	The added files shown in the diff are source/test files with normal headers, e.g. dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:1 and…
6	NEEDS-INFO	No diff-citable evidence shows that pre-commit was run locally and passed before push.
7	NEEDS-INFO	GitHub bot/Copilot review-thread state, fixes, won't-fix justifications, and any deferred-ticket IDs are not present in the diff.
8	N/A	This PR is on a users/... branch rather than a stacked feature branch, so the stacked-feature-branch bot-run caveat does not apply.
9	N/A	The diff does not show a split or trim made solely to satisfy a PR-size cap.
10	PASS	The added tests exercise concrete behavior, such as decode routing assertions in dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:83 and GPU numeric scenarios in…
11	NEEDS-INFO	No diff-citable evidence shows the PR title/body were read back from the GitHub API field after editing.

§G graded 34/34

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:211-245 tests C++/Python byte-identity separately from on-GPU correctness in dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_numeric.py:89-105.
2	PASS	dnn-providers/hip-kernel-provider/rocke/library/builders/mla/verify_mla_decode.py:116-124 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_numeric.py:89-105 implement the on-silicon numeric gate against fp64 reference with bf16 tolerances, and…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:130-136 and 193-245 gate all modern LLVM flavors (llvm22, llvm23, llvm24) and C++/Python byte-identity, alongside CPU suites in…
4	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:94-243 and 349-535 add behavioral tests for decode routing, grid limits, and launch checks that fail without the new code.
5	NEEDS-INFO	Several test files in the PR (test_mla_decode.py, test_mla_decode_reference.py, test_mla_decode_runner.py, test_device_max_grid_size.py) are omitted from the diff budget, so reachability and falsifiability of all tests cannot be verified.
6	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md, dnn-providers/hip-kernel-provider/rocke/library/examples/gfx950/mla_decode/README.md, and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py are omitted from the diff, preventing verification of prose claims and resource tables against code/tests.
7	N/A	This PR does not add or modify a design doc (`DESIGN.md` is not in the changed files list).
8	PASS	dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:195-210 and 539 loudly raise ValueError on non-bf16 tensors, and dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:61 filters strictly to bf16.
9	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py (which contains TestResources per mla_decode.py:124) and the example READMEs are omitted from the diff, so tool-sourced VGPR/LDS assertions cannot be inspected.
10	N/A	The PR explicitly makes no performance claims and is a functional bring-up.
11	PASS	dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:79 and dnn-providers/hip-kernel-provider/rocke/library/builders/mla/verify_mla_decode.py:582 route every shape through `MLARequest -> dispatch_mla` (MLA is documented as…
12	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:211-245 verifies C++ and Python lowering parity across all 14 decode cases and all served LLVM flavors, while documenting that MLA has no C++ instance mirror.
13	N/A	The PR uses existing DSL builder methods (`IRBuilder`) and does not introduce or recommend raw LLVM intrinsic mangled symbols.
14	N/A	The PR is a functional bring-up with no bottleneck classification or performance claims.
15	PASS	All graded files and hunks come directly from the provided PR diff for commit `c41875b` (e.g., dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:1-100).
16	N/A	This is a completed code PR, not a pre-implementation plan.
17	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_benchmark_shapes.py:67-70 runs `benchmark_mla_decode_live.main` end-to-end in CPU resolve mode, and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_numeric.py:59-105 runs…
18	NEEDS-INFO	The tests for the verification guards and runner checks (`library/tests/test_mla_decode.py`, `library/tests/test_mla_decode_runner.py`, `library/tests/test_mla_decode_reference.py`) are omitted from the diff, so we cannot verify whether each guard is tested to fail for its named reason and documents its blind spots.
19	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/builders/mla/verify_mla_decode.py:246-272`, `make_problem` allocates `len(k_lens) * max_blocks` shuffled pages, poisons all unowned cache rows (including partial last pages) with `NaN`, and sets unused block_table…
20	N/A	This PR does not introduce or modify a detector/linter for banned performance claims.
21	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:160-291` and `dnn-providers/hip-kernel-provider/rocke/library/builders/mla/verify_mla_decode.py:494-554`, all CLI flags (`--warmup`, `--iterations`, `--repeats`…
22	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:181-190`, `test_mla_decode_golden_records_no_build_errors` asserts that no golden entry contains an `error` key, and…
23	N/A	This PR is a functional bring-up of the MLA decode kernel and makes no performance or routing speedup claims.
24	NEEDS-INFO	The diff for `dnn-providers/hip-kernel-provider/rocke/library/tests/golden/mla_prefill_gfx942_ir_sha256.json` (+2/-2 lines) and `dnn-providers/hip-kernel-provider/rocke/library/tests/golden/mla_decode_ir_sha256.json` is omitted, so we cannot inspect what changed in the existing prefill golden fixture.
25	N/A	This PR introduces new MLA decode kernels rather than applying a codegen micro-optimization to an existing kernel.
26	N/A	This PR makes no benchmark speedup claims against a baseline or external library.
27	N/A	This PR does not report a performance A/B study toggling optimization knobs.
28	N/A	This PR adds a new MLA decode kernel path and does not present a performance delta table or counterfactual comparison.
29	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:213-228`, `exp2_fast` is called on `fsub(ml[0], big_m)` where `big_m = max_s(m_s)`, so the argument is provably `<= 0` (and in `mla_decode.py:865-875` `exp2_fast` is used for online…
30	NEEDS-INFO	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py` and the example READMEs (`library/examples/gfx942/mla_decode/README.md`, `library/examples/gfx950/mla_decode/README.md`) are omitted from the diff, though `test_mla_prefill_gfx942.py:474-476` shows `analyze_hsaco` is used for prefill resource checks.
31	N/A	This PR is a new kernel bring-up and does not claim to fix a performance defect against a baseline.
32	NEEDS-INFO	The documentation files recording the compiler and resource tables (`dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md` and `gfx950/mla_decode/README.md`) and `library/tests/test_mla_decode.py` are omitted from the diff.
33	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:333-359`, `_resolve_num_cus` was hoisted from `dispatch/mla/gfx942.py:95-113` to `resolve_num_cus` with identical logic for positive `req.num_cus` and live-device arch matching (narrowing…
34	N/A	This PR does not add an opt-in output flag; `lse` is always a required output of the MLA decode kernel signature…

§H graded 5/5

#	Verdict	Evidence
1	FAIL	dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:424
2	FAIL	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:678
3	FAIL	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:617
4	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:629
5	NEEDS-INFO	`test_mla_decode.py` (+730 lines) and `test_mla_decode_runner.py` (+484 lines) are omitted from the truncated diff, so negative test coverage for every launch input (`qa`, `c_kv`, `k_rope`, `cu_seqlens_q`, `cu_seqlens_k`, `block_table`, `num_seg`) cannot be verified.

§I graded 6/6

#	Verdict	Evidence
1	N/A	This PR is a functional bring-up of a single MLA decode layout per architecture without a multi-candidate tuning space or routing…
2	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:160 clamps `num_seg` to `mla_decode_max_num_seg` and `MAX_GRID_DIM` (repair), while lines 266-269 and 299-314 loudly reject `llvm20`, missing `absorbed_query=True`, or unsupported head counts…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_benchmark_shapes.py:48 tests curated model shapes from `mla_shapes.json`, and `dnn-providers/hip-kernel-provider/rocke/library/builders/mla/verify_mla_decode.py:141` exercises generated ragged…
4	N/A	This PR does not add a multi-candidate routing benchmark or tuning table comparing shipped vs best-achievable selections.
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:28 documents the checked-in IR SHA256 golden snapshot and provides the explicit `--write` regeneration path (`test_mla_decode_golden.py:265`).
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:591 uses existing general DSL/layout primitives (`TransposeLdsReader` and `query_reduce` in `_mla_body.py`) without introducing any new instruction-named DSL op.

§J graded 7/7

#	Verdict	Evidence
1	NEEDS-INFO	The Markdown hunks for the new and changed README/dsl_docs files are omitted, so the equations, operand annotations, and notation-table orientation needed for type-checking are missing.
2	NEEDS-INFO	The Markdown hunks are omitted, so the notation tables, symbol definitions, domains, and log-base statements cannot be verified.
3	NEEDS-INFO	The changed Markdown prose is omitted, so it is not possible to tell whether the new docs restate an algorithm owned by a sibling document or only describe the delta.
4	NEEDS-INFO	The novelty claims in the Markdown files are not visible, and the needed origin/develop comparison for those claims is not provided in the truncated diff.
5	NEEDS-INFO	The Markdown hunks are omitted, so any scaling or complexity claims in the docs and their repeated sites cannot be checked.
6	NEEDS-INFO	The model tables, constants, tuned band edges, and provenance text in the Markdown documentation are not visible in the truncated diff.
7	NEEDS-INFO	The new README Markdown contents are omitted, so the sibling naming/ownership convention and declared entry point cannot be verified.

§K graded 13/13

#	Verdict	Evidence
1	N/A	No tuned candidate or KnobSpace-backed candidate with multiple named configurations per request is added; MLA decode is a…
2	N/A	The diff adds no KnobAxis or tunable knob fields; the new MLA spec fields are not exposed as tuning axes.
3	N/A	No tuning_id/config_key payload or tuning-store identity is changed because MLA has no accepted non-auto tuning ids.
4	N/A	No KnobSpace.defaults() or recorded() implementation is added or changed.
5	N/A	No knob inertness, base-value, refusal, or known-wrong knob policy is added or changed.
6	N/A	No KnobSpace.validate path, tuning_specs policy-free constructor, or tuned sweep/pin legality gate is added or changed.
7	N/A	No KnobSpace.base(req) or tuning-problem base resolver is added or changed.
8	N/A	No accepted knob-value boundary is added; MLA does not accept tuning knob values in this diff.
9	N/A	The added benchmark resolves and times live dispatch selections but does not record or replay tuning configurations.
10	N/A	No bare tuning_id lookup or full-space tuning-id resolution path is added or changed.
11	N/A	No tuned sweep entry point, sampling entry point, pin_to_spec storage path, or sweep-level iterator contract is added or changed.
12	PASS	The routing change is explicitly pinned: decode selects at one query per sequence in dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:111, prefill wins at two queries in…
13	PASS	The benchmark sends shapes through dispatch_mla in dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:74 and only classifies unsupported shapes after dispatch raises in…

§L graded 10/10

#	Verdict	Evidence
1	N/A	This PR does not add or modify a gfx target in arch_specs.json or backend registries.
2	N/A	This PR does not add or modify any MMA op_id in the hardware catalog or lowering.
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:82-85 and :474-476 use the catalog helper wrappers mfma_16x16x16_for_dtype and mfma_16x16x32_for_dtype without hard-coding intrinsic names.
4	N/A	This PR only targets gfx942 and gfx950 (both wave64) and does not alter ArchTarget.wave_size or add wave32/gfx1250 support.
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:165-171 sets SCORE_K=16 and PV_TRANSPOSE_READ=False on MlaDecodeSpec (gfx942), while dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_gfx950.py:44-49 isolates SCORE_K=32…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:32-60 and dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:72-95 keep kernel schedule and dispatch logic in library/kernels and library/dispatch without mutating ArchTarget…
7	N/A	This PR operates exclusively on bf16 and f32 tensors and does not add or modify sub-byte/FP8 dtypes or packing formats.
8	N/A	This PR does not use or modify scaled MMA operations.
9	N/A	This PR does not add a new gfx architecture target to the platform.
10	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:154-168 defines flavor_honors_agpr_alloc by delegating to _is_modern_flavor, and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:130-136 derives served flavors…

§M graded 10/10

#	Verdict	Evidence
1	NEEDS-INFO	The diff truncates dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py before the end of the new kernel body; the full hunk is needed to verify no device SSA Value is used in host-time Python control flow and no scoped SSA values escape.
2	N/A	The diff does not touch legalize_for_target or target legalization code.
3	N/A	The diff does not change ck.dsl.ir/v1 serialization, print_ir interchange behavior, or serialization versioning.
4	N/A	The diff does not change the portable IR, recipe, bundle schema, importer, recorder coverage, or magic-division implementation.
5	N/A	The new kernels do not use AsyncTileLoader for async DRAM-to-LDS copies.
6	NEEDS-INFO	The diff truncates dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py in the new kernel body; the full hunk is needed to verify all barriers, pointer loads, masking, and cross-lane operations obey the hazards rule.
7	PASS	The new buffer path constructs resources through the IR builder rather than hard-coding resource word 3 in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:484.
8	PASS	The shared MLA flavor gate imports and delegates to _is_modern_flavor in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:165 and returns that check at dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:167.
9	N/A	The diff does not change public C entry points, C headers under cpp/include/rocke, or HIP launch behavior in the codegen-only C++…
10	N/A	The diff does not emit !dbg metadata or change debug-info/code-object consumption paths.

§N graded 18/18

#	Verdict	Evidence
1	PASS	`MAX_GRID_DIM` in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264` is a non-SSA host constant shared across GEMM/wgrad (`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/split_k.py:33,168`) and MLA decode…
2	PASS	The MLA decode kernels reuse shared helpers (`mfma_16x16x16_for_dtype`, `mfma_16x16x32_for_dtype`, `safe_inv_l`, `TransposeLdsReader`, `SignatureBuilder`, `read_ahead`, `bf16x4_rne_finite`, and `_mla_body` helpers) in…
3	N/A	This PR does not modify or use `SoftwarePipeline`.
4	N/A	This PR does not add or modify `EpilogueOp`, `_PATTERN_TABLE`, or graph/epilogue fusion kernels.
5	N/A	This PR does not add or modify any API in `rocke.helpers.tiling`.
6	N/A	This PR does not add or modify front-door/toolbox tiling helpers or `LayoutStyle` classes in `rocke.helpers.tiling`.
7	N/A	This PR does not use `rocke.helpers.tiling` or custom tiling layout overrides.
8	N/A	This PR does not use or modify `rocke.helpers.tiling` memory verbs (`lds_tile_alloc`, `lds_store`, `lds_read`).
9	N/A	This PR does not use or modify `rocke.helpers.tiling` N-D tensor axis roles or `FlowStage` transforms.
10	N/A	This PR does not change `rocke.helpers.tiling` emission, LDS staging, or MMA issue.
11	N/A	This PR does not use or modify `rocke.helpers.tiling` (`TileMma`, `Tiling`, `lds_swizzle`, or `derive_sched_group_counts`).
12	N/A	This PR makes no LDS bank-conflict or coalescing claims under `rocke.helpers.tiling`.
13	N/A	This PR does not build or modify `WarpDistributionEncoding`, `make_tile_desc`, or `RegisterMapper` in `rocke.helpers.tiling`.
14	N/A	This PR does not use or modify `rocke.helpers.tiling` data-movement verbs or tiling demo pipelines.
15	N/A	This PR does not modify `rocke.helpers.tiling` allocation or bounds checks.
16	N/A	This PR does not use or modify `optimize_layout` or `evaluate_transform` in `rocke.helpers.tiling`.
17	N/A	This PR does not touch the `rocke.helpers.tiling` recorder, round-trip analysis, or visualization modules.
18	N/A	This PR does not touch `InterleavedStyle` or demo kernels under `rocke.helpers.tiling`.

§O graded 16/16

#	Verdict	Evidence
1	NEEDS-INFO	The new hand-built IR in library/kernels/mla/mla_decode.py is truncated before the end of build_mla_decode, so the verifier-relevant emitted IR cannot be fully inspected.
2	NEEDS-INFO	The scf.for/scf.if bodies and yields in library/kernels/mla/mla_decode.py are partly after the truncation point, so loop/yield/terminator structure cannot be fully checked.
3	PASS	Deterministic IR bytes are pinned by the new golden hash test that hashes lowered LLVM IR for each case and flavor in dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:145 and compares it to the fixture in…
4	N/A	The diff does not change IR type equality, dtype conversion, SmemType naming, or smem exclusive parsing.
5	N/A	The diff does not edit core/passes.py, cpp/core/passes.cpp, op purity tables, or optimize_ir behavior.
6	N/A	The diff does not change constant folding logic for arithmetic operations.
7	PASS	The new decode kernel uses structured kernel attrs, including DECODE_AGPR_ALLOC in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:113 and max_workgroup_size/agpr_alloc assignment in…
8	N/A	The diff does not add scheduler strategy handling or raw compile_kernel scheduler options.
9	NEEDS-INFO	The new barrier-heavy MLA decode builder is truncated, so the complete emitted wait/barrier pattern for gfx942 and gfx950 cannot be inspected.
10	N/A	The diff does not change lowerer attr validation, dpp_ctrl_attr handling, or pointer parameter LLVM type lowering.
11	FAIL	The new live-device architecture check compares target strings directly with split(':') in dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:337 instead of going through ArchTarget/base-arch helpers.
12	PASS	The new grid cap is centralized as MAX_GRID_DIM in dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:262 and split_k aliases it instead of keeping its own literal in…
13	N/A	The diff does not touch public C structs, public C enums, ROCKE_ABI_VERSION, or the Python BINARY_ABI mirror.
14	NEEDS-INFO	platform/CMakeLists.txt is listed as changed but its hunk is omitted, so new warning flags, test registration, and CMake behavior cannot be checked.
15	PASS	The new parity test forces strict C++ lowering and fails on backend coverage gaps in dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:205 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:233.
16	NEEDS-INFO	platform/dsl_docs/development/known_gaps.md is changed but its hunk is omitted, so it cannot be checked for widened or overstated debt.

§P graded 14/14

#	Verdict	Evidence
1	PASS	Builders take `(spec, *, arch)`: `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:514`, `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_gfx950.py:52`…
2	PASS	Signatures are constructed via `SignatureBuilder` with canonical pointer/scalar types (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:415`, `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:149`), and no…
3	PASS	Harnesses take the launch grid directly from the dispatch result or `MlaDecodeLaunch` (`dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:103`, `dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:398`)…
4	PASS	The LDS gate in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:367` charges the exact `kv_lds + qr_lds` allocation (`_lds_bytes` at line 292) against `attention_lds_capacity_bytes(spec.ARCH)` rather than a literal; MLA is a Python-only…
5	N/A	This PR does not touch WMMA/RDNA GEMM, `GemmMultiD`, or `GemmMultiAbd`.
6	NEEDS-INFO	The split-KV workspace (`part_acc_ptr`, `part_ml_ptr`) does not require caller-zeroing because empty segments write `(NEG_INIT, 0)` to `part_ml` and the combine kernel skips reading `part_acc` when `l <= 0` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:235`), which is tested by poisoning the workspace with NaN in…
7	N/A	This PR does not touch Stream-K, persistent GEMM, or `GroupedGemmLauncher`.
8	N/A	This PR implements bf16 MLA decode and does not perform quantization or MX scale decoding.
9	N/A	This PR does not add or modify numeric emulation or quantized decode helpers (the truncation in `build_mla_decode` is truncated in…
10	N/A	This PR does not add or modify any manifest `kind` or touch `helpers/manifest.py`.
11	PASS	Launches go through `KernelLauncher` (`dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:111`), workspace buffers are cached on the runner instance…
12	PASS	`get_device_max_grid_size` in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/hip_module.py:351` queries `_hipDeviceGetAttribute` via ctypes without importing torch or loading a second LLVM.
13	PASS	The host-side constant `MAX_GRID_DIM = 65535` in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264` is value-identical to the previous `MAX_GRID_DIM_Z = 65535` in…
14	PASS	No `expected.json` files or `tflops_lower_bound` values are added, `benchmark_mla_decode_live.py:30` explicitly avoids recording absolute performance numbers, and `MAX_GRID_DIM` in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264` is…

§Q graded 7/7

#	Verdict	Evidence
1	NEEDS-INFO	The diffs for dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt and dnn-providers/hip-kernel-provider/rocke/platform/tests/runtime/test_device_max_grid_size.py are omitted, so CMake registration and installed-layout test portability cannot be verified from the provided diff.
2	PASS	MLA kernels live in library/kernels/mla/ (Python-only builders, no platform/cpp/instances mirror) and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:212 (test_mla_decode_cpp_python_byte_identity) lowers all 14 cases across gfx942 and…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:242 fails on BackendCoverageGap or zero compared cases, dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:31 fails when ROCKE_REQUIRE_COMPILER_TOOLS=1 and tools…
4	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:130 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:193 record and verify golden hashes across every served flavor in LLVM_FLAVORS (llvm20 is explicitly…
5	NEEDS-INFO	While visible tests such as dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_numeric.py:28 and dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:328 are portable across devices and flavors without torch or a GPU, the diffs for dnn-providers/hip-kernel-provider/rocke/platform/tests/runtime/test_device_max_grid_size.py and…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:31 preserves the exact AssertionError vs unittest.SkipTest semantics when factoring out the resource-gate helper, and…
7	N/A	This PR does not modify cross-engine backend resolution or ROCKE_BACKEND / ROCKE_CPP_STRICT rejection contracts.

§R graded 11/11

#	Verdict	Evidence
1	N/A	This PR is a functional bring-up of a previously missing operator regime (MLA decode, `max_seqlen_q == 1`) with no prior in-tree…
2	PASS	Arch-specific recipes (`SCORE_K = 32` for `mfma_f32_16x16x32_bf16`, `PV_TRANSPOSE_READ = True` for `ds_read_tr16_b64`, and per-arch `DECODE_SCORE_PREFETCH`/`DECODE_SCORE_FENCE_EVERY`) are gated by `spec.ARCH` in…
3	PASS	The kernel follows the utilities skills for LDS padding (`WT_PAD` vs `V8_PAD` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:252`), read-ahead prefetching (`read_ahead` with `sched_barrier` in…
4	NEEDS-INFO	The PR body and `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:135` cite `TestResources` in `dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py` for static VGPR, scratch, and ISA verification, but `test_mla_decode.py` (+730 lines) is omitted from the truncated diff.
5	NEEDS-INFO	Global memory accesses are added in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:573-664` (and in the truncated epilogue after line 800) and `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:209-275`, but the epilogue is truncated and no `analyze_coalescing` / `assert_asm_backed` report is present in the visible diff.
6	PASS	The kernels are built with raw `IRBuilder` outside `rocke.helpers.tiling` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:505`) and have exact/fp64-reference and GPU-oracle verification tests covering the LDS/MMA layouts…
7	N/A	The PR makes no bottleneck classification claims and states in the description that this is functional bring-up with no…
8	PASS	`b.sched_barrier(0)` and `DECODE_SCORE_PREFETCH` / `DECODE_SCORE_FENCE_EVERY` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:124-145` and `725-731` document the exact hot-loop shape and register-spill rationale per arch, and `b.sync()`…
9	N/A	This PR is not a performance regression fix and does not retune an existing kernel across shapes.
10	PASS	Compiler hazards around `agpr_alloc` on `llvm20` are gated via `flavor_honors_agpr_alloc` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:154-169`), pointer alias semantics (`noalias`, `readonly`, `writeonly`, `align`) are set on kernel…
11	N/A	No new optimization failure mode, heuristic study, or `dsl_probes` tool is introduced, and `measured_results.md` is untouched.

§S graded 8/8

#	Verdict	Evidence
1	NEEDS-INFO	The diffs for `platform/CMakeLists.txt`, `platform/tests/runtime/test_device_max_grid_size.py`, `library/tests/test_mla_decode.py`, `library/tests/test_mla_decode_runner.py`, and `library/tests/test_mla_decode_reference.py` are omitted from the prompt budget, so we cannot verify all new symbols' direct unit tests and CMake CI lane registrations from the visible diff.
2	PASS	`MLARequest` in `library/dispatch/mla/common.py:137` adds required `max_seqlen_q: int` and changes `op` default to `MLA_OP` (`mla_fwd`); all callers and tests are updated in `library/benchmarks/common/mla_shapes.py:130`…
3	PASS	`dispatch.mla.gfx942._resolve_num_cus` was moved/renamed to `dispatch.mla.common.resolve_num_cus` (`library/dispatch/mla/common.py:346`) and all callers/tests were updated (`library/dispatch/mla/gfx942.py:68`…
4	PASS	`MAX_GRID_DIM` is defined in `platform/python/rocke/helpers/spec.py:264` and added to `__all__` in `platform/python/rocke/helpers/spec.py:33` and `platform/python/rocke/helpers/__init__.py:694`; `_HIP_ATTR_MAX_GRID_DIMS` and `_max_grid_size` in…
5	NEEDS-INFO	While `library/tests/dispatch/mla/test_mla_decode_wiring.py:125-237` and `library/tests/test_mla_decode_golden.py:193-209` assert exact IR hashes and negative `ValueError` matches, the diffs for `platform/tests/runtime/test_device_max_grid_size.py`, `library/tests/test_mla_decode.py`, and `library/tests/test_mla_decode_runner.py` are omitted, so negative tests and behavioral assertions for…
6	PASS	CPU-checkable contracts (IR golden hashes, C++/Python byte-identity, spec/request validation, and torch-free device grid limit checks) run on CPU without torch or a GPU in `library/tests/test_mla_decode_golden.py:193-246` and…
7	PASS	`library/tests/test_mla_decode_golden.py:212-246` (`test_mla_decode_cpp_python_byte_identity`) lowers all 14 MLA decode cases through both `backend="python"` and `backend="cpp"` with `ROCKE_CPP_STRICT=1` across all served LLVM flavors and asserts `py == cpp`.
8	FAIL	The PR body's 'Test plan' section lists pytest/suite counts and GPU archs (`gfx942`, `gfx950`), and `platform/python/rocke/helpers/spec.py:264` / `platform/python/rocke/runtime/hip_module.py:351` do not break README usage, but the PR description does not state that it ran the README entry points `PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py` and `python tests/run_all.py` from `rocke/platform/`.

yraparti · 2026-11-03T16:16:42Z

+    num_warps: int = 4
+    dtype: str = "bf16"
+
+    def __post_init__(self) -> None:


[HIGH] §B.4 ★ ★ At a trust boundary, REJECT — never NORMALIZE into validity. Masking (& 0xFF), clamping, coercion or getattr(x, field, default) on a meaningful field rewrites malformed input into DIFFERENT valid input, turning a loud failure into a silent wrong answer. Name what each such expression…

dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:194-224 and dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:81-91 check value <= 0 on integer fields (num_heads, num_seg, etc.) without checking type(value) is int; for direct-spec callers bypassing dispatch, MlaDecodeSpec(num_heads=64, num_seg=True) or num_seg=2.0 is accepted (True > 0 is False for split_kv, 2.0 > 1 enables the split-KV ABI).

Fix: In MlaDecodeSpec.__post_init__ (mla_decode.py:201-214) and MlaDecodeCombineSpec.__post_init__ (mla_decode_combine.py:86-91), validate if type(value) is not int: (including num_warps) before range checks and raise ValueError.

yraparti · 2026-11-03T16:16:42Z

+
+def _tools_required() -> bool:
+    """Whether this run supplies the compiler tools, and so requires them."""
+    return os.environ.get(REQUIRE_ENV) == "1"


[HIGH] §D.1 ★ ★ Knobs/features on the KernelSpec (validator + selector + name tag + cache-key), NOT os.environ (env only as a selector-level diagnostic override folded onto the spec). Reviewers may reject env vars outright — even a reproducibility/pin one (#9739: "please dont use env variables", on a…

dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:28 reads the environment

Fix: Put the knob on the KernelSpec and delete the environment read.

yraparti · 2026-11-03T16:16:42Z

+            "scale": float(scale),
+            "num_seqs": int(self.request.num_seqs),
+            # The kernel's row pitch: the stride, which a column-sliced view
+            # does not share with its shape.
+            "block_table_stride": int(block_table.stride(0)),
+            "max_k": int(self.request.max_seqlen_k),
+        }
+        config = LaunchConfig(grid=self.grid, block=self.block, stream=stream)
+        if self._combine is None:
+            vals.update(out_ptr=out, lse_ptr=lse)
+            return [(self._latent, vals, config)], out, lse
+        part_acc, part_ml = self.workspace(qa.device)
+        vals.update(part_acc_ptr=part_acc, part_ml_ptr=part_ml, num_seg=spec.num_seg)
+        combine_vals = {
+            "out_ptr": out,
+            "lse_ptr": lse,
+            "part_acc_ptr": part_acc,
+            "part_ml_ptr": part_ml,
+            "num_seg": spec.num_seg,
+        }
+        combine_config = LaunchConfig(
+            grid=self.launch.combine_grid(),
+            block=self.launch.combine_block(),
+            stream=stream,
+        )
+        return (
+            [
+                (self._latent, vals, config),
+                (self._combine, combine_vals, combine_config),
+            ],
+            out,
+            lse,
+        )
+
+    def latent(
+        self,
+        *,
+        qa: torch.Tensor,
+        c_kv: torch.Tensor,
+        k_rope: torch.Tensor,
+        cu_seqlens_q: torch.Tensor,
+        cu_seqlens_k: torch.Tensor,
+        block_table: torch.Tensor,
+        scale: float,
+        out: torch.Tensor | None = None,
+        lse: torch.Tensor | None = None,
+        stream: int = 0,
+        check_contents: bool = False,


2 checklist findings on this snippet

[HIGH] §H.1 ★ (line 424) ★ A hardware/resource bound derived from SPEC fields (e.g. an AMDGPU buffer-SRD num_records sized from num_kv_blocks*block_size*num_kv_heads*head_size) MUST be matched by the launch tensors — host-validate their shape/size against the spec that sizes the bound BEFORE compile/launch, else the…

dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:424

Fix: Change check_contents: bool = False to check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward (leaving bind_latent with check_contents=False for the hot/timed path).

[MEDIUM] §E.15 (line 377) No dead parameters — and a parameter ACCEPTED then IGNORED is worse than one absent: the caller believes it was used, and in measurement code it silently discards a real observation. Grep each parameter name in the body. (#12070: substitute_share(total_armA_us=…) never read it — the measured…

dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:591-593 declares scalar kernel parameters num_seqs and max_k (b.param("num_seqs", I32) and b.param("max_k", I32)) whose return handles are discarded and never read in the kernel body, while dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:377-381 passes req.num_seqs and req.max_seqlen_k into them on every launch.

Fix: Remove the unused num_seqs and max_k scalar parameters from mla_decode_signature and build_mla_decode (and pass them only in run_prefill_oracle when invoking the prefill kernel), or use them in-kernel.

yraparti · 2026-11-03T16:16:42Z

+        b.ret()
+
+    cu_k_start = b.global_load_i32(cu_k, seq_idx)
+    cu_k_stop = b.global_load_i32(cu_k, b.add(seq_idx, b.const_i32(1)))
+    s_k = b.sub(cu_k_stop, cu_k_start)
+    n_k_tiles = b.div(b.add(s_k, b.const_i32(BK - 1)), b.const_i32(BK))
+
+    def _min(x, y):
+        return b.select(b.cmp_lt(x, y), x, y)
+
+    if SPLIT:
+        seg = b.block_id_z()
+        per_seg = b.div(b.add(n_k_tiles, b.sub(num_seg_p, b.const_i32(1))), num_seg_p)
+        t_lo = _min(b.mul(seg, per_seg), n_k_tiles)
+        t_hi = _min(b.add(t_lo, per_seg), n_k_tiles)
+    else:
+        t_lo = b.const_i32(0)
+        t_hi = n_k_tiles
+    # Tiles wholly inside S_k run unmasked; at most one tail tile follows.
+    t_full = _min(b.div(s_k, b.const_i32(BK)), t_hi)
+    t_full = b.select(b.cmp_lt(t_full, t_lo), t_lo, t_full)
+
+    # ---- absorbed query, read straight into the score's B layout ------------
+    # SW columns at SK t + lane_ks of head lane % 16: exactly the B operand of
+    # score k-step t, so there is no transform at all.
+    q_base = b.add(
+        b.mul(b.add(b.mul(token, b.const_i32(H)), head), b.const_i32(QA_COLS)),
+        lane_ks,
+    )
+    qa_frags = [
+        b.global_load_vN(qa_ptr, shift(b, t * SK, q_base), dtype, SW)
+        for t in range(N_LAT_STEPS)
+    ]
+    rope_frags = [
+        b.global_load_vN(qa_ptr, shift(b, R_KV + j * SK, q_base), dtype, SW)
+        for j in range(ROPE_K_ITERS)
+    ]
+    # Q_rope is parked in LDS and re-read every key tile, as in the prefill:
+    # the VGPRs it would pin are what keeps the loop inside 256. Each lane
+    # reads back exactly what it wrote, so no barrier is involved.
+    qr_lds = b.smem_alloc(dtype, [BLOCK_H, D_ROPE + PAD], name_hint="qr_lds")
+    qr_row = b.add(b.mul(wave, b.const_i32(HPW)), lane_row)
+    for j in range(ROPE_K_ITERS):
+        b.smem_store_vN(qr_lds, [qr_row, shift(b, j * SK, lane_ks)], rope_frags[j], SW)
+
+    scale_log2 = b.fmul(scale_p, b.const_f32(LOG2E))
+    neg_init = b.const_f32(NEG_INIT)
+    zero_f = b.const_f32(0.0)
+
+    iter_args = [("m", neg_init), ("l", zero_f)]
+    iter_args += [(f"acc{t}", b.zero_vec_f32(4)) for t in range(N_R_TILES)]
+
+    if TR_PV:
+        # ds_read_tr16_b64 lane addressing for a K = 16 (one key tile) operand.
+        tr = TransposeLdsReader(K=MFMA_K).bind(b, lane)
+    else:
+        lane_row4 = b.mul(lane_row, b.const_i32(4))
+    kv_lds = b.smem_alloc(dtype, [BK, QA_COLS + PAD], name_hint="kv_lds")
+
+    c_chunks = (BK * R_KV) // (THREADS * C_STAGE_W)
+    c_chunks_per_row = R_KV // C_STAGE_W
+    KR_W = _kr_stage_width(spec)


2 checklist findings on this snippet

[HIGH] §H.2 ★ (line 678) ★ Guard indirection VALUES, not just indices — a masked/bounds-guarded table read protects the lookup INDEX (e.g. logical page < n_pages), NOT the value it returns. Any value used as a memory index (block-table physical id) must itself be range-checked, or shown contained by a HW bounds-check —…

dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:678

Fix: Default check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward (or range-check page against num_pages on device before b.zext(b.readfirstlane(page), I64)), because folding page64 * stride into the 64-bit base pointer bypasses num_records and check_contents is currently off by default.

[LOW] §H.3 (line 617) Hot-loop div/mod by a compile-time constant lowers to shift/mask ONLY when the constant is a power of two (else a magic-number mul+shift). AMDGPU has no integer-divide so it is never a runtime divide, but enforce power-of-two in the spec (production page/tile sizes are 16/32) or benchmark the…

dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:617

Fix: Enforce spec.block_k & (spec.block_k - 1) == 0 in supports_mla_decode (or MlaDecodeSpec.__post_init__), since b.div(..., b.const_i32(BK)) is used in the kernel and supports_mla_decode currently only checks spec.block_k % MFMA_N == 0.

yraparti · 2026-11-03T16:16:42Z

+    """Whether the build box's device is ``arch`` (base arch, target-ID
+    features ignored): only then may dispatch read a live device property, so
+    a cross-compile never takes another part's."""
+    try:


[HIGH] §O.11 ★ ★ Hardware decisions query ArchTarget (from_gfx, fits_lds, lds_capacity_bytes, supports_dtype_combo, async_lds_max_dwords, has_*, the MMA catalog), never gfx.startswith(…) or string family compares. Resolve the gfx string once, never mutate the lru_cached ArchTarget or…

The new live-device architecture check compares target strings directly with split(':') in dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:337 instead of going through ArchTarget/base-arch helpers.

Fix: Replace the string split comparison in _live_device_is with the existing target-id/base-arch resolution helper or ArchTarget-based comparison, and use the same resolved base arch for both the live device and request.

yraparti · 2026-11-03T16:16:42Z

+def _time(
+    shapes, *, arch: str, warmup: int, iters: int, repeats: int, seed: int
+) -> int:
+    import torch


[MEDIUM] §E.2 Library doesn't import torch for device queries — use rocke.runtime.hip_module (prefer hipDeviceGetAttribute).

dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:169 adds import torch

Fix: Query the device through rocke.runtime.hip_module instead of importing torch.

yraparti · 2026-11-03T16:16:42Z

+_max_grid_size: Dict[int, Optional[Tuple[int, int, int]]] = {}
+
+
+def get_device_max_grid_size(device: int = 0) -> tuple[int, int, int] | None:


[LOW] §S.8 The PR states it ran the README's entry points from rocke/platform/: PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py (relative-path guard, cross-engine byte-identity, pytest, and ctest when built); a manifest or runtime change…

The PR body's 'Test plan' section lists pytest/suite counts and GPU archs (gfx942, gfx950), and platform/python/rocke/helpers/spec.py:264 / platform/python/rocke/runtime/hip_module.py:351 do not break README usage, but the PR description does not state that it ran the README entry points PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py from rocke/platform/.

Fix: Run PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py from rocke/platform/ and record them in the PR test plan.

…k-1504-mla-decode

Review follow-up for the MLA decode PR (checklist review of the previous head). Integer fields. MlaDecodeSpec and MlaDecodeCombineSpec checked their counts with "<= 0" only, so a direct caller could pass True or 2.0: True passed the check, and num_seg=2.0 picked the split-KV variant. Both specs now reject any count whose type is not exactly int, num_warps included, before the range checks. Power-of-two block_k. The kernel divides by block_k. Dispatch only admits a page size of 16, so the IR does not change, but a direct spec with another block_k is now refused instead of compiling a magic-number divide. Base arch. _live_device_is stripped target-ID features by splitting on ":"; it now resolves both the live device and the request through base_arch_from_target_id, as does the decode benchmark's default --arch. A device that reports no name supplies nothing. Tests cover a bool and a float for each spec, a non-power-of-two block_k, and an unnamed live device. Decode IR goldens and C++/Python parity are unchanged at llvm22, llvm23 and llvm24. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…k-1504-mla-decode

aosewski · 2026-11-04T14:30:15Z

Re the rocKE checklist review of c41875b. Develop is merged; fixes are in 934f5bb.

§B.4 fixed: both decode specs reject counts that are not exactly int (True, 2.0), num_warps included.
§O.11 fixed: _live_device_is and the benchmark's default --arch use base_arch_from_target_id.
§H.3 fixed: block_k must be a power of two (IR unchanged; dispatch admits only 16).
§H.1/§H.2 kept: page ids are caller input; check_contents checks them on the host and stays opt-in because it syncs. The verify harness always enables it.
§E.15 deferred: dropping num_seqs/max_k changes the signature (golden re-bless + GPU runs on both arches); it goes into the follow-up that moves decode onto _mla_body.
§D.1 not applicable: ROCKE_REQUIRE_COMPILER_TOOLS is the CI lane's skip-to-fail switch (see TESTING.md), not a kernel knob.
§E.2 kept: the benchmark times the whole op on torch tensors; the import is lazy and the arch comes from hip_module.
§S.8 done: tests/test_rocke.py and tests/run_all.py results are in the PR test plan.

cgmillette · 2026-11-04T14:45:03Z

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

PR #13923 — add MLA decode kernel for gfx942 and gfx950 (aosewski, AICK-1504) — 🟠 Follow-up: flavor rule; duplicate gate helper

Prior item	Status
Copilot (10-19) ×7: block-table i32 bounds; `cu_seqlens_q` span; projection shapes; `max_seqlen_q` / `num_seg` truncation; non-ROCm torch; `HipError`	🟢 FIXED, unchanged • fix sites untouched by `c41875b4`, `934f5bb4` (read: diff) • runner tests with torch: 43 ✓ at head (ran); CI skips them (no torch) • 7/7 threads resolved
CongMa13 (10-22) APPROVED at `fa8e6110`	🟢 APPROVED `fa8e6110` • older head; still counts on develop
cgmillette (10-28) APPROVED at `87fb6ff2`	🟢 APPROVED `87fb6ff2` • older head; still counts on develop • since: llvm24 serving, review refactors, spec checks (author's)
Our 11-02 N1: decode re-inlines MLA family emitters	🟡 DEFERRED, mostly • author: tail needs re-bless + device runs; `_mla_body` still moving: holds, #14502 reworks it (read) • `lane_ids`, `softmax_consts`, LSE helper are IR-identical: land them now (ran) • file a ticket under AICK-1504 for the rest, after #14502; note it in `known_gaps.md`
Our 11-02 N2–N4: flavor list; gate rule copy; selection reads device	🟢 FIXED `c41875b4` • one flavor rule; a new served flavor fails the golden (ran: mutation); see N1 • one gate helper; undoing its fail branch fails a test (ran: mutation); see N2 • selection reads no device; launch check is torch-free, ran in CI (read)
yraparti (11-03) §B.4, §H.3, §O.11: loose ints; `block_k`; arch split	🟢 FIXED `934f5bb4` · threads open • `True`, `1.9`, `2.0` refused; new tests catch removal (ran: mutation); see N3 • live-device check uses `base_arch_from_target_id`; same answers (ran) • `split(":")` left: `verify_mla_decode.py:660`, both kernel gates, numeric test (Rule 9)
yraparti (11-03) §H.1, §H.2: page ids unchecked unless `check_contents`	🟢 Answered, holds · threads open • host check opt-in (it syncs); verify harness opts in (read) • family contract: develop's prefill trusts page ids too (read)
yraparti (11-03) §E.15: unused `num_seqs`, `max_k` kernel args	🟡 DEFERRED · thread open • author: dropping them changes the signature, needs re-bless + GPU runs • holds: pack mirrors prefill's oracle kernel, `mla_decode.py:405` (read)
yraparti (11-03) §D.1, §E.2: env read in test gate; benchmark imports torch	🟢 N/A · threads open • repo Rule 3 exempts test-harness env reads (read) • benchmark needs torch for tensors and streams, not device queries (read)
yraparti (11-03) §S.8: README entry points not in test plan	🟢 ADDRESSED (body) · thread open • `test_rocke.py`, `run_all.py` results now recorded (read) • 150 local failures called host-only; CI `rocke_pytest` green at `c41875b4` (read)
Bots / CI	🟡 Running at `047d7fde` • running: math-libs builds, hipDNN Linux; 22 ✓ · 4 running · 0 ✗ • `c41875b4` green: MLA lane 469 ✓ / 3 skip, gate 4/4; numerics skip, no torch (read: CI log) • merge: blocked: 7 threads open (answered in a comment); no rocke-lib owner approval
Our checks	🟢 Clean at `047d7fde` • merge vs `77098d63`: clean, 3 behind; own delta since our checks: `934f5bb4` (ran) • build + ctest 45 ✓ / 1 skip; parity, goldens ✓ ×4 flavors; installed 14/14 files (ran at `c41875b4`; carries) • MLA lane at llvm24, tools required: 480 ✓ / 3 skip; C++ parity ✓ (ran) • rules: 5 size hits are shape math or a tested budget; 2 encoding hits, ASCII (ran)
Code quality	🟠 Should-fix • flavor rule keys on datalayout; new lane refused silently: N1 (ran) • gate helper is a third home; owner #14218: N2 (read) • our 11-02 N2–N4 fixed (ran); N1 deferral sound once tracked (read) • minor: N3; runner's grid check repeats dispatch's `_check_grids` (read)

#	Sev	file:line	Finding
N1	🟠	`kernels/mla/_mla_body.py:168`	The one flavor rule asks the datalayout, not the compiler lane • impact: a new-generation lane is refused with every test green (ran) • `c41875b4` also dropped the unserved-lane guard test (read) • fix: `_flavor_at_least(flavor, "llvm22")`: same lanes today (ran)
N2	🟠	`tests/_mla_resource_gate.py:31`	Third home for `require_tools_or_skip`; #14218 owns it • impact: copies drift: a refused option skips here (`tests/test_mla_decode.py:730`), fails there (ran) • our 11-02 N3 fix named this module; we missed #14218's • fix: this PR lands first: take the owner's names (`tests/_compiler_test_utils.py`, `REQUIRE_COMPILER_TOOLS_ENV`); send that case through it
N3	🟡	`kernels/mla/mla_decode.py:218`	New int checks raise `ValueError`; the family raises `TypeError` • impact: one mistake, two exception types, three texts across MLA specs (ran) • fix: one `_mla_body` check raising `TypeError` (PYTHON_STYLE §13.3); drop `with_num_seg`'s

Utility: 🟢 Positive — adds MLA decode on gfx942/gfx950; flavor rule and gate helper need rework.

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: the delta since our last review (c41875b4, 934f5bb4, two develop merges), merge with develop, build + ctest, goldens, C++ parity, the MLA lane at LLVM 24, runner tests with torch, probes and mutations for each fix, a simulated new compiler lane, CI logs at c41875b4, code quality. Flags whether a closer look is needed; not an approval or rejection.

cgmillette · 2026-11-04T15:47:49Z

Automated overlap notice (rocKE merge monitor)

Hi @aosewski, this PR, #14218 and #14502 each move develop's private MLA skip-or-fail helper (_require_tools_or_skip, tests/test_mla_prefill_gfx942.py:43 on develop) into a test module of their own. Row N2 of our 4 November review has the evidence. Paths are under dnn-providers/hip-kernel-provider/rocke/, in library/ unless they start with platform/:

Here: require_tools_or_skip in tests/_mla_resource_gate.py:31-40, with REQUIRE_ENV (:23), imported at tests/test_mla_decode.py:27 and tests/test_mla_prefill_gfx942.py:23.
fix(rocke): never pass a compiler option the loaded compiler does not know #14218: the same function in tests/_compiler_test_utils.py:44-49, with REQUIRE_COMPILER_TOOLS_ENV (:41) and skip_where_compiler_refuses_option (:52-64), which the attention backward tests use too. refactor(rocke): MLA prefill on shared emitters; odd head counts on the packed layout #14502's copy (tests/_compiler_tools.py:44-48) has its own notice.

The two PRs already treat one case differently: a compile option the loaded comgr refuses. Your decode gate skips it even under ROCKE_REQUIRE_COMPILER_TOOLS=1 (tests/test_mla_decode.py:562-570, truth-table row :730); #14218 fails it there and skips it otherwise (tests/test_compiler_test_utils.py:68-84). Decode compiles with such an option: its agpr_alloc = (0, 0) (kernels/mla/mla_decode.py:120) adds -mllvm -amdgpu-mfma-vgpr-form.

Owner: #14218. Its module is the shared home of the compile-test skip policy, where this PR's is named and documented for the two MLA gates, and #14218 wrote it first (2 November; yours came on 3 November). Resolution: the merge plan lands this PR first, so take the owner's shape now, and whichever PR lands second keeps one module:

git mv tests/_mla_resource_gate.py tests/_compiler_test_utils.py, holding fix(rocke): never pass a compiler option the loaded compiler does not know #14218's text without what needs its option table (skip_where_compiler_refuses_option and its imports).
Import it as fix(rocke): never pass a compiler option the loaded compiler does not know #14218 and develop's test helpers do, after the absolute imports: from ._compiler_test_utils import REQUIRE_COMPILER_TOOLS_ENV, require_tools_or_skip in test_mla_decode.py (renaming REQUIRE_ENV), and require_tools_or_skip alone in test_mla_prefill_gfx942.py.
Take the owner's rule: return "tools", not "skip", at test_mla_decode.py:570, so the refusal goes through require_tools_or_skip. test_comgr_rejecting_a_compile_option_is_a_skip (:654) and row :730 then expect the tools verdict and the failure; update the docstrings at :545-553, :587-595 and :625-631 to match. CI's compiler accepts the option, so CI sees no change.

Suggested order: this PR first (wave 1 of the merge plan), then #14218 (wave 2). #14218 then rebases onto it: it keeps its own _compiler_test_utils.py (the add/add conflict resolves to its file), resolves test_mla_prefill_gfx942.py and platform/CMakeLists.txt, and does step 4 below, because its compile path refuses the option before comgr sees it. Nothing more for this PR. If #14218 lands first, this PR drops its copy and imports the owner's.

Rebase, if #14218 lands first (checked on develop 961dd76390 merged with #14218 at 45876faa7f), in your merge of develop, from dnn-providers/hip-kernel-provider/rocke/library:

git checkout origin/develop -- tests/_compiler_test_utils.py.
tests/test_mla_prefill_gfx942.py: take develop's side of the imports and of _OccupancyGuard.setUpClass; keep your flavor_honors_agpr_alloc check in TestAgprAccResources.
platform/CMakeLists.txt, rocke_library_mla_pytest: keep your six decode files and develop's test_compiler_test_utils.py, and merge the two comments.
In _RESOURCE_PROBE, import CompilerOptionNotAccepted from rocke.runtime beside compile_kernel, and after except ImportError (test_mla_decode.py:510-512) add an except CompilerOptionNotAccepted as exc: branch that prints f"loaded compiler cannot build this kernel: {exc}" and exits 78.

From your current head instead, make the test_mla_decode.py edits above, delete tests/_mla_resource_gate.py rather than moving it, and in step 2 also delete the old import the merge keeps at test_mla_prefill_gfx942.py:23.

On develop 961dd76390 with this PR and the edits above, test_mla_decode.py and test_mla_prefill_gfx942.py passed (156 tests) with and without ROCKE_REQUIRE_COMPILER_TOOLS=1; with an unknown -mllvm option injected into the decode probe, the gate failed under the variable and skipped without it (at your head it skipped both ways). With #14218 merged first and steps 1-4, those files and test_compiler_test_utils.py passed (160) both ways; with a simulated LLVM 20 compiler the gate failed under the variable and skipped without it, and without step 4 it failed both ways. Merged the other way, #14218 onto your edits, the same three files conflict, and with step 4 the rest is the same tree.

If you think this PR should own it instead, reply here and on #14218, and the reviews will follow the new decision.

Automated overlap check: one owner per shared piece of code, coordinated on the PRs. Nothing else changed.

yraparti

rocKE checklist re-review of the fixes since `c41875b42c05`

ROCm/rocm-libraries#13923 · 047d7fde0a25 · model gemini-4-argon, fallback gpt-5.5

PASS · graded 238/238 · BLOCKER 0, HIGH 3, MEDIUM 2, LOW 1, NEEDS-INFO 34

The diff was truncated before grading. Items that needed a missing hunk are NEEDS-INFO.

Since the last review (`c41875b42c05`)

Resolved 3 · still open 6 · new 0

✅ §B.4 was FAIL, now PASS — had asked: In MlaDecodeSpec.__post_init__ (mla_decode.py:201-214) and MlaDecodeCombineSpec.__post_init__ (mla_decode_combine.py:86-91), validate if type(value) is not int: (including num_warps) before range checks and raise ValueError.
✅ §H.3 was FAIL, now PASS — had asked: Enforce spec.block_k & (spec.block_k - 1) == 0 in supports_mla_decode (or MlaDecodeSpec.__post_init__), since b.div(..., b.const_i32(BK)) is used in the kernel and supports_mla_decode currently only checks spec.block_k % MFMA_N == 0.
✅ §O.11 was FAIL, now PASS — had asked: Replace the string split comparison in _live_device_is with the existing target-id/base-arch resolution helper or ArchTarget-based comparison, and use the same resolved base arch for both the live device and request.
⏳ §D.1 still FAIL (HIGH): Put the knob on the KernelSpec and delete the environment read.
⏳ §E.2 still FAIL (MEDIUM): Query the device through rocke.runtime.hip_module instead of importing torch.
⏳ §E.15 still FAIL (MEDIUM): Remove the unused num_seqs and max_k scalar parameters from mla_decode_signature and build_mla_decode (and pass them only in run_prefill_oracle when invoking the prefill kernel), or use them in-kernel.
⏳ §H.1 still FAIL (HIGH): Change check_contents: bool = False to check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward (keeping check_contents=False in bind_latent).
⏳ §H.2 still FAIL (HIGH): Default check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward, or range-check page against num_pages on device before forming the 64-bit pointer offset.
⏳ §S.8 still FAIL (LOW): Run PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py from rocke/platform/ and record them in the PR test plan.

Findings

HIGH §D · dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:28 · Put the knob on the KernelSpec and delete the environment read.
HIGH §H · dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:424 · Change check_contents: bool = False to check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward (keeping check_contents=False in bind_latent).
HIGH §H · dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:686 · Default check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward, or range-check page against num_pages on device before forming the 64-bit pointer offset.
MEDIUM §E · dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:169 · Query the device through rocke.runtime.hip_module instead of importing torch.
MEDIUM §E · dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:599 · Remove the unused num_seqs and max_k scalar parameters from mla_decode_signature and build_mla_decode (and pass them only in run_prefill_oracle when invoking the prefill kernel), or use them in-kernel.
LOW §S · platform/python/rocke/helpers/spec.py:264 · Run PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py from rocke/platform/ and record them in the PR test plan.

Needs info

§A.5 The diff for dnn-providers/hip-kernel-provider/rocke/library/tests/golden/mla_prefill_gfx942_ir_sha256.json (+2 -2 lines) is still omitted from the prompt, so the exact change in the existing prefill golden fixture cannot be verified by line number.
§C.9 Bug fixes and gates are applied across sibling sites: flavor_honors_agpr_alloc unifies decode and prefill (dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:266, dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_prefill_gfx942.py:748), exact type(value) is not int validation is applied in both MlaDecodeSpec (dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:216) and MlaDecodeCombineSpec (dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:90), and require_tools_or_skip unifies both resource gates (dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:31). — no…
§D.11 mla_decode_gfx942 and mla_decode_gfx950 are registered for the disjoint max_seqlen_q == 1 decode regime where no prior MLA candidate existed (dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:296-300, dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:327). — no diff line supports this verdict
§E.7 Not resolved in the visible diff: the diffs for dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md, dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/attention.md, and dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/index.md are omitted from the prompt budget, and dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:23 introduces ROCKE_REQUIRE_COMPILER_TOOLS while dsl_docs/reference/env_flags.md is not in the PR file list.
§F.6 The new commits since the previous review do not resolve this item because local pre-commit execution before push is not recorded in the diff.
§F.7 The new commits since the previous review touch code and tests (e.g., dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:217), but GitHub bot/Copilot review thread states, won't-fix justifications, and deferred ticket IDs are not visible in the diff.
§F.11 The new commits since the previous review do not resolve this item because verifying the PR title/body via the GitHub API is an external action not visible in the diff.
§G.5 While dnn-providers/hip-kernel-provider/rocke/platform/tests/runtime/test_device_max_grid_size.py:1-79 is now visible in the diff, test_mla_decode.py (+750 -0), test_mla_decode_reference.py (+560 -0), and test_mla_decode_runner.py (+484 -0) remain omitted from the diff budget, so reachability and falsifiability of all tests in those files cannot be verified.

Section notes

§B: All 6 items pass; the previous FAIL on item 4 (integer type validation in MlaDecodeSpec.post_init and MlaDecodeCombineSpec.post_init) was resolved in commit 25.
§D: Re-review of items 1–11 in Section D. Item 1 was a false positive on a CTest test-tool gate (ROCKE_REQUIRE_COMPILER_TOOLS); all kernel specs, selectors, and builders are free of os.environ reads.
§D: Items 12-21 of Section D re-reviewed against head 047d7fd.
§G: Items 18–34 of Section G (Verification honesty).
§H: Trace of device-side masks & bounds in mla_decode.py: 1. Query early-exit: extent source = cu_seqlens_q[seq_idx..seq_idx+1] (token, q_stop); visited range = 1 query per sequence (token); agrees when q_lens[seq] in {0, 1}, host-validated via MLARequest.max_seqlen_q == 1 and check_contents. 2. Key tile loop bounds (t_lo, t_full, t_hi): extent source = s_k = cu_seqlens_k[seq_idx+1] - cu_seqlens_k[seq_idx], n_k_tiles = ceil(s_k / BK); visited range = [t_lo, t_hi) dynamically bounded by n_k_tiles (no compile-time Skv loop). Agrees with runtime s_k. 3. Page table lookup (_load_page): extent source = block_table[seq_idx * bt_stride + tile] with tile clamped to last_tile = max(0, t_hi - 1) and guarded by has_tiles = (t_lo < t_hi); visited range = [0, ceil(s_k / BK)); agrees with host check block_table.shape[1] >= ceil(max_seqlen_k / page_block_size) provided s_k <= max_seqlen_k (checked in check_contents). 4. Per-page buffer SRD (_stage_loads): extent source = C_PAGE_BYTES (PAGE * R_KV *…
§J: Shapes table: | doc/hunk | operands/inner dimensions/result | status | | Markdown hunks for new/changed README/ALGORITHM/DESIGN files | not visible in the provided diff | cannot type-check | Symbols table: | doc/hunk | symbol/domain definition | status | | Markdown hunks for new/changed README/ALGORITHM/DESIGN files | notation tables and domain bridges not visible | cannot verify | Novelty table: | doc/hunk | claim | status | | Markdown hunks for new/changed README/ALGORITHM/DESIGN files | added/new/shared claims not visible | cannot compare to develop |

Graded with the rocKE PR checklist, one model pass per section. FAIL rows that sit on a diff line are inline comments.

§A graded 14/14

#	Verdict	Evidence
1	PASS	The diff now includes `dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md:203`, which resolves the previous NEEDS-INFO by…
2	N/A	No core lowering-op or op-handler files (`core/lower_*`, `core/ir`) are…
3	PASS	Changes in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:1` and `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_d…
4	N/A	No C++ kernel builder or helper twin exists or is modified by this PR.
5	NEEDS-INFO	The diff for `dnn-providers/hip-kernel-provider/rocke/library/tests/golden/mla_prefill_gfx942_ir_sha256.json` (+2 -2 lines) is still omitted from the prompt, so the exact change in the existing prefill golden fixture cannot be verified by…
6	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:212` (`test_mla_decode_cpp_python_byte_identity`) explicitly tests…
7	PASS	Existing prefill kernel logic in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_prefill_gfx942.py:745` only refactors the host admission…
8	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:212` provides `test_mla_decode_cpp_python_byte_identity` and the PR test…
9	N/A	This PR does not contain a shared op-handler fix or lowerer-only arch…
10	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:119` defines golden hash cases across both arches (`gfx942`, `gfx950`)…
11	N/A	No new IR op is introduced in this PR.
12	N/A	`core/backend.py` and `ROCKE_BACKEND` fallback logic are not touched in…
13	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:138` lowers kernels via `_lower_kernel_to_llvm_python` for deterministic…
14	PASS	Changes in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264` and…

§B graded 6/6

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:81-87 documents the head-count scope (`DECODE_NUM_HEADS = (64, 128)`) with a TODO…
2	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:271-275 calls `spec_for(req)` inside `support(req)` and catches `ValueError` from…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:327 sets `ShapeRange("max_seqlen_q", min=2)` on prefill so `max_seqlen_q == 1` is…
4	PASS	Resolved since the previous review: dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:201-223 and…
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:87-216 scopes admission to `DECODE_NUM_HEADS = (64, 128)` and…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:311-381 (`supports_mla_decode`) and…

§C graded 10/10

#	Verdict	Evidence
1	PASS	Shared predicates and logic are extracted into single named helpers: `flavor_honors_agpr_alloc` in…
2	PASS	One reusable builder per kernel family (`build_mla_decode(spec)` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:520` and…
3	PASS	Builders and `kernel_name()` are named by operation and codegen knobs, never workload sizes (`num_seg` is excluded from `kernel_name()`, keeping only…
4	PASS	The shared MLA decode body in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:165` uses spec class attributes (`ARCH`…
5	PASS	Every touched file belongs directly to the MLA decode kernel, dispatch, verification, tests, or the shared `MAX_GRID_DIM` / `get_device_max_grid_size`…
6	PASS	Shared helpers (`resolve_num_cus` in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:351`, `flavor_honors_agpr_alloc` in…
7	N/A	This PR introduces a new MLA decode feature rather than performing a…
8	PASS	Reuses existing helpers from `develop` including `read_ahead` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:86`)…
9	NEEDS-INFO	Bug fixes and gates are applied across sibling sites: `flavor_honors_agpr_alloc` unifies decode and prefill (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:266`…
10	PASS	Widening `MLARequest` with `max_seqlen_q` and `MLA_OP = 'mla_fwd'` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:54`…

§D graded 21/21

#	Verdict	Evidence
1	FAIL	dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:28 reads the environment
2	PASS	`make_mla_decode_candidate` declares `Capability(arches=(arch,), ...)` where `arch = spec_cls.ARCH` (`gfx942` or `gfx950`), and `DECODE_NUM_HEADS = (64…
3	PASS	`support` in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:248-276` runs cheap request/selector/flavor checks before `spec_for`…
4	PASS	`decode_spec_for` and `_check_grids` take `arch` from `req` and only query device CU count when `req.num_cus == 0` via `resolve_num_cus(req)`…
5	PASS	`MlaDecodeSpec.identity()` removes launch-only `num_seg` and includes `split_kv` plus class emission attributes (`arch`, `score_k`, `pv_transpose_read`)…
6	PASS	Single-pass vs split-KV routing keeps `decode_spec_for`, `mla_decode_grid`, `mla_decode_signature`, `MlaDecodeLaunch`, and `MlaDecodeSpec.identity()` in…
7	N/A	This PR does not introduce machine-regenerable heuristics or hand-pinned…
8	PASS	The LLVM flavor cohort rule `flavor_honors_agpr_alloc` is defined once in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:154-169`…
9	PASS	`make_mla_decode_candidate` registers `mla_decode_gfx942` and `mla_decode_gfx950` on `MLA_REGISTRY` with `Capability`, `build`, `grid`, `block`, and…
10	PASS	The new `MLARequest.max_seqlen_q` field (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:137, 228`) and `MlaDecodeSpec.num_seg`…
11	NEEDS-INFO	`mla_decode_gfx942` and `mla_decode_gfx950` are registered for the disjoint `max_seqlen_q == 1` decode regime where no prior MLA candidate existed (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:296-300`…
12	PASS	Priority `DECODE_PRIORITY = 5` is justified against prefill (10) in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:73-75` and…
13	PASS	`dispatch_mla` delegates directly to `MLA_REGISTRY.select` without catching or retrying on `PinRefused` in…
14	PASS	`MlaDecodeSpec.identity()` omits the problem-dependent launch argument `num_seg` and keeps only `split_kv` and arch emission choices so `compile_key` does…
15	PASS	The PR intentionally partitions the MLA regime by `max_seqlen_q` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:325-327`), which…
16	PASS	The `mla` family exposes `registered_mla_combos`, `mla_sweep_space`, and `dispatch_mla_all` with `_request_errors` short-circuiting…
17	PASS	CPU routing tests in `dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:94-245` check cohort selection across…
18	PASS	`MLA_REGISTRY` keeps `require_build=True` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/__init__.py:48-50`) and no ratchets or exempt sets…
19	N/A	This PR adds MLA decode kernels and dispatch under `dispatch/mla/` and does…
20	PASS	The PR explicitly documents in `dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md:203-214` and…
21	N/A	This PR introduces a new MLA decode kernel and its dispatcher registration…

§E graded 17/17

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/hip_module.py:343 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/he…
2	FAIL	dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:169 adds `import torch`
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_numeric.py:23 gates torch with pytest.importorskip('torch'), and…
4	PASS	Layering is respected downward: kernels (dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:67) -> dispatch…
5	PASS	dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/hip_module.py:351 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/he…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:212 verifies C++/Python byte-identity and IR SHA256 goldens across all…
7	NEEDS-INFO	Not resolved in the visible diff: the diffs for dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md, dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/attention.md, and…
8	PASS	Visible diffs (e.g. dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:1 and…
9	PASS	All touched files in the PR diff (e.g. dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:1…
10	PASS	The PR body notes explicit coordination with #12464 and merge simulation with develop (with develop merged into the branch in the latest commits), and…
11	PASS	Only deliverable source, tests, shape JSONs, and required instance/example READMEs are committed (e.g.…
12	PASS	Both gfx942 and gfx950 are covered: dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:385 and…
13	PASS	Renamed/hoisted helpers (_resolve_num_cus -> resolve_num_cus in dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:351…
14	PASS	Removed lines in modified files (e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:95…
15	FAIL	Not resolved in the new commits: dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:599-601 still declares scalar kernel parameters `num_seqs` and `max_k` (`b.param("num_seqs", I32)` and `b.param("max_k", I32)`) whose return handles are discarded and never read in the kernel body, while…
16	N/A	This item applies to how review feedback is written rather than the code…
17	PASS	No absolute software-achieved performance figures appear in added lines (e.g. dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_ml…

§F graded 11/11

#	Verdict	Evidence
1	PASS	Branch `users/aosewski/aick-1504-mla-decode` uses the recognized `users/<user>/` prefix; code is added in…
2	PASS	title `feat(rocke): add MLA decode kernel for gfx942 and gfx950` matches type(scope): summary and is 56 characters
3	PASS	description is at least 30 characters and contains an issue reference
4	PASS	test file `dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py` is in the diff
5	PASS	All changed files in the diff are source, test, config, or doc files (e.g., `dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla…
6	NEEDS-INFO	The new commits since the previous review do not resolve this item because local `pre-commit` execution before push is not recorded in the diff.
7	NEEDS-INFO	The new commits since the previous review touch code and tests (e.g., `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:217`), but GitHub bot/Copilot review thread states, won't-fix justifications, and deferred ticket…
8	N/A	This PR targets `develop` directly from a `users/...` branch rather than…
9	N/A	The diff does not show any artificial splitting or trimming done to meet a…
10	PASS	The added and updated tests verify real invariants, such as spec field validation in `dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/tes…
11	NEEDS-INFO	The new commits since the previous review do not resolve this item because verifying the PR title/body via the GitHub API is an external action not visible in the diff.

§G graded 34/34

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:211-245 checks C++/Python lowering byte-identity separately from on-GPU…
2	PASS	dnn-providers/hip-kernel-provider/rocke/library/builders/mla/verify_mla_decode.py:116-124 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:130-136 and 193-245 gate all served modern LLVM flavors (llvm22, llvm23…
4	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:94-243 and 349-560 add behavioral tests for decode routing…
5	NEEDS-INFO	While dnn-providers/hip-kernel-provider/rocke/platform/tests/runtime/test_device_max_grid_size.py:1-79 is now visible in the diff, test_mla_decode.py (+750 -0), test_mla_decode_reference.py (+560 -0), and test_mla_decode_runner.py (+484 -0)…
6	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md (+295 -0), dnn-providers/hip-kernel-provider/rocke/library/examples/gfx950/mla_decode/README.md (+129 -0), and…
7	N/A	This PR does not add or modify a design doc (`DESIGN.md` is not in the PR's…
8	PASS	dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:195-210 and 539 loudly raise ValueError on non-bf16 tensors, and…
9	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py (+750 -0, which contains TestResources per mla_decode.py:124) and the example READMEs remain omitted from the diff budget, so tool-sourced VGPR/LDS assertions cannot be…
10	N/A	The PR explicitly makes no performance claims and is a functional bring-up.
11	PASS	dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:79 and…
12	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:211-245 verifies C++ and Python lowering parity across all 14 decode cases…
13	N/A	The PR uses existing DSL builder methods (`IRBuilder`) and does not…
14	N/A	The PR is a functional bring-up with no bottleneck classification or…
15	PASS	All graded files and hunks come directly from the provided PR diff for commit `047d7fd` (e.g.…
16	N/A	This is a completed code PR, not a pre-implementation plan.
17	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_benchmark_shapes.py:67-70 runs `benchmark_mla_decode_live.main` end-to-end in CPU…
18	NEEDS-INFO	The unit test files for the verification guards and runner checks (`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py`, `test_mla_decode_runner.py`, and `test_mla_decode_reference.py`) remain omitted from the diff, so…
19	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/builders/mla/verify_mla_decode.py:246-272`, `make_problem` allocates `len(k_lens) * max_blocks`…
20	N/A	This PR does not introduce or modify a detector or linter for banned…
21	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:160-291` and…
22	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:181-190`, `test_mla_decode_golden_records_no_build_errors` asserts that…
23	N/A	This PR is a functional bring-up of the MLA decode kernel and makes no…
24	NEEDS-INFO	The diffs for `dnn-providers/hip-kernel-provider/rocke/library/tests/golden/mla_prefill_gfx942_ir_sha256.json` (+2/-2 lines) and `dnn-providers/hip-kernel-provider/rocke/library/tests/golden/mla_decode_ir_sha256.json` remain omitted from the…
25	N/A	This PR introduces new MLA decode kernels rather than applying a codegen…
26	N/A	This PR makes no benchmark speedup claims against a baseline or external…
27	N/A	This PR does not report a performance A/B study toggling optimization knobs.
28	N/A	This PR adds a new MLA decode kernel path and does not present a…
29	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:233`, `exp2_fast` is called on `fsub(ml[0], big_m)` where `big_m =…
30	NEEDS-INFO	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py` and the example READMEs (`library/examples/gfx942/mla_decode/README.md`, `library/examples/gfx950/mla_decode/README.md`) remain omitted from the diff, so we cannot…
31	N/A	This PR is a new kernel bring-up and does not claim to fix a performance…
32	NEEDS-INFO	The documentation files recording the compiler and resource tables (`dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md` and `gfx950/mla_decode/README.md`) and `library/tests/test_mla_decode.py` remain…
33	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:333-364`, `_resolve_num_cus` was hoisted from `dispatch/mla/gfx942.py:95-113` to…
34	N/A	This PR does not add an opt-in output flag; `lse` is always a required…

§H graded 5/5

#	Verdict	Evidence
1	FAIL	In dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:424 and :510, `check_contents: bool = False` was not updated since the previous review and remains default-off in both `MlaDecodeRunner.latent` and `MlaDecodeRunner.forward`, whereas the checklist requires contents checks that…
2	FAIL	In dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:686-699, the physical page ID from `block_table` is zero-extended to 64-bit (`page64 = b.zext(b.readfirstlane(page), I64)`) and multiplied by `stride` into the 64-bit base pointer (`b.global_ptr_add(cache, offset)`), which bypasses the…
3	PASS	Resolved in commit 25: `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:216-218` now enforces `if self.block_k & (self.block_k -…
4	PASS	In dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:617-631 and :669-676, the kernel loops over the runtime tile range `[t_lo…
5	NEEDS-INFO	The test files `dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py` (+750 lines) and `dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_runner.py` (+484 lines) remain omitted from the truncated diff…

§I graded 6/6

#	Verdict	Evidence
1	N/A	This PR is a functional bring-up of a single MLA decode layout per…
2	PASS	`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:160` clamps `num_seg` to `mla_decode_max_num_seg` and `MAX_GRID_DIM` (repair)…
3	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_benchmark_shapes.py:48` tests curated model shapes from `mla_shapes.json`, and…
4	N/A	This PR does not add a multi-candidate routing benchmark or tuning table…
5	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:28` documents the checked-in IR SHA256 golden snapshot and provides the…
6	PASS	`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:599` uses existing general DSL/layout primitives (`TransposeLdsReader` and…

§J graded 7/7

#	Verdict	Evidence
1	NEEDS-INFO	The previous NEEDS-INFO is not resolved: the Markdown hunks for the new and changed README/ALGORITHM/DESIGN files are omitted, so operand annotations, notation-table orientation, and equation result shapes cannot be checked.
2	NEEDS-INFO	The previous NEEDS-INFO is not resolved: the Markdown notation tables, symbol definitions, domain statements, and log-base mappings are not visible in the provided diff.
3	NEEDS-INFO	The previous NEEDS-INFO is not resolved: the changed Markdown prose for the new and modified docs is omitted, so it cannot be determined whether the docs restate an algorithm owned by a sibling document or only describe deltas.
4	NEEDS-INFO	The previous NEEDS-INFO is not resolved: novelty claims in the Markdown files are not visible, and no origin/develop comparison for those claims is provided in the diff.
5	NEEDS-INFO	The previous NEEDS-INFO is not resolved: the Markdown hunks that may contain scaling or complexity claims and repeated sites are omitted, so the aggregate and consistency cannot be checked.
6	NEEDS-INFO	The previous NEEDS-INFO is not resolved: any model tables, constants, tuned band edges, provenance text, and out-of-scope rows in the Markdown documentation are not visible in the provided diff.
7	NEEDS-INFO	The previous NEEDS-INFO is not resolved: the new README Markdown contents are omitted, so sibling naming/ownership conventions, declared entry points, and overlap with peer docs cannot be verified.

§K graded 13/13

#	Verdict	Evidence
1	N/A	No tuned candidate or KnobSpace-backed candidate with multiple named…
2	N/A	The PR adds no KnobAxis or tunable knob fields; the new MLA spec fields are…
3	N/A	No tuning_id or config_key payload or tuning-store identity is changed…
4	N/A	No KnobSpace.defaults() or recorded() implementation is added or changed.
5	N/A	No knob inertness, base-value, refusal, or known-wrong knob policy is added…
6	N/A	No KnobSpace.validate path, tuning_specs policy-free constructor, or tuned…
7	N/A	No KnobSpace.base(req) or tuning-problem base resolver is added or changed.
8	N/A	No accepted knob-value boundary is added; MLA refuses all non-empty…
9	N/A	The added benchmark resolves and times live dispatch selections in memory…
10	N/A	No bare tuning_id lookup or full-space tuning-id resolution path is added…
11	N/A	No tuned sweep entry point, sampling entry point, pin_to_spec storage path…
12	PASS	The routing change between decode and prefill is explicitly pinned: decode selects at max_seqlen_q == 1 in…
13	PASS	The benchmark sends shapes through dispatch_mla in dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:79 and…

§L graded 10/10

#	Verdict	Evidence
1	N/A	This PR does not add or modify a gfx target in arch_specs.json or backend…
2	N/A	This PR does not add or modify any MMA op_id in the hardware catalog or…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:82-85 and :482-484 use the catalog helper wrappers mfma_16x16x16_for_dtype and…
4	N/A	This PR only targets gfx942 and gfx950 (both wave64) and does not alter…
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:165-171 sets SCORE_K=16 and PV_TRANSPOSE_READ=False on MlaDecodeSpec (gfx942)…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:32-60 and dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:7…
7	N/A	This PR operates exclusively on bf16 and f32 tensors and does not add or…
8	N/A	This PR does not use or modify scaled MMA operations.
9	N/A	This PR does not add a new gfx architecture target to the platform.
10	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:154-168 defines flavor_honors_agpr_alloc by delegating to _is_modern_flavor, and…

§M graded 10/10

#	Verdict	Evidence
1	NEEDS-INFO	The diff still truncates dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py at line 790 before the end of build_mla_decode; the full hunk is needed to verify no device SSA Value is used in host-time Python control flow…
2	N/A	The diff does not touch legalize_for_target or target legalization code.
3	N/A	The diff does not change ck.dsl.ir/v1 serialization, print_ir interchange…
4	N/A	The diff does not change the portable IR, recipe, bundle schema, importer…
5	N/A	The new kernels do not use AsyncTileLoader for async DRAM-to-LDS copies.
6	NEEDS-INFO	The diff still truncates dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py at line 790 in the new kernel body; the full hunk is needed to verify all barriers, pointer loads, masking, and cross-lane operations in the…
7	PASS	The new buffer path constructs resources through the IR builder rather than hard-coding resource word 3 in…
8	PASS	The shared MLA flavor gate imports and delegates to _is_modern_flavor in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:166 and…
9	N/A	The diff does not change public C entry points, C headers under…
10	N/A	The diff does not emit !dbg metadata or change debug-info/code-object…

§N graded 18/18

#	Verdict	Evidence
1	PASS	`MAX_GRID_DIM` in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264` is a non-SSA host constant shared across GEMM/wgrad…
2	PASS	The MLA decode kernels reuse shared helpers (`mfma_16x16x16_for_dtype`, `mfma_16x16x32_for_dtype`, `safe_inv_l`, `TransposeLdsReader`, `SignatureBuilder`…
3	N/A	This PR does not modify or use `SoftwarePipeline`.
4	N/A	This PR does not add or modify `EpilogueOp`, `_PATTERN_TABLE`, or…
5	N/A	This PR does not add or modify any API in `rocke.helpers.tiling`.
6	N/A	This PR does not add or modify front-door/toolbox tiling helpers or…
7	N/A	This PR does not use `rocke.helpers.tiling` or custom tiling layout…
8	N/A	This PR does not use or modify `rocke.helpers.tiling` memory verbs…
9	N/A	This PR does not use or modify `rocke.helpers.tiling` N-D tensor axis roles…
10	N/A	This PR does not change `rocke.helpers.tiling` emission, LDS staging, or…
11	N/A	This PR does not use or modify `rocke.helpers.tiling` (`TileMma`, `Tiling`…
12	N/A	This PR makes no LDS bank-conflict or coalescing claims under…
13	N/A	This PR does not build or modify `WarpDistributionEncoding`…
14	N/A	This PR does not use or modify `rocke.helpers.tiling` data-movement verbs…
15	N/A	This PR does not modify `rocke.helpers.tiling` allocation or bounds checks.
16	N/A	This PR does not use or modify `optimize_layout` or `evaluate_transform` in…
17	N/A	This PR does not touch the `rocke.helpers.tiling` recorder, round-trip…
18	N/A	This PR does not touch `InterleavedStyle` or demo kernels under…

§O graded 16/16

#	Verdict	Evidence
1	NEEDS-INFO	The hand-built IR in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py is still truncated after line 850 before the end of build_mla_decode, so the complete emitted IR cannot be inspected.
2	NEEDS-INFO	While scf_for_iter and scf_yield in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:205-265 are complete, the loop bodies and yields in…
3	PASS	Deterministic IR bytes are pinned by the golden hash test in dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:138 and…
4	N/A	The diff does not modify IR type equality, dtype conversion, SmemType…
5	N/A	The diff does not touch core/passes.py, cpp/core/passes.cpp, op purity…
6	N/A	The diff does not modify constant folding logic in either engine.
7	PASS	The decode kernel sets structured kernel attributes DECODE_AGPR_ALLOC = (0, 0) in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py…
8	N/A	The diff does not set or modify CodegenPolicy or SchedulerStrategy…
9	NEEDS-INFO	The barrier and wait schedule inside build_mla_decode in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py is truncated after line 850, so the full loop body and epilogue synchronization cannot be verified.
10	N/A	The diff does not modify lowerer attribute validation or pointer parameter…
11	PASS	Resolved since the previous review: _live_device_is now resolves both target IDs with base_arch_from_target_id in…
12	PASS	The per-axis grid cap is centralized as MAX_GRID_DIM in dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264 and aliased by…
13	N/A	The PR's files under review do not modify public C structs, enums…
14	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt is modified in the PR but its diff hunk is omitted, so CMake flags and target changes cannot be inspected.
15	PASS	The C++/Python lowering parity test enforces ROCKE_CPP_STRICT=1 and checks all served flavors in…
16	PASS	Now visible in the diff: dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md:203-214 accurately records that…

§P graded 14/14

#	Verdict	Evidence
1	PASS	Builders take `(spec, *, arch)`: `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:522`…
2	PASS	Signatures are constructed via `SignatureBuilder` with canonical pointer/scalar types (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_deco…
3	PASS	Harnesses take the launch grid directly from the dispatch result or `MlaDecodeLaunch` (`dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_dec…
4	PASS	The LDS gate in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:375` charges the exact `kv_lds + qr_lds` allocation…
5	N/A	This PR does not touch WMMA/RDNA GEMM, `GemmMultiD`, or `GemmMultiAbd`.
6	NEEDS-INFO	Resolved against the current diff: the split-KV workspace (`part_acc_ptr`, `part_ml_ptr`) does not require caller-zeroing because empty segments write `(NEG_INIT, 0)` to `part_ml` and the combine kernel skips reading `part_acc` when `l <= 0`…
7	N/A	This PR does not touch Stream-K, persistent GEMM, or `GroupedGemmLauncher`.
8	N/A	This PR implements bf16 MLA decode and does not perform quantization or MX…
9	N/A	This PR does not add or modify numeric emulation or quantized decode helpers.
10	N/A	This PR does not add or modify any manifest `kind` or touch…
11	PASS	Launches go through `KernelLauncher` (`dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:111`), workspace buffers are…
12	PASS	`get_device_max_grid_size` in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/hip_module.py:351` queries `_hipDeviceGetAttribute`…
13	PASS	The host-side constant `MAX_GRID_DIM = 65535` in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264` is value-identical to…
14	PASS	No `expected.json` files or `tflops_lower_bound` values are added, `dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_l…

§Q graded 7/7

#	Verdict	Evidence
1	NEEDS-INFO	While dnn-providers/hip-kernel-provider/rocke/platform/tests/runtime/test_device_max_grid_size.py:16 is now visible and imports only from the installed rocke package, the diff hunk for…
2	PASS	MLA is documented as Python-only in dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md:203 (no platform/cpp/instances…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:242 fails on BackendCoverageGap or zero compared cases…
4	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:130 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decod…
5	PASS	Resolved from previous review: dnn-providers/hip-kernel-provider/rocke/platform/tests/runtime/test_device_max_grid_size.py:24 is now visible and tests…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:31 preserves the exact AssertionError vs unittest.SkipTest semantics when…
7	N/A	This PR does not modify cross-engine backend resolution or ROCKE_BACKEND /…

§R graded 11/11

#	Verdict	Evidence
1	N/A	This PR is a functional bring-up of a previously missing operator regime…
2	PASS	Arch-specific recipes (`SCORE_K = 32` for `mfma_f32_16x16x32_bf16`, `PV_TRANSPOSE_READ = True` for `ds_read_tr16_b64`, and per-arch…
3	PASS	The kernel follows the utilities skills for LDS padding (`WT_PAD` vs `V8_PAD` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py…
4	NEEDS-INFO	Unresolved since previous review: `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:135` cites `TestResources` in `dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py` for static VGPR, scratch…
5	NEEDS-INFO	Unresolved since previous review: global memory accesses are added in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:581-672` (and in the truncated epilogue after line 808) and…
6	PASS	The kernels are built with raw `IRBuilder` outside `rocke.helpers.tiling` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:513`)…
7	N/A	The PR makes no bottleneck classification claims and states in the…
8	PASS	`b.sched_barrier(0)` and `DECODE_SCORE_PREFETCH` / `DECODE_SCORE_FENCE_EVERY` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py…
9	N/A	This PR is not a performance regression fix and does not retune an existing…
10	PASS	Compiler hazards around `agpr_alloc` on `llvm20` are gated via `flavor_honors_agpr_alloc` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla…
11	N/A	No new optimization failure mode, heuristic study, or `dsl_probes` tool is…

§S graded 8/8

#	Verdict	Evidence
1	NEEDS-INFO	`platform/tests/runtime/test_device_max_grid_size.py:38-79` is now visible and directly tests `get_device_max_grid_size`, and `library/tests/dispatch/mla/test_mla_decode_wiring.py:348-425` exercises `MAX_GRID_DIM` and `MlaDecodeLaunch`…
2	PASS	`MLARequest` in `library/dispatch/mla/common.py:137` adds required `max_seqlen_q: int` and changes `op` default to `MLA_OP` (`mla_fwd`); all callers and…
3	PASS	`dispatch.mla.gfx942._resolve_num_cus` was moved/renamed to `dispatch.mla.common.resolve_num_cus` (`library/dispatch/mla/common.py:351`) and all…
4	PASS	`MAX_GRID_DIM` is defined in `platform/python/rocke/helpers/spec.py:264` and added to `__all__` in `platform/python/rocke/helpers/spec.py:33` and…
5	NEEDS-INFO	`platform/tests/runtime/test_device_max_grid_size.py:38-79` is now visible and asserts positive and negative cases for `get_device_max_grid_size`, and `library/tests/dispatch/mla/test_mla_decode_wiring.py:125-237` /…
6	PASS	CPU-checkable contracts (IR golden hashes, C++/Python byte-identity, spec/request validation, and torch-free device grid limit checks) run on CPU without…
7	PASS	`library/tests/test_mla_decode_golden.py:212-246` (`test_mla_decode_cpp_python_byte_identity`) lowers all 14 MLA decode cases through both…
8	FAIL	Not resolved since the previous review: the PR body's 'Test plan' section lists pytest/suite counts and GPU archs (`gfx942`), and `platform/python/rocke/helpers/spec.py:264` / `platform/python/rocke/runtime/hip_module.py:351` do not break README usage, but the PR description still does not state that it ran the…

yraparti · 2026-11-04T16:25:41Z

+
+def _tools_required() -> bool:
+    """Whether this run supplies the compiler tools, and so requires them."""
+    return os.environ.get(REQUIRE_ENV) == "1"


[HIGH] §D.1 ★ ★ Knobs/features on the KernelSpec (validator + selector + name tag + cache-key), NOT os.environ (env only as a selector-level diagnostic override folded onto the spec). Reviewers may reject env vars outright — even a reproducibility/pin one (#9739: "please dont use env variables", on a…

dnn-providers/hip-kernel-provider/rocke/library/tests/_mla_resource_gate.py:28 reads the environment

Fix: Put the knob on the KernelSpec and delete the environment read.

yraparti · 2026-11-04T16:25:41Z

+            "scale": float(scale),
+            "num_seqs": int(self.request.num_seqs),
+            # The kernel's row pitch: the stride, which a column-sliced view
+            # does not share with its shape.
+            "block_table_stride": int(block_table.stride(0)),
+            "max_k": int(self.request.max_seqlen_k),
+        }
+        config = LaunchConfig(grid=self.grid, block=self.block, stream=stream)
+        if self._combine is None:
+            vals.update(out_ptr=out, lse_ptr=lse)
+            return [(self._latent, vals, config)], out, lse
+        part_acc, part_ml = self.workspace(qa.device)
+        vals.update(part_acc_ptr=part_acc, part_ml_ptr=part_ml, num_seg=spec.num_seg)
+        combine_vals = {
+            "out_ptr": out,
+            "lse_ptr": lse,
+            "part_acc_ptr": part_acc,
+            "part_ml_ptr": part_ml,
+            "num_seg": spec.num_seg,
+        }
+        combine_config = LaunchConfig(
+            grid=self.launch.combine_grid(),
+            block=self.launch.combine_block(),
+            stream=stream,
+        )
+        return (
+            [
+                (self._latent, vals, config),
+                (self._combine, combine_vals, combine_config),
+            ],
+            out,
+            lse,
+        )
+
+    def latent(
+        self,
+        *,
+        qa: torch.Tensor,
+        c_kv: torch.Tensor,
+        k_rope: torch.Tensor,
+        cu_seqlens_q: torch.Tensor,
+        cu_seqlens_k: torch.Tensor,
+        block_table: torch.Tensor,
+        scale: float,
+        out: torch.Tensor | None = None,
+        lse: torch.Tensor | None = None,
+        stream: int = 0,
+        check_contents: bool = False,


2 checklist findings on this snippet

[HIGH] §H.1 ★ (line 424) ★ A hardware/resource bound derived from SPEC fields (e.g. an AMDGPU buffer-SRD num_records sized from num_kv_blocks*block_size*num_kv_heads*head_size) MUST be matched by the launch tensors — host-validate their shape/size against the spec that sizes the bound BEFORE compile/launch, else the…

In dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:424 and :510, check_contents: bool = False was not updated since the previous review and remains default-off in both MlaDecodeRunner.latent and MlaDecodeRunner.forward, whereas the checklist requires contents checks that force a device-to-host sync to be gated behind a default-on flag (off-able on the hot/graph-captured path, which bind_latent already provides).

Fix: Change check_contents: bool = False to check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward (keeping check_contents=False in bind_latent).

[MEDIUM] §E.15 (line 377) No dead parameters — and a parameter ACCEPTED then IGNORED is worse than one absent: the caller believes it was used, and in measurement code it silently discards a real observation. Grep each parameter name in the body. (#12070: substitute_share(total_armA_us=…) never read it — the measured…

Not resolved in the new commits: dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:599-601 still declares scalar kernel parameters num_seqs and max_k (b.param("num_seqs", I32) and b.param("max_k", I32)) whose return handles are discarded and never read in the kernel body, while dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:377-381 passes req.num_seqs and req.max_seqlen_k into them on every launch.

Fix: Remove the unused num_seqs and max_k scalar parameters from mla_decode_signature and build_mla_decode (and pass them only in run_prefill_oracle when invoking the prefill kernel), or use them in-kernel.

yraparti · 2026-11-04T16:25:41Z

+
+    c_chunks = (BK * R_KV) // (THREADS * C_STAGE_W)
+    c_chunks_per_row = R_KV // C_STAGE_W
+    KR_W = _kr_stage_width(spec)


[HIGH] §H.2 ★ ★ Guard indirection VALUES, not just indices — a masked/bounds-guarded table read protects the lookup INDEX (e.g. logical page < n_pages), NOT the value it returns. Any value used as a memory index (block-table physical id) must itself be range-checked, or shown contained by a HW bounds-check —…

In dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:686-699, the physical page ID from block_table is zero-extended to 64-bit (page64 = b.zext(b.readfirstlane(page), I64)) and multiplied by stride into the 64-bit base pointer (b.global_ptr_add(cache, offset)), which bypasses the SRD bound check for out-of-range page IDs (and negative page IDs like -1 zero-extend to 0xFFFFFFFF), while host-side validation in MlaDecodeRunner.latent (mla_decode_runner.py:424) still defaults to check_contents: bool = False (unresolved from the previous review).

Fix: Default check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward, or range-check page against num_pages on device before forming the 64-bit pointer offset.

yraparti · 2026-11-04T16:25:41Z

+def _time(
+    shapes, *, arch: str, warmup: int, iters: int, repeats: int, seed: int
+) -> int:
+    import torch


[MEDIUM] §E.2 Library doesn't import torch for device queries — use rocke.runtime.hip_module (prefer hipDeviceGetAttribute).

dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:169 adds import torch

Fix: Query the device through rocke.runtime.hip_module instead of importing torch.

yraparti · 2026-11-04T16:25:42Z

+_max_grid_size: Dict[int, Optional[Tuple[int, int, int]]] = {}
+
+
+def get_device_max_grid_size(device: int = 0) -> tuple[int, int, int] | None:


[LOW] §S.8 The PR states it ran the README's entry points from rocke/platform/: PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py (relative-path guard, cross-engine byte-identity, pytest, and ctest when built); a manifest or runtime change…

Not resolved since the previous review: the PR body's 'Test plan' section lists pytest/suite counts and GPU archs (gfx942), and platform/python/rocke/helpers/spec.py:264 / platform/python/rocke/runtime/hip_module.py:351 do not break README usage, but the PR description still does not state that it ran the README entry points PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py from rocke/platform/.

Fix: Run PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py from rocke/platform/ and record them in the PR test plan.

cgmillette · 2026-11-04T16:29:34Z

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

PR #13923 — add MLA decode kernel for gfx942 and gfx950 (aosewski, AICK-1504) — 🔴 Follow-up: new spill test fails in CI

CI update to our follow-up 4 at 047d7fde, which went out while CI ran. Both Linux gfx94X hipkernelprovider shards now fail on one of this PR's new tests.

#	Sev	file:line	Finding
N5	🔴	`tests/test_mla_decode.py:599`	`TestResources.test_no_spills_two_workgroups_per_cu[4-gfx942]` fails in CI • the decode resource probe reports 8 bytes of scratch at `num_seg` 4 on gfx942: `assert 8 == 0`, on both shards (read: CI logs) • `c41875b4` passed both shards; since then: `934f5bb4` and two develop merges. Our MLA lane at llvm24 passes it at `047d7fde` (ran, follow-up 4), so the shards' compiler spills where ours does not • fix: find the spill with the shards' compiler and keep `num_seg` 4 spill-free; the test is new here, so CI turns red if this lands as is

Bots / CI	Status
`047d7fde`	🔴 Linux gfx94X hipkernelprovider ✗ on both shards: N5; 85 of 86 ctests ✓; `rocke_library_mla_pytest` 479 ✓, 1 ✗, 3 skipped (read) • still running: the Windows gfx110X tests

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

Follow-up CI update, verified by the orchestrator (Claude Opus 5.5) from the Linux gfx94X hipkernelprovider job logs at 047d7fde. Flags whether a closer look is needed; not an approval or rejection.

Copilot

🟡 Changes recommended

Unresolved critical defects in input validation, device and stream handling, and kernel arithmetic can permit unsafe GPU launches and incorrect decode results.

10 open issues

Reject cache pointers that are not 16-byte aligned · New
Reject overlapping inputs and outputs · New
Widen offsets before validating differences · New
Reject tensors on a device different from the runner · New
Reject CPU tensors before launching kernels · New
Run content validation on the supplied stream · New
Allocate outputs and workspace on the launch stream · New
Guard key-length tile arithmetic against int32 overflow · New
Handle finite scores below the empty sentinel · New
Resolve duplicate section numbering and clarify lowering coverage · New

7 resolved since last review

_{Review effort: Max}

+    the kernel gate, and the combine gate when split. Dispatch's own choice and
+    a forced split count (:meth:`MlaDecodeLaunch.with_num_seg`) pass the same
+    gate."""
+    ok, why = check_mla_decode_extents(spec, total_q=req.total_q)


+#: the ``m`` an empty segment reports. Finite on purpose: with every segment
+#: empty the merge's ``M`` equals it too, so ``2**(m_s - M)`` is 1 rather than
+#: ``-inf - -inf = NaN``, and ``l = 0`` then zeroes the row.
+NEG_INIT = -1e30


+## 5. Library MLA — no C++ instance mirror, no emitter parity pair
+
+**Subsystem:** `library/kernels/mla/` (`mla_prefill_fwd`, `mla_decode`,
+`mla_decode_combine`)
+
+The MLA kernels are built by their Python builders only: there is no
+`platform/cpp/instances` counterpart and no `library/tests/parity/mla_*_emit.{py,c}`
+pair. Lowering is still dual-engine for decode
+(`library/tests/test_mla_decode_golden.py::test_mla_decode_cpp_python_byte_identity`
+lowers every case through both engines), and the Python-lowered IR of both
+kernels is pinned by the `library/tests/golden/mla_*_ir_sha256.json` goldens.


CI's LLVM 24 compiler moved to a newer build, and under it the gfx942 split-KV decode kernel spilled one VGPR (8 bytes of scratch), failing TestResources at num_seg 4 on both gfx94X shards. The spilled value was the lane's head index: the prologue computes h0 + lane_row for the query load, and the epilogue's workspace row reused it, so it stayed live across the whole key loop. No score read-ahead or fence setting avoids the spill; the single-pass kernel was not affected. The split epilogue now builds the row from its wave-uniform part, (token * num_seg + seg) * H + h0, passed through readfirstlane, plus lane_row, which is live through the loop anyway. readfirstlane keeps the compiler from reassociating the sum back into the prologue's head index; the value is wave-uniform, so the result is unchanged. Under both CI LLVM 24 builds and LLVM 22, all decode kernels have zero scratch on gfx942 and gfx950 (the gfx950 split kernel also drops one VGPR). Only the split cases' IR moves: their goldens are re-blessed at llvm22, llvm23 and llvm24, single-pass and combine are unchanged, and C++/Python parity holds. The example READMEs' resource tables are re-measured. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Review follow-up for the MLA decode PR: one flavor rule by lane, one integer check, the runner's remaining input checks, and the shared compile-test skip policy in its owner's shape. Flavor rule. flavor_honors_agpr_alloc compared the datalayout a lane emits, so a new compiler lane would be refused with every test green. It now ranks lanes (_flavor_at_least against AGPR_ALLOC_MIN_FLAVOR, llvm22): the same lanes today, and a newer lane is served. Tests pin a simulated newer lane and an unknown name. Integer fields. require_int_fields in _mla_body is the one exact-int check of the MLA specs and raises TypeError, as prefill already did; the decode, combine and prefill specs use it, and with_num_seg relies on the spec's check instead of its own ValueError. Key length. Dispatch refuses max_seqlen_k past what the kernel's i32 tile count (s_k + block_k - 1) holds; check_mla_decode_extents takes max_seqlen_k. Runner. MlaDecodeRunner refuses tensors that are not on a GPU or not on the device its kernels were loaded on, pointers short of the alignment the kernels declare, and outputs that share storage with an input or with each other (the kernels declare them noalias). check_contents widens the offsets to int64 before subtracting, so wrapped int32 offsets are rejected, and reads them back on the launch stream; outputs and the split workspace are allocated on the launch stream. Skip policy. tests/_mla_resource_gate.py becomes tests/_compiler_test_utils.py with the name and text of the module an open PR owns (REQUIRE_COMPILER_TOOLS_ENV, require_tools_or_skip), so whichever lands second keeps one module. A compile option the loaded comgr refuses now follows that rule: it fails where the run requires the tools and skips elsewhere. The rule's own test fails, not skips, if the helper skips where it must fail. Smaller: target IDs resolve through base_arch_from_target_id in both kernel gates, the verify harness and the numeric test; the combine's finite running-max floor is documented; known_gaps.md describes the MLA families in one section (the decode block had reused number 5) and notes that decode is not yet on the shared _mla_body emitters. No kernel IR changes: the goldens and C++/Python parity are unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…n decode Decode emitted its own copies of three shared pieces. It now calls them: lane_ids and softmax_consts from kernels/mla/_mla_body.py for the lane setup and the log2-domain scale and sentinels, and natural_lse_from_log2_stats for the LSE in the latent and combine kernels. The combine's NEG_INIT is the family's NEG_INF, which the latent kernel's softmax constants carry, so the two kernels cannot disagree on the empty-segment max. Each swap emits the same instructions in the same order: every decode and combine IR golden is unchanged at llvm22, llvm23 and llvm24, and C++/Python parity holds. The emitters whose reuse would move the IR (the masked key tail among them) stay for the follow-up noted in known_gaps.md. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…k-1504-mla-decode

aosewski · 2026-11-05T12:55:04Z

@cgmillette thanks for the follow-up, the CI update and the overlap notice. All addressed at 1b26101 (develop merged again):

N5, the gfx942 spill: fixed in 58123e2. CI's LLVM 24 moved to a newer build, which spilled one VGPR: the lane's head index, kept live from the prologue to the split epilogue. No read-ahead or fence setting avoids it. The epilogue now builds the workspace row from its wave-uniform part through readfirstlane, plus lane_row. Zero scratch on both arches under both CI builds and LLVM 22; the 18 split goldens are re-blessed; device numerics pass on gfx942 and gfx950, also with the new CI compiler's binaries.
N1, the flavor rule: fixed in f48c4f3. flavor_honors_agpr_alloc ranks lanes (_flavor_at_least(flavor, "llvm22")); a test serves a simulated newer lane, and the old datalayout rule fails it.
N2, the gate helper: done as the overlap notice suggests. tests/_compiler_test_utils.py with fix(rocke): never pass a compiler option the loaded compiler does not know #14218's names and text, minus skip_where_compiler_refuses_option; a refused compile option is now the tools verdict (fails under the variable, skips otherwise). fix(rocke): never pass a compiler option the loaded compiler does not know #14218 keeps its file when it lands second.
N3: fixed. One require_int_fields in _mla_body raises TypeError for every MLA spec, prefill included; with_num_seg relies on it. The leftover split(":") sites now use base_arch_from_target_id.
The 11-02 N1 deferral: lane_ids, softmax_consts and the LSE helper are reused in 6808bc4, IR unchanged. The rest (masked tail, softmax row, kernel-argument table, the two unused arguments) is tracked in AICK-3039 under AICK-1504, after refactor(rocke): MLA prefill on shared emitters; odd head counts on the packed layout #14502, and noted in known_gaps.md section 6.
Runner grid check vs dispatch's _check_grids: kept both on purpose. Dispatch checks the platform's MAX_GRID_DIM so selection stays host-independent; the runner checks the device's own limit at launch.

Copilot's ten comments are in f48c4f3 too, except the -1e30 floor, which is documented rather than changed (it would move both MLA families' IR).

yraparti

rocKE checklist re-review of the fixes since `047d7fde0a25`

ROCm/rocm-libraries#13923 · 1b26101cf685 · model gemini-4-argon, fallback gpt-5.5

PASS · graded 238/238 · BLOCKER 0, HIGH 3, MEDIUM 2, LOW 0, NEEDS-INFO 45

The diff was truncated before grading. Items that needed a missing hunk are NEEDS-INFO.

Since the last review (`047d7fde0a25`)

Resolved 1 · still open 5 · new 0

✅ §S.8 was FAIL, now PASS — had asked: Run PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py from rocke/platform/ and record them in the PR test plan.
⏳ §D.1 still FAIL (HIGH): Put the knob on the KernelSpec and delete the environment read.
⏳ §E.2 still FAIL (MEDIUM): Query the device through rocke.runtime.hip_module instead of importing torch.
⏳ §E.15 still FAIL (MEDIUM): Remove the unused num_seqs and max_k scalar parameters from mla_decode_signature and build_mla_decode (and pass them only in run_prefill_oracle when invoking the prefill kernel), or use them in-kernel.
⏳ §H.1 still FAIL (HIGH): Change check_contents: bool = False to check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward (keeping check_contents=False in bind_latent).
⏳ §H.2 still FAIL (HIGH): Default check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward, or range-check page against num_pages on device before forming the 64-bit pointer offset.

Findings

HIGH §D · dnn-providers/hip-kernel-provider/rocke/library/tests/_compiler_test_utils.py:40 · Put the knob on the KernelSpec and delete the environment read.
HIGH §H · dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:541 · Change check_contents: bool = False to check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward (keeping check_contents=False in bind_latent).
HIGH §H · dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:693 · Default check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward, or range-check page against num_pages on device before forming the 64-bit pointer offset.
MEDIUM §E · dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:169 · Query the device through rocke.runtime.hip_module instead of importing torch.
MEDIUM §E · dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:581 · Remove the unused num_seqs and max_k scalar parameters from mla_decode_signature and build_mla_decode (and pass them only in run_prefill_oracle when invoking the prefill kernel), or use them in-kernel.

Needs info

§A.1 The current platform/dsl_docs/development/known_gaps.md hunk is omitted, so I cannot verify that the MLA Python-only family entry still exists after the latest edits.
§A.2 The current diff lists changes under both Python and C++ lowering/core files, but the relevant platform/python/rocke/core/lower_* and platform/cpp/core/lower_* hunks are omitted, so mirrored behavior cannot be verified.
§A.3 The current diff includes kernel-level attention/conv/MLA selector or spec-builder changes, but several corresponding C++ instance/selector hunks are omitted, so selector/spec parity cannot be verified.
§A.4 The current diff includes C++ twin changes for conv/attention/platform instances, but the emitted-op ordering in the corresponding Python and C++ hunks is omitted, so operand/order parity cannot be checked.
§A.5 The previous NEEDS-INFO is not resolved in this prompt: the hunk for library/tests/golden/mla_prefill_gfx942_ir_sha256.json is still omitted, and additional representative golden hunks are also omitted, so intentional versus accidental re-blessing cannot be verified.
§A.7 The current diff adds and changes hot-path kernel logic in multiple omitted conv/attention/platform hunks, so I cannot verify that covered byte-identical shapes remain guarded from unconditional IR changes.
§A.9 The current diff lists new parity mini-pair files and platform lowering changes, but the actual mini-pair hunks are omitted, so I cannot verify that the shared op/lowerer changes are proven by the required <op>_emit.{py,c} pairs.
§A.10 dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:119 defines MLA decode golden cases, but the current diff also adds other rocKE IR-changing features whose golden JSON hunks are omitted, so full flag-on golden coverage cannot be verified.

Section notes

§A: The current diff is heavily truncated; items that depend on omitted rocKE hunks are marked NEEDS-INFO rather than guessed.
§E: Re-review at head 1b26101: items 2 and 15 remain FAIL (not resolved in the new commits), and item 7 remains NEEDS-INFO because the doc diffs are omitted and ROCKE_REQUIRE_COMPILER_TOOLS is not added to env_flags.md.
§G: Re-review of head 1b26101. Items 5, 6, and 9 remain NEEDS-INFO because test_mla_decode.py, test_mla_decode_reference.py, test_mla_decode_runner.py, and the example READMEs are still omitted from the diff budget.
§G: Graded items 18-34 of section G against current head 1b26101.
§H: Trace of device-side masks and bounds in build_mla_decode (kernels/mla/mla_decode.py) and build_mla_decode_combine (kernels/mla/mla_decode_combine.py): 1. Query sequence empty guard: extent source is runtime cu_seqlens_q[seq_idx] and cu_seqlens_q[seq_idx+1] (token >= q_stop, mla_decode.py:610-614); kernel exits early before any global writes. 2. Key tile loop range [t_lo, t_hi) and unmasked prefix [t_lo, t_full): extent source is runtime s_k = cu_seqlens_k[seq_idx+1] - cu_seqlens_k[seq_idx] and num_seg (mla_decode.py:617-635). The kernel visits [t_lo, t_full) unmasked and [t_full, t_hi) masked; they agree. 3. Page SRD bounds c_rsrc and kr_rsrc (mla_decode.py:694-710): extent source is compile-time C_PAGE_BYTES = PAGE * R_KV * 2 and KR_PAGE_BYTES = PAGE * D_ROPE * 2 (or 0 when live = has_tiles = (t_lo < t_hi) is false). The kernel visits c_chunks of C_STAGE_W and kr_chunks of KR_W within a single page, which matches the SRD bound for that…
§J: Shapes: | Equation / Op | Operands | Inner dim | Result vs Declared | |---|---|---|---| | Markdown hunks omitted | library/examples/{gfx942,gfx950}/mla_decode/README.md, library/examples/gfx942/mla_prefill/README.md, platform/dsl_docs/... | (omitted from diff) | NEEDS-INFO (Markdown hunks omitted to fit model budget) | Symbols: | Symbol | Defined in | Domain | |---|---|---| | Markdown hunks omitted | library/examples/{gfx942,gfx950}/mla_decode/README.md, platform/dsl_docs/... | NEEDS-INFO (Markdown hunks omitted to fit model budget) | Novelty: | Claim | Symbol / Path | Status | |---|---|---| | Markdown hunks omitted | library/examples/{gfx942,gfx950}/mla_decode/README.md, platform/dsl_docs/development/known_gaps.md, platform/dsl_docs/instances/attention.md | NEEDS-INFO (Markdown hunks omitted to fit model budget) |

Graded with the rocKE PR checklist, one model pass per section. FAIL rows that sit on a diff line are inline comments.

§A graded 14/14

#	Verdict	Evidence
1	NEEDS-INFO	The current `platform/dsl_docs/development/known_gaps.md` hunk is omitted, so I cannot verify that the MLA Python-only family entry still exists after the latest edits.
2	NEEDS-INFO	The current diff lists changes under both Python and C++ lowering/core files, but the relevant `platform/python/rocke/core/lower_` and `platform/cpp/core/lower_` hunks are omitted, so mirrored behavior cannot be verified.
3	NEEDS-INFO	The current diff includes kernel-level attention/conv/MLA selector or spec-builder changes, but several corresponding C++ instance/selector hunks are omitted, so selector/spec parity cannot be verified.
4	NEEDS-INFO	The current diff includes C++ twin changes for conv/attention/platform instances, but the emitted-op ordering in the corresponding Python and C++ hunks is omitted, so operand/order parity cannot be checked.
5	NEEDS-INFO	The previous NEEDS-INFO is not resolved in this prompt: the hunk for `library/tests/golden/mla_prefill_gfx942_ir_sha256.json` is still omitted, and additional representative golden hunks are also omitted, so intentional versus accidental…
6	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:212` adds `test_mla_decode_cpp_python_byte_identity`, a cross-engine…
7	NEEDS-INFO	The current diff adds and changes hot-path kernel logic in multiple omitted conv/attention/platform hunks, so I cannot verify that covered byte-identical shapes remain guarded from unconditional IR changes.
8	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:212` explicitly lowers every MLA decode golden case through the C++ engine…
9	NEEDS-INFO	The current diff lists new parity mini-pair files and platform lowering changes, but the actual mini-pair hunks are omitted, so I cannot verify that the shared op/lowerer changes are proven by the required `<op>_emit.{py,c}` pairs.
10	NEEDS-INFO	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:119` defines MLA decode golden cases, but the current diff also adds other rocKE IR-changing features whose golden JSON hunks are omitted, so full flag-on golden…
11	NEEDS-INFO	The current diff lists possible IR-op touchpoint changes in C++ IR, lowerers, docs, and tests, but the relevant core IR/lowering hunks are omitted, so I cannot determine whether a new IR op hits every required table and lowerer.
12	NEEDS-INFO	The current diff modifies `platform/python/rocke/core/backend.py`, but that hunk is omitted, so I cannot verify that `ROCKE_BACKEND=both` still avoids Python substitution on C++ failures and that `CPP_UNPORTED_ARCHES` remains empty.
13	NEEDS-INFO	The current diff modifies LLVM lowering files, but the lowerer hunks are omitted, so I cannot verify that emitted `.ll` remains free of build IDs, timestamps, hashes, or debug locations.
14	NEEDS-INFO	`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:257` adds `MAX_GRID_DIM` in a helper with a C++ twin, but the corresponding `platform/cpp/helpers/spec.cpp` and header hunks are omitted, so the twin update cannot…

§B graded 6/6

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:81-87 documents the head-count scope (`DECODE_NUM_HEADS = (64, 128)`) with a TODO…
2	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:272-276 calls `spec_for(req)` inside `support(req)` and catches `(TypeError…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:327 sets `ShapeRange("max_seqlen_q", min=2)` on prefill so `max_seqlen_q == 1` is…
4	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:180-190 defines `require_int_fields` checking `if type(value) is not int:`, called…
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:87-218 scopes admission to `DECODE_NUM_HEADS = (64, 128)` and…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:316-386 (`supports_mla_decode`) and…

§C graded 10/10

#	Verdict	Evidence
1	NEEDS-INFO	Shared predicates and logic are factored into single named helpers: `flavor_honors_agpr_alloc` and `require_int_fields` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:160` and…
2	PASS	One reusable builder per kernel family (`build_mla_decode(spec)` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:527` and…
3	PASS	Builders and `kernel_name()` are named by operation and codegen knobs, never workload sizes (`num_seg` is excluded from `kernel_name()`, keeping only…
4	PASS	The shared MLA decode body in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:167` uses spec class attributes (`ARCH`…
5	PASS	Every touched file in the PR belongs directly to the MLA decode kernel, dispatch, verification, tests, or the shared `MAX_GRID_DIM` /…
6	PASS	Shared helpers (`resolve_num_cus` in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:351`, `flavor_honors_agpr_alloc` and…
7	N/A	This PR introduces a new MLA decode feature rather than performing a…
8	PASS	Reuses existing helpers from `develop` including `read_ahead` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:85`)…
9	PASS	Resolved since previous review: `require_int_fields` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:180` now unifies…
10	PASS	Widening `MLARequest` with `max_seqlen_q` and `MLA_OP = 'mla_fwd'` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:54`…

§D graded 21/21

#	Verdict	Evidence
1	FAIL	dnn-providers/hip-kernel-provider/rocke/library/tests/_compiler_test_utils.py:40 reads the environment
2	PASS	`make_mla_decode_candidate` declares `Capability(arches=(arch,), ...)` where `arch = spec_cls.ARCH` (`gfx942` or `gfx950`), and `DECODE_NUM_HEADS = (64…
3	PASS	`support` in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:250-277` runs cheap request/selector/flavor checks before `spec_for`…
4	PASS	`decode_spec_for` and `_check_grids` take `arch` from `req` and only query device CU count when `req.num_cus == 0` via `resolve_num_cus(req)`…
5	PASS	`MlaDecodeSpec.identity()` removes launch-only `num_seg` and includes `split_kv` plus class emission attributes (`arch`, `score_k`, `pv_transpose_read`)…
6	PASS	Single-pass vs split-KV routing keeps `decode_spec_for`, `mla_decode_grid`, `mla_decode_signature`, `MlaDecodeLaunch`, and `MlaDecodeSpec.identity()` in…
7	N/A	This PR does not introduce machine-regenerable heuristics or hand-pinned…
8	PASS	The LLVM flavor cohort rule `flavor_honors_agpr_alloc` is defined once in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:160-177`…
9	PASS	`make_mla_decode_candidate` registers `mla_decode_gfx942` and `mla_decode_gfx950` on `MLA_REGISTRY` with `Capability`, `build`, `grid`, `block`, and…
10	PASS	The new `MLARequest.max_seqlen_q` field (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:137, 228`) and `MlaDecodeSpec.num_seg`…
11	PASS	The new production candidates `mla_decode_gfx942` and `mla_decode_gfx950` serve the disjoint `max_seqlen_q == 1` decode regime where prefill explicitly…
12	PASS	Priority `DECODE_PRIORITY = 5` is justified against prefill (10) in `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:73-75` and…
13	PASS	`dispatch_mla` delegates directly to `MLA_REGISTRY.select` without catching or retrying on `PinRefused` in…
14	PASS	`MlaDecodeSpec.identity()` omits the problem-dependent launch argument `num_seg` and keeps only `split_kv` and arch emission choices so `compile_key` does…
15	PASS	The PR intentionally partitions the MLA regime by `max_seqlen_q` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:325-327`), which…
16	PASS	The `mla` family exposes `registered_mla_combos`, `mla_sweep_space`, and `dispatch_mla_all` with `_request_errors` short-circuiting…
17	PASS	CPU routing tests in `dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:94-245` check cohort selection across…
18	PASS	`MLA_REGISTRY` keeps `require_build=True` (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/__init__.py:48-50`) and no ratchets or exempt sets…
19	N/A	This PR adds MLA decode kernels and dispatch under `dispatch/mla/` and does…
20	PASS	The PR explicitly documents in `dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:14-18` that MLA has no C++ instance mirror…
21	N/A	This PR introduces a new MLA decode kernel and its dispatcher registration…

§E graded 17/17

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/hip_module.py:351 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/he…
2	FAIL	dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:169 adds `import torch`
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/builders/mla/ref_mla_decode.py:24 uses numpy only, and…
4	PASS	Layering is respected downward: kernels (dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:67) -> dispatch…
5	PASS	dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/hip_module.py:351 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/he…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:212 verifies C++/Python byte-identity and IR SHA256 goldens across all…
7	NEEDS-INFO	Not resolved in the visible diff: the diffs for dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md, dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/attention.md, and…
8	PASS	Visible diffs (e.g. dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:1 and…
9	PASS	All PR-specific touched files in the diff (e.g. dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:1…
10	PASS	The PR body notes explicit coordination with #12464 and merge simulation with develop (with develop merged into the branch in the latest commits), and…
11	PASS	Only deliverable source, tests, shape JSONs, and required instance/example READMEs are committed (e.g.…
12	PASS	Both gfx942 and gfx950 are covered: dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:385 and…
13	PASS	Renamed/hoisted helpers (`_resolve_num_cus` -> `resolve_num_cus` in dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:351…
14	PASS	Removed lines in modified files (e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/gfx942.py:95…
15	FAIL	Not resolved in the new commits: dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:581-583 still declares scalar kernel parameters `num_seqs` and `max_k` (`b.param("num_seqs", I32)` and `b.param("max_k", I32)`) whose return handles are discarded and never read in the kernel body, while…
16	N/A	This item applies to how review feedback is written rather than the code…
17	PASS	No absolute software-achieved performance figures appear in added lines (e.g. dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_ml…

§F graded 11/11

#	Verdict	Evidence
1	PASS	Branch `users/aosewski/aick-1504-mla-decode` uses the recognized `users/<user>/` prefix; code is added in…
2	PASS	title `feat(rocke): add MLA decode kernel for gfx942 and gfx950` matches type(scope): summary and is 56 characters
3	PASS	description is at least 30 characters and contains an issue reference
4	PASS	test file `dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py` is in the diff
5	PASS	All changed files in the diff are source, test, config, or doc files (e.g., `dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla…
6	NEEDS-INFO	The new commits since the previous review do not resolve this item because local `pre-commit` execution before push is not recorded in the diff.
7	NEEDS-INFO	The new commits since the previous review update code and tests (e.g., `dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:50` and `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:201`)…
8	N/A	This PR targets `develop` directly from a `users/...` branch rather than…
9	N/A	The diff does not show any artificial splitting or trimming done to meet a…
10	PASS	The added and updated tests verify real invariants, such as forced split-count type validation in…
11	NEEDS-INFO	The new commits since the previous review do not resolve this item because verifying the PR title and body via the GitHub API is an external action not visible in the diff.

§G graded 34/34

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:211-245 checks C++/Python lowering byte-identity separately from on-GPU…
2	PASS	dnn-providers/hip-kernel-provider/rocke/library/builders/mla/verify_mla_decode.py:116-124 and 495-556 implement the on-silicon numeric gate against the…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:130-136 and 193-245 gate all served modern LLVM flavors (llvm22, llvm23…
4	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:94-243 and 365-608 add behavioral tests for decode routing…
5	NEEDS-INFO	The diffs for dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py (+791 -0), dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_reference.py (+560 -0)…
6	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md (+295 -0), dnn-providers/hip-kernel-provider/rocke/library/examples/gfx950/mla_decode/README.md (+129 -0), and…
7	N/A	This PR does not add or modify a design doc (`DESIGN.md` is not in the PR's…
8	PASS	dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:249-263 and 656 loudly raise ValueError on non-bf16 tensors, and…
9	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py (+791 -0, which contains TestResources per mla_decode.py:132) and the example READMEs remain omitted from the diff budget, so tool-sourced VGPR/LDS assertions cannot be…
10	N/A	The PR explicitly makes no performance claims and is a functional bring-up.
11	PASS	dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:79 and…
12	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:211-245 verifies C++ and Python lowering parity across all 14 decode cases…
13	N/A	The PR uses existing DSL builder methods (`IRBuilder`) and does not…
14	N/A	The PR is a functional bring-up with no bottleneck classification or…
15	PASS	All graded files and hunks come directly from the provided PR diff for commit `1b26101` (e.g.…
16	N/A	This is a completed code PR, not a pre-implementation plan.
17	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_benchmark_shapes.py:67-70 runs `benchmark_mla_decode_live.main` end-to-end in CPU…
18	NEEDS-INFO	The unit test files for the verification guards and runner checks (`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py`, `test_mla_decode_runner.py`, and `test_mla_decode_reference.py`) remain omitted from the diff, so…
19	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/builders/mla/verify_mla_decode.py:246-272`, `make_problem` allocates `len(k_lens) * max_blocks`…
20	N/A	This PR does not introduce or modify a detector or linter for banned…
21	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:160-291` and…
22	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:181-190`, `test_mla_decode_golden_records_no_build_errors` asserts that…
23	N/A	This PR is a functional bring-up of the MLA decode kernel and makes no…
24	NEEDS-INFO	The diffs for `dnn-providers/hip-kernel-provider/rocke/library/tests/golden/mla_prefill_gfx942_ir_sha256.json` (+2/-2 lines) and `dnn-providers/hip-kernel-provider/rocke/library/tests/golden/mla_decode_ir_sha256.json` remain omitted from the…
25	N/A	This PR introduces new MLA decode kernels rather than applying a codegen…
26	N/A	This PR makes no benchmark speedup claims against a baseline or external…
27	N/A	This PR does not report a performance A/B study toggling optimization knobs.
28	N/A	This PR adds a new MLA decode kernel path and does not present a…
29	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:240`, `exp2_fast` is called on `fsub(b.vec_extract(ml, 0), big_m)`…
30	NEEDS-INFO	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py` and the example READMEs (`library/examples/gfx942/mla_decode/README.md`, `library/examples/gfx950/mla_decode/README.md`) remain omitted from the diff, so we cannot…
31	N/A	This PR is a new kernel bring-up and does not claim to fix a performance…
32	NEEDS-INFO	The documentation files recording the compiler and resource tables (`dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md` and `gfx950/mla_decode/README.md`) and `library/tests/test_mla_decode.py` remain…
33	PASS	In `dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:333-364`, `_resolve_num_cus` was hoisted from `dispatch/mla/gfx942.py:95-113` to…
34	N/A	This PR does not add an opt-in output flag; `lse` is always a required…

§H graded 5/5

#	Verdict	Evidence
1	FAIL	Not resolved in the new commits: in dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:541 and :626, `check_contents: bool = False` remains default-off in both `MlaDecodeRunner.latent` and `MlaDecodeRunner.forward`, whereas the checklist requires contents checks that force a…
2	FAIL	Not resolved in the new commits: in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:693-706, the physical page ID from `block_table` is zero-extended to 64-bit (`page64 = b.zext(b.readfirstlane(page), I64)`) and multiplied by `stride` into the 64-bit base pointer…
3	PASS	In dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:222-224, `MlaDecodeSpec.__post_init__` enforces that `block_k` is a power of…
4	PASS	In dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:617-635 and :756-768, the kernel loops over the runtime tile range `[t_lo…
5	NEEDS-INFO	While `dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:415-426` adds negative tests for `num_seg` (`with_num_seg`), the test files…

§I graded 6/6

#	Verdict	Evidence
1	N/A	This PR is a functional bring-up of a single MLA decode layout per…
2	PASS	`dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:160` clamps `num_seg` to `mla_decode_max_num_seg` and `MAX_GRID_DIM` (repair)…
3	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_benchmark_shapes.py:48` tests curated model shapes from `mla_shapes.json`, and…
4	N/A	This PR does not add a multi-candidate routing benchmark or tuning table…
5	PASS	`dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:28` documents the checked-in IR SHA256 golden snapshot and provides the…
6	PASS	`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:614` uses existing general DSL/layout primitives (`TransposeLdsReader` and…

§J graded 7/7

#	Verdict	Evidence
1	NEEDS-INFO	The previous NEEDS-INFO is not resolved in the visible diff: the hunks for `dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md`…
2	NEEDS-INFO	The previous NEEDS-INFO is not resolved in the visible diff: the Markdown hunks for `dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md`…
3	NEEDS-INFO	The previous NEEDS-INFO is not resolved in the visible diff: the contents of `dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md`…
4	NEEDS-INFO	The previous NEEDS-INFO is not resolved in the visible diff: the Markdown hunks for `dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md`…
5	NEEDS-INFO	The previous NEEDS-INFO is not resolved in the visible diff: the Markdown hunks for `dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md`…
6	NEEDS-INFO	The previous NEEDS-INFO is not resolved in the visible diff: the hunks for `dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md`…
7	NEEDS-INFO	The previous NEEDS-INFO is not resolved in the visible diff: the hunks for `dnn-providers/hip-kernel-provider/rocke/library/examples/gfx942/mla_decode/README.md`…

§K graded 13/13

#	Verdict	Evidence
1	N/A	No tuned candidate or KnobSpace-backed candidate with multiple named…
2	N/A	The PR adds no KnobAxis or tunable knob fields; the new MLA spec fields are…
3	N/A	No tuning_id or config_key payload or tuning-store identity is changed…
4	N/A	No KnobSpace.defaults() or recorded() implementation is added or changed.
5	N/A	No knob inertness, base-value, refusal, or known-wrong knob policy is added…
6	N/A	No KnobSpace.validate path, tuning_specs policy-free constructor, or tuned…
7	N/A	No KnobSpace.base(req) or tuning-problem base resolver is added or changed.
8	N/A	No accepted knob-value boundary is added; MLA refuses all non-empty…
9	N/A	The added benchmark resolves and times live dispatch selections in memory…
10	N/A	No bare tuning_id lookup or full-space tuning-id resolution path is added…
11	N/A	No tuned sweep entry point, sampling entry point, pin_to_spec storage path…
12	PASS	The routing change between decode and prefill is explicitly pinned: decode selects at max_seqlen_q == 1 in…
13	PASS	The benchmark sends shapes through dispatch_mla in dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:79 and…

§L graded 10/10

#	Verdict	Evidence
1	N/A	This PR does not add or modify a gfx target in arch_specs.json or backend…
2	N/A	This PR does not add or modify any MMA op_id in the hardware catalog or…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:82-85 and :497-499 use the catalog helper wrappers mfma_16x16x16_for_dtype and…
4	N/A	This PR only targets gfx942 and gfx950 (both wave64) and does not alter…
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:165-171 sets SCORE_K=16 and PV_TRANSPOSE_READ=False on MlaDecodeSpec (gfx942)…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/decode.py:32-60 and dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:7…
7	N/A	This PR operates exclusively on bf16 and f32 tensors and does not add or…
8	N/A	This PR does not use or modify scaled MMA operations.
9	N/A	This PR does not add a new gfx architecture target to the platform.
10	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body.py:157-177 defines flavor_honors_agpr_alloc by delegating to _flavor_at_least, and…

§M graded 10/10

#	Verdict	Evidence
1	NEEDS-INFO	The diff still truncates dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py at line 805 before the end of build_mla_decode; the full hunk is needed to verify no device SSA Value is used in host-time Python control flow…
2	N/A	The diff does not touch legalize_for_target or target legalization code.
3	N/A	The diff does not change ck.dsl.ir/v1 serialization, print_ir interchange…
4	N/A	The diff does not change the portable IR, recipe, bundle schema, importer…
5	N/A	The new kernels do not use AsyncTileLoader for async DRAM-to-LDS copies.
6	NEEDS-INFO	The diff still truncates dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py at line 805 in the new kernel body; the full hunk is needed to verify all barriers, pointer loads, masking, and cross-lane operations in the…
7	PASS	The new buffer path constructs resources through the IR builder rather than hard-coding resource word 3 in…
8	PASS	The shared MLA flavor gate delegates to _flavor_at_least with AGPR_ALLOC_MIN_FLAVOR in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla_body…
9	N/A	The diff does not change public C entry points, C headers under…
10	N/A	The diff does not emit !dbg metadata or change debug-info/code-object…

§N graded 18/18

#	Verdict	Evidence
1	PASS	`MAX_GRID_DIM` in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264` is a non-SSA host constant shared across GEMM/wgrad…
2	PASS	The MLA decode kernels reuse shared helpers (`mfma_16x16x16_for_dtype`, `mfma_16x16x32_for_dtype`, `safe_inv_l`, `natural_lse_from_log2_stats`…
3	N/A	This PR does not modify or use `SoftwarePipeline`.
4	N/A	This PR does not add or modify `EpilogueOp`, `_PATTERN_TABLE`, or…
5	N/A	This PR does not add or modify any API in `rocke.helpers.tiling`.
6	N/A	This PR does not add or modify front-door/toolbox tiling helpers or…
7	N/A	This PR does not use `rocke.helpers.tiling` or custom tiling layout…
8	N/A	This PR does not use or modify `rocke.helpers.tiling` memory verbs…
9	N/A	This PR does not use or modify `rocke.helpers.tiling` N-D tensor axis roles…
10	N/A	This PR does not change `rocke.helpers.tiling` emission, LDS staging, or…
11	N/A	This PR does not use or modify `rocke.helpers.tiling` (`TileMma`, `Tiling`…
12	N/A	This PR makes no LDS bank-conflict or coalescing claims under…
13	N/A	This PR does not build or modify `WarpDistributionEncoding`…
14	N/A	This PR does not use or modify `rocke.helpers.tiling` data-movement verbs…
15	N/A	This PR does not modify `rocke.helpers.tiling` allocation or bounds checks.
16	N/A	This PR does not use or modify `optimize_layout` or `evaluate_transform` in…
17	N/A	This PR does not touch the `rocke.helpers.tiling` recorder, round-trip…
18	N/A	This PR does not touch `InterleavedStyle` or demo kernels under…

§O graded 16/16

#	Verdict	Evidence
1	NEEDS-INFO	The hand-built IR in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py is still truncated after line 865 before the end of build_mla_decode, so the complete emitted IR cannot be inspected.
2	NEEDS-INFO	While scf_for_iter and scf_yield in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode_combine.py:212-272 are complete, the loop bodies and yields in…
3	PASS	Deterministic IR bytes are pinned by the golden hash test in dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:138 and…
4	N/A	The PR's files under review do not modify IR type equality, dtype…
5	N/A	The PR's files under review do not touch core/passes.py…
6	N/A	The PR's files under review do not modify constant folding logic in either…
7	PASS	The decode kernel sets structured kernel attributes DECODE_AGPR_ALLOC = (0, 0) in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py…
8	N/A	The PR's files under review do not set or modify CodegenPolicy or…
9	NEEDS-INFO	The barrier and wait schedule inside build_mla_decode in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py is truncated after line 865, so the full loop body and epilogue synchronization cannot be verified.
10	N/A	The PR's files under review do not modify lowerer attribute validation or…
11	PASS	Architecture target IDs are resolved via base_arch_from_target_id in dnn-providers/hip-kernel-provider/rocke/library/dispatch/mla/common.py:337-345…
12	PASS	The per-axis grid cap is centralized as MAX_GRID_DIM in dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264 and aliased by…
13	N/A	The PR's files under review do not modify public C structs, enums…
14	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt is modified in the PR but its diff hunk is omitted, so CMake flags and target changes cannot be inspected.
15	PASS	The C++/Python lowering parity test enforces ROCKE_CPP_STRICT=1 and checks all served flavors in…
16	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md is modified in the PR but its diff hunk is omitted in this run, so the updated documentation cannot be inspected.

§P graded 14/14

#	Verdict	Evidence
1	PASS	Builders take `(spec, *, arch)`: `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:534`…
2	PASS	Signatures are constructed via `SignatureBuilder` with canonical pointer/scalar types (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_deco…
3	PASS	Harnesses take the launch grid directly from the dispatch result or `MlaDecodeLaunch` (`dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_dec…
4	PASS	The LDS gate in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:375` charges the exact `kv_lds + qr_lds` allocation…
5	N/A	This PR does not touch WMMA/RDNA GEMM, `GemmMultiD`, or `GemmMultiAbd`.
6	NEEDS-INFO	Resolved against the current diff: the split-KV workspace (`part_acc_ptr`, `part_ml_ptr`) does not require caller-zeroing because empty segments write `(NEG_INIT, 0)` to `part_ml` and the combine kernel skips reading `part_acc` when `l <= 0`…
7	N/A	This PR does not touch Stream-K, persistent GEMM, or `GroupedGemmLauncher`.
8	N/A	This PR implements bf16 MLA decode and does not perform quantization or MX…
9	N/A	This PR does not add or modify numeric emulation or quantized decode helpers.
10	N/A	This PR does not add or modify any manifest `kind` or touch…
11	PASS	Launches go through `KernelLauncher` (`dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:163`), workspace buffers are…
12	PASS	`get_device_max_grid_size` in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/hip_module.py:351` queries `_hipDeviceGetAttribute`…
13	PASS	The host-side constant `MAX_GRID_DIM = 65535` in `dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/spec.py:264` is value-identical to…
14	PASS	No `expected.json` files or `tflops_lower_bound` values are added, `dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_l…

§Q graded 7/7

#	Verdict	Evidence
1	NEEDS-INFO	The diff hunk for dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt (+19 -9 lines) is still omitted in the provided diff, so CMake test registration for the new tests cannot be verified.
2	PASS	MLA has no C++ instance mirror (documented as Python-only in known_gaps.md) and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden…
3	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:242 fails on BackendCoverageGap or zero compared cases…
4	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode_golden.py:130 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decod…
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/mla/test_mla_decode_wiring.py:470 tests _check_grids_fit_device without torch or a GPU…
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/_compiler_test_utils.py:37 preserves the exact AssertionError vs unittest.SkipTest semantics when…
7	N/A	This PR does not modify cross-engine backend resolution or ROCKE_BACKEND /…

§R graded 11/11

#	Verdict	Evidence
1	N/A	This PR is a functional bring-up of a previously missing operator regime…
2	PASS	Arch-specific recipes (`SCORE_K = 32` for `mfma_f32_16x16x32_bf16`, `PV_TRANSPOSE_READ = True` for `ds_read_tr16_b64`, and per-arch…
3	PASS	The kernel follows the utilities skills for LDS padding (`WT_PAD` vs `V8_PAD` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py…
4	NEEDS-INFO	Unresolved since previous review: `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:135` cites `TestResources` in `dnn-providers/hip-kernel-provider/rocke/library/tests/test_mla_decode.py` for static VGPR, scratch…
5	NEEDS-INFO	Unresolved since previous review: global memory accesses are added in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:581-672` (and in the truncated epilogue after line 808) and…
6	PASS	The kernels are built with raw `IRBuilder` outside `rocke.helpers.tiling` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:513`)…
7	N/A	The PR makes no bottleneck classification claims and states in the…
8	PASS	`b.sched_barrier(0)` and `DECODE_SCORE_PREFETCH` / `DECODE_SCORE_FENCE_EVERY` in `dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py…
9	N/A	This PR is not a performance regression fix and does not retune an existing…
10	PASS	Compiler hazards around `agpr_alloc` on `llvm20` are gated via `flavor_honors_agpr_alloc` (`dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/_mla…
11	N/A	No new optimization failure mode, heuristic study, or `dsl_probes` tool is…

§S graded 8/8

#	Verdict	Evidence
1	NEEDS-INFO	Still NEEDS-INFO: `library/tests/dispatch/mla/test_mla_decode_wiring.py:285-526` and `library/tests/test_mla_decode_golden.py:77-246` exercise `MAX_GRID_DIM`, `MlaDecodeLaunch`, `mla_decode_launch`, `resolve_num_cus`, and the decode/combine…
2	PASS	`MLARequest` in `library/dispatch/mla/common.py:137-158` adds required `max_seqlen_q: int` and changes `op` default to `MLA_OP` (`mla_fwd`); all callers…
3	PASS	`dispatch.mla.gfx942._resolve_num_cus` was moved/renamed to `dispatch.mla.common.resolve_num_cus` (`library/dispatch/mla/common.py:351`) and all…
4	PASS	`MAX_GRID_DIM` is defined in `platform/python/rocke/helpers/spec.py:264` and added to `__all__` in `platform/python/rocke/helpers/spec.py:33` and…
5	NEEDS-INFO	Still NEEDS-INFO: `library/tests/dispatch/mla/test_mla_decode_wiring.py:119-237,429-526` and `library/tests/test_mla_decode_golden.py:193-209` assert exact IR hashes and negative `ValueError`/`TypeError` matches for dispatch and…
6	PASS	CPU-checkable contracts (IR golden hashes, C++/Python byte-identity, spec/request validation, and torch-free device grid limit checks) run on CPU without…
7	PASS	`library/tests/test_mla_decode_golden.py:212-246` (`test_mla_decode_cpp_python_byte_identity`) lowers all 14 MLA decode cases through both…
8	PASS	Resolved since the previous review: the PR body's 'Test plan' section now explicitly states 'From `rocke/platform/`: `tests/test_rocke.py` -- 368 tests OK…

yraparti · 2026-11-05T16:15:53Z

+def require_tools_or_skip(reason: str) -> NoReturn:
+    """Skip the running test for ``reason``, or fail it with ``AssertionError``
+    when ``ROCKE_REQUIRE_COMPILER_TOOLS=1`` says the tools must be there."""
+    if os.environ.get(REQUIRE_COMPILER_TOOLS_ENV) == "1":


[HIGH] §D.1 ★ ★ Knobs/features on the KernelSpec (validator + selector + name tag + cache-key), NOT os.environ (env only as a selector-level diagnostic override folded onto the spec). Reviewers may reject env vars outright — even a reproducibility/pin one (#9739: "please dont use env variables", on a…

dnn-providers/hip-kernel-provider/rocke/library/tests/_compiler_test_utils.py:40 reads the environment

Fix: Put the knob on the KernelSpec and delete the environment read.

yraparti · 2026-11-05T16:15:53Z

+            "num_seqs": int(self.request.num_seqs),
+            # The kernel's row pitch: the stride, which a column-sliced view
+            # does not share with its shape.
+            "block_table_stride": int(block_table.stride(0)),
+            "max_k": int(self.request.max_seqlen_k),
+        }
+        config = LaunchConfig(grid=self.grid, block=self.block, stream=stream)
+        if self._combine is None:
+            vals.update(out_ptr=out, lse_ptr=lse)
+            return [(self._latent, vals, config)], out, lse
+        with _on_stream(stream, qa.device):
+            part_acc, part_ml = self.workspace(qa.device)
+        vals.update(part_acc_ptr=part_acc, part_ml_ptr=part_ml, num_seg=spec.num_seg)
+        combine_vals = {
+            "out_ptr": out,
+            "lse_ptr": lse,
+            "part_acc_ptr": part_acc,
+            "part_ml_ptr": part_ml,
+            "num_seg": spec.num_seg,
+        }
+        combine_config = LaunchConfig(
+            grid=self.launch.combine_grid(),
+            block=self.launch.combine_block(),
+            stream=stream,
+        )
+        return (
+            [
+                (self._latent, vals, config),
+                (self._combine, combine_vals, combine_config),
+            ],
+            out,
+            lse,
+        )
+
+    def latent(
+        self,
+        *,
+        qa: torch.Tensor,
+        c_kv: torch.Tensor,
+        k_rope: torch.Tensor,
+        cu_seqlens_q: torch.Tensor,
+        cu_seqlens_k: torch.Tensor,
+        block_table: torch.Tensor,
+        scale: float,
+        out: torch.Tensor | None = None,
+        lse: torch.Tensor | None = None,
+        stream: int = 0,


2 checklist findings on this snippet

[HIGH] §H.1 ★ (line 541) ★ A hardware/resource bound derived from SPEC fields (e.g. an AMDGPU buffer-SRD num_records sized from num_kv_blocks*block_size*num_kv_heads*head_size) MUST be matched by the launch tensors — host-validate their shape/size against the spec that sizes the bound BEFORE compile/launch, else the…

Not resolved in the new commits: in dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:541 and :626, check_contents: bool = False remains default-off in both MlaDecodeRunner.latent and MlaDecodeRunner.forward, whereas the checklist requires contents checks that force a device-to-host sync to be gated behind a default-on flag (off-able on the hot/graph-captured path, which bind_latent at line 568 already provides).

Fix: Change check_contents: bool = False to check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward (keeping check_contents=False in bind_latent).

[MEDIUM] §E.15 (line 495) No dead parameters — and a parameter ACCEPTED then IGNORED is worse than one absent: the caller believes it was used, and in measurement code it silently discards a real observation. Grep each parameter name in the body. (#12070: substitute_share(total_armA_us=…) never read it — the measured…

Not resolved in the new commits: dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:581-583 still declares scalar kernel parameters num_seqs and max_k (b.param("num_seqs", I32) and b.param("max_k", I32)) whose return handles are discarded and never read in the kernel body, while dnn-providers/hip-kernel-provider/rocke/library/builders/mla/mla_decode_runner.py:495-499 passes self.request.num_seqs and self.request.max_seqlen_k into them on every launch.

Fix: Remove the unused num_seqs and max_k scalar parameters from mla_decode_signature and build_mla_decode (and pass them only in run_prefill_oracle when invoking the prefill kernel), or use them in-kernel.

yraparti · 2026-11-05T16:15:53Z

+    c_chunks = (BK * R_KV) // (THREADS * C_STAGE_W)
+    c_chunks_per_row = R_KV // C_STAGE_W
+    KR_W = _kr_stage_width(spec)
+    kr_chunks = (BK * D_ROPE) // (THREADS * KR_W)


[HIGH] §H.2 ★ ★ Guard indirection VALUES, not just indices — a masked/bounds-guarded table read protects the lookup INDEX (e.g. logical page < n_pages), NOT the value it returns. Any value used as a memory index (block-table physical id) must itself be range-checked, or shown contained by a HW bounds-check —…

Not resolved in the new commits: in dnn-providers/hip-kernel-provider/rocke/library/kernels/mla/mla_decode.py:693-706, the physical page ID from block_table is zero-extended to 64-bit (page64 = b.zext(b.readfirstlane(page), I64)) and multiplied by stride into the 64-bit base pointer (b.global_ptr_add(cache, offset)), which bypasses the SRD bound check for out-of-range page IDs (and negative page IDs like -1 zero-extend to 0xFFFFFFFF), while host-side validation in MlaDecodeRunner.latent (mla_decode_runner.py:541) and MlaDecodeRunner.forward (mla_decode_runner.py:626) still defaults to check_contents: bool = False.

Fix: Default check_contents: bool = True in MlaDecodeRunner.latent and MlaDecodeRunner.forward, or range-check page against num_pages on device before forming the 64-bit pointer offset.

yraparti · 2026-11-05T16:15:53Z

+def _time(
+    shapes, *, arch: str, warmup: int, iters: int, repeats: int, seed: int
+) -> int:
+    import torch


[MEDIUM] §E.2 Library doesn't import torch for device queries — use rocke.runtime.hip_module (prefer hipDeviceGetAttribute).

dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_mla_decode_live.py:169 adds import torch

Fix: Query the device through rocke.runtime.hip_module instead of importing torch.

cgmillette · 2026-11-06T15:00:25Z

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

PR #13923 — add MLA decode kernel for gfx942 and gfx950 (aosewski, AICK-1504) — 🟠 Follow-up: fixes hold; alignment table restated

Prior item	Status
Ours (11-04) N5: `TestResources…[4-gfx942]` fails in CI	🟢 FIXED `58123e26` • ✓ on both gfx94X shards at head, CI's compiler (read: CI logs) • wave-uniform row via `readfirstlane`; only the 18 split golden entries moved (ran)
Ours (11-04) N1: flavor rule asks the datalayout	🟢 FIXED `f48c4f35` • `_flavor_at_least(flavor, "llvm22")`: same lanes today (ran) • a simulated newer lane is served; its golden fails until blessed (ran: probe)
Ours (11-04) N2: third home for `require_tools_or_skip`	🟢 FIXED `f48c4f35` • `tests/_compiler_test_utils.py` is 14218's text minus its context manager (ran: diff) • both gates call it; a refused option fails under the variable (ran: probe; reproduces at `047d7fde`) • follows the recorded owner; 14218 keeps its module when it lands
Ours (11-04) N3: int checks raise `ValueError`	🟢 FIXED `f48c4f35` • one `require_int_fields`: `TypeError`, one text, for decode, combine, prefill (ran) • `True`, `2.0`, `1.9` refused; `with_num_seg` relies on it (ran)
Ours (11-04) minor: runner grid check repeats `_check_grids`	🟢 Answered, holds • both read one grid source (read) • platform limit at selection, device limit at launch (ARCHITECTURE.md §6.2) (read)
Ours (11-02) N1: decode re-inlines MLA family emitters	🟡 DEFERRED, tracked • `lane_ids`, `softmax_consts`, LSE helper reused in `6808bc45`; goldens untouched, green (ran) • rest in AICK-3039 under AICK-1504, after 14502; `known_gaps.md` §6 (read) • holds: the rest moves the IR (re-bless, device runs)
Copilot (11-05) ×10: alignment, aliasing, int32 wrap, device, CPU tensors, streams ×2, key length, `-1e30`, §5 heading	🟢 FIXED `f48c4f35` ×9 • Copilot's examples refused at head, accepted at `047d7fde` (ran: GPU tensors) • device, stream items: runner tests 61 ✓ with torch (ran); CI skips that file (no torch) • `-1e30` floor documented, family-wide: holds · 10 threads open: aosewski resolves
Copilot (10-19) ×7	🟢 FIXED, hold • same runner tests 61 ✓ at head (ran); 7/7 threads resolved
yraparti (11-03) §B.4, §H.3, §O.11, §S.8	🟢 FIXED `934f5bb4`, `f48c4f35` · threads open • §B.4, §S.8 PASS in their head re-review (read) • last `split(":")` sites now `base_arch_from_target_id` (ran: grep)
yraparti (11-03→11-05) §H.1, §H.2: page ids unchecked by default	🟢 Answered, holds · threads open • host check opt-in (it syncs); verify harness opts in (read) • develop's prefill trusts page ids too; their re-review keeps FAIL: yraparti, aosewski settle
yraparti (11-03→11-05) §E.2: torch for device queries	🟡 OPEN, in part · thread open • ⚪ our 11-04 N/A: the benchmark's device-name print asks torch (`:186`) (read) • new in `f48c4f35`: runner's device index asks torch (`mla_decode_runner.py:186`), not `hip_module.current_device()` (read) • fix: `hip_module` for both; torch stays for tensors, streams
yraparti (11-03→11-05) §D.1: env read in test helper	🟢 N/A · thread open • repo Rule 3 exempts test-harness env reads (read) • the variable is develop's (TESTING.md §5.1.4); the module is 14218's (read)
yraparti (11-03) §E.15: unused `num_seqs`, `max_k`	🟡 DEFERRED · thread open • in AICK-3039: drop both, re-bless, device runs (read) • holds
CongMa13 (10-22) APPROVED	🟢 APPROVED `fa8e6110` • older head; still counts on develop
cgmillette (10-28) APPROVED	🟢 APPROVED `87fb6ff2` • older head; still counts on develop • since: llvm24 serving, spill fix, runner checks (author's)
Bots / CI	🟡 Red, infra • fails: Windows gfx110X shard 1/2, both attempts — `rocke_pytest` timeouts on the replaced Azure runners, as on develop today; the PR's test there ✓ (read) • 45 ✓ · 2 ✗ (that shard, its Summary) · 50 skipped; gfx94X shards ✓, MLA lane 489 ✓ / 3 skip (read) • merge: blocked: Multi-Arch Summary until a re-run; 26/33 threads open; no rocke-lib approval
Our checks	🟢 Clean • merge vs `614e241a12` (adds 14448, 14629; CI ran before both): clean, 59 behind; lane registrations, grid query, MLA tests 449 ✓, runner 61 ✓ (ran) • at `60b43a5a52`: build, ctest 51 ✓ / 1 skip; parity ×4, goldens, differential ✓; the PR's tests 565 ✓ / 5 skip; installed 455 ✓ / 5 skip (ran) • whole `rocke_pytest` there 6599 ✓, 2 ✗ on our install's missing example data, ✓ in CI; decode byte identity ✓; rules, ABI clean (ran)
Code quality	🟠 Should-fix • runner restates the kernels' pointer alignment, untested: N6 (ran) • our N1–N3 fixed; the old flavor rule fails two tests (ran) • minor: N7; third `_rope_tables` (home: `ref_mla_attn.py`); `forward`'s own stream switch beside `_on_stream` (read)

#	Sev	file:line	Finding
N6	🟠	`builders/mla/mla_decode_runner.py:50`	`_POINTER_ALIGN` restates the kernels' `align=`; no test ties them • impact: equal today (ran); widen a kernel's `align` and the runner and every MLA decode test stay silent (ran: probe) • develop since 14448: `KernelLauncher.from_artifact` checks each declared `align` at every launch and `bind`; launchers built from bytes, as here, check nothing (read) • fix: merge develop, build both launchers with `from_artifact`, drop `_POINTER_ALIGN`; they then follow the declaration (ran: emulated)
N7	🟡	`kernels/mla/mla_decode_combine.py:78`	`NEG_INIT = NEG_INF` keeps a second name alive as an alias • impact: callers split: the latent kernel uses `NEG_INF`; the combine (`:216`) and a test use `NEG_INIT` (read) • fix: `NEG_INF` at both; drop `NEG_INIT` and its `__all__` entry (API_DESIGN_STANDARD Part 1 rule 5); IR unchanged

Utility: 🟢 Positive — MLA decode on two arches; review fixes verified; alignment table open

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: the delta since 047d7fde (three commits, one develop merge), merge with develop (build + ctest, parity, goldens, differential golden, the PR's tests, installed lanes, whole rocke_pytest at 60b43a5a52; lane registrations and the PR's tests at 614e241a12), saved probes, Copilot's examples on GPU tensors, runner tests with torch, an alignment-drift probe, CI logs at head (both attempts), the follow-up ticket, threads, code quality. Flags whether a closer look is needed; not an approval or rejection.

Copilot

🟡 Changes recommended

The supported r_kv=1024 combine path emits an unsupported 16-element bf16 store and assumes stronger alignment than its output pointer guarantees, causing build failure or undefined behavior.

6 open issues

7 resolved since last review

_{Review effort: Max}

+    out_row = b.add(
+        b.mul(b.add(b.mul(token, b.const_i32(H)), head), b.const_i32(R_KV)), col
+    )
+    b.global_store_vN(out, out_row, b.vec_trunc_f32_to_bf16(vals), CPL, align=2 * CPL)


+    out_row = b.add(
+        b.mul(b.add(b.mul(token, b.const_i32(H)), head), b.const_i32(R_KV)), col
+    )
+    b.global_store_vN(out, out_row, b.vec_trunc_f32_to_bf16(vals), CPL, align=2 * CPL)


+        errors = _request_errors(req)
+        if errors:
+            return False, "; ".join(errors)


aosewski requested review from a team as code owners October 19, 2026 10:34

Copilot AI balanced review requested due to automatic review settings October 19, 2026 10:34

github-actions Bot added documentation project: hip-kernel-provider rocke: platform Touches rocKE platform labels Oct 19, 2026

Copilot started reviewing on behalf of aosewski October 19, 2026 10:35

View session

assistant-librarian Bot added the organization: ROCm label Oct 19, 2026

Copilot AI reviewed Oct 19, 2026

View reviewed changes

AnaghaRaoAMD force-pushed the users/anarao/mla-prefill branch from 9d1511e to b26bc7f Compare October 20, 2026 00:49

aosewski marked this pull request as draft October 20, 2026 12:48

AnaghaRaoAMD force-pushed the users/anarao/mla-prefill branch from b26bc7f to 666a141 Compare October 20, 2026 16:55

Base automatically changed from users/anarao/mla-prefill to develop October 20, 2026 19:53

aosewski and others added 13 commits October 21, 2026 04:40

fix(rocke): let gather_paged_kv handle an empty key range (AICK-1504)

94d1b98

reshape(-1) cannot infer an axis from a zero-element tensor, so a sequence with S_k = 0 crashed the shared MLA reference. Pass the latent and rope widths explicitly. Co-Authored-By: Claude <noreply@anthropic.com>

docs(rocke): document MLA decode and its correctness records (AICK-1504)

125dd25

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

build(rocke): register MLA decode tests in CTest

c6c1f8a

Adds the decode CPU tests to the library MLA lane and the numeric test to the GPU lane, so installed-test runs cover decode like prefill. Co-Authored-By: Claude <noreply@anthropic.com>

Copilot finished work on behalf of aosewski November 2, 2026 12:49

Copilot started work to fix failing Actions workflows on behalf of aosewski November 2, 2026 12:53

View session

Copilot finished work on behalf of aosewski November 2, 2026 13:09

aosewski unassigned Copilot Nov 2, 2026

yraparti reviewed Nov 3, 2026

View reviewed changes

aosewski and others added 3 commits November 4, 2026 07:58

Merge remote-tracking branch 'origin/develop' into users/aosewski/aic…

0b643fe

…k-1504-mla-decode

Merge remote-tracking branch 'origin/develop' into users/aosewski/aic…

047d7fd

…k-1504-mla-decode

yraparti reviewed Nov 4, 2026

View reviewed changes

tenpercent requested a max review from Copilot November 5, 2026 00:24

Copilot started reviewing on behalf of tenpercent November 5, 2026 00:24

View session

Copilot AI reviewed Nov 5, 2026

View reviewed changes

aosewski and others added 4 commits November 5, 2026 06:37

Merge remote-tracking branch 'origin/develop' into users/aosewski/aic…

1b26101

…k-1504-mla-decode

yraparti reviewed Nov 5, 2026

View reviewed changes

tenpercent requested a max review from Copilot November 6, 2026 19:54

Copilot started reviewing on behalf of tenpercent November 6, 2026 20:04

View session

Copilot AI reviewed Nov 6, 2026

View reviewed changes

		_max_grid_size: Dict[int, Optional[Tuple[int, int, int]]] = {}


		def get_device_max_grid_size(device: int = 0) -> tuple[int, int, int] \| None:

Conversation

aosewski commented Oct 19, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Problem

Fix

Test plan

Known limits and follow-ups

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

therock-pr-bot Bot commented Oct 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Oct 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

cgmillette commented Nov 2, 2026

PR #13923 — add MLA decode kernel for gfx942 and gfx950 (aosewski, AICK-1504) — 🟠 Follow-up: decode duplicates MLA family code

Uh oh!

aosewski commented Nov 3, 2026

Uh oh!

yraparti left a comment

Choose a reason for hiding this comment

rocKE checklist review

Findings

Needs info

Section notes

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

aosewski commented Nov 4, 2026

Uh oh!

cgmillette commented Nov 4, 2026

PR #13923 — add MLA decode kernel for gfx942 and gfx950 (aosewski, AICK-1504) — 🟠 Follow-up: flavor rule; duplicate gate helper

Uh oh!

cgmillette commented Nov 4, 2026

Uh oh!

yraparti left a comment

Choose a reason for hiding this comment

rocKE checklist re-review of the fixes since c41875b42c05

Since the last review (c41875b42c05)

Findings

Needs info

Section notes

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

cgmillette commented Nov 4, 2026

PR #13923 — add MLA decode kernel for gfx942 and gfx950 (aosewski, AICK-1504) — 🔴 Follow-up: new spill test fails in CI

Uh oh!

aosewski commented Oct 19, 2026 •

edited

Loading

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

edited

Loading

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

edited

Loading

rocKE checklist re-review of the fixes since `c41875b42c05`

Since the last review (`c41875b42c05`)

rocKE checklist re-review of the fixes since `047d7fde0a25`

Since the last review (`047d7fde0a25`)