Skip to content

feat(hip-kernel-provider): paged/ragged SDPA engine for gfx942/gfx950 - #14247

Open
dpeabody-amd wants to merge 9 commits into
developfrom
users/dpeabody/hipdnn-cdna-paged-attn
Open

dpeabody-amd wants to merge 9 commits into
developfrom
users/dpeabody/hipdnn-cdna-paged-attn

Conversation

@dpeabody-amd

@dpeabody-amd dpeabody-amd commented Oct 22, 2026 •

Copy link
Copy Markdown
Contributor

ISSUE ID : AICK-1847

Draft. Based on develop (1702865), which includes #15024 (merged as 6ea2887), #14448, #14259, #15080, #14629, #14822, #14907, #14908 and #11642. It builds on the runtime head-count, window and batch kernel arguments #15024 added to the gfx942/gfx950 tiled 2D kernels. The switch of #14254's gfx1151 paged matcher to the shared paged checks needs #14254 (not merged yet) and lands as a follow-up commit after it (see Shared paged-graph checks). The engine follows the paged SDPA contract of #14358 (merged) and is checked against hipDNN's CPU SDPA reference. It serves cuDNN's FLOAT sink_token, which hipDNN's frontend admits, and also a sink in the Q data type (see Attention sinks).

Summary

Adds hipkernel:AttentionPagedUnified, a kernel-ingestor native engine for gfx942 and gfx950. It serves the paged K/V SDPA forward graph of #14358 with one launch of the rocKE tiled 2D unified attention kernel, in a new runtime_layout mode. It covers the two shapes a serving framework needs from one plan:

  • uniform decode or spec-decode: a dense Q over a paged K/V cache;
  • mixed prefill/decode batches: a ragged Q/O over a paged K/V cache.

Lengths, block ids and ragged offsets are read on the device. The engine needs no workspace and makes no host sync or allocation in execute, so the plan can be captured into a HIP graph and replayed with new lengths, block ids and data.

It relies on the engine plugin API 1.4.0 that develop already reports (since #14202), so hipDNN offers the provider graphs with ragged tensors and token-unit ragged offsets.

What changes

  • rocKE runtime_layout (kernels/{gfx942,gfx950}/attention_tiled_2d.py), a new spec field:

    • It builds on perf(rocke): runtime heads, window and batch for gfx942/gfx950 tiled attention #15024's runtime kernel arguments (query and KV head counts, sliding window, and the batch trip count of the ragged query-block search). On top of them the mode makes the Q/O/K/V strides and whether Q is dense or ragged kernel arguments too, and appends 16 kernel arguments after perf(rocke): runtime heads, window and batch for gfx942/gfx950 tiled attention #15024's, the last being sink_is_f32 (see Attention sinks). Its window mask is always on, with perf(rocke): runtime heads, window and batch for gfx942/gfx950 tiled attention #15024's sliding_window argument as its size, and a dense Q needs no query-block search: it runs the search with zero iterations and reads no offset.
    • The page tables and sequence lengths are read on the device, and a grid-stride loop covers every query block.
    • K pages through block_tables_ptr and V through a new block_tables_v_ptr argument. Both share block_table_stride, and they may be the same pointer, so a graph with separate K and V page tables (cuDNN's form) is served as well as vLLM's one shared table.
    • K/V use one 64-bit buffer resource per page, built once per tile load from one page-table read (a runtime_layout tile never spans more than one page). On gfx942 this sits beside the per-page i64 loader of perf(rocke): build one K/V resource per page on the gfx942 2D i64 path #14349 and uses the same mechanism; on gfx950 it replaces a page-table read per copy.
    • Page-table and K/V reads stop at seq_len_kv. Rows with no visible key write zeros, and so do a dense Q's rows past seq_len_q (up to S_q), as feat(hipdnn): SDPA CPU and GPU references cover the full SDPA graph API #14358 requires. A ragged Q writes only its live rows, the first seq_len_q[b] rows from its origin: under RFC 0014 as feat(hipdnn): ragged tensors with per-batch origins, lengths and strides (ALMIOPEN-3151) #15080 rewrote it, the rows of an offset block past seq_len_q[b] belong to no sequence and the buffer may end at the last live row, so the kernel leaves them alone and never writes past it. A query block before the first origin (ragged_offset[0] > 0) finds no sequence; it reads no entry at index -1 and writes nothing.
    • The spec checks refuse a num_warps outside {1, 2, 4, 8} before any arithmetic that divides by it, in Python and in C++.
    • LDS: on gfx942 the support gate (supports_tiled_2d_spec) and both builders, Python and C++ (rocke_gfx942_attention_tiled_2d_runtime_layout_lds_bytes), charge a runtime_layout kernel's LDS against ArchTarget lds_capacity_bytes and refuse a spec over it, with the same text. On gfx950 the general tiled 2D LDS gate, which the C++ support check already mirrors, applies.
    • runtime_layout_catalog_spec() builds the catalog cells.
    • The rocKE C++ twin (platform/cpp/instances/gfx9{42,50}/attention_tiled_2d_*.cpp) mirrors the mode op for op: the runtime_layout spec field (ROCKE_ABI_VERSION and BINARY_ABI 13, develop being 12 since refactor(rocke): model indexed MMA operands #11642; feat(hip-kernel-provider): paged/ragged SDPA in Gfx1151WmmaAttention #14254 also bumps it, so whichever of the two lands second takes 14), the appended kernel arguments (shared in the new instance_attention_tiled_2d_rtl_internal.h), the dense-or-ragged q-block mapping and grid-stride loop, the device-read page tables and lengths, the per-page 64-bit buffer resource with V through block_tables_v_ptr, the dense padded-row zero stores and live-rows-only ragged stores, sinks in both forms, the rtl kernel-name token and both architectures' runtime-layout spec checks. The rocke_engine binding's fastKV builder now reads runtime_layout too, so it refuses a fastKV runtime_layout spec as Python does. The parity emitters gain runtime-layout configs (gfx942 indices 20 to 38, after fix(rocke): give the gfx942 cfvst V loader an i64 path #14629's cfvst configs 18 and 19, and gfx950 indices 72 to 90, after perf(rocke): runtime heads, window and batch for gfx942/gfx950 tiled attention #15024's D192 config 71: every catalog head size and block size cell, both dtypes, GQA 1 to 16, num_kv_heads > 1, with and without sinks), so check_byte_identity.py gates the mode, and the differential golden (llvm_gfx_all.json, .ll for llvm20/22/23/24 and canonical IR) records them, so CI pins their IR. The prove_parity_binding mirror (attention_tiled_parity_configs.py) also gains fix(rocke): give the gfx942 cfvst V loader an i64 path #14629's gfx942 configs 18 and 19, which develop added to the emitters only, so the mirror stays index for index with the emitters.
    • The representative-IR golden (platform/tests/golden/rocke_representative_ir_sha256/attention.json, refactor(rocke): split the representative-IR golden into per-family shards #14679's per-family shard) pins six runtime_layout objects at llvm20/22/23/24, three per architecture: one without sinks and two with them (D128, and D256, which on gfx942 is the 2-warp body). Dense or ragged Q and the sink form are launch arguments, so these objects cover both layouts and both sink forms.
    • Docs: rocKE's SUPPORT_MATRIX.md gains a tiled 2D runtime_layout=True row and dsl_docs/instances/attention.md a "Runtime layout" subsection.
    • runtime_layout=False emits unchanged IR.
    • Register pressure: the runtime strides live in SGPRs, and most runtime_layout kernels spill some of them into VGPR lanes (never to scratch). The engine README's "Register pressure" section gives the counts over the shipped catalog: 264 of 288 gfx942 and 232 of 288 gfx950 kernels spill (at most 57 and 76 SGPRs), the lane moves are a median 1.7% and 1.5% of a kernel's instructions, and the fixed-layout twins spill none on gfx942 and 1 SGPR in 4 kernels on gfx950 (2 kernels and 6 SGPRs before perf(rocke): runtime heads, window and batch for gfx942/gfx950 tiled attention #15024), so the spills come from the runtime strides, not from perf(rocke): runtime heads, window and batch for gfx942/gfx950 tiled attention #15024. Removing them means fewer runtime arguments, so the cost is accepted. The runtime cost is not measured: no fixed-layout kernel serves the same paged graph.
    • On gfx950, additive bias and the longest-first query-block order (use_qblock_reverse, whose reversal does not match the runtime_layout grid-stride range) cannot be combined with runtime_layout, and the runtime-layout checks reject them, in Python and in the C++ twin. On both architectures they also reject emit_lse (the optional LSE output of feat(rocke): optional LSE output on unified tiled 2D attention #14616). Their head-count check is perf(rocke): runtime heads, window and batch for gfx942/gfx950 tiled attention #15024's shared rule (check_head_counts, rocke_attention_check_head_counts in C++), with its text. The gfx942 4-warp GQA builder refuses a runtime_layout spec instead of building an _rtl kernel that ignores the runtime layout.
    • Dispatch never sets the flag. It is a semantic field, so no tuning sweep varies it, and while it is off it stays out of the tuning spec's cache key and config defaults, so every stored spec_hash, config_key and pin is unchanged.
    • New tests: test_gfx942_attention_runtime_layout.py and test_attention_tiled_2d_runtime_layout_gfx950.py. On the device they include scenarios where V sits in other blocks than K through its own table, and sequences with 0 < seq_len_kv < seq_len_q, whose leading rows see no key. On the host they check every catalog cell and the rejection matrix of both architectures, in Python and, through the rocke_engine binding, in the C++ twin. The device cases use rocke.runtime.hip_module and numpy only, with no torch.
    • CTest: a new rocke_library_attention_runtime_layout_pytest entry (labels host, quick) runs the host checks of both files with -m "not gpu", and rocke_library_gpu_pytest runs both files on GPU runners. Neither file needs torch, so TheRock's gfx942 runners run the gfx942 device cases (TheRock has no gfx950 job, see Known integration work). ROCKE_ENGINE_test_categories_external.yaml puts the new entry in the quick and ffm-quick tiers.
  • Engine hipkernel:AttentionPagedUnified (packs/AttentionPagedUnifiedNative.cpp, descriptors under descriptors/rocKE/attention_paged_unified/, generated from IngestorGenerator/configs/attention_paged_unified.yaml). There is one KDP per architecture with 288 objects each, sink variants included. The UED sdk_version is 1.4.0. It serves:

    • fp16 / bf16;
    • head size 64, 128 or 256 and block size 16, 32 or 64;
    • GQA/MQA;
    • bottom-right causal with an optional left bound (cuDNN meaning: left_bound = L keeps the L keys ending at the diagonal, so a vLLM window W is left_bound = W);
    • INT32 page tables [B, 1, M, 1]: one tensor for K and V, or separate page_table_k and page_table_v tensors with the same width M and batch stride;
    • seq_len_q and seq_len_kv, both required, as for every paged graph;
    • max_seq_len_kv set or unset (the kernel reads only seq_len_kv, which the inferred extent bounds the same way);
    • a dense Q, or a ragged Q/O sharing one token-unit offset tensor in vLLM's query_start_loc layout (see Known integration work);
    • K/V containers with any equal strides, including an interleaved K|V cache.

    Everything else is declined, and the cause is logged at INFO. That includes graphs feat(hipdnn): SDPA CPU and GPU references cover the full SDPA graph API #14358 makes valid but this kernel cannot read: one-sided paging (one_sided_paging), K and V page tables of different widths or batch strides and packed page tables (page_table), and differing K and V block sizes (block_size). A ragged Q/O whose ragged_offset_multiplier is not its token stride is declined too (ragged): RFC 0014 allows offsets in other units, but the kernel reads stored offsets as token counts. The ASM SDPA engine declines it the same way, and every known producer (the vLLM binding, the SdpaRaggedFprop sample, the golden bundles) uses the token stride. A declared tensor alignment (engine API 1.3.0) below what the kernels load is declined with cause alignment: Q, K, V, O and sink_token must declare a multiple of 16 bytes, and the INT32 page tables, seq_len_q, seq_len_kv and ragged offsets a multiple of 4. Because the UED declares sdk_version 1.4.0, ragged and aligned graphs reach this matcher rather than stopping at the ingestor's schema gate, so the engine declines what it cannot serve itself. Other declines: softcap, ALiBi, FP8 scales, stats, dropout, attn_mask, non-paged or ragged K/V, other masks, left_bound = 0, unsupported widths, ratios or strides, a zero ragged multiplier, and graphs past the kernel's index limits. The engine README lists the causes and carries the engine-specific paged notes (preconditions, one plan for every step, vLLM binding).

  • Shared paged-graph checks (packs/SdpaGraphMatch.hpp, from feat(hip-kernel-provider): add gfx950 unified attention 2D/3D ingestor engines #14242). The engine uses the header's SDPA match helpers (findTensor, sdpaNode, hasSdpaForwardNode, describeOperand, sdpa_graph::DeclinedBy, the axis constants) instead of private copies. The stricter operand checks it needs (usesOverrideShapes, hasNonNegativeStrides, declaresAlignment, isInt32Vector) and the cuDNN paged-K/V graph checks (SEM-009) move into the header too, so every pack serving a paged graph applies one set of rules with one cause vocabulary, logged with the calling engine's name:

    • pagesKAndV: one_sided_paging;
    • pagedContainersAgree: block_size, shape (one block size, one block count; head counts and head sizes stay the engine's to check);
    • pageTablesFor: page_table (INT32 [B, 1, M, 1], unit entry stride, not packed; separate K and V tables share M and the batch stride);
    • sequenceLengthsFor: seq_len_kv, seq_len_q (both required, INT32 [B, 1, 1, 1]);
    • raggedQueryFor, countsOffsetsInTokens: ragged (Q and O through one [B+1, 1, 1, 1] offset table counted in tokens; K/V never ragged);
    • declaresServedAlignment: alignment (vector operands at the engine's vector width, INT32 operands at 4 bytes);
    • sinkFor: sinks (FLOAT or the Q data type, [1, Hq, 1, 1], unit head stride; it reports which form), and alignment below the engine's own minimum sink alignment;
    • takesNoAttentionMask: attn_mask.

    The reasons they log are the texts feat(hip-kernel-provider): paged/ragged SDPA in Gfx1151WmmaAttention #14254's gfx1151 paged matcher logs for the same rules, so that engine can call them without changing what it reports; where feat(hip-kernel-provider): paged/ragged SDPA in Gfx1151WmmaAttention #14254 now logs a different text (its sink and two ragged-offset reasons), its switch keeps those checks private. The engine keeps its own kernel limits and the order of its declines. The header's checkedProduct is the one overflow-checked product that admits zero factors (an extent of 1 contributes (1 - 1) * stride), unlike pack_features::checkedProduct; this engine and the gfx942 unified 3D pack use it instead of local copies. FlydslSdpaShared.hpp, which already takes computesInFloat from the header (feat(hip-kernel-provider): add gfx942 dense attention AOT ingestor engine #14908), now takes checkedProduct and declaresAlignment from it too instead of its identical copies. The commit that switches Gfx1151WmmaAttentionNative.cpp (feat(hip-kernel-provider): paged/ragged SDPA in Gfx1151WmmaAttention #14254) to these checks, keeping its cause keys, reasons, sink alignment rules and gfx1151 kernel limits, touches only that file. It cannot build on this branch without feat(hip-kernel-provider): paged/ragged SDPA in Gfx1151WmmaAttention #14254, so it is kept on a combined test branch and lands as a follow-up commit once feat(hip-kernel-provider): paged/ragged SDPA in Gfx1151WmmaAttention #14254 and this PR are in develop.

  • Saved plans. The engine supports the ingestor's plan saving and restoring. Its prepared dispatch is an IngestorPreparedDispatch; the handler saves the bound uids, every by-value kernel argument (sink_is_f32 included) and the CTA size as typed launch values under hipkernel.attention_paged_unified.dispatch.v1 (packs/AttentionPagedUnifiedLaunchValues.hpp), registered next to the descriptor's dispatch name for the same handler. restoreLaunch rebuilds them through restoreIngestorLaunch, and prepare() and the restore share the function that applies the launch geometry. graph_match and the reader share one scale rule (any finite float), so every plan the engine builds also loads, zero and negative scales included. The reader refuses a value of another type, a missing or extra value, and one outside what the kernel takes. The README has a "Saved plans" section.

  • Attention sinks (sink_token, SEM-022):

    • Sinks are one logit per query head, [1, Hq, 1, 1], FLOAT (cuDNN's form, the one hipDNN's frontend admits) or in the Q data type (vLLM keeps sinks in the model dtype). They are passed to rocKE attention_unified's sink input.
    • The rocKE runtime_layout mode gains sink support. It previously rejected use_sinks, and its sink init bounded the head by the compile-time GQA ratio, which would have dropped the sink for every KV head after the first.
    • The launch argument sink_is_f32 tells the kernel which form sink_ptr holds (nonzero: Hq fp32 logits; 0: Hq logits in the Q data type). Each head's sink is read with masked loads of both forms, so neither reads past the other's buffer, and the same objects serve both forms.
    • The catalog adds sink variants: 288 objects per architecture, up from 144.
    • A sink in any other data type is declined with cause sinks. Graphs with a Q-dtype sink still need a frontend that accepts that data type (being discussed with the hipDNN team); until then those integration cases skip with the reason, while the FLOAT cases run.
    • The sink buffer must be 16-byte aligned; a graph declaring less for it is declined with cause alignment.
  • Tests:

    • matcher unit tests: accept cases (including separate K and V page tables, FLOAT and Q-dtype sinks, and every finite scale), plus decline cases with the logged cause read back (including the new paged declines, page tables of different widths or batch strides, a missing seq_len_q, left_bound = 0, an unset max_seq_len_kv, a ragged multiplier other than the token stride, ragged K or V, an offsets tensor of B entries instead of B+1, a K/V block span past the buffer range, a sink in a data type other than FLOAT or Q's, and declared alignments below what the kernels load: DeclinesOperandsDeclaringLessThanTheKernelsLoad and DeclinesASinkDeclaringLessThanSixteenBytes);

    • the generated descriptor census for the gfx942 and gfx950 shards;

    • TestGfx950AttentionDenseSchemaGate, a hand-written census suite over the gfx950 product shard, in its own commit. hipkernel:Gfx950AttentionDense's matcher checks neither ragged operands nor a declared alignment, and its UED declares sdk_version 1.0.0, so with the plugin at API 1.4.0 the ingestor's schema gate (GenericPlanBuilder::understandsGraph) is what keeps those graphs away from it. The suite builds the engine from the shipped descriptor set, offers it a ragged SDPA graph and one whose Q declares an 8-byte alignment (both stamped 1.3.0, as GraphDescriptor stamps them; the plugin's 1.4.0 admits them), and requires the schema-gate decline. A third case checks that the same graph without either feature passes the gate.

    • saved plans: TestAttentionPagedUnifiedLaunchValues (round trips of the dense, ragged, separate K/V table, Q-dtype and FLOAT sink cases and of zero and negative scales, and the refusals, sharing the common reader refusals with the other packs through IngestorPlanTestUtilities.hpp), a checked-in byte sample of the v1 contract in TestIngestorPlanLaunchContractGoldens (AttentionPagedUnifiedV1SampleStillReadsAsItsContract, like the other contracts), TestIngestorDispatchAliases, and the GPU suite IntegrationGpuAttentionPagedUnifiedSaveLoad, which saves a plan in both forms, loads it into a fresh handle and requires the output to match the original plan's bit for bit;

    • IngestorGenerator/tests/test_gfx950_attention_dense_schema_gate.py pins the sdk_version of hipkernel:Gfx950AttentionDense's authored config, engine manifest and committed UED below 1.3.0, so every build that runs the IngestorGenerator suite checks the gate's precondition, not only builds that ship that engine;

    • GPU integration tests against a gather-then-dense reference, which is checked against hipDNN's CPU SDPA reference (CpuFpReferenceSdpa over gatherPagedKv, with a FLOAT copy of the sink values). They cover decode, spec-decode with padded rows (which must be zero), sequences with 0 < seq_len_kv < seq_len_q whose leading rows see no key (which must be zero; dense and ragged), an unset max_seq_len_kv, D64/D256, ragged mixed batches with gaps (the rows past seq_len_q, one gap after the last sequence included, must keep their contents; the Q/O storage holds S_q rows from every origin, so a stray write lands where the check sees it), a first origin past 0, a fused-QKV token stride, window plus GQA, separate K and V page tables (V in other blocks than K, dense decode and ragged with a window), FLOAT and Q-dtype sinks, and HIP-graph capture/replay with changed lengths, block ids and data.

    • support claims: the five quick paged bundles the engine serves (bf16/hd128_causal_paged/Decode, fp16/hd64_causal_paged/SpecDecode, fp16/hd64_window_ragged_paged/Mixed, and the FLOAT-sink bf16/hd64_causal_paged_sinks/Decode and bf16/hd64_window_ragged_paged_sinks/Mixed) claim hipkernel:AttentionPagedUnified on gfx942 and gfx950 (linux). Only the git-tracked .support.json sidecars change; the bundle data is feat(hipdnn): SDPA CPU and GPU references cover the full SDPA graph API #14358's, unchanged. SUPPORT_MATRIX_OVERVIEW.md is not regenerated here: since ci(integration-tests): refresh the support matrix overview apart from claim PRs #15293 it is refreshed in a PR of its own.

    • a gfx942 and a gfx950 row in HKP_PINNED_ATTENTION_ENGINES (src/CMakeLists.txt, the table feat(hip-kernel-provider): add gfx942 dense attention AOT ingestor engine #14908 introduced), which registers for each shipped architecture an engine-pinned external integration lane (hip-kernel-provider-{gfx942,gfx950}-attention-paged-unified-external-integration_*, categories in hipkernel_AttentionPagedUnified_test_categories_integration.yaml, filter HKP_SDPA_FWD_ALL_GTEST_FILTER), a GPU-reference target and the Packs census, routed into that architecture's shard. Each row's TOML (config/hipkernel_AttentionPagedUnified_{gfx942,gfx950}.toml) skips the other architecture, so a package carrying both shards runs the bundles once. The census table now collects a multi-architecture engine's rows into one Packs registration (a suite may be declared once), and a _hkp_packs_cases_<engine> list gives this engine's hand-written Packs cases; the schema-gate suite registers as a companion suite of the dense engine's row. The integration suite is a row of _knob_suites, so test_attention_bundle_filters.py checks it with the others. integration-tests/docs/file-formats.md now lists every per-engine categories file, the new one included. It runs the SdpaFwd bundles against the engine with the claims enforced, so a claimed bundle the engine declines or fails breaks the lane.

Behaviour and API changes

  • New engine name hipkernel:AttentionPagedUnified. Dense (non-paged) graphs never match it.

Known integration work

  • A ragged Q must use vLLM's query_start_loc layout: non-decreasing origins with ragged_offset[b+1] - ragged_offset[b] >= seq_len_q[b], and ragged_offset[B] the end of the last block (the first origin need not be 0). The kernel finds each query block's sequence by a binary search over the offsets and sizes its grid from ragged_offset[B]. RFC 0014, as feat(hipdnn): ragged tensors with per-batch origins, lengths and strides (ALMIOPEN-3151) #15080 rewrote it, also allows origins in any order and an unused ragged_offset[B] when seq_len_q is bound. Those graphs are valid, but nothing in the graph shows the offset values, so the engine cannot decline them at match time, and its output for them is wrong. What it can check at match time, it still checks: seq_len_q must be bound and the offset table must be INT32 [B+1, 1, 1, 1]. A grid over (B, blocks per sequence) with each sequence's own origin removes the requirement and was tried, but on a mixed batch (63 decode rows plus one 2048-token prefill, bf16, Hq 64, Hkv 8, D 128) it made the kernel 54% slower on gfx942 and 129% slower on gfx950, because almost every query block is empty. The follow-up is a packed numbering driven by seq_len_q: each query block finds its sequence from a wave prefix sum over ceil(min(seq_len_q[b], S_q) / BLOCK_Q), so the offsets are read only as origins.
  • Split-KV (3D) decode needs kernel slots (RFC 0021), so this PR has one launch only.
  • The non-paged ragged case (encoder / prefill with ragged K/V) is not served: the kernel addresses only paged K/V.
  • Four of feat(hipdnn): SDPA CPU and GPU references cover the full SDPA graph API #14358's paged bundles (bf16/hd128_causal_ragged_paged/Mixed, fp16/hd64_causal_paged_sinks/Decode, standard bf16/hd128_causal_paged/LongDecode and standard bf16/hd64_window_ragged_paged_sinks/Long) also ask for the softmax statistics output, which this engine does not produce, so it declines them with cause stats. feat(rocke): optional LSE output on unified tiled 2D attention #14616 added an LSE output to the tiled 2D kernel, but only for the fixed layout ([total_q, Hq, 1], compile-time head count); serving stats here needs that output in runtime_layout mode (runtime stats strides and head count, -inf for padded rows) and LSE catalog variants, so it is left for a follow-up.
  • TheRock CI has no pull_request, push or nightly job that tests gfx950 (the gfx950-dcgpu family runs only on submodule bumps, and the gfx94X job fetches only the gfx942 shard). So TheRock does not test the gfx950 KDP, kernel changes, census entry or integration cases; gfx950 coverage comes from the author's runs below. The same holds for TestGfx950AttentionDenseSchemaGate, which, like the existing hipkernel:Gfx950AttentionDense census, registers only in the gfx950 shard of a build that ships that engine (its kernels are a DVC artifact).
  • Overlap with hip-kernel-provider: enable minimal single-token attention decode #13108 (hipkernel:AttentionDecode): both build on the tiled 2D kernels. hip-kernel-provider: enable minimal single-token attention decode #13108 builds an identity page table and query offsets in workspace, while this engine binds the graph's own page table, lengths and offsets with no workspace.
  • Overlap with fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967 (OOB reads of K/V and the block table in unified attention): both bound the tiled 2D block-table lookups by the sequence's own ceil(seq_len / BS) entries and use an out-of-range buffer offset for K/V reads past seq_len. fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967 owns both. If it lands first, this PR switches its runtime_layout paths to fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967's block_table_max_idx and buffer_offset_in_sequence at its rebase; if this PR lands first, fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967 makes those steps.

Validation (Linux, Slurm compute nodes)

Head 3bed02a (9 commits) on develop 1702865; git merge-tree --write-tree against develop 1702865 (fetched 2026-11-07, still its tip) is clean. Round 4 changed no emitted IR of a valid spec. It added the gfx942 runtime_layout LDS refusal to both builders, the representative-IR golden entries and the docs, and rebased onto three develop commits that do not touch rocKE. So the GPU results below are from the previous head fe41d27 and were not rerun. Checks on 3bed02a (built as its identical tree 2ec59d3a81d):

  • CI-config build on a CPU node (GPU_TARGETS=gfx942;gfx950, clang-tidy 20, Flash2 and ASM SDPA on) builds with no compiler or clang-tidy diagnostics.
  • rocKE ctest: 60/60.
  • Byte identity with Python is green for LLVM 20, 22, 23 and 24 (gfx950 tiled 2D 91/91, gfx942 tiled 2D stress 51/51), and the differential goldens report GOLDEN OK.
  • Representative-IR golden: rocke_ir_parity_harness.py --check tests/golden/rocke_representative_ir_sha256 reports IR parity OK at llvm20/22/23/24, on the head and at 6114fe0, the commit that adds the entries.
  • runtime_layout host pytest with a standalone rocke_engine build: 187 passed, covering the native reject and LDS cases. With the C LDS gate disabled, test_native_builder_rejects_lds_over_budget fails.
  • Unit binary: 1772 passed, 140 skipped, and the same 30 kpack-fixture failures as before.
  • Census: 25/25.
  • rocKE dispatch pytest: 1836 passed.
  • Descriptor regeneration: none of the 8 files differ (584 ids reused).
  • test_abi.py:
    • Against a librocke.so built from the head: rocke_abi_version() is 13; 27 passed and 1 failed (the COMGR-less online-lowering case, as before).
    • In the provider build: 21 passed, 7 skipped.
  • The remaining nonzero CTest groups match the previous head's job exactly.
  • rocKE README entry points, run from rocke/platform/ on a CPU node in the same environment for this head and for develop 1702865:
    • PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py: OK, 368 tests (142 skipped), on both.
    • python tests/run_all.py: identical results on both.
      • Byte-identity gate: green, 100 families.
      • Arch-domain lane: green.
      • pytest: 10297 passed, 5497 skipped.
      • ctest: 60/60.
      • The ROCKE_BACKEND=both pass is skipped (no rocke_engine in that environment).
      • Two gates fail on both trees, so the overall result is RED on both: the relative-path guard (CMakeLists.txt:428, the integration-bundles path from 58ff0d0), and the emitted-IR validity gate, on the same 22 cases (7 gfx950 scalar attention, 15 gfx1250 Winograd conv; LLVM 23 "Operand has incorrect register class").
  • pre-commit on the changed files: all hooks pass.

Previous head fe41d27 on develop 477ab5b. Everything below is on fe41d27 unless labelled otherwise.

  • Node-local GPU builds (GPU_TARGETS = the node's architecture, tidy off):

    Check gfx942 (MI300X) gfx950 (MI355X)
    IntegrationGpuAttentionPagedUnified + IntegrationGpuAttentionPagedUnifiedSaveLoad 21 passed, 3 Q-dtype sink cases skipped 21 passed, 3 Q-dtype sink cases skipped
    Unit tests (matcher, plugin, compiled-in engines) 66/66 63 passed, 3 skipped (TestHipFlash2Engine gfx942-only cases)
    Whole hip_kernel_provider_tests binary 1901 passed, 11 skipped, 30 failed 1860 passed, 52 skipped, 30 failed
    Census 10/10 (20/20 with feat(hip-kernel-provider): add gfx942 dense attention AOT ingestor engine #14908's gfx942 dense artifacts pulled) 15/15
    rocKE runtime_layout pytest 222 passed, 11 skipped 715 passed, 7 skipped
    perf(rocke): runtime heads, window and batch for gfx942/gfx950 tiled attention #15024's test_attention_tiled_runtime_shape_numeric.py 25 passed, 10 skipped (gfx950 cases) 28 passed, 7 skipped (gfx942 cases)
    Claimed bundles, claims enforced (--verification-mode gpu --validator cpu) 5 pass, 5 claims confirmed, no claim failures; 4 declined stats 5 pass, 5 claims confirmed, no claim failures; 4 declined stats
    --write-support-claims on a copy with the claims cleared writes exactly this arch's entry for the 5 claimed bundles same
    Engine-pinned external lane (quick + ffm-quick), from the new table rows 2/2 2/2
    GPU-reference target hip_kernel_provider_gfx942_attention_paged_unified_gpu_ref_integration_tests passed not run
    • The 30 unit failures on each architecture are TestAttentionDenseSignature, TestAttentionDenseDispatch, TestGfx942AttentionDense{Signature,Dispatch} and TestGfx950AttentionDenseSignature cases whose kpack fixture path does not exist in this environment (library 'there-is-no-archive-here.kpack' resolves to ''), the same cause as the 11 TestGfx950AttentionDense* failures before feat(hip-kernel-provider): add gfx942 dense attention AOT ingestor engine #14908 parametrised those suites over both architectures. They fail the same way with feat(hip-kernel-provider): add gfx942 dense attention AOT ingestor engine #14908's gfx942 dense DVC artifacts pulled.
    • feat(hip-kernel-provider): add gfx942 dense attention AOT ingestor engine #14908's own GPU tests on gfx942 (its DVC artifacts pulled): IntegrationGpuGfx942AttentionDense* 9/9, its engine-pinned quick + ffm-quick lane 2/2, its GPU-reference target passed.
    • TheRock-style installed ctest -L quick: gfx942 fails the same 19 entries as the previous head (fixture failures in the unit suites and the rocKE Python entries without a toolchain in the install prefix); gfx950 fails 4 (the unit suites, rocke_library_mla_pytest, and rocke_pytest, which hit its 300 s limit).
  • CI-config build on a CPU node (GPU_TARGETS=gfx942;gfx950, clang-tidy 20 via CMAKE_CXX_CLANG_TIDY, Flash2 and ASM SDPA on):

    • Builds with no compiler or clang-tidy diagnostics. The only build warnings are develop's gfx942 unified 3D UHD heuristic not being shipped because its DVC model file is not pulled.
    • rocKE build-tree ctest: 60/60. rocKE ABI tests: 21 passed, 7 skipped (the C-engine cases skip because the provider build links no shared librocke.so).
    • Byte identity with Python for LLVM 20, 22, 23 and 24: green (gfx950 tiled 2D 91/91, gfx942 tiled 2D stress 51/51). Differential goldens (.ll at four flavors and canonical IR): GOLDEN OK.
    • Descriptor regeneration: 8 files compared, none differ (584 ids reused, none minted).
    • feature_matrix.py --check reports SUPPORT_MATRIX_OVERVIEW.md behind; it is behind on develop too (it lacks feat(hip-kernel-provider): add gfx942 dense attention AOT ingestor engine #14908's claims), and since ci(integration-tests): refresh the support matrix overview apart from claim PRs #15293 no check requires it and claim PRs leave it alone. The renderer passes.
    • rocKE tests/library/tests/dispatch: 1836 passed, 257 skipped. test_attention_bundle_filters.py passes locally (18 tests) with the engine's two table rows.
    • The CTest groups that return nonzero fail the same way on develop 477ab5b, built and run with the same configure and commands (tidy off): hipdnn_golden_sdpa_bwd_generator_python_tests, uhd_gen_tests, hip-kernel-provider_test_name_validation, the installed hip_kernel_provider_tests_*_suite entries (fixture failures above), rocke_library_mla_pytest, and rocke_pytest (26 failures without a toolchain on this head; on develop it hit its 300 s limit).
    • IngestorGenerator pytest: 1873 passed, 70 skipped, 2 failed, 23 errors; develop fails identically (2 provenance tests need a .git, 23 conv dgrad AOT errors need unpulled DVC files). descriptor-packaging pytest: 1198 passed, 7 skipped, 1 failed (test_worker_inherits_the_in_tree_rocke_paths, harness PYTHONPATH); develop fails the same test.
    • Whole unit binary on the CPU node: 1772 passed, 140 skipped and the same 30 fixture failures as on the GPU nodes. Census: 25/25.
    • test_abi.py against a librocke.so built from this head (online.build_lib()): rocke_abi_version() returns 13; 27 passed, 1 failed, the failure being an online-lowering case that cannot load COMGR in that job's environment (same on the combined tree).
  • Combined tree with feat(hip-kernel-provider): paged/ragged SDPA in Gfx1151WmmaAttention #14254 (tmp/v3-with-14254 @ fccef3e0457): this head, feat(hip-kernel-provider): paged/ragged SDPA in Gfx1151WmmaAttention #14254's head a891f662ada picked on top with its ABI entry at 14, and the commit that switches Gfx1151WmmaAttention to the shared checks. Built for gfx942;gfx950;gfx1100;gfx1151;gfx1201 (tidy off) on a CPU node: whole unit binary 1802 passed, 139 skipped and the same 30 fixture failures; census 85/85; rocKE ctest 60/60; librocke.so reports ABI 14; test_attention_bundle_filters.py passes (18 tests).

  • pre-commit on the files changed by the rebase: all hooks pass.

ASIC coverage

gfx942 and gfx950, one KDP each (288 objects per architecture). No other architecture matches the engine.

@dpeabody-amd
dpeabody-amd force-pushed the users/dpeabody/hipdnn-cdna-paged-attn branch 2 times, most recently from 76fa8e2 to 1f27e15 Compare October 23, 2026 01:48
dpeabody-amd added a commit that referenced this pull request Oct 23, 2026
…inks (#14246)

ISSUE ID : AICK-1847

Small correctness fix. Looking for review from someone who knows the ASM
SDPA kernels.

## Summary
`ASM_SDPA_ENGINE` and `HIP_FLASH2_ENGINE` claim SDPA forward graphs that
set `sink_token` (attention sinks), then run them as if no sink were
present. Neither engine's kernels take a sink input, so these graphs
silently produce wrong results instead of declining.

A sink is a per-head logit that enters the softmax denominator
(`exp(sink - m)`) without contributing a value row. cuDNN models it with
the same attribute (`sink_token`). Ignoring it changes every output row.

Found while adding sink support to the paged SDPA engines, where the
sink graphs reached ASM SDPA. Related PRs: #14245 (paged K/V SDPA
contract) and #14247 (paged/ragged SDPA engine). This PR does not depend
on either.

## What changes
- `SdpaFwdPlanBuilder::isApplicable` (ASM SDPA) declines a graph with
`sink_token_tensor_uid` set.
- `HipFlash2FwdPlanBuilder::isApplicable` (HIP flash2, off by default)
declines it too.
- New tests `TestSdpaFwdPlanBuilder.IsApplicableRejectsSinkToken` and
`TestHipFlash2FwdPlanBuilder.RejectsSinkToken`: the same graph is
accepted without the sink and declined with it. Both build the sink
graph with one shared helper, `addSinkToken`, which gives the sink
tensor uid 0 so the tests check that the field is present, not that its
value is nonzero.

### Behaviour and API changes
- SDPA forward graphs with a sink no longer run on ASM SDPA or HIP
flash2. They now decline, so another engine can claim them or the caller
falls back (the frontend reports `GRAPH_NOT_SUPPORTED` when no engine
applies).
- No API change.

## Questions for an ASM reviewer
- Do any of the vendored `fmha_v3_fwd` kernels have a sink input that
this engine could pass through instead of declining? I found none in the
kernel argument struct (`asm/SdpaFwdKernelArgs.hpp`).

## Validation
Built and run on an MI300X (gfx942) with `HIPDNN_ENABLE_SDPA=ON`,
`ENABLE_ASM_SDPA_ENGINE=ON` and `ENABLE_HIP_FLASH2_ENGINE=ON`, with
clang-tidy 20 enabled as in the CI configuration (no diagnostics):
- `TestSdpaFwdPlanBuilder.*`, `TestHipFlash2FwdPlanBuilder.*` and
`TestAsmSdpaEngine.*`: 38 of 38 pass. Both new tests run (not skipped).
- Full `hip_kernel_provider_tests`: 893 pass, 7 skipped (unrelated
single-device and gfx950-arch gates), 0 fail.
- Mutation: deleting the ASM decline fails
`IsApplicableRejectsSinkToken`; deleting the flash2 decline fails
`RejectsSinkToken`.

Not run:
- gfx950. The new ASM check is plain `isApplicable` code that does not
depend on the architecture, but I have not executed it there.
- CI does not cover the new tests: its runners have no GPU, so they
skip, and `ENABLE_HIP_FLASH2_ENGINE` is off in every CI configuration,
so CI does not compile the flash2 change.
- No integration test builds a sink graph, so the fallback described
above is not exercised end to end.

## ASIC coverage
ASM SDPA: gfx942 and gfx950. HIP flash2: gfx942 only (its plan builder
declines other architectures).
adickin-amd pushed a commit that referenced this pull request Nov 2, 2026
…PI (#14358)

## Summary

This PR brings the hipDNN SDPA **CPU and GPU reference implementations**
up to everything the hipDNN SDPA graph API can express, forward and
backward. It also adds the tests and docs that let the references serve
as the oracle every attention engine is compared against.

cuDNN's semantics are the contract. Each place where hipDNN deliberately
differs from cuDNN is recorded in
`projects/hipdnn/docs/SdpaSemanticsRegister.md` (SEM-xxx). The design,
the full support table and the testing strategy are in
**`projects/hipdnn/docs/SdpaReferenceImplementations.md`**.

Builds on #14483 (cuDNN's `left_bound` meaning: L keeps the L keys
ending at the diagonal) and #14312 (SdpaBackward O and LSE derived from
the CPU forward), both merged.

JIRA ID : ALMIOPEN-2863
Related: ALMIOPEN-2928, AICK-1847

## What the references support

Both references (CPU `CpuFpReferenceSdpa`; GPU `GpuFpReferenceSdpa`,
hipRTC kernels) and both graph executors support:

**Forward (`SdpaAttributes`)**

| Feature | Notes |
|---|---|
| dtypes | FP32, FP16, BF16; FP8 E4M3 / E4M3 FNUZ / E5M2 / E5M2 FNUZ
inputs with BF16 or same-format fp8 output |
| Layout, GQA/MQA, `D_qk != D_v` | any strides; independent `H_k` /
`H_v` |
| Additive bias | FLOAT/HALF/BF16, rank 1-4 broadcast |
| Diagonal band | `left_bound` / `right_bound` / alignment and the
deprecated causal flags |
| Padding | `seq_len_q/kv`; padded rows give O = 0, LSE = -inf |
| Paged K/V | K, V or both, each with its own block size; extent
inferred in cuDNN's order when unset |
| Ragged (RFC 0014) | native packed forward (BF16, FP8 E4M3); every
other ragged graph, including backward and ragged Q over paged K/V,
through a ragged-as-dense rewrite |
| Attention sink | FLOAT `[1, H, 1, 1]` |
| ALiBi | cuDNN slopes, top-left causal |
| Block-sparse mask | cuDNN's UINT8 bitmask, one bit per 128x128 tile |
| Dropout | explicit mask; Philox through the engine's `rng_dump` |
| fp8 scaling | Q/K/V descales (scalar or per KV head); `scale_s` /
`descale_s` / `scale_o`; `amax_s` / `amax_o` |
| Outputs | LSE, Max, Sum_exp |

**Backward (`SdpaBackwardAttributes`)**: a new GPU backward reference.
Both references support FP32/FP16/BF16, GQA, `D_qk != D_v`, LSE or
recompute, bias + dBias, band, padding, ALiBi, explicit-mask dropout,
and ragged through the rewrite.

**Reference-only**: logit softcap, and an opt-in GPU P-rounding mode
that models provider kernels.

**Not supported because the graph API lacks it**: backward Philox (no
backward `rng_dump`), backward sinks / dSink, the fp8 backward node, and
`score_mod`.

**Structure**: each direction is one pipeline (score: scale, softcap,
bias, ALiBi, band, block mask; then softmax with sink; then P @ V),
configured by a named options struct. fp8, paging and ragged are thin
layers around that core. The CPU and GPU references share no code, only
the semantics. The graph-translation layer (`PlanUtils.hpp`, the ragged
rewrite) is shared by both executors, and the dense CPU and GPU plans
use one applicability gate per direction.

## Behavior changes

- **Frontend validation follows cuDNN (user-visible).** `SdpaFwdNode`
now rejects: `right_bound < -1`; `max_seq_len_kv` without paging; paging
without both `seq_len_q` and `seq_len_kv`; page tables that are not
INT32 `[B, 1, M, 1]`; M != `ceil(max_seq_len_kv / block_size)` when the
extent is set; `seq_len_*` not INT32 `[B, 1, 1, 1]`; `sink_token` not
FLOAT `[1, H_q, 1, 1]`. One-sided paging, per-side block sizes, an
inferred extent and packed page tables stay allowed. Each rule cites
cuDNN and has a decline test. In-repo callers that built such graphs are
fixed.
- **Breaking `test_sdk` change (deliberate).** The positional
`CpuFpReferenceSdpa::forward` / `backward` overloads are removed in
favor of `SdpaFwdOptions` / `SdpaBwdOptions` (`.setScale`, `.setMask`,
`.setBand(left, right, topLeft)`; causal is `.setBand(-1, 0, topLeft)`).
There is no forwarding shim on purpose: `left_bound` changed meaning
(SEM-001), so a shim would silently change old callers' results. Every
in-repo caller is migrated.
- **gfx1151 WMMA attention (`Gfx1151WmmaAttentionNative.cpp`
`maskFor`)**: the deprecated `causal_mask` / `causal_mask_bottom_right`
flags compose with explicit bounds and alignment, as the references do,
instead of overriding them. Graphs that set only the flags bind as
before.

## Absorbed PRs (authorship kept)

- **#14004** (EwanC, ALMIOPEN-2863): SDPA backward golden bundles and
generator, extended with GQA, padding, dBias, explicit dropout and an
FP64 autograd cross-check.
- **#13935** (scerzh, ALMIOPEN-2928): CPU ragged SDPA forward plan.
- **#14245** (dpeabody-amd, AICK-1847): the frontend paged / sink /
seq_len validation, the GROUP-mode ASM test fix, and the paged golden
generator. Its separate paged plans are not taken: paging goes through
the dense reference path, so every other feature composes with it.
Stacked work on #14245 (#14247, #14254) may need to follow the stricter
`seq_len_q` rule.

## Testing

How correctness is established (detail in
`SdpaReferenceImplementations.md` §3):

1. **Simple by construction**: FP32, no tiling, three-pass softmax, one
code path per feature.
2. **Closed-form oracles and properties**: uniform scores, single key,
dead rows; GQA = repeated heads, band and block mask = -inf bias,
padding = truncated batches, paged = dense, sink = extra zero-value key,
ragged = per-batch dense.
3. **Finite-difference gradient checks** of the backward.
4. **CPU vs GPU parity** (`hipdnn_gpu_ref_tests`, `TestGpuSdpa*Plan`):
every feature at small shapes with NaN-filled padding; Standard-tier
medium sweeps up to 1100 keys; fp8 device codecs bit-exact against the
host.
5. **PyTorch golden bundles** (`hipdnn_golden_data_tests` checks every
bundle against both references, with no engine involved). The backward
generator is FP64-autograd checked; the fp8-scaling and block-mask/ALiBi
paths reproduce cuDNN's own reference code. Golden data is 78 DVC
pointers (26.8 MiB, every bundle under 1 MiB, pushed to
`hipdnn/golden-data`).
6. **Mutation testing**: injected reference and gate bugs each fail a
test.
7. **Applicability**: every schema field is gated, with one decline test
per rule; a graph is declined, never accepted and then thrown on.

**Results on gfx942** (MI300X, `-Werror`; the current head is covered by
CI):

| Suite | Passed |
|---|---|
| test_sdk / frontend / public_frontend | 2609 (1 skipped, pre-existing)
/ 2247 / 487 |
| gpu_ref | 13836 |
| integration-tests unit | 1149 |
| golden (every bundle with golden data, both references) | 220 |
| ASM_SDPA_ENGINE SdpaBackward bundles under #14312's tolerances | 31
(41 graphs the engine does not accept) |
| hip-kernel-provider / its SDPA integration / ASM SDPA external quick
suite | 1393 / 45 / 26 |
| flatbuffers_sdk / plugin_sdk / generator pytest | 272 / 1293 / 146 |

**Devices**: gfx942 runs everything above. gfx950: no run. gfx1151 has
host matcher tests only (no usable device node; see Known issues).
Windows and the other families are covered by TheRock CI. GPU reference
kernels are hipRTC source, so there is no `GPU_TARGETS`-dependent host
configuration.

**Test time**: the new medium sweeps and long bundles are Standard tier
(~18 s). The GPU reference forward is faster than on `develop` (masked
keys skip the dot product), with byte-identical outputs.

## Known issues

- **fp8 scaling over 128 keys (SEM-024)**: cuDNN quantizes P per 128-key
tile against a running max; the references are exact (and bit-exact
against cuDNN's `fp8_ref`) up to 128 keys, and the bundles stay within
that.
- **gfx1151 (accepted risk)**: the `maskFor` composition change has host
matcher tests only; a 13-graph device check is prepared for when a
gfx1151 node is available.
- Once this lands, #14004, #13935 and #14245 are closed as superseded.

## Pre-merge

- [ ] CI green on the current head

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

---------

Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: Ewan Crawford <ewan.crawford@streamhpc.com>
Co-authored-by: Sergei Grigoriev <102019268+scerzh@users.noreply.github.com>
Co-authored-by: dpeabody-amd <dpeabody-amd@users.noreply.github.com>
@AnaghaRaoAMD

Copy link
Copy Markdown
Contributor

Interaction between the plugin API bump to 1.4.0 and the rocKE dense attention engines (AICK-2919).

Gfx950AttentionDense's matcher (gfx950AttentionDenseGraphMatches) never checks Q/K/V/O for ragged_offset_tensor_uid or the declared alignment, and its kernel takes no offset table. A ragged graph whose dims and strides look dense would be read as one dense batch: wrong values, no error. Today the plugin's 1.2.0 API keeps those graphs away from it entirely.

After this PR the only guard left is the ingestor schema gate: GenericPlanBuilder::understandsGraph declines any graph whose min_required_engine_api_version is above the engine descriptor's sdk_version. Ragged and non-default-alignment graphs are stamped 1.3.0, and every rocKE CDNA attention descriptor declares sdk_version: 1.0.0 (gfx950 dense on develop, and #14908, #14242, #14259). So those graphs are still declined before the dense matcher runs, and the bump is safe as written. But the protection is one descriptor field: regenerating any of those engines with sdk_version >= 1.3.0 would silently expose the bug.

Two asks:

  1. Please add a test here that pins this: a ragged SDPA graph, and one with a non-default alignment, offered to hipkernel:Gfx950AttentionDense with the plugin at 1.4.0, is declined with the schema-gate cause.
  2. Merge order for the rest of the stack: the explicit ragged and alignment declines from AICK-2919 must land before any rocKE attention engine descriptor moves to sdk_version >= 1.3.0. That covers gfx950 dense and feat(hip-kernel-provider): add gfx942 dense attention AOT ingestor engine #14908, feat(hip-kernel-provider): add gfx950 unified attention 2D/3D ingestor engines #14242, feat(hipdnn): gfx942 unified 3D split-KV decode as a hipDNN ingestor engine #14259. feat(hip-kernel-provider): add gfx950 unified attention 2D/3D ingestor engines #14242 already declines ragged operands; none of them check alignment yet.

@dpeabody-amd
dpeabody-amd force-pushed the users/dpeabody/hipdnn-cdna-paged-attn branch from 1f27e15 to aa80f2c Compare November 3, 2026 16:55
@dpeabody-amd
dpeabody-amd changed the base branch from users/dpeabody/hipdnn-paged-sdpa-contract to develop November 3, 2026 16:55
@therock-pr-bot

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

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

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

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

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

🙋 Wish to Override Policy?

@therock-pr-bot

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

Copy link
Copy Markdown

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

@dpeabody-amd dpeabody-amd reopened this Nov 3, 2026
@dpeabody-amd
dpeabody-amd force-pushed the users/dpeabody/hipdnn-cdna-paged-attn branch 2 times, most recently from 7e1191c to d49477e Compare November 4, 2026 00:51
@dpeabody-amd
dpeabody-amd marked this pull request as ready for review November 4, 2026 03:15
@dpeabody-amd
dpeabody-amd requested review from a team as code owners November 4, 2026 03:15
@dpeabody-amd
dpeabody-amd force-pushed the users/dpeabody/hipdnn-cdna-paged-attn branch from d49477e to 8e16db7 Compare November 4, 2026 05:32
@dpeabody-amd
dpeabody-amd force-pushed the users/dpeabody/hipdnn-cdna-paged-attn branch from 8e16db7 to 922a229 Compare November 5, 2026 18:31
@dpeabody-amd

Copy link
Copy Markdown
Contributor Author

Pushed 922a229: rebuilt on develop, which now includes #15024. Replies to Brian's inline threads are on each thread.

@cgmillette, first-round review

Thanks for the review. Since it, the branch is rebuilt on #15024 and now on develop (see above), so the line numbers below are the new head's.

  • Prior item, AnaghaRaoAMD Ask 1 (gate pin never runs): see N3.
  • Prior item, BrianHarrisonAMD 2e-2 notes and plan saving: answered in the inline threads below.
  • N1 (conflicts): resolved by the rebuild. TestGfx950AttentionDenseSchemaGate registers under if(_hkp_gfx950_attention_dense) after develop's 2D/3D handoff block, with no else(); the integration CMake keeps develop's block and this PR's paged block, each under its own guard. ROCKE_ABI_VERSION and BINARY_ABI are 10: develop is 9. feat(hip-kernel-provider): paged/ragged SDPA in Gfx1151WmmaAttention #14254 also bumps them, so whichever of the two lands second takes the next number. The history text now says the bool lands in tail padding, so sizeof stays 152. (51e3dc3, fec3540)
  • N2 (private copies of the match helpers): the engine now uses packs/SdpaGraphMatch.hpp (findTensor, sdpaNode, hasSdpaForwardNode, describeOperand, sdpa_graph::DeclinedBy, the axis constants). The stricter checks move into that header (usesOverrideShapes, hasNonNegativeStrides, declaresAlignment, isInt32Vector), and so do the cuDNN paged-graph checks, with one cause vocabulary for every paged engine. The local checkedProduct is renamed spanProduct: pack_features::checkedProduct refuses a zero factor, which every extent of 1 gives, so it could not replace it. (19f309b)
  • N3 (the gate pin runs in no CI lane): IngestorGenerator/tests/test_gfx950_attention_dense_schema_gate.py pins the dense engine's sdk_version below 1.3.0 in the authored config, the engine manifest and the committed UED, so every build that runs the IngestorGenerator suite checks it. The census control now asserts that the engine serves the baseline graph (EXPECT_TRUE(offer(g))). Setting any one of the three versions to 1.3.0 fails exactly that source's case. (fec3540)
  • N4 (SGPR spills): measured again on the new head, over the shipped catalog (288 kernels per architecture) against the same specs with runtime_layout off: 264 of 288 gfx942 kernels and 232 of 288 gfx950 kernels spill SGPRs (at most 57 and 76), all into VGPR lanes, never to scratch. The fixed-layout twins spill none on gfx942 and 1 SGPR in 4 kernels on gfx950, and perf(rocke): runtime heads, window and batch for gfx942/gfx950 tiled attention #15024 barely changes that (2 kernels and 6 SGPRs before it), so the spills come from the runtime strides. The lane moves are a median 1.7% (gfx942) and 1.5% (gfx950) of a kernel's instructions, at most 7.4% and 6.3%; the lanes they hold lower the VGPR-limited occupancy estimate for 70 of 192 gfx942 and 209 of 288 gfx950 kernels. Trimming them means fewer runtime arguments, which is what lets one plan serve every layout, so I kept them and recorded the counts and the cost in the engine README ("Register pressure"). I have not measured the runtime cost: no fixed-layout kernel serves the same paged graph. (59de321)

Overlap with #14254

#14254 already logs this PR's cause keys. This PR moves the shared paged checks into SdpaGraphMatch.hpp with the reasons #14254 logs for the same rules, so its tests keep passing unchanged. The commit that switches #14254's gfx1151 paged matcher to those checks is ready, but it needs #14254 underneath, so it goes onto this branch when #14254 is in develop (see the PR body). On a throwaway branch with #14254's commits on this head, it applies unchanged and the combined unit tests pass.

dpeabody-amd added a commit that referenced this pull request Nov 5, 2026
…1151WmmaAttention

hipDNN's paged K/V contract (SdpaSemanticsRegister SEM-008/009, #14358)
differs from the one this engine was written against:

- Every paged graph carries seq_len_q and seq_len_kv, also with a dense Q.
  The matcher now requires both and always binds seq_len_q. The dense-Q
  form without seq_len_q is gone: GraphLayout keeps PAGED_DENSE_Q (dense Q,
  live rows seq_len_q) and PAGED_RAGGED_Q, and an object without seqused_q
  serves no graph.
- Rows past seq_len_q are padding with O = 0 and LSE = -inf. The kernel
  now writes them (previous commit); the GPU test expects that rather than
  untouched rows.
- Separate K and V page tables (page_table_k, page_table_v) are valid
  graphs and are now served: the matcher binds both, and the launch passes
  them as the kernel's block_table and block_table_v (one table bound to both
  sides is passed twice). The two tables must have one width and, with more
  than one sequence, one row stride.
- One-sided paging, per-side block sizes or page counts, K and V tables of
  different widths or row strides, and packed page tables are valid graphs
  the kernels cannot read, so each is declined with a logged cause.
- An unset max_seq_len_kv is valid. The kernels use it only as the bound on
  seq_len_kv, so it is inferred as the table width times the block size, as
  cuDNN does.
- A graph may declare a device-pointer alignment below the default 16 bytes
  (engine API 1.3.0, which the plugin's 1.4.0 admits). The kernels load
  Q / K / V / O rows as 16-byte vectors, INT32 operands as 4-byte words and
  the sink with 16-bit loads, so the matcher declines a graph declaring less
  for any of them (cause 'alignment', dense and paged) where hipDNN can
  still pick another engine, rather than the plan failing at execute.
- left_bound follows cuDNN (L keys ending at the diagonal); maskFor already
  translates it to the kernel's window_left = L - 1. The GPU test's windows
  follow it (32-, 16- and 128-key windows).

Catalog: the 48 paged objects with dense_q but without seqused_q (24 plain,
24 with sinks) could serve no hipDNN graph and are dropped; 96 paged
objects remain (dense Q and ragged Q, each with and without sinks), 144
prefill objects in all. Retained descriptor ids are unchanged. The census
is regenerated.

Decline logs: the paged checks log the cause keys #14247 uses for the same
contract rules, so the shared checks it moves into SdpaGraphMatch.hpp keep
every key this engine logs: one_sided_paging; page_table (a packed table,
K and V tables of different widths or row strides, a table that is not
INT32 [B, 1, M, 1] with a unit entry stride); block_size (different K and V
block sizes, a page size that is not a power of two, and one no paged object
ships for, which used to find no object without a log); shape (different K
and V page counts); seq_len_q, seq_len_kv and ragged (a missing or malformed
sequence-length or offset tensor, declined silently before); attn_mask
(additive bias on a paged graph). Each reason names the operand and values
involved. The dense layout logs 'padding' for seq_len_q, seq_len_kv,
max_seq_len_kv or padding_mask. Each line is headed by the declining
engine's name (sdpa_graph::DeclinedBy): the decode engine runs the shared
dense matcher too, and used to log under the prefill engine's name; it now
also logs why it declines a paged graph.

Tests: the matcher scaffolding both unit tests share (engine symbols, the
graph fixture, object metadata, the decline-log reader, the ABI helpers)
moves into Gfx1151WmmaAttentionTestSupport.hpp. The unit tests bind
separate K and V tables (dense and ragged Q) and a V container with its own
strides, and assert the engine, cause and reason of each decline: every
paged check in isolation, the unshipped page sizes, alignment (sink
included), FP8, dropout, a forced implementation, the other unserved
features and ALiBi; the dense declines are asserted for both engines. The
GPU test checks against test_sdk's CpuFpReferenceSdpa (gatherPagedKv per
page table, sequence lengths, band, sink and LSE) instead of its own
reference, adds two cases whose V pages sit in other blocks than their K
pages, through a second page table, two whose V page, head and token strides
all differ from K's, and runs the 96 and 256 buckets at pages 32 and 64 and
the 128 bucket at page 64.

Sinks: the engine still accepts a sink_token in the Q data type and
declines a FLOAT one. hipDNN's frontend accepts only a FLOAT sink, so the
GPU test's sink cases now skip with a message when the frontend rejects a
Q-dtype sink; the paged-rules probe that skipped every case is removed, and
the other cases run.

Re-pin the gfx1151_wmma_attention.yaml digest in the IngestorGenerator
fixture unslotted_render_golden.json (test_slots.py): the contract
change alters that config's render by design.
@dpeabody-amd
dpeabody-amd force-pushed the users/dpeabody/hipdnn-cdna-paged-attn branch from 922a229 to 33842d2 Compare November 5, 2026 23:55
@cgmillette

Copy link
Copy Markdown
Contributor

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

PR #14247 — paged/ragged SDPA engine for gfx942/gfx950 (dpeabody-amd, AICK-1847) — 🟠 Follow-up: N1–N4 fixed; conflicts, golden open

Prior item Status
AnaghaRaoAMD (11-03) Asks 1–2: gate pin; merge order 🟢 FIXED 7da10089
• Ask 1: the pin now runs in CI (N3)
• Ask 2: 14259 landed at sdk_version 1.0.0, as agreed (read)
BrianHarrisonAMD (11-04) ×3: 2e-2 notes ×2; plan saving 🟢 FIXED aeeb1c1f, 33842d27
• notes at both bounds; threads resolved (read)
• CI: launch-value round trips ✓, GPU save-load 4 ✓, sink case skips (read: logs)
AnaghaRaoAMD (11-04) rtl rows missing from the golden 🟠 OPEN, no reply yet
• merge: 19 + 19 rtl configs unblessed per flavor, develop 0; still GOLDEN OK (ran)
• fix: record both families, as asked
AnaghaRaoAMD (11-04) FLOAT sinks declined; sink objects unreachable 🟠 OPEN, no reply yet
• engine serves Q-dtype sinks only; develop's frontend admits FLOAT only (read)
• all 3 GPU sink cases skip in CI (read: logs); body: hipDNN decision pending
AnaghaRaoAMD (11-04) gfx950 dense mode runs 17 search loads 🟢 FIXED aeeb1c1f
• dense: zero iterations; ragged: runtime trip count from num_seqs (read)
• thread still unresolved
AnaghaRaoAMD (11-04) ×2: CMake _others; fkv_build 🟡 OPEN, no reply yet
• _others still from this build's archs (src/CMakeLists.txt:556) (read)
• fastkv binding ignores runtime_layout: C++ builds non-rtl IR, Python raises (ran)
Our 11-04 N1: six files conflict with develop 🟢 FIXED aeeb1c1f, 7da10089
• all six resolved in the rebuild; probe passes at head and on the merge (ran)
• 3 new conflicts with 14259: Our checks
Our 11-04 N2: 7th private copy of SDPA match helpers 🟢 FIXED cc28936b
• engine calls sdpa_graph::; stricter and paged checks moved there (read)
• spanProduct kept: shared checkedProduct refuses zero (read); see N6
Our 11-04 N3: schema-gate pin runs in no CI lane 🟢 FIXED 7da10089
• IngestorGenerator pin 3 ✓; each 1.3.0 mutant fails its own case (ran)
• both hipDNN superbuilds ran that suite at head (read: logs); control asserts offer()
Our 11-04 N4: rtl kernels spill SGPRs 🟢 FIXED 506b3374
• cost recorded in the engine README, kept by design (read)
• 576-kernel grid recompiled on the merge: no scratch, no failures (ran)
Bots / CI 🟢 Passed at head
• 52 ✓, 53 skipped; required ✓ (read 06:17Z)
• attempt 1: Windows rocke_pytest timed out on a slow runner; re-run ✓
• merge ref predates 14259; dirty, so no new runs
Our checks 🟠 Dirty
• merge vs 54a60371: 3 conflicts, all 14259's (ran); keep both census blocks, add its lane to the docs list, rerun feature_matrix.py
• resolved: ctest 51 ✓ / 1 skip; byte identity GREEN ×4; goldens pass; rtl off = develop IR, 3000 specs (ran)
• tests 310 ✓ / 656 skip (GPU) (ran); ABI: develop is 9, takes develop + 1 when it lands; hipDNN C++ not built here
Code quality 🟠 Should-fix
• N2, N3 fixed; each pin mutant fails its own case (ran)
• save/restore reuses develop's restore helpers; scale rule split: N5 (read)
• N6; minor: tests redo expectCommonReadRefusals; pageTablesFor needs pagesKAndV() first (read)
# Sev file:line Finding
N5 🟠 packs/AttentionPagedUnifiedNative.cpp:1320 A plan for a non-positive attn_scale saves, then never loads
• impact: graph_match serves any finite scale (:707); the reader refuses <= 0 (:1320), and a test pins that (TestAttentionPagedUnifiedLaunchValues.cpp:329) (read)
• fix: one served-scale predicate both call, like the dense pack's isSupportedAttentionScale; add a save-then-load case at a negative scale
N6 🟡 packs/AttentionPagedUnifiedNative.cpp:868 (+ :209, :141) Second homes for facts develop or the PR already holds
• impact: kernelSignature() restates the 21 arguments of unified2DKernelSignature() (Gfx950AttentionUnified2DNative.cpp:464); spanProduct is develop's zero-admitting checkedProduct a 4th time; NO_UID restates the header's (read)
• fix: share the base signature from a header and append the rtl tail; one spanProduct in SdpaGraphMatch.hpp; NO_UID as ATTENTION_PAGED_UNIFIED_NO_UID

Utility: 🟢 Positive — N1–N4 fixed with CI evidence; golden and sinks still 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: changes since 8e16db72 (a rebuild on develop, then a rebase), CI logs at head (4 rocKE shards, both hipDNN superbuilds, attempt 1), a resolved merge with develop (build + ctest, byte identity, goldens, IR sweep vs develop, grid compile, tests), the IngestorGenerator pin and its mutants, a binding probe, code quality. Flags whether a closer look is needed; not an approval or rejection.

@cgmillette

Copy link
Copy Markdown
Contributor

Automated overlap notice (rocKE merge monitor)

Hi @dpeabody-amd, this PR and #14967 both bound the tiled 2D block-table lookups by the sequence's own ceil(seq_len / BS) entries, and both write the out-of-range buffer offset that makes a K/V read past seq_len load zero. Rows N3 and N5 of our 6 November follow-up review of #14967 flagged them. This notice records the owner of both and adds rebase steps checked on the merge.

The copies, here at 33842d270d and in #14967 at c3f4ca5292 (paths under dnn-providers/hip-kernel-provider/rocke/):

  • Bound, here: the runtime layout sets block_table_max_idx to seq_base + ceil(seq_len / BS) in library/kernels/gfx950/attention_tiled_2d.py:2462-2473 and in the gfx950 C++ preloop (platform/cpp/instances/gfx950/attention_tiled_2d_kv_descriptors_loads.cpp:399-406); gfx942's _rtl_kv_block checks the same entries as lblk * BS < seq_len (library/kernels/gfx942/attention_tiled_2d.py:3184; C++ gfx942/attention_tiled_2d_kv_descriptors_loads.cpp:707-708).
  • Bound, fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967: every layout bounds the lookups by min(seq_base + ceil(seq_len / BS), num_seqs * bt_stride): block_table_max_idx in library/kernels/gfx942/attention_tiled_2d.py:2626-2634 and library/kernels/gfx950/attention_tiled_2d.py:2170-2178, and rocke__block_table_max_idx in both C++ loaders (gfx942 :143-158, gfx950 :138-153).
  • Offset, here: _RTL_KV_OOB_VOFF = 0x7FFFFFF0 (library/kernels/gfx942/attention_tiled_2d.py:1156; C++ ctx->rtl_kv_oob, gfx942/attention_tiled_2d_kv_descriptors_loads.cpp:416) is the voffset of the masked block-table read (:3185-3187; C++ :712) and of K/V tokens past seq_len (:3203-3204; C++ :748-749). gfx950's _rtl_kv_voff selects the page span instead (library/kernels/gfx950/attention_tiled_2d.py:2552; C++ :637-638).
  • Offset, fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967: OUT_OF_RANGE_OFFSET = 0x7FFFFFFF and buffer_offset_in_sequence in library/kernels/common/_attention_kv_seq_tail.py:22 and :33-38, mirrored once in platform/cpp/include/rocke/instance_attention_kv_seq_tail_internal.h:15 and :29-37.

Owner: #14967 for both, as our review proposed and the monitor recorded. Its bound covers every layout, its offset is one named constant per engine in shared homes (kernels/common and a common C header), and it merges with develop without a conflict, so it lands first. This PR's copies came first (21 October). Resolution: this PR drops its runtime-layout bound in both engines and the gfx942 lblk check for #14967's block_table_max_idx, and _RTL_KV_OOB_VOFF and the gfx950 span select for #14967's buffer_offset_in_sequence.

Order: the merge plan on develop c24ffe52d2 lists neither PR among its 35 near-merge PRs; this PR appears only in its binary ABI lane. #14967 merges with develop without a conflict; this PR conflicts with develop in 3 files outside rocKE (hip-kernel-provider/src/tests/CMakeLists.txt, integration-tests/SUPPORT_MATRIX_OVERVIEW.md, integration-tests/docs/file-formats.md) and was last pushed on 5 November.

Rebase onto develop with #14967, which adds 6 conflicting files (both tiled 2D kernels, their C++ loaders and their .c parity emitters):

  1. gfx950/attention_tiled_2d.py (5 hunks): in the bound hunk keep fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967's block_table_max_idx and drop the if RTL: branch; in the four loader hunks keep both branches, yours first. The runtime layout keeps a tile inside one block (tile_size == block_size on gfx950, :1122; a divisor of it on gfx942, :1200), so fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967's LANE_FREE_PAGES is off there and the order changes no kernel.
  2. gfx950/attention_tiled_2d_kv_descriptors_loads.cpp (6 hunks): keep both new helper blocks and both loader branches, yours first. In the V loader keep fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967's linear_half = ... assignment, not a second rocke_value_t* linear_half declaration: fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967 declares it before the branches for its V-tail check, and a shadowing declaration still compiles but makes that check rebuild it, so the C++ IR drifts from Python's on partial-row copies. Then delete the if(ctx->RTL) block in rocke_gfx950_attn2d_emit_preloop (:399-406 here): it merges without a conflict and, kept, leaves the C++ engine on the uncapped bound while Python takes fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967's. CI checks only the Python engine's IR, so only a byte-identity run of both engines would show it.
  3. gfx942 (attention_tiled_2d.py, 3 hunks; its C++ loader, 2): keep both sides, with fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967's linear_half = None (C++: = NULL) before your RTL branch. In _rtl_kv_block / rocke__rtl_kv_block, bound the table read by block_table_max_idx (the index seq_base + lblk below it: the same entries as lblk * BS < seq_len, plus the table end).
  4. Build the masked voffsets of _rtl_kv_block, _rtl_kv_call and _rtl_kv_voff with buffer_offset_in_sequence (C++ rocke_kv_buffer_offset_in_sequence), and drop _RTL_KV_OOB_VOFF and ctx->rtl_kv_oob.
  5. V rows: fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967's V-tail guard now also runs on your runtime-layout V copies, whose voffset already bounds tokens by seq_len. Keep one of the two: skip the guard on the runtime-layout branch, or drop your V select and let the guard bound them.
  6. Parity indices: both PRs add configs at gfx942 index 18 and gfx950 indices 72-73. The .c emitters conflict; the .py emitters merge without a conflict, and your _CONFIGS[18], [72] and [73] silently replace fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967's, while platform/cpp/bindings/attention_tiled_parity_configs.py already lists your rows after fix(rocke): stop OOB reads of K/V and the block table in unified attention #14967's. Renumber your runtime-layout rows to gfx942 19-37 and gfx950 74-92 in both emitters and their idx comments.
  7. The out-of-range values and the gfx942 compare change the runtime-layout IR: re-run the engines' byte identity and your runtime-layout numeric tests on gfx942 and gfx950.

Checked with git merge-tree on develop c24ffe52d2 (no build): #14967 merges cleanly; this PR conflicts in 3 files, and in 9 once #14967 is in. In that trial merge the gfx950 C++ keeps both bounds and the .py emitters keep both sets of indices, without a conflict. The copies come from both heads' sources.

If you think this PR should own them instead, reply here and on #14967, 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.

@dpeabody-amd
dpeabody-amd force-pushed the users/dpeabody/hipdnn-cdna-paged-attn branch from 33842d2 to f1d0d68 Compare November 6, 2026 17:38
@tenpercent
tenpercent requested a max review from Copilot November 6, 2026 17:53

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

🟡 Changes recommended

Ragged-query mapping can read out of bounds or leave outputs unwritten, valid zero/negative-scale plans cannot reload, and zero-warps inputs can trigger modulo by zero.

4 open issues
What changed in this PR

Adds hipkernel:AttentionPagedUnified to serve paged SDPA with dense or ragged queries on gfx942 and gfx950 through packaged rocKE tiled-attention kernels.

Changes:

  • Registers the engine, descriptor metadata, and versioned saved-launch contract.
  • Wires runtime-layout support through C++ kernels, bindings, ABI versioning, and parity configurations.
  • Adds schema-gate tests, installed test registrations, engine-pinned integration lanes, and support claims.
File Description
projects/​hipdnn/​tools/​IngestorGenerator/​tests/​test_gfx950_attention_dense_schema_gate.py Updated as part of this pull request.
projects/​hipdnn/​tools/​IngestorGenerator/​README.md Updated as part of this pull request.
dnn-providers/​integration-tests/​SUPPORT_MATRIX_OVERVIEW.md Updated as part of this pull request.
dnn-providers/​integration-tests/​integration-test-bundles/​quick/​SdpaFwd/​bshd/​fp16/​hd64_window_ragged_paged/​Mixed/​Mixed.support.json Updated as part of this pull request.
dnn-providers/​integration-tests/​integration-test-bundles/​quick/​SdpaFwd/​bshd/​fp16/​hd64_causal_paged/​SpecDecode/​SpecDecode.support.json Updated as part of this pull request.
dnn-providers/​integration-tests/​integration-test-bundles/​quick/​SdpaFwd/​bshd/​bf16/​hd128_causal_paged/​Decode/​Decode.support.json Updated as part of this pull request.
dnn-providers/​integration-tests/​docs/​file-formats.md Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​TestKernelIngestorEngine.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​packs/​TestIngestorDispatchAliases.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​packs/​TestGfx950AttentionDenseSchemaGate.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​engines/​kernel_ingestor_engine/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​tests/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​integration_tests/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​packs/​AttentionPagedUnifiedLaunchValues.hpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​IngestorPacks.hpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​IngestorPacks.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​attention_paged_unified/​README.md Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​attention_paged_unified/​operation_is_gfx950.umd.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​attention_paged_unified/​operation_is_gfx942.umd.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​attention_paged_unified/​kernel_dtype_matches_graph.umd.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​attention_paged_unified/​attention_paged_unified.ued.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​attention_paged_unified/​attention_paged_unified.udd.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​descriptors/​rocKE/​attention_paged_unified/​attention_paged_unified.kmd.json Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​engines/​kernel_ingestor_engine/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​src/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​python/​rocke/​portable_ir/​src/​abi.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx950/​attention_tiled_2d_q_load_gather.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx950/​attention_tiled_2d_public_entry_glue.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx950/​attention_tiled_2d_kv_descriptors_loads.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx950/​attention_tiled_2d_kv_body_qk_softmax.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx950/​attention_tiled_2d_kv_body_pv_epilogue.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx950/​attention_tiled_2d_fastkv_regp_helpers.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx942/​attention_tiled_2d.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx942/​attention_tiled_2d_q_load_gather.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx942/​attention_tiled_2d_public_entry_glue.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx942/​attention_tiled_2d_loop_scaffold.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx942/​attention_tiled_2d_kv_descriptors_loads.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx942/​attention_tiled_2d_kv_body_qk_softmax.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​gfx942/​attention_tiled_2d_epilogue.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​instance_gfx950_attention_tiled_2d_internal.h Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​instance_gfx942_attention_tiled_2d_internal.h Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​instance_attention_tiled_2d_rtl_internal.h Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​helper_helper_rocke.instances.gfx942.attention_tiled_2d.h Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​abi.h Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​bindings/​rocke_engine_attention.cpp Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​bindings/​attention_tiled_parity_configs.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​CMakeLists.txt Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​gfx950_attention_tiled_2d_emit.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​gfx950_attention_tiled_2d_emit.c Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​gfx942_attention_tiled_2d_emit.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​gfx942_attention_tiled_2d_emit.c Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​dispatch/​attention/​unified_rules.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​dispatch/​attention/​tuning_specs.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​dispatch/​attention/​common.py Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​ROCKE_ENGINE_test_categories_external.yaml Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​hipkernel_AttentionPagedUnified_test_categories_integration.yaml Updated as part of this pull request.

Review effort: Max

Comment on lines +424 to +427
rocke_value_t* rows_ragged = rocke_b_sub(b, off_stop, off_start);
rocke_value_t* owned = rocke_b_select(b, dense, q_len_u, rows_ragged);
rocke_value_t* live = rocke_b_smin(b, owned, slq);
ctx->cur_batch_q_len = rocke_b_select(b, has_slq, live, owned);

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Agreed that these are valid graphs and that the engine gives wrong output for them. I tried the fix you suggest, mapping q-blocks per batch (grid of B × ceil(S/BLOCK_Q) q-blocks, each batch's own origin and seq_len_q), in all four builders. It is correct for any origin order, but it is much slower on the serving case it exists for, a mixed batch: bf16, Hq 64, Hkv 8, D 128, block 16, 63 decode rows (KV 4096) plus one 2048-token prefill, median of 100 runs:

gfx942 gfx950
packed numbering (this PR) 2844 µs 954 µs
per-batch grid 4378 µs (+54%) 2186 µs (+129%)
per-batch, grid capped at the packed size 3627 µs (+27%) 1281 µs (+34%)

Decode-only batches were unchanged or faster. The cost is the empty q-blocks: with S = 2048, 63 of 64 sequences get 255 empty blocks each (about 128k of 131k CTAs do nothing), and testing liveness first or numbering sequences fastest did not recover it. So the PR keeps the packed numbering and documents the requirement in the engine README and the kernels: origins non-decreasing, off[b+1] - off[b] >= seq_len_q[b], and off[B] the end of the last block, which vLLM's query_start_loc meets. What the engine can check at match time it still checks (seq_len_q bound, an INT32 [B+1,1,1,1] offset table). The follow-up that removes the requirement without empty blocks is a packed numbering driven by seq_len_q: each q-block finds its sequence from a wave prefix sum over ceil(min(seq_len_q[b], S) / BLOCK_Q), so the offsets are only read as origins.

@yraparti yraparti left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

rocKE checklist review

ROCm/rocm-libraries#14247 · f1d0d688696e · model gemini-4-argon, fallback gpt-5.5

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

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

Findings

  • HIGH §A · dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:503 · Add representative golden hash entries for the runtime-layout flag-on IR for each affected architecture, LLVM flavor, and emitted path.
  • HIGH §P · dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1105 · Add a C++ runtime-layout LDS byte calculation and capacity check matching the Python gate, and cover it with a parity/rejection test case.
  • MEDIUM §E · dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:528 · Add the required rocKE support/reference documentation for the new runtime_layout attention feature in the same PR.
  • LOW §S · dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h:109 · Run PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py from rocke/platform/ and record the results in the PR description.

Needs info

  • §A.4 Missing canonical IR or byte-identity output showing that the new Python and C++ runtime-layout emitter order is CANON_EQUAL for the added cases.
  • §A.5 Missing representative-IR golden status showing existing byte-identical shapes remained unchanged and any intentional runtime-layout IR changes were handled intentionally.
  • §A.8 Missing evidence that the C++ attention twin was reviewed or checked with the cross-engine byte-identity gate, rather than relying on CI being green.
  • §B.4 runtime_layout: bool = False in public specs UnifiedAttention2DTiledSpec (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:526-535 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:543-567) changes the kernel ABI by appending 15 kernel arguments (RUNTIME_LAYOUT_EXTRA_ARGS / _RTL_PARAMS), but __post_init__ only checks if type(self.emit_lse) is not bool: and omits if type(self.runtime_layout) is not bool:, allowing truthy non-bools (like "false" or 1) to coerce into an enabled runtime-layout ABI. — no diff line supports this verdict
  • §B.6 _runtime_layout_reject_reason and _supports_runtime_layout (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1183-1333 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1061-1135) check only functional/architectural invariants (pipeline flags, head divisibility, tile/DMA copy divisibility, LDS capacity) and report hardware/algorithm invariant reasons. — no diff line supports this verdict
  • §C.2 No bespoke shape-class builder is added; runtime-layout support is integrated via spec.runtime_layout inside build_unified_attention_2d_tiled in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1863 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1366. — no diff line supports this verdict
  • §C.3 kernel_name() appends the rtl codegen mode token ("rtl" if self.runtime_layout else "") without any workload sizes in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1108 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1032. — no diff line supports this verdict
  • §D.3 Kernel-level _supports_runtime_layout and _runtime_layout_reject_reason return specific invariant reasons for each check (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1183, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1281, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1061). — no diff line supports this verdict

Section notes

  • §D: Items 1-11 of Section D (Dispatcher: registration, routing & spec integration). The PR adds the runtime_layout spec field for hipDNN's AttentionPagedUnified ingestor engine; Python attention dispatch does not route to runtime_layout.
  • §E: Diff is truncated; items that require checking omitted hunks are marked NEEDS-INFO.
  • §G: Items 1-17 of 34 in Section G.
  • §G: Items 18-34 of Section G graded. Several test and documentation files were omitted from the truncated diff.
  • §H: Trace of device-side masks and bounds in runtime_layout (RTL): 1. Page-table lookup mask: - gfx942: blk_ok = lblk_tok < seq_len where lblk = (kv_tile_idx * T) / BS (kernels/gfx942/attention_tiled_2d.py:3197-3202). Extent source: runtime seq_len (seq_len_kv[seq_idx]). Range visited: kv_tile_idx in [0, num_tiles), where num_tiles = ceil(min(context_len + qb_start_pos + bm1_div_nqk + 1, seq_len) / T) (:2640-2642). Since T divides BS, lblk_tok <= tok0 < num_tiles * T, and when lblk_tok >= seq_len, bt_voff is set to _RTL_KV_OOB_VOFF (0x7FFFFFF0 > big_bytes = 0x7FFF0000), returning 0 from HW buffer load. - gfx950: idx < block_table_max_idx where block_table_max_idx = seq_base + ceil(seq_len / BS) (kernels/gfx950/attention_tiled_2d.py:2471-2476, :2529-2537). Extent source: runtime seq_len. Visited range: kv_tile_idx < num_tiles <= ceil(seq_len / T) with T == BS. Agree. 2. Per-page KV buffer resource & token voffset mask: - Each page builds…
  • §J: Shapes: | Equation / Operand | Inner Dims | Result vs Declared | |---|---|---| | None in visible diff (descriptors/rocKE/attention_paged_unified/README.md omitted) | N/A | N/A | Symbols: | Symbol | Defined At | Domain | |---|---|---| | <ENGINE>_test_categories_integration.yaml | dnn-providers/integration-tests/docs/file-formats.md:599 | Filename pattern | | (All symbols in descriptors/rocKE/attention_paged_unified/README.md omitted from diff) | N/A | N/A | Novelty: | Claim | Type | Verified Against | |---|---|---| | hipkernel_AttentionPagedUnified added to integration test categories table | Added | dnn-providers/integration-tests/docs/file-formats.md:599 | | Claims in descriptors/rocKE/attention_paged_unified/README.md (155 lines) | Unknown | Omitted from diff |

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 new Python attention runtime_layout spec field is mirrored in the C++ attention spec/twin (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:503…
2 N/A No core lowering-op or op-handler files under core/lower_* or core/ir are changed in the shown diff.
3 PASS The kernel-level runtime_layout knob is carried through the C++ binding and parity configuration rather than only the Python dispatch layer (dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp:824…
4 NEEDS-INFO Missing canonical IR or byte-identity output showing that the new Python and C++ runtime-layout emitter order is CANON_EQUAL for the added cases.
5 NEEDS-INFO Missing representative-IR golden status showing existing byte-identical shapes remained unchanged and any intentional runtime-layout IR changes were handled intentionally.
6 PASS The diff adds cross-engine C/Python parity emit cases for the new runtime-layout configurations rather than relying only on golden files (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.c:347…
7 PASS The new emitted runtime-layout logic is gated on spec.runtime_layout/ctx->RTL, with appended runtime-layout params declared only in that branch (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1359…
8 NEEDS-INFO Missing evidence that the C++ attention twin was reviewed or checked with the cross-engine byte-identity gate, rather than relying on CI being green.
9 N/A The shown diff does not make a shared op-handler or lowerer-only arch-divergence fix that would require a mini-pair proof.
10 FAIL A new flag-on IR feature is introduced via runtime_layout (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:503) and only parity/runtime-layout cases are visible (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.py:334), with no representative golden hash entries for the new flag-on IR in the file list/diff.
11 N/A No new IR opcode is introduced in the shown diff.
12 N/A The shown diff does not change ROCKE_BACKEND=both behavior or add any architecture to CPP_UNPORTED_ARCHES.
13 N/A No emitted .ll stamping or debug-location gate behavior is changed in the shown diff.
14 N/A The shown diff does not change a classic helper under platform/cpp/helpers/; the touched attention instance helper/header code…
§B graded 6/6
# Verdict Evidence
1 PASS New rejections in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1183-1206, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1061-1135, and…
2 PASS Preflight gates _supports_tiled_2d_spec (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:7354-7355) and supports_tiled_2d_spec (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:5447-5450)…
3 PASS runtime_layout defaults to False and is excluded from rocKE dispatch (dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:990-998, tuning_specs.py:183-185), while the hipDNN ingestor pack logs every decline reason at INFO…
4 NEEDS-INFO runtime_layout: bool = False in public specs UnifiedAttention2DTiledSpec (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:526-535 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:543-567) changes the kernel ABI by appending 15 kernel arguments (RUNTIME_LAYOUT_EXTRA_ARGS / _RTL_PARAMS), but __post_init__ only checks `if…
5 PASS The newly opened runtime_layout gate scope (fp16/bf16, head sizes 64/128/256, block sizes 16/32/64, GQA ratios 1..16, sinks, sub-block tiles on gfx942) is covered in parity emitters…
6 NEEDS-INFO _runtime_layout_reject_reason and _supports_runtime_layout (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1183-1333 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1061-1135) check only functional/architectural invariants (pipeline flags, head divisibility, tile/DMA copy divisibility, LDS capacity) and report hardware/algorithm…
§C graded 10/10
# Verdict Evidence
1 PASS Repeated predicates and parameter declarations are factored into shared helpers such as SW_ACTIVE = SLIDING_WINDOW > 0 or RTL in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1369, rocke_attn2d_rtl_declare_params in…
2 NEEDS-INFO No bespoke shape-class builder is added; runtime-layout support is integrated via spec.runtime_layout inside build_unified_attention_2d_tiled in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1863 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1366. — no diff line supports this verdict
3 NEEDS-INFO kernel_name() appends the rtl codegen mode token ("rtl" if self.runtime_layout else "") without any workload sizes in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1108 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1032. — no diff line supports this verdict
4 PASS Arch-specific kernel and instance builders stay in their respective gfx942 and gfx950 modules (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:526…
5 PASS All touched files correspond to the inherent fan-out across Python/C++ engines, gfx942/gfx950 architectures, parity configs (dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/attention_tiled_parity_configs.py:333), ABI bump…
6 PASS The change cleanly builds on prior refactoring (check_head_counts, runtime kernargs, and shared SdpaGraphMatch.hpp helpers in dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/packs/SdpaGraphMatch.hpp:261) while keeping runtime_layout=False…
7 N/A This PR introduces a new feature (runtime_layout mode and hipkernel:AttentionPagedUnified engine) rather than a standalone…
8 PASS The implementation reuses existing primitives and helpers on develop such as check_head_counts / rocke_attention_check_head_counts (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1196…
9 PASS Invariants and rejection guards are applied across all sibling sites: Python gfx942 (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1183), Python gfx950…
10 PASS Widening UnifiedAttention2DTiledSpec with runtime_layout is handled across all consumers: excluded from tuning spec hashes and defaults when off (dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:996…
§D graded 21/21
# Verdict Evidence
1 PASS no added os.environ or getenv call
2 N/A This PR does not add or modify any KernelCandidate or Capability declarations in the Python dispatcher registry.
3 NEEDS-INFO Kernel-level _supports_runtime_layout and _runtime_layout_reject_reason return specific invariant reasons for each check (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1183, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1281, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1061). — no diff line…
4 PASS The new catalog/geometry spec helpers runtime_layout_geometry and runtime_layout_catalog_spec take arch explicitly as a parameter and do not probe the host device (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1226…
5 PASS runtime_layout is included in the tuning spec identity when True and omitted when False to preserve existing spec_hash/config_key values, is marked as a semantic field in _semantic_fields, and appends the 'rtl' token to kernel_name()…
6 N/A This PR does not add or modify any cohort routing in the Python dispatcher.
7 N/A This PR does not add or modify any hand-pinned overrides or heuristic tables in the dispatcher.
8 N/A This PR does not define or modify any dispatcher cohort predicates.
9 N/A This PR does not register a new KernelCandidate in ATTENTION_ROUTE_REGISTRY or ATTENTION_EXECUTION_REGISTRY; runtime_layout is…
10 PASS runtime_layout is an ingestor-engine catalog mode rather than a Python AttentionRequest feature, which is explicitly documented in the PR body and code comments and wired end-to-end through the hipDNN ingestor engine and C++/Python bindings…
11 N/A This PR does not register or promote any candidate in the Python dispatcher registries.
12 N/A This PR does not add or reprioritize a Python CandidateRegistry candidate, so dispatcher priority ordering and name tie-breaking…
13 N/A This PR does not add a CandidateRegistry pin wrapper or pinned dispatcher entry point, so pinned-request fallback behavior does…
14 PASS The disabled runtime_layout field is excluded from the spec hash/cache identity at dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:991, and the C/Python ABI versions are bumped for the new spec field at…
15 N/A This PR does not add a new Python dispatcher candidate or arch module in the CandidateRegistry sense, so additive registration…
16 N/A This PR adds a kernel-ingestor engine and runtime-layout kernel mode, not a new Python dispatcher operator family requiring…
17 N/A No new CandidateRegistry candidate is added, so the CPU routing test requirements through admits() do not apply.
18 N/A The diff does not modify dispatcher registry ratchets or exempt sets, so require_build/require_binding/require_torch_binding…
19 PASS The diff explicitly avoids claiming dispatcher auto-routing for the new runtime_layout mode: dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:991 says dispatch never sets it, and…
20 PASS The Python kernel spec adds runtime_layout at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:526, and the C++ side mirrors it in the public spec/binding at…
21 N/A This is a new runtime-layout kernel mode and ingestor engine, not a selection-layer bug fix where the chooser should change…
§E graded 17/17
# Verdict Evidence
1 NEEDS-INFO The diff omits platform hunks and new test files needed to verify that platform code/tests do not import library code.
2 PASS no added torch import under library/ outside tests
3 NEEDS-INFO The omitted test hunks are needed to confirm any torch use is gated with pytest.importorskip and not used to locate ROCm libraries.
4 NEEDS-INFO The omitted library test hunks and full import sections are needed to verify no upward imports and that known violation lists only shrink.
5 NEEDS-INFO The platform CMake/test hunks are omitted, so absolute or repo-escaping paths and helper-script hygiene under platform/ cannot be fully checked.
6 NEEDS-INFO Emitter trees are modified, but the diff does not provide process evidence that no lint autofix removed intentional handles or that byte identity was rerun after formatting/linting.
7 FAIL dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:528 adds the new public runtime_layout feature, but the changed file list does not include the required rocKE reference/support docs such as instances/SUPPORT_MATRIX.md or dsl_docs/instances entries.
8 NEEDS-INFO Several documentation and generated descriptor hunks are omitted, so Jira IDs, PR numbers, reviewer handles, or tracker URLs in committed artifacts cannot be ruled out.
9 NEEDS-INFO The diff is too truncated to determine whether all touched areas are part of the same deliverable or include unrelated fixes.
10 NEEDS-INFO No trusted diff artifact provides the required merge-simulation evidence: overlapping open PRs, agreed land order, and rerun affected goldens/suites on develop plus those PRs.
11 NEEDS-INFO Large documentation, descriptor, generated, and test hunks are omitted, so scratch notes or throwaway design files cannot be ruled out.
12 PASS Both target arches are covered by added parity/runtime-layout rows for gfx942 at dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/attention_tiled_parity_configs.py:326 and for gfx950 at…
13 N/A No dead-code removal or concept rename is evident from the provided diff.
14 N/A No rebase, cherry-pick, or port operation is evident from the provided diff.
15 NEEDS-INFO Large new source and test hunks are omitted, so every newly accepted parameter cannot be checked for use in its function body.
16 N/A No reviewer-facing severity labels or review comments are part of this code diff.
17 PASS no absolute throughput, bandwidth, or latency figure in added lines
§F graded 11/11
# Verdict Evidence
1 PASS branch users/dpeabody/hipdnn-cdna-paged-attn uses the users// prefix
2 PASS title feat(hip-kernel-provider): paged/ragged SDPA engine for gfx942/gfx950 matches type(scope): summary and is 69 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/test_attention_tiled_2d_runtime_layout_gfx950.py is in the diff
5 PASS The supplied file list contains no key/cert/secret paths, and visible additions are source/test files such as dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/packs/AttentionPagedUnifiedNative.cpp:1.
6 NEEDS-INFO No citable diff evidence or CI artifact shows that pre-commit was run locally and green before push.
7 NEEDS-INFO Bot/Copilot review-thread state and any fix, won't-fix justification, or deferred-ticket IDs are not included in the diff.
8 N/A The PR targets develop, so the stacked-PR-on-a-feature-branch bot-skipping case does not apply.
9 N/A The supplied diff and metadata do not indicate that changes were split or trimmed for a PR-size cap.
10 PASS The added tests exercise concrete behavior rather than being empty bot-green checks, e.g. the ragged-graph schema-gate test at dnn-providers/hip-kernel-provider/src/tests/engines/kernel_ingestor_engine/packs/TestGfx950AttentionDenseSchemaGate.cpp:197.
11 NEEDS-INFO No citable diff evidence shows that the title/body were verified from the PR API fields after editing.
§G graded 34/34
# Verdict Evidence
1 NEEDS-INFO Parity configs are added in dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.py:372 and gfx950_attention_tiled_2d_emit.py:1194, while device numerical tests in test_gfx942_attention_runtime_layout.py and test_attention_tiled_2d_runtime_layout_gfx950.py are omitted from the truncated diff.
2 NEEDS-INFO The GPU test files (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, IntegrationGpuAttentionPagedUnified.cpp) are omitted from the diff, and on-silicon execution results are not shown.
3 NEEDS-INFO Local non-GPU gate sweep output (tools/run_checks.py, check_byte_identity.py across llvm20/22/23, check_ir_validity.py, pytest) cannot be verified from the diff alone.
4 NEEDS-INFO TestGfx950AttentionDenseSchemaGate.cpp:201-244 has falsifiable tests, but the new runtime-layout and pack unit test files (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, TestAttentionPagedUnified*.cpp) are omitted from the diff.
5 NEEDS-INFO TestGfx950AttentionDenseSchemaGate.cpp:171-182 tests through makeEngine/isApplicable, but the remaining new test files (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, TestAttentionPagedUnified*.cpp) are omitted from the diff.
6 NEEDS-INFO Stale kernel-argument count in C++ comments: dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_tiled_2d_internal.h:188, dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d_ctx_prologue.cpp:538, dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx950_attention_tiled_2d_internal.h:187, and…
7 NEEDS-INFO The engine README (dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md) and device test files are omitted from the truncated diff.
8 NEEDS-INFO The test harnesses (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, IntegrationGpuAttentionPagedUnified.cpp) are omitted from the truncated diff.
9 NEEDS-INFO The PR body cites SGPR spill and instruction counts from descriptors/rocKE/attention_paged_unified/README.md, which is omitted from the truncated diff.
10 N/A This PR introduces functional support for a new paged/ragged SDPA engine and makes no latency or throughput performance claims.
11 NEEDS-INFO AttentionPagedUnifiedNative.cpp is truncated at line 600 before the dispatch registration and handler implementation, and pack/integration tests are omitted from the diff.
12 PASS Both Python and C++ rocKE builders are updated in lockstep (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:526, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:543…
13 N/A No new LLVM intrinsic overload classes or mangled intrinsic symbols are introduced or modified in this PR.
14 N/A This PR does not make hardware bottleneck classification claims.
15 PASS The PR diff is inspected directly from head commit f1d0d68 (e.g., dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h:109).
16 PASS Symbols and shared helpers match existing codebase conventions across layers (e.g., check_head_counts / rocke_attention_check_head_counts in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d.cpp:38 and 220).
17 NEEDS-INFO Emitter scripts are updated in dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.py:337 and gfx950_attention_tiled_2d_emit.py:1163, but end-to-end execution output is not present in the diff.
18 PASS dnn-providers/hip-kernel-provider/src/tests/engines/kernel_ingestor_engine/packs/TestGfx950AttentionDenseSchemaGate.cpp:198 proves the schema guard fires on ragged and non-default-alignment graphs (lines 198, 214) and includes a control case where the fault cannot…
19 NEEDS-INFO The test fixture files (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, IntegrationGpuAttentionPagedUnified.cpp, TestAttentionPagedUnifiedMatchers.cpp) are omitted from the diff.
20 N/A This PR does not add or modify a claim-ban detector or linter.
21 N/A This PR does not add benchmarking or performance measurement code.
22 NEEDS-INFO The new test files (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py) and golden contract test (TestIngestorPlanLaunchContractGoldens.cpp) are omitted from the diff, so skip conditions and golden entry formats cannot be verified.
23 N/A This PR adds a new paged/ragged SDPA engine and runtime_layout kernel mode without making speedup or routing A/B claims.
24 N/A This PR does not re-bless existing IR goldens; runtime_layout=False leaves existing emitted IR unchanged.
25 N/A This PR is a feature addition (runtime_layout mode and hipkernel:AttentionPagedUnified engine) rather than a codegen…
26 N/A This PR does not include performance benchmarks or speedup claims.
27 N/A This PR does not report a performance A/B comparison.
28 NEEDS-INFO dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:525 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:542 keep runtime_layout=False byte-identical (0-delta) and do not table counterfactual comparisons. — no diff line supports this verdict
29 N/A This PR does not swap guarded math primitives for fast approximations.
30 NEEDS-INFO The PR description cites SGPR-to-VGPR lane spill and zero scratch usage from descriptors/rocKE/attention_paged_unified/README.md, which is omitted from the diff.
31 N/A This PR adds a new engine and kernel mode rather than claiming a fix or speedup over a baseline defect.
32 NEEDS-INFO The static register-pressure and instruction-share measurements cited in the PR body live in descriptors/rocKE/attention_paged_unified/README.md, which is omitted from the diff.
33 PASS dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:995, tuning_specs.py:185, and unified_rules.py:401 explicitly exclude runtime_layout when False so existing dispatch/tuning decision functions and cache keys remain 0-delta on develop.
34 N/A runtime_layout is a layout/execution mode rather than an opt-in extra output flag, and it explicitly rejects emit_lse at…
§H graded 5/5
# Verdict Evidence
1 NEEDS-INFO Host-side shape/stride/span validation for kv_block_span_bytes and MAX_PAGE_TABLE_ENTRIES is visible in AttentionPagedUnifiedNative.cpp:477-504, but the dispatch/prepare/execute sections of AttentionPagedUnifiedNative.cpp (truncated after line 735) and the test files are omitted from the diff.
2 NEEDS-INFO In rocke/library/kernels/gfx942/attention_tiled_2d.py:3205-3206 and rocke/library/kernels/gfx950/attention_tiled_2d.py:2538-2541, the physical page id phys read from the block table is multiplied in 64-bit (b.mul(b.zext(phys, I64), rtl_kv_page_stride_b)) and added to cache to form a per-page buffer resource base, which bypasses num_records containment across the KV container (num_blocks); because…
3 NEEDS-INFO Power-of-two constants are enforced in the spec: head_size in (64, 128, 256), block_size in (16, 32, 64), num_warps in (1, 2, 4, 8), and tile_size dividing block_size as a multiple of 16 (rocke/library/kernels/gfx942/attention_tiled_2d.py:1199-1205, 1299-1314; rocke/library/kernels/gfx950/attention_tiled_2d.py:1123-1135, 138-149), and in AttentionPagedUnifiedNative.cpp:179-181. — no diff line…
4 NEEDS-INFO Both num_tiles and the per-page/per-token masks are derived dynamically on-device from the same runtime seq_len (seq_len_kv[seq_idx]), and softmax masks out col_abs >= max_seq_prefix_len where max_seq_prefix_len = min(context_len + qb_start_pos + bm1_div_nqk + 1, seq_len) (rocke/library/kernels/gfx942/attention_tiled_2d.py:2640-2647, 3199-3219, 5321-5322…
5 NEEDS-INFO Host validations in graphMatches and SdpaGraphMatch.hpp are visible (AttentionPagedUnifiedNative.cpp:324-605, SdpaGraphMatch.hpp:261-632), but the launch-time checks in AttentionPagedUnifiedNative.cpp (truncated at line 735) and all unit/negative test files (TestAttentionPagedUnifiedMatchers.cpp, TestAttentionPagedUnifiedLaunchValues.cpp, TestAttentionPagedUnifiedPacks.cpp…
§I graded 6/6
# Verdict Evidence
1 N/A This PR adds the runtime_layout kernel mode and the hipkernel:AttentionPagedUnified engine with a single deterministic catalog…
2 PASS Unsupported spec combinations, builders, and graph attributes are loudly rejected with explicit reasons in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1183…
3 NEEDS-INFO The diffs for dnn-providers/hip-kernel-provider/rocke/library/tests/test_gfx942_attention_runtime_layout.py, dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_tiled_2d_runtime_layout_gfx950.py, and the C++ pack tests are omitted from the truncated diff.
4 N/A This PR does not add or update performance benchmark or routing comparison tables with shipped vs achievable columns.
5 NEEDS-INFO The diffs for dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md and projects/hipdnn/tools/IngestorGenerator/configs/attention_paged_unified.yaml are omitted, so the regeneration path for the generated descriptors cannot be verified from the diff.
6 N/A This PR does not introduce any new DSL ops or hardware-instruction-named primitives.
§J graded 7/7
# Verdict Evidence
1 NEEDS-INFO The new Markdown document dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md (+155 lines) is omitted from the truncated diff, so any equations or shape annotations inside it cannot be verified.
2 NEEDS-INFO The contents of dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md (+155 lines) are omitted from the diff, preventing verification of symbol definitions and domains.
3 NEEDS-INFO The contents of dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md (+155 lines) are omitted from the diff, so we cannot check whether it restates algorithms owned by another document.
4 NEEDS-INFO The contents of dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md (+155 lines) and dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md (+8 -2 lines) are omitted from the diff.
5 NEEDS-INFO The contents of dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md (+155 lines) are omitted from the diff, so any complexity or scaling claims in it cannot be checked.
6 NEEDS-INFO The contents of dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md (+155 lines, which the PR body notes contains register-pressure and catalog constants) are omitted from the diff.
7 NEEDS-INFO dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md is added (+155 lines), but its diff hunk is omitted, so its entry-point declaration and scope cannot be verified.
§K graded 13/13
# Verdict Evidence
1 N/A This PR does not add or modify any tuned dispatch candidates.
2 NEEDS-INFO The new runtime_layout spec field (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:526, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:543) is explicitly classified in _semantic_fields at dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/tuning_specs.py:185. — no diff line supports this verdict
3 PASS runtime_layout is omitted from spec_hash items when False (dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:996) and excluded from defaults()…
4 PASS KnobSpace.defaults() in dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/unified_rules.py:401 excludes runtime_layout alongside emit_lse so problem-Independent defaults and portability are preserved.
5 N/A This PR does not add or modify any tuning knobs, axes, or drop/refuse rules.
6 NEEDS-INFO dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/tuning_specs.py:185 classifies runtime_layout as a semantic field without invoking production heuristics, and spec legality is gated in __post_init__ / _supports_tiled_2d_spec (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:533…
7 N/A This PR does not modify KnobSpace.base(req) or request normalization.
8 N/A This PR does not touch knob normalization or axis type coercion.
9 N/A This PR does not touch benchmark recording, replay, or PinRefused handling.
10 N/A This PR does not modify KnobSpace.find or tuning_id resolution.
11 N/A This PR does not touch sweep entry points or core.pin_to_spec.
12 N/A This PR does not reorder production stacks, change accept_default, or alter any dispatch variant's base spec.
13 N/A This PR does not modify any benchmark or sweep harness.
§L graded 10/10
# Verdict Evidence
1 N/A This PR does not add or modify a gfx target in the architecture specs or ISA backend registries.
2 N/A This PR does not add or modify any MMA op_id or fragment layout info.
3 N/A This PR does not change MMA atom selection, accumulator layouts, or intrinsic emission.
4 NEEDS-INFO The gfx942 and gfx950 kernel builders use the target's wave size (64 lanes) and reject non-target architectures via validate_tiled_attention_arch in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1304 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1126. — no diff line supports this verdict
5 NEEDS-INFO gfx942 and gfx950 runtime_layout implementations remain separate and enforce their own arch-specific pipeline constraints (16x16x16 on gfx942 vs 16x16x32 on gfx950) in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1120 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1061. — no diff line supports this verdict
6 N/A This PR does not touch ArchTarget, core/arch/, or family policy files, nor does it import CK Tile's dispatcher.
7 NEEDS-INFO The runtime_layout mode restricts supported dtypes to fp16/bf16 and rejects kv_storage_dtype (FP8) in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1185 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1082. — no diff line supports this verdict
8 N/A This PR does not add or modify scaled-MMA operations or scale contracts.
9 N/A This PR does not add a new gfx architecture target to the platform.
10 N/A This PR does not touch intrinsic arch-domain artifacts or LLVM flavor membership lists.
§M graded 10/10
# Verdict Evidence
1 NEEDS-INFO Host-time Python if statements test only spec booleans (if RTL:), while runtime SSA values use b.cmp_* with b.select, b.scf_if, and b.scf_for without escaping across regions (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1895-1982, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1572-1661). — no diff line supports this verdict
2 N/A The diff does not touch legalize_for_target or any target legalization passes.
3 N/A The diff does not modify print_ir or ck.dsl.ir/v1 IR serialization (core/ir_serialize.py or cpp/core/ir/serialize.cpp).
4 N/A The diff does not modify the portable IR schema, importer, recipe expansion, or recorder coverage (only bumping BINARY_ABI in…
5 NEEDS-INFO The runtime_layout async K/V loads keep the existing wave-uniform LDS base (K_wave_base/V_wave_base) and only supply the per-page buffer resource and masked voff (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:3242-3251, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:2741-2764), with spec validation enforcing whole async-DMA…
6 NEEDS-INFO The grid-stride loop and rtl_rows_if/rtl_live_block conditions are CTA-uniform, b.sync() is placed after exiting the conditional before the next loop iteration (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6376-6382, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:5419-5424), and out-of-bounds/conditional reads use…
7 PASS Buffer resources for runtime_layout are constructed via b.buffer_rsrc / rocke_b_buffer_rsrc without hard-coding word 3 or using illegal temporal hints (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:3189-3206…
8 N/A The diff does not add any new LLVM intrinsic declarations or modify backend engine selection precedence.
9 PASS Existing public headers under cpp/include/rocke/ are not renamed, ROCKE_ABI_VERSION and BINARY_ABI are bumped in lockstep from 10 to 11 (dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h:109…
10 N/A The diff does not touch !dbg emission, LLVM module debug flags, or DWARF consumers.
§N graded 18/18
# Verdict Evidence
1 N/A The PR does not add or modify any module under platform/python/rocke/helpers/.
2 PASS The runtime-layout paths reuse shared helpers binary_search_seq_idx (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1927, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1607) and…
3 N/A The PR does not use or modify SoftwarePipeline.
4 N/A The PR does not add or modify fused GEMM/epilogue kernels or fusion helpers.
5 N/A The PR does not add or modify any rocke.helpers.tiling API.
6 N/A The PR does not touch front-door/toolbox tiling helpers or LayoutStyle.
7 N/A The PR does not use or modify helpers/tiling/ layout overrides or MMA drivers.
8 N/A The PR does not use or modify helpers/tiling/ LDS or global memory verbs.
9 N/A The PR does not use or modify helpers/tiling/ axis roles, batch reductions, or FlowStage.
10 N/A The PR does not change helpers/tiling/ emit, LDS staging, or MMA issue.
11 N/A The PR does not use or modify helpers/tiling/ traits, TileMma, lds_swizzle, or derive_sched_group_counts.
12 N/A The PR makes no LDS bank-conflict or global coalescing claims under helpers/tiling/.
13 N/A The PR does not touch WarpDistributionEncoding, make_tile_desc, or RegisterMapper.
14 N/A The PR does not use helpers/tiling/ data-movement verbs or add a tiling pipeline variant.
15 N/A The PR does not modify helpers/tiling/ allocation or bounds checks on transformed descriptors.
16 N/A The PR does not use or modify optimize_layout or evaluate_transform.
17 N/A The PR does not touch the tiling recorder, round-trip analysis, or tiling visualization.
18 N/A The PR does not touch InterleavedStyle or demo kernels under helpers/tiling/kernels/.
§O graded 16/16
# Verdict Evidence
1 N/A The diff does not modify core/verify.py, cpp/core/verify.cpp, or introduce new address spaces or required IR attributes.
2 PASS In dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1918 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d_ctx_prologue.cpp:135, the new scf.for loop has matching i32 lower, upper, and step…
3 NEEDS-INFO In dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1867 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1541, parameters and IR ops are emitted in fixed tuple order without relying on id(), set iteration, or unsorted dicts. — no diff line supports this verdict
4 PASS In dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1870 and 1984, module singletons I32 and I64 are used directly and SmemType is not modified.
5 N/A The diff does not touch core/passes.py, cpp/core/passes.cpp, or op purity tables.
6 N/A The diff does not touch constant folding in core/passes.py or cpp/core/passes.cpp.
7 N/A The diff does not modify kernel attributes, COMGR options, or cluster_dims validation.
8 N/A The diff does not touch CodegenPolicy, SchedulerStrategy, or compile_kernel options.
9 N/A The diff does not touch waitcnt encoding, named/split barriers, or optimization_barrier lowering.
10 PASS In dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1869 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_attention_tiled_2d_rtl_internal.h:68, new pointer params use standard global PtrType with…
11 PASS In dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1237 and 1299, hardware LDS capacity and architecture validation query attention_lds_capacity_bytes(arch) and validate_tiled_attention_arch(arch) rather than string prefix checks.
12 N/A The diff does not touch core target catalogs or lowerer vector-load / async-LDS width caps.
13 PASS Growing rocke_attention_tiled_2d_spec_t in dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_helper_rocke.instances.gfx942.attention_tiled_2d.h:192 is accompanied by bumping ROCKE_ABI_VERSION to 11 in…
14 PASS New header dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_attention_tiled_2d_rtl_internal.h:9 uses #ifndef include guards (no #pragma once), rocke_* snake_case naming, and bindings changes in…
15 N/A The diff does not modify engine fallback logic, LLVM flavor resolution, or debug lowerers.
16 PASS In dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d.cpp:162 and 242, C++ spec validation faithfully mirrors Python rejections without widening known_gaps.md debt.
§P graded 14/14
# Verdict Evidence
1 PASS dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:525 adds runtime_layout as a defaulted spec field, and dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp:824 reads it from the spec dict rather…
2 PASS dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_attention_tiled_2d_rtl_internal.h:1 defines the runtime-layout kernarg order shared by the C++ twins, and dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h:101 bumps…
3 N/A This PR changes attention runtime-layout kernels and an ingestor engine, not GEMM, implicit-GEMM convolution, gfx950 block-scaled…
4 FAIL dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1105 adds runtime-layout LDS sizing/geometry logic, but the C++ runtime-layout validator added in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d.cpp:130 mirrors feature/head/tile checks without a corresponding LDS-capacity gate.
5 N/A The diff does not change WMMA/RDNA GEMM, GemmMultiD, or GemmMultiAbd narrow-path admission rules.
6 N/A The diff does not add or modify Stream-K, MoE sorting, fused MoE, or split-K caller-zeroed workspaces.
7 N/A The diff does not change Stream-K, persistent Stream-K, GroupedGemmLauncher, or the GEMM support matrix limits.
8 N/A The diff does not change quantization helpers, MX GEMM, quantized clamps, or e8m0 decode logic.
9 N/A The diff does not add numeric emulation or decode-path arithmetic that needs overflow, Inf, NaN, denormal, or rounding-edge tests.
10 N/A The diff adds ingestor descriptors and packs, not a new rocKE runtime manifest kind or manifest schema consumer.
11 NEEDS-INFO The hunk containing the new AttentionPagedUnifiedNative.cpp kernel launch/execute path is truncated, so buffer retention, stream synchronization, fencing, and timing behavior cannot be verified.
12 N/A The diff does not change COMGR or HIP library discovery/loading rules or claim to fix a HIP launch error.
13 NEEDS-INFO The PR adds host launch sizing and dispatch/cache-related logic, but the relevant test hunks and any trusted diff-side list of numeric suites run for the touched attention operator are omitted or unavailable.
14 PASS New runtime-layout kernel logic is placed in the attention kernel/instance paths, including dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:525 and…
§Q graded 7/7
# Verdict Evidence
1 NEEDS-INFO The diffs for dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt, dnn-providers/hip-kernel-provider/ROCKE_ENGINE_test_categories_external.yaml, and the two new test files (test_gfx942_attention_runtime_layout.py and test_attention_tiled_2d_runtime_layout_gfx950.py) are omitted from the truncated diff, so registration and installed-layout path usage cannot be verified.
2 PASS The new runtime_layout kernel variant adds 19 parity configs per architecture to the C/Python parity emitters and binding config mirror: dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.c:350…
3 NEEDS-INFO The PR does not weaken any gate criteria or claim numerics from IR parity; it extends the parity emitters (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.c:350, dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.c:1233) for cross-engine agreement alongside device numerical tests. — no diff line supports this verdict
4 N/A This PR does not modify or re-bless any platform LLVM IR goldens (runtime_layout=False leaves existing emitted IR unchanged) and…
5 NEEDS-INFO The contents of the new test files (dnn-providers/hip-kernel-provider/rocke/library/tests/test_gfx942_attention_runtime_layout.py, dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_tiled_2d_runtime_layout_gfx950.py, and projects/hipdnn/tools/IngestorGenerator/tests/test_gfx950_attention_dense_schema_gate.py) are omitted from the diff, so device/HIP/clang skip guards and target portability…
6 N/A This PR does not refactor any test harness or verification gate.
7 NEEDS-INFO The PR adds new spec rejection rules in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:579 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1061 (and C++ twins in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d.cpp:162), but the diffs for test_gfx942_attention_runtime_layout.py and…
§R graded 11/11
# Verdict Evidence
1 N/A runtime_layout is a new semantic mode for hipDNN's paged/ragged SDPA contract (classified in _semantic_fields at…
2 NEEDS-INFO The runtime_layout additions are separated per architecture in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1226-1246 (LDS capacity budget and sub-block tile_size) and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1138-1149 (whole-block tiles, 4/2 warps), and incompatible pipeline knobs are explicitly rejected in…
3 NEEDS-INFO LDS sizing accounts for non-aliasable live allocations across the grid-stride loop and bank-conflict padding in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1209-1223, and per-page 64-bit buffer resources follow the existing i64 page-resource pattern (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:3164-3207…
4 NEEDS-INFO The PR body summarizes SGPR-to-VGPR lane spills and zero scratch spills from dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md (diff omitted), and changes global load/store and barrier primitives in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:3192-3219 and…
5 PASS Global Q/O accesses keep global_load_vN / global_store_vN with width 8 (align=16) and unit dim stride (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1993-2002, 2610, 6358…
6 N/A The diff does not change any LDS or MMA layout; runtime_layout reuses the default 16x16x16 (gfx942) and 16x16x32 (gfx950)…
7 N/A The PR makes no bottleneck classification claims (memory-bound, LDS-bound, barrier-bound, or latency-bound).
8 NEEDS-INFO No waits or barriers are removed, and the single added b.sync() at the end of each grid-stride loop iteration (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6376-6382, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:5419-5424) documents the exact LDS reuse hazard between consecutive q-blocks (Q staging aliasing K_lds and Acc_lds…
9 N/A This PR is not a performance regression fix and does not change any dispatch heuristic or knob affecting existing shapes…
10 N/A The PR makes no benchmark timing claims, occupancy-from-wall-clock conclusions, hand-written assembly parity claims, or DSL…
11 N/A The PR introduces no new optimization failure mode, heuristic, case study, or probe/analysis tool.
§S graded 8/8
# Verdict Evidence
1 NEEDS-INFO The PR adds public symbols (UnifiedAttention2DTiledSpec.runtime_layout, runtime_layout_lds_bytes, runtime_layout_geometry, runtime_layout_num_warps, runtime_layout_catalog_spec, rocke_gfx942_attention_tiled_2d_runtime_layout_reject_reason, rocke_gfx950_attention_tiled_2d_runtime_layout_reject_reason), but the diffs for…
2 NEEDS-INFO Existing APIs (UnifiedAttention2DTiledSpec, build_unified_attention_2d_tiled, build_gfx942_4warp_gqa, supports_tiled_2d_spec, ROCKE_ABI_VERSION / BINARY_ABI) are updated in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:526 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:543, but the diffs for…
3 N/A The diff does not remove or rename any public rocke platform or library API.
4 N/A The diff does not modify __all__ in any module, and new internal helpers in attention_tiled_2d.py are underscore-prefixed…
5 NEEDS-INFO New fail-fast ValueError and NotImplementedError raises are added in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:535, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1261, dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6739…
6 NEEDS-INFO Parity emitter configs are added in dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.py:372 and dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.py:1194, but the diffs for test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, and…
7 NEEDS-INFO The C++ twin and binding are updated in dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp:827, dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/attention_tiled_parity_configs.py:333, dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.c:353, and…
8 FAIL The PR touches dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h:109 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/portable_ir/src/abi.py:84, but the PR description does not state that PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py were run from rocke/platform/.

@cgmillette

Copy link
Copy Markdown
Contributor

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

PR #14247 — paged/ragged SDPA engine for gfx942/gfx950 (dpeabody-amd, AICK-1847) — 🟠 Follow-up: conflicts; golden, ragged mapping open

Prior item Status
AnaghaRaoAMD (11-04), yraparti (11-07 §A.10): rtl rows missing from the golden 🟠 OPEN, no reply yet
• the PR still adds no golden entry; merge reads GOLDEN OK, rtl configs unblessed (ran)
• byte identity alone lets both engines move together: record both families
AnaghaRaoAMD (11-04) FLOAT sinks declined; sink objects unreachable 🟠 OPEN, no reply yet
• engine still declines FLOAT sinks (AttentionPagedUnifiedNative.cpp:85); develop admits FLOAT only (read)
• body: waits on a hipDNN decision
AnaghaRaoAMD (11-04) ×2: CMake _others; fkv_build 🟡 OPEN, no reply yet
• _others still from this build's archs (src/CMakeLists.txt:609) (read)
• fkv_build still ignores runtime_layout (rocke_engine_attention.cpp:1004) (read)
Copilot (11-06) ×2: ragged origin past 0; offsets out of order 🟠 OPEN, no reply yet
• first origin past block 0: the search returns −1, then reads cu_q[-1] and seq_lens[-1] (read)
• other RFC 0014 orders: the engine README itself says the output is wrong (read)
• fix: map query blocks per sequence, as the README proposes
Copilot (11-06) num_warps=0 reaches % 0 🟡 OPEN, no reply yet
• the rtl reject helper divides before any warp check (gfx942/attention_tiled_2d.cpp:319) (read)
• fix: check the warp count first, in both engines
yraparti (11-07) §P.4: rtl LDS gate 🟡 OPEN, no reply yet
• both engines' builders build rtl specs past the LDS budget; clang refuses one (ran)
• the gate lives only in the support check (gfx942/attention_tiled_2d.py:1322), so the engines agree
• fix: move it into both engines' rtl reject helpers
yraparti (11-07) §E.7 reference docs; §S.8 README runs 🟡 OPEN, no reply yet
• no rocKE reference doc names runtime_layout (read)
• body lists ctest runs, not test_rocke.py or run_all.py (read)
CongMa13 (11-06) replies owed on 11 items 🟡 OPEN
• no author reply since; 13 of 16 threads open (read)
Our 11-06 N5 (+ Copilot 11-06): non-positive attn_scale plan never loads 🟠 OPEN, no reply yet
• matcher serves any finite scale (AttentionPagedUnifiedNative.cpp:711); reader refuses <= 0 (:1323) (read)
• fix as posted: one served-scale predicate both call, like isSupportedAttentionScale
Our 11-06 N6: second homes for facts 🟠 OPEN, and wider
• sdpa_graph::declaresAlignment is develop's FlydslSdpaShared.hpp:92, body for body (ran)
• integerMetadata/stringMetadata (:786, :797) restate develop's UnifiedAttentionGraphMatch.hpp:157,175 (ran); spanProduct, NO_UID, kernelSignature() unchanged
• fix: call develop's helpers; move declaresAlignment and the zero-admitting product down into one shared header
Our 11-06 overlap notice: 14967 owns the K/V bound and offset 🟠 OPEN, no reply yet
• 14967 still open; it lands first, then this PR drops its copies (read)
Earlier items: BrianHarrisonAMD ×3; AnaghaRaoAMD Asks 1–2, dense search loads; our N1–N4 🟢 FIXED a259510d, f1d0d688, c75c1ec9, d9c6f0df, 9f3ebe52
• all hold after the history rewrite into 8 commits (read)
• N1's probe passes at head; the new conflicts are N7 (ran)
Bots / CI 🟢 Passed at head
• 54 ✓, 52 skipped; 1 duplicate PR Bot run cancelled (read)
• runs predate 14822, 14908 and 11642 on develop; dirty, so no new runs
Our checks 🟠 Dirty
• merge vs 477ab5b9: 13 files conflict (ran): N7
• ABI-resolved merge: ctest 59 ✓ / 1 skip; byte identity GREEN ×4 (gfx942 39, gfx950 91 configs); goldens pass (ran)
• tests 334 ✓ / 656 skip (GPU) (ran); ABI: develop is 12, takes develop + 1 when it lands; hipDNN C++ not built here
Code quality 🟠 Should-fix
• N6 wider: develop's exported helpers copied again (ran)
• paged checks have one home; save/restore reuses develop's restore helpers (ran)
• live-rows fix in both engines and tested; its rule sits in 4 copies: N8 (ran)
# Sev file:line Finding
N7 🟠 rocke/platform/cpp/include/rocke/abi.h:109 (+ 12 files) 13 files conflict with develop
• impact: no merge and no new CI; develop moved under 11642, 14908, 14822, 15102 and 15293
• fix: keep both sides in the pack registry, CMake lists and launch-contract samples; take develop's support-matrix overview (15293 stops per-PR regeneration); ABI: develop is 12, takes develop + 1 when it lands (that side built and gated here); update the body's ABI text and drop its "Draft." line
N8 🟡 instances/gfx942/attention_tiled_2d_ctx_prologue.cpp:165 (+ gfx950 attention_tiled_2d_public_entry_glue.cpp:424) The ragged-row rule is copied per arch in each engine; this round's fix edited all four copies
• impact: the C++ owned/live rows and q-block total repeat per arch beside the shared instance_attention_tiled_2d_rtl_internal.h; Python repeats them in both kernels (ran)
• fix: one static inline helper in that header, and one in rocke.helpers.attention beside binary_search_seq_idx

Utility: 🟢 Positive — the delta keeps both engines byte-identical; conflicts and ragged mapping open


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

Follow-up review: first pass by qwen3.6 (local), verified by the orchestrator (Claude Opus 5.5). Checks: changes since 33842d27 (history rewritten into 8 commits), a merge with develop (13 conflicts; ABI side resolved: build + ctest, byte identity, goldens, tests), CI at head, an LDS probe in both engines, code quality. Flags whether a closer look is needed; not an approval or rejection.

@dpeabody-amd
dpeabody-amd force-pushed the users/dpeabody/hipdnn-cdna-paged-attn branch from f1d0d68 to fe41d27 Compare November 7, 2026 07:00
@dpeabody-amd

Copy link
Copy Markdown
Contributor Author

Pushed fe41d27.

The branch is rebased on develop 477ab5b. It conflicted with #11642's ABI bump and with #14908 (gfx942 dense attention engine) in SdpaGraphMatch.hpp, the test CMake lists, the launch-value tests and goldens, and file-formats.md. ROCKE_ABI_VERSION and BINARY_ABI are now 13 (develop is 12 since #11642); #14254 also bumps it, so whichever of the two lands second takes 14. #14908 also moved the attention engines' lanes into one CMake table (HKP_PINNED_ATTENTION_ENGINES); this engine is now two rows of it (gfx942 and gfx950), so it gets the same external lane, GPU-reference target and census as the other attention engines, and test_attention_bundle_filters.py checks its filter against its categories YAML. Following #15293, the branch no longer regenerates SUPPORT_MATRIX_OVERVIEW.md.

Commits on the new head:

  • c93a883 feat(rocke): runtime_layout mode for the gfx942/gfx950 tiled 2D unified attention
  • a517dfc feat(hip-kernel-provider): hipkernel:AttentionPagedUnified paged/ragged SDPA for gfx942/gfx950
  • c3dc31c feat(rocke): attention sinks in the gfx942/gfx950 tiled 2D runtime_layout mode
  • 2e50b48 feat(hip-kernel-provider): AttentionPagedUnified sink objects
  • b8eb84e feat(hip-kernel-provider): attention sinks in hipkernel:AttentionPagedUnified
  • 2361623 test(hip-kernel-provider): pin the schema gate in front of hipkernel:Gfx950AttentionDense
  • 5c7a85b refactor(hip-kernel-provider): share the cuDNN paged-graph checks in SdpaGraphMatch.hpp
  • 96ae6d8 feat(hip-kernel-provider): save and restore hipkernel:AttentionPagedUnified plans
  • fe41d27 feat(hip-kernel-provider): serve FLOAT attention sinks in hipkernel:AttentionPagedUnified (new)

Fixes are folded into the commit that introduced the code; FLOAT sinks are a new commit because they change what the engine serves.

Copilot review summary

Covered by the four threads above: the shifted first origin and the saved scale are fixed, num_warps = 0 is refused before the modulo, and the origin-order case is documented with the measurements that ruled out the per-batch grid.

@cgmillette, follow-up review

Thanks. Against the table:

  • Golden rows (Anagha): recorded per commit; see thread 4473823510.
  • FLOAT sinks (Anagha): the engine now serves them, and the sink integration cases run on develop; see thread 4473823528.
  • gfx950 dense search loads: fixed last round; replied on the thread.
  • CMake _others, fkv_build: fixed; see the threads.
  • Conflicts: rebased on develop 477ab5b; git merge-tree --write-tree against develop 1702865 (fetched after the rebase) is clean.
  • N5 (non-positive scale saves but never loads): fixed with one predicate shared by graph_match and the reader, and a save-then-load test at zero and negative scales (96ae6d8).
  • N6 (second homes):
    • The 21 base arguments live once, in packs/AttentionTiled2DSignature.hpp (tiled2DKernelSignature()); the Unified2D pack uses it and this engine appends its runtime_layout tail (a517dfc).
    • spanProduct is gone: sdpa_graph::checkedProduct in SdpaGraphMatch.hpp is the one zero-admitting product, and the gfx942 unified 3D pack's local copy uses it too; FlydslSdpaShared.hpp now takes it and declaresAlignment from the header instead of its identical copies (5c7a85b).
    • The file-local NO_UID is replaced by ATTENTION_PAGED_UNIFIED_NO_UID (96ae6d8).
    • expectReadRefusal and expectCommonReadRefusals moved to IngestorPlanTestUtilities.hpp, and both launch-value test files use them (96ae6d8).
    • pageTablesFor calls pagesKAndV itself, so it has no precondition (5c7a85b).

@cgmillette, overlap notice with #14967

Agreed: #14967 owns the block-table bound (block_table_max_idx) and the out-of-range buffer offset (buffer_offset_in_sequence), since it covers every layout and keeps one named constant per engine in shared homes. Whichever lands first, the other follows the steps in your notice:

@CongMa13

Thanks. Every thread you listed now has a reply: eight are fixed and can be resolved (golden rows, FLOAT sinks, the gfx950 dense search, CMake _others, fkv_build, the shifted first origin, the saved scale, num_warps = 0), and the origin-order thread is answered with the measurements and the documented limitation, left open for review. Both of cgmillette's comments have replies above.

…ed attention


Adds UnifiedAttention2DTiledSpec.runtime_layout on top of the runtime shape
kernargs the tiled 2D bodies already take (num_query_heads, num_kv_heads,
sliding_window, and the cu_q search trip count derived from num_seqs). With
it, the Q/O/K/V strides and a dense-or-ragged Q are kernel arguments too;
the page tables and seq lens are read on device; a grid-stride loop covers
every q block; page-table and K/V reads stop at seq_len; rows with no
visible key write zeros, and so do the query rows a sequence owns past
seq_lens_q (padding), as the hipDNN contract requires. The mode appends 15
kernel arguments after the runtime shape kernargs. Its window mask is always
on, with the sliding_window kernarg as its size, and its ragged q-block
search uses the shared runtime trip count (no iteration for a dense Q).

K pages through block_tables_ptr and V through the appended
block_tables_v_ptr, which shares block_table_stride and may be the same
pointer, so cuDNN-style graphs with separate K and V page tables are served.

K/V use one 64-bit buffer resource per page, built once per tile load from
one page-table read: a runtime_layout tile never spans more than one page.
On gfx942 this is the per-page mechanism of the i64 path (#14349), which the
runtime_layout loaders sit beside. runtime_layout=False emits unchanged IR.

Additive bias (gfx950), the optional LSE output (emit_lse) and the gfx950
longest-first query-block order (use_qblock_reverse) are not combinable with
runtime_layout and are rejected by the runtime-layout checks; the gfx942
4-warp GQA builder refuses a runtime_layout spec. On gfx942 the support
gate and both builders (Python and C++) charge runtime_layout_lds_bytes
against ArchTarget lds_capacity_bytes and refuse a spec over it.
The runtime-layout checks take their head-count rule and text from the
shared check_head_counts (rocke_attention_check_head_counts in C++).

The rocKE C++ twin mirrors the mode op for op. rocke_attention_tiled_2d_spec_t
gains runtime_layout, so ROCKE_ABI_VERSION and BINARY_ABI move to 13.

Dispatch never sets runtime_layout: it is a semantic field (no sweep varies
it), and while it is off it stays out of the tuning spec's cache key and
config defaults, so every spec_hash, config_key and stored pin is unchanged.

The CPU checks of the two runtime_layout test files run in the new
rocke_library_attention_runtime_layout_pytest CTest entry (-m "not gpu");
their GPU numerics run in rocke_library_gpu_pytest. The scenarios include
sequences with 0 < seq_len < q_len, whose leading rows see no key.
…ed SDPA for gfx942/gfx950


Native ingestor engine over rocKE tiled 2D unified attention (runtime_layout):
paged K/V with arbitrary strides (incl. vLLM's interleaved K|V view), INT32
page tables (one tensor for K and V, or separate page_table_k and
page_table_v tensors with the same width and batch stride), seq_len_q and
seq_len_kv, dense or ragged Q, bottom-right causal with optional left bound
(cuDNN meaning: L keys up to the diagonal), GQA, fp16/bf16, head size
64/128/256, block size 16/32/64. Query rows past seq_len_q come back zero.
max_seq_len_kv may be unset: the kernel reads only seq_len_kv, which the
inferred extent bounds the same way.

Declines everything else with a logged cause, including valid paged graphs
the kernel cannot read: one-sided paging, K and V page tables of different
widths or batch strides, differing K and V block sizes, packed page
tables, a ragged Q/O whose ragged_offset_multiplier is not its token stride
(the kernel reads stored offsets as token counts), and a declared tensor
alignment below what the kernels load (cause alignment: 16 bytes for
Q/K/V/O, 4 for the INT32 page tables, sequence lengths and ragged offsets).
UED sdk_version 1.4.0, so these graphs reach the matcher rather than stopping
at the schema gate.

Adds matcher unit tests, the descriptor census, and GPU integration tests
against a gather-then-dense reference that is itself checked against
hipDNN's CPU SDPA reference (CpuFpReferenceSdpa over gatherPagedKv), including
padded query rows, rows with no visible key (0 < seq_len_kv < seq_len_q), an
unset max_seq_len_kv, separate K and V page tables and HIP-graph
capture/replay.

The three quick paged bundles the engine serves claim it on gfx942 and gfx950
(linux), and one engine-pinned external integration lane per shipped
architecture runs the SdpaFwd bundles against it with those claims enforced.
The engine takes a gfx942 and a gfx950 row of HKP_PINNED_ATTENTION_ENGINES,
which registers those lanes, a GPU-reference target per architecture and the
Packs census (the census table now collects a multi-arch engine's rows into
one registration); each row's TOML skips the other architecture, so a package
carrying both shards runs the bundles once.
@cgmillette

Copy link
Copy Markdown
Contributor

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

PR #14247 — paged/ragged SDPA engine for gfx942/gfx950 (dpeabody-amd, AICK-1847) — 🟠 Follow-up: conflicts fixed; ragged order open

Prior item Status
AnaghaRaoAMD (11-04), yraparti (11-07 §A.10): rtl rows missing from the golden 🟢 FIXED c93a883a, c3dc31c1, fe41d279
• 19 + 19 rtl rows recorded at all four flavors and canonical (ran)
• differential gate: no unblessed change (ran); yraparti's thread has no reply yet
AnaghaRaoAMD (11-04) FLOAT sinks declined; sink objects unreachable 🟢 FIXED fe41d279
• sinkFor admits FLOAT or the Q dtype; sink_is_f32 saved with plans (read)
• 9 of 19 rtl parity rows per arch use sinks: byte-identical, golden-recorded (ran)
• GPU integration cases not run here
AnaghaRaoAMD (11-04) ×2: CMake _others; fkv_build 🟢 FIXED a517dfc1, c93a883a
• committed per-arch TOMLs skip the other arch by name (read)
• fkv_build reads runtime_layout (rocke_engine_attention.cpp:1034) (read)
Copilot (11-06) ×2: ragged origin past 0; num_warps=0 reaches % 0 🟢 FIXED c93a883a
• origin: seq_ok gates the rows in all four builders (read)
• warps: refused before the % (gfx942/attention_tiled_2d.cpp:324); host tests pass (ran)
Copilot (11-06) ragged offsets out of order 🟠 DEFERRED
• author: a per-sequence grid is correct but much slower on mixed batches
• fair for vLLM's layout; other valid RFC 0014 orders still match and silently give wrong output
• fix: a hipDNN layout contract the matcher requires, or a device order check that poisons the output
yraparti (11-07) §P.4: rtl LDS gate 🟡 OPEN, no reply yet
• both builders still build rtl specs past the LDS estimate (ran)
• fix: the gate in both engines' rtl reject helpers
yraparti (11-07) §E.7 reference docs; §S.8 README runs 🟡 OPEN, no reply yet
• no rocKE reference doc names runtime_layout (ran)
• body still lists no test_rocke.py or run_all.py run (read)
CongMa13 (11-06) replies owed on 11 items 🟢 FIXED
• author replied on every listed thread (read)
Our 11-06 N5 (+ Copilot 11-06): non-positive attn_scale plan never loads 🟢 FIXED 96ae6d86
• one predicate, isServedAttentionScale, for matcher and reader (read)
• save-load test over zero and negative scales (read; hipDNN C++ not built here)
Our 11-06 N6: second homes for facts 🟡 OPEN, narrowed
• fixed: declaresAlignment, checkedProduct, NO_UID; tiled2DKernelSignature() keeps develop's 21 args in order (ran)
• left: integerMetadata/stringMetadata (AttentionPagedUnifiedNative.cpp:785, :796) restate unified_attention:: (read)
• fix: using unified_attention::integerMetadata, as the Unified2D pack does
Our 11-06 overlap notice: 14967 owns the K/V bound and offset 🟢 AGREED
• author accepts 14967 as owner, with steps for either land order (read)
• 14967 still open; the second to land follows them
Our 11-07 N7: 13 files conflict with develop 🟢 FIXED (rebase onto 477ab5b9)
• merge vs 1702865a: 0 conflicts (ran)
• body's opening "Draft." line remains: N9
Our 11-07 N8: ragged-row rule copied per arch 🟡 OPEN, and grown; no reply yet
• seq_ok/owned/live repeat in gfx942 ctx_prologue.cpp:152 and gfx950 public_entry_glue.cpp:416, and per Python kernel (read)
• new: Python _sink_logit copied per arch (gfx942/attention_tiled_2d.py:2791, gfx950 :2281); C++ shares one (ran)
• fix: one helper each in rocke.helpers.attention and instance_attention_tiled_2d_rtl_internal.h
Earlier items: BrianHarrisonAMD ×3; AnaghaRaoAMD Asks 1–2, dense search loads; our N1–N4 🟢 FIXED
• all hold after the rebase (read)
Bots / CI 🟡 Running
• 38 ✓, 44 skipped, 0 failed (read)
• running: gfx94X hipkernelprovider shards 1–2, miopenprovider; Windows gfx110X math-libs
• merge: clean; blocked on review
Our checks 🟢 Clean
• merge vs 1702865a: clean (ran)
• ctest 59 ✓ / 1 skip; byte identity GREEN ×4 (gfx942 39, gfx950 91 configs); goldens pass (ran)
• installed tests 340 ✓ / 803 skip (GPU); compile: no new spills (ran); ABI: develop is 12, takes develop + 1 when it lands
Code quality 🟡 Minor
• N6 narrowed to the metadata pair; rest fixed (ran)
• FLOAT sink logit shared in C++, copied in Python: N8 (ran)
• save/restore, paged checks, test helpers: one home each (ran)
# Sev file:line Finding
N9 🟡 PR body, line 3 The body still opens "Draft." on an older base
• impact: the squash commits it: "Draft.", base 886472ee; ABI text names only 14254 as the other bump
• fix: drop the line; state base 477ab5b9; ABI: develop is 12, takes develop + 1 when it lands

Utility: 🟢 Positive — rebase clears conflicts; golden, sinks and saved plans fixed; ragged order deferred


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

Follow-up review: first pass by qwen3.6 (local), verified by the orchestrator (Claude Opus 5.5). Checks: changes since f1d0d688 (rebase onto develop, one new commit), a merge with develop (clean: build + ctest, byte identity, goldens, installed tests, compile), CI at head, an LDS probe in both engines, code quality. Flags whether a closer look is needed; not an approval or rejection.

…yout mode


runtime_layout now composes with use_sinks. The sink init bounds the query
head with the runtime num_query_heads kernarg, as every tiled 2D body does, so
heads of KV head > 0 keep their sink; the head index is the same
kv_head * qpk + row % qpk mapping the Q/O accesses use.
runtime_layout_catalog_spec takes use_sinks (default False). Non-RTL emission
is unchanged. Adds spec-acceptance tests and GPU numerics with sinks and two
KV heads against an fp32 reference.

The representative-IR golden (rocke_representative_ir_sha256/attention.json)
pins three runtime_layout objects per architecture at llvm20/22/23/24: one
without sinks and two with them (D128 and D256, the latter the gfx942
2-warp body).
Every shipped (arch, dtype, head size, block size, GQA ratio) object now also
ships with use_sinks (name suffix _sinks), distinguished by a new KMD field
use_sinks (int, default 0). 144 -> 288 objects per arch (288 -> 576 total).
Descriptors regenerated with IngestorGenerator; retained documents and kernel
entries keep their ids. The pack census lists the new objects.
…dUnified


graph_match accepts sink_token [1, Hq, 1, 1] in the Q data type with a unit
head stride, declines FLOAT sinks, other shapes and strided heads with cause
sinks, and declines a sink declaring less than a 16-byte alignment with cause
alignment; kernel_match selects the use_sinks objects exactly when the graph
has a sink; execute passes the sink buffer as sink_ptr. The Q-dtype sink is a hipDNN
extension of cuDNN's FLOAT sink_token: vLLM keeps sinks in the model dtype.

Matcher unit tests for the accept/select/decline cases, GPU integration tests
with sinks (dense GQA decode, ragged Q with a window) against the
gather-then-dense reference and hipDNN's CPU SDPA reference (FLOAT sink of the
same values), README. A frontend that accepts only FLOAT sinks rejects these
graphs, so the sink integration tests skip with that reason there.
…Gfx950AttentionDense


Gfx950AttentionDense's graph_match checks neither a ragged_offset on Q/K/V/O nor
a declared tensor alignment. Since engine plugin API 1.4.0, the only thing
keeping such graphs away from it is GenericPlanBuilder::understandsGraph(),
which declines a graph whose min_required_engine_api_version (1.3.0 for both
features) is above the descriptor's sdk_version (1.0.0).

A census suite over the gfx950 product shard builds the engine from the loaded
descriptor set, as the provider does, stamps a ragged graph and an 8-byte
aligned graph the way hipDNN's GraphDescriptor does, checks the plugin's API
version admits them, and requires the schema-gate decline. The control offers
the baseline graph and requires the engine to serve it. The suite fails if the
gate goes or the descriptor's sdk_version rises to 1.3.0 before the matcher
declines both features itself. It registers as a companion suite
of the dense engine's HKP_PINNED_ATTENTION_ENGINES row.

That census only registers where a build ships the engine for gfx950, so an
IngestorGenerator test also pins the sdk_version of the authored config, its
engine manifest and the committed UED below 1.3.0 in every build that runs the
IngestorGenerator suite.
…SdpaGraphMatch.hpp


hipkernel:AttentionPagedUnified's graph_match carried its own copies of the
SDPA match helpers that packs/SdpaGraphMatch.hpp now shares (findTensor,
sdpaNode, hasSdpaForwardNode, describeOperand, the decline head and the axis
constants). It now uses the header's, with sdpa_graph::DeclinedBy. The
overflow-checked product it shares with hipkernel:Gfx942AttentionUnified3d
moves into the header as sdpa_graph::checkedProduct: unlike
pack_features::checkedProduct it takes zero factors, which every extent of 1
gives. FlydslSdpaShared.hpp, which
already takes computesInFloat from the header, now takes checkedProduct and
declaresAlignment from it too instead of keeping identical copies.

The stricter operand checks it needs move into the header too
(usesOverrideShapes, hasNonNegativeStrides, declaresAlignment, isInt32Vector),
together with the cuDNN paged-K/V graph contract checks (SEM-009), so every
pack serving a paged graph applies one set of rules with one cause vocabulary,
logged with the calling engine's name:
- pagesKAndV: one_sided_paging;
- pagedContainersAgree: block_size, shape (one block size, one block count;
  head counts and head sizes stay the engine's to check);
- pageTablesFor: page_table (INT32 [B, 1, M, 1], unit entry stride, not
  packed; separate K and V tables share M and the batch stride);
- sequenceLengthsFor: seq_len_kv, seq_len_q (both required, INT32 [B,1,1,1]);
- raggedQueryFor, countsOffsetsInTokens: ragged (Q and O through one
  [B+1,1,1,1] offset table counted in tokens; K/V never ragged);
- declaresServedAlignment: alignment (vector operands at the engine's vector
  width, INT32 operands at 4 bytes);
- sinkFor: sinks, and alignment below the engine's own minimum sink
  alignment (16 bytes for this engine's kernels);
- takesNoAttentionMask: attn_mask.

The reasons they log are the texts hipkernel:Gfx1151WmmaAttention's paged
matcher (#14254) logs for the same rules, so that engine can call these checks
without changing what it reports. The decline lines go through
HIPDNN_PLUGIN_LOG_INFO((DeclinedBy{...}) << ...): the braces' comma would
otherwise split the macro argument.

The engine keeps its kernel limits (head and block sizes, GQA ratios, strides
and address ranges) and the order of its declines; it checks itself that V has
K's head count and head size. A malformed page-table tensor now logs page_table
rather than operand. The header's file comment covers paged graphs.
…nified plans


The engine is one kpack launch with no slots and no workspace; lengths, page
tables, offsets and sinks are device buffers bound at execute. The prepared
dispatch is now an IngestorPreparedDispatch, and the handler saves the bound
uids, every by-value kernel argument and the CTA size as typed launch values
under hipkernel.attention_paged_unified.dispatch.v1 (registered next to the
UDD's unversioned name, same handler). restoreLaunch rebuilds them through
restoreIngestorLaunch; prepare() and restore share one function that applies
the launch geometry. The reader refuses values outside what the kernel takes.

Tests: CPU launch-value round trips for dense, ragged, separate K/V tables and
sinks plus refusals; the alias table; a checked-in byte sample of the v1
contract in TestIngestorPlanLaunchContractGoldens (a ragged Q with a sink and
separate K and V tables, every value distinct), read back field by field; GPU
save/load round trips in a fresh handle, both save forms, bit-identical to the
original plan's output.
…ttentionPagedUnified


hipDNN's frontend admits only cuDNN's FLOAT sink_token, so the Q-dtype sinks
the engine read could not reach it on develop. The rocKE tiled 2D
runtime_layout kernels gain a sink_is_f32 launch argument (the 16th of the
runtime_layout tail): nonzero means sink_ptr holds Hq fp32 logits, zero means
Hq logits in the Q data type. Each head's sink is read with masked loads of
both forms, so neither reads past the other's buffer, and the same use_sinks
objects serve both; there are no new objects. Python and the C++ twin build
the same IR.

The engine accepts a FLOAT sink_token as well as one in the Q data type
(sdpa_graph::sinkFor), binds sink_is_f32 from it, passes it at launch and
saves it as the launch value sink_is_f32. The FLOAT integration cases run
on develop's frontend; the Q-dtype ones still skip on its rejection. Two
FLOAT-sink quick bundles of #14358 are now claimed.

The representative-IR golden's runtime_layout entries are re-recorded for
the new argument, and rocKE's SUPPORT_MATRIX.md and dsl_docs/instances/attention.md
describe the runtime_layout mode.
@cgmillette

Copy link
Copy Markdown
Contributor

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

PR #14247 — paged/ragged SDPA engine for gfx942/gfx950 (dpeabody-amd, AICK-1847) — 🟠 Follow-up: gfx942 LDS fixed; ragged open

Prior item Status
Copilot (11-06) ragged offsets out of order 🟠 DEFERRED
• unchanged since 11-07; the one open thread
• other valid RFC 0014 orders still match and silently give wrong output
• fix: a hipDNN layout contract the matcher requires, or a device order check that poisons the output
yraparti (11-07) §P.4: rtl LDS gate 🟡 PARTLY FIXED 95848d08
• gfx942: gate and both builders refuse over-budget rtl specs (ran)
• gfx950: no LDS check in the gate or either builder: N10 (ran)
yraparti (11-07) §A.10 golden; §E.7 docs; §S.8 README runs 🟢 FIXED 6114fe07, 3bed02a4
• 6 rtl representative cases at four flavors; golden gate passes (ran)
• attention.md Runtime layout section; SUPPORT_MATRIX row (read)
• test_rocke.py, run_all.py runs in the body, head and develop (read)
Our 11-06 N6: second homes for facts 🟡 OPEN, unchanged
• integerMetadata/stringMetadata (AttentionPagedUnifiedNative.cpp:785, :796) restate unified_attention:: (ran)
• fix: using unified_attention::integerMetadata, as the Unified2D pack does
Our 11-07 N8: ragged-row rule copied per arch 🟡 OPEN, unchanged
• seq_ok/owned/live per arch: gfx942 ctx_prologue.cpp:152, gfx950 public_entry_glue.cpp:416 (ran)
• Python _sink_logit per arch: gfx942 :2806, gfx950 :2281 (ran)
• fix: one helper each in rocke.helpers.attention and instance_attention_tiled_2d_rtl_internal.h
Our 11-07 N9: body opens "Draft." on an older base 🟡 OPEN, narrowed
• base now 1702865a (fixed); still opens "Draft." (read)
• ABI text names only 14254 and a fixed number; 14707 also bumps 12 → 13 (ran)
• fix: drop the line; ABI: develop is 12, takes develop + 1 when it lands
Our 11-06 overlap notice: 14967 owns the K/V bound and offset 🟢 AGREED
• author's steps for either land order stand (read)
• 14967 still open; the second to land follows them
Earlier items: BrianHarrisonAMD ×3; AnaghaRaoAMD ×7; Copilot ×3; CongMa13; our N1–N5, N7 🟢 FIXED
• all hold: the delta since fe41d279 touches only the LDS gate, goldens and docs (ran)
• merge vs 1702865a: still 0 conflicts (ran)
Bots / CI 🟡 Running
• 21 ✓, 26 skipped, 0 failed (read)
• running: math-libs gfx94X, gfx950, gfx125X, Windows gfx110X; hipDNN Linux superbuild; no gfx94X rocKE shard yet at this head
• merge: clean; blocked on review
Our checks 🟢 Clean
• merge vs 1702865a: clean, 0 behind (ran)
• ctest 59 ✓ / 1 skip; byte identity GREEN ×4 (gfx942 39, gfx950 91 configs); goldens pass (ran)
• installed tests 342 ✓ / 805 skip (GPU); compile: no new spills; LDS probe, both arches and engines (ran)
• ABI: develop is 12, takes develop + 1 when it lands
Code quality 🟡 Minor
• gfx942 LDS: one estimate per engine, used by gate and both builders; tests in both engines (native case skips in CI) (ran)
• Python builder reads capacity past attention_lds_capacity_bytes and repeats the gate's text: one reject helper (read)
• gfx950 LDS rule lives only in dispatch: N10; N6, N8 unchanged
# Sev file:line Finding
N10 🟡 library/kernels/gfx950/attention_tiled_2d.py:5471 gfx950 rtl specs past the LDS budget pass the gate and both builders
• impact: D256 with 8 warps (block 32 or 64) is admitted and built in Python and C++, then clang refuses it for local memory (ran); shipped objects use 2 or 4 warps
• the reply and attention.md:447 say a general gate applies; only dispatch's _gfx950_2d_legality checks LDS
• fix: move _gfx950_2d_lds_bytes beside the kernel, with a C twin; gate, both builders and dispatch call it

Utility: 🟢 Positive — gfx942 LDS gate fixed in both engines; gfx950 still ungated


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

Follow-up review: first pass by qwen3.6 (local), verified by the orchestrator (Claude Opus 5.5). Checks: changes since fe41d279 (rebase onto develop; gfx942 LDS gate, representative goldens, docs), a merge with develop (clean: build + ctest, byte identity, goldens, installed tests, compile), CI at head, an LDS probe on both arches in both engines, code quality. Flags whether a closer look is needed; not an approval or rejection.

dpeabody-amd added a commit that referenced this pull request Nov 7, 2026
…1151WmmaAttention

hipDNN's paged K/V contract (SdpaSemanticsRegister SEM-008/009, #14358)
differs from the one this engine was written against:

- Every paged graph carries seq_len_q and seq_len_kv, also with a dense Q.
  The matcher now requires both and always binds seq_len_q. The dense-Q
  form without seq_len_q is gone: GraphLayout keeps PAGED_DENSE_Q (dense Q,
  live rows seq_len_q) and PAGED_RAGGED_Q, and an object without seqused_q
  serves no graph.
- A dense Q's rows past seq_len_q, up to Sq, are padding with O = 0 and
  LSE = -inf. The kernel now writes them (an earlier commit); the GPU test
  expects that rather than untouched rows.
- A ragged Q follows RFC 0014 (#15080): origins may come in any order,
  ragged_offset[B] is neither used nor validated, and the rows of an offset
  block past seq_len_q belong to no sequence. The kernel reads only each
  sequence's origin and writes only its live rows; the GPU test lays the
  sequences out in reverse order, poisons ragged_offset[B] and expects
  those rows untouched.
- Separate K and V page tables (page_table_k, page_table_v) are valid
  graphs and are now served: the matcher binds both, and the launch passes
  them as the kernel's block_table and block_table_v (one table bound to both
  sides is passed twice). The two tables must have one width and, with more
  than one sequence, one row stride.
- One-sided paging, per-side block sizes or page counts, K and V tables of
  different widths or row strides, and packed page tables are valid graphs
  the kernels cannot read, so each is declined with a logged cause.
- An unset max_seq_len_kv is valid. The kernels use it only as the bound on
  seq_len_kv, so it is inferred as the table width times the block size, as
  cuDNN does.
- A graph may declare a device-pointer alignment below the default 16 bytes
  (engine API 1.3.0, which the plugin's 1.4.0 admits). The kernels load
  Q / K / V / O rows as 16-byte vectors, INT32 operands as 4-byte words and
  the sink with 16-bit loads, so the matcher declines a graph declaring less
  for any of them (cause 'alignment', dense and paged) where hipDNN can
  still pick another engine, rather than the plan failing at execute.
- left_bound follows cuDNN (L keys ending at the diagonal); maskFor already
  translates it to the kernel's window_left = L - 1. The GPU test's windows
  follow it (32-, 16- and 128-key windows).

Catalog: the 48 paged objects with dense_q but without seqused_q (24 plain,
24 with sinks) could serve no hipDNN graph and are dropped; 96 paged
objects remain (dense Q and ragged Q, each with and without sinks), 144
prefill objects in all. Retained descriptor ids are unchanged. The census
is regenerated.

Decline logs: the paged checks log the cause keys #14247 uses for the same
contract rules, so the shared checks it moves into SdpaGraphMatch.hpp keep
every key this engine logs: one_sided_paging; page_table (a packed table,
K and V tables of different widths or row strides, a table that is not
INT32 [B, 1, M, 1] with a unit entry stride); block_size (different K and V
block sizes, a page size that is not a power of two, and one no paged object
ships for, which used to find no object without a log); shape (different K
and V page counts); seq_len_q, seq_len_kv and ragged (a missing or malformed
sequence-length or offset tensor, declined silently before); attn_mask
(additive bias on a paged graph). Each reason names the operand and values
involved. The dense layout logs 'padding' for seq_len_q, seq_len_kv,
max_seq_len_kv or padding_mask. Each line is headed by the declining
engine's name (sdpa_graph::DeclinedBy): the decode engine runs the shared
dense matcher too, and used to log under the prefill engine's name; it now
also logs why it declines a paged graph.

Tests: the matcher scaffolding both unit tests share (engine symbols, the
graph fixture, object metadata, the decline-log reader, the ABI helpers)
moves into Gfx1151WmmaAttentionTestSupport.hpp. The unit tests bind
separate K and V tables (dense and ragged Q) and a V container with its own
strides, and assert the engine, cause and reason of each decline: every
paged check in isolation, the unshipped page sizes, alignment (sink
included), FP8, dropout, a forced implementation, the other unserved
features and ALiBi; the dense declines are asserted for both engines. The
GPU test checks against test_sdk's CpuFpReferenceSdpa (gatherPagedKv per
page table, sequence lengths, band, sink and LSE) instead of its own
reference, adds two cases whose V pages sit in other blocks than their K
pages, through a second page table, two whose V page, head and token strides
all differ from K's, and runs the 96 and 256 buckets at pages 32 and 64 and
the 128 bucket at page 64.

Sinks: the engine still accepts a sink_token in the Q data type and
declines a FLOAT one. hipDNN's frontend accepts only a FLOAT sink, so the
GPU test's sink cases now skip with a message when the frontend rejects a
Q-dtype sink; the paged-rules probe that skipped every case is removed, and
the other cases run.

Re-pin the gfx1151_wmma_attention.yaml digest in the IngestorGenerator
fixture unslotted_render_golden.json (test_slots.py): the contract
change alters that config's render by design.

@yraparti yraparti left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

rocKE checklist re-review of the fixes since f1d0d688696e

ROCm/rocm-libraries#14247 · 3bed02a4f10a · model gemini-4-argon, fallback gpt-5.5

FAIL · graded 238/238 · BLOCKER 1, HIGH 3, MEDIUM 2, LOW 1, NEEDS-INFO 73

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

Since the last review (f1d0d688696e)

Resolved 3 · still open 1 · new 6

  • ✅ §A.10 was FAIL, now NEEDS-INFO — had asked: Add representative golden hash entries for the runtime-layout flag-on IR for each affected architecture, LLVM flavor, and emitted path.
  • ✅ §E.7 was FAIL, now NEEDS-INFO — had asked: Add the required rocKE support/reference documentation for the new runtime_layout attention feature in the same PR.
  • ✅ §P.4 was FAIL, now PASS — had asked: Add a C++ runtime-layout LDS byte calculation and capacity check matching the Python gate, and cover it with a parity/rejection test case.
  • ⏳ §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 the results in the PR description.
  • 🆕 §C.1 new FAIL (MEDIUM): Define a named SW_ACTIVE = SLIDING_WINDOW > 0 or RTL in the gfx942 builder and replace all repeated SLIDING_WINDOW > 0 or RTL sites with it.
  • 🆕 §C.10 new FAIL (HIGH): Correct all runtime-layout kernarg-count comments to the actual 16 appended arguments and grep the repo for stale 15 appended or 17 appended wording.
  • 🆕 §G.6 new FAIL (HIGH): Update the five C++ comments in helper_helper_rocke.instances.gfx942.attention_tiled_2d.h:187, instance_gfx942_attention_tiled_2d_internal.h:188, attention_tiled_2d_ctx_prologue.cpp:570, instance_gfx950_attention_tiled_2d_internal.h:187, and attention_tiled_2d_public_entry_glue.cpp:855 to say '16 appended' to match ROCKE_ATTN2D_RTL_NUM_PARAMS and RUNTIME_LAYOUT_EXTRA_ARGS.
  • 🆕 §H.2 new FAIL (BLOCKER): Add a default-on host contents validation (with an explicit opt-out for graph-captured hot paths) that checks every K and V page-table physical id used by the launch is in range of the K/V container block count before launch.
  • 🆕 §H.5 new FAIL (HIGH): Add host validation for every launch input that feeds a device bound or mask, including page-table physical ids, sequence lengths, ragged offsets, strides and counts, and add negative tests that exercise each rejected input.
  • 🆕 §L.4 new FAIL (MEDIUM): Replace hard-coded wave-lane constants in the new runtime-layout geometry and launch logic with the exact target wave size, and mirror that in the C++ runtime-layout/pack code.

Findings

  • BLOCKER §H · dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:3254 · Add a default-on host contents validation (with an explicit opt-out for graph-captured hot paths) that checks every K and V page-table physical id used by the launch is in range of the K/V container block count before launch.
  • HIGH §C · dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_helper_rocke.instances.gfx942.attention_tiled_2d.h:185 · Correct all runtime-layout kernarg-count comments to the actual 16 appended arguments and grep the repo for stale 15 appended or 17 appended wording.
  • HIGH §G · dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_attention_tiled_2d_rtl_internal.h:6 · Update the five C++ comments in helper_helper_rocke.instances.gfx942.attention_tiled_2d.h:187, instance_gfx942_attention_tiled_2d_internal.h:188, attention_tiled_2d_ctx_prologue.cpp:570, instance_gfx950_attention_tiled_2d_internal.h:187, and attention_tiled_2d_public_entry_glue.cpp:855 to say '16 appended' to match ROCKE_ATTN2D_RTL_NUM_PARAMS and RUNTIME_LAYOUT_EXTRA_ARGS.
  • HIGH §H · dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/packs/AttentionPagedUnifiedNative.cpp:57 · Add host validation for every launch input that feeds a device bound or mask, including page-table physical ids, sequence lengths, ragged offsets, strides and counts, and add negative tests that exercise each rejected input.
  • MEDIUM §C · dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:2678 · Define a named SW_ACTIVE = SLIDING_WINDOW > 0 or RTL in the gfx942 builder and replace all repeated SLIDING_WINDOW > 0 or RTL sites with it.
  • MEDIUM §L · dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1133 · Replace hard-coded wave-lane constants in the new runtime-layout geometry and launch logic with the exact target wave size, and mirror that in the C++ runtime-layout/pack code.
  • LOW §S · dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h:116 · Run PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py from rocke/platform/ and record the results in the PR description.

Needs info

  • §A.5 dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention.json (+192 -0 lines) and dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json (+195 -5 lines) are in the file list, but their diffs are omitted to fit the model budget; need those hunks to verify existing representative shapes are untouched and why 5 lines changed in llvm_gfx_all.json.
  • §A.6 The PR adds cross-engine C and Python parity emit configurations (indices 20-38 on gfx942 and 72-90 on gfx950) rather than relying solely on Python-drift goldens (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.c:353; dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.py:372; dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.c:1233; dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.py:1194). — no diff line supports this verdict
  • §A.8 While the C++ twin was hand-reviewed in the diff (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d_ctx_prologue.cpp:108) and parity cases were added (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.c:353), the diff for dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/rocke_ir_parity_harness.py (+65 -0 lines) is omitted and no check_byte_identity.py run output is included.
  • §A.10 Since the previous FAIL, the author updated dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention.json (+192 -0 lines) and dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json (+195 -5 lines), which addresses the missing representative golden hash entries for runtime_layout (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:528), but the diff hunks for both JSON files are omitted to fit the model budget so the exact per-LLVM-flavor and per-arch entries cannot be verified.
  • §B.4 Not resolved since previous review: runtime_layout: bool = False in public specs UnifiedAttention2DTiledSpec (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:528-537 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:544-568) changes the kernel ABI by appending 16 kernel arguments (RUNTIME_LAYOUT_EXTRA_ARGS / _RTL_PARAMS), but __post_init__ only validates if type(self.emit_lse) is not bool: and omits if type(self.runtime_layout) is not bool:, allowing truthy non-bools (such as "false" or 1) to coerce into an enabled runtime-layout ABI. — no diff line supports this verdict
  • §B.5 The newly opened runtime_layout gate scope (fp16/bf16, head sizes 64/128/256, block sizes 16/32/64, GQA ratios 1..16, sinks, and sub-block tiles on gfx942) is covered in parity emitters (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.py:372-390 and dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.py:1194-1212) and dedicated test suites (test_gfx942_attention_runtime_layout.py and test_attention_tiled_2d_runtime_layout_gfx950.py). — no diff line supports this verdict
  • §B.6 Resolved from previous review with diff lines: _runtime_layout_reject_reason and _supports_runtime_layout in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1190-1340 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1068-1148 check only functional/architectural invariants (incompatible pipeline flags, head count divisibility, tile/DMA copy divisibility, and LDS capacity) and report hardware/algorithm invariant reasons. — no diff line supports this verdict
  • §C.2 The runtime-layout path is integrated into the existing tiled 2D builders via if spec.runtime_layout in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1650 and the existing gfx950 builder sets RTL = spec.runtime_layout in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1372, with no visible bespoke shape-class build_* script. — no diff line supports this verdict

Section notes

  • §A: Re-review against current head 3bed02a. Item 4 is PASS after inspecting the C++ temporary-binding emit order in the diff; items 5, 8, and 10 are NEEDS-INFO because the golden and parity harness hunks (attention.json, llvm_gfx_all.json, rocke_ir_parity_harness.py) are omitted from the diff budget, though attention.json is now present in the modified file list with +192 -0 lines.
  • §C: Re-review against the visible current diff; omitted hunks were not used as evidence.
  • §E: The supplied diff is still truncated in multiple hunks; items that require omitted files or omitted hunks are marked NEEDS-INFO rather than guessed.
  • §G: Re-review of head 3bed02a (items 1-17 of 34). Stale appended-argument counts (15 and 17 vs 16) remain in C++ header and builder comments (item 6).
  • §G: Items 18–34 of 34 in section G.
  • §H: Trace: K/V per-page buffer bound extent comes from runtime kernarg kv_block_span_bytes and the kernel visits one compile-time KV tile T/BS within a page; token lanes past seq_len use an OOB voffset, but the physical page-id value can move the 64-bit base outside the allocation. Block-table lookup bound extent comes from seq_len as ceil(seq_len/BS) or lblk_tok < seq_len while the kernel visits tiles up to num_tiles derived from max_seq_prefix_len; out-of-range logical lookups default to page 0 and are masked. Softmax/key mask extent comes from runtime seq_len via max_seq_prefix_len/num_tiles while each compile-time tile visits T columns; columns/tokens past seq_len are masked. Q/O row bounds come from q_uniform_len, seq_lens_q, and ragged offsets while each q-block visits BLOCK_Q-derived rows; dense padding stores zero and ragged only stores live rows. Sink bound comes from num_query_heads and sink_is_f32; loads are masked to index 0 when qh is out of range.

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 new Python attention runtime_layout spec field and kernel body logic are mirrored in the C++ attention spec and instance twins…
2 N/A No core lowering-op or op-handler files under core/lower_* or core/ir…
3 PASS The kernel-level runtime_layout knob and validation rules are mirrored across the C++ bindings (t2d_build and fkv_build), C++ spec validation, and…
4 PASS Resolved by inspecting the C++ emitters in the diff: C++ helpers bind operands to temporaries left-to-right one builder call per statement in Python order…
5 NEEDS-INFO dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention.json (+192 -0 lines) and dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json (+195…
6 NEEDS-INFO The PR adds cross-engine C and Python parity emit configurations (indices 20-38 on gfx942 and 72-90 on gfx950) rather than relying solely on Python-drift goldens…
7 PASS All new emitted IR is guarded behind if RTL: / if (ctx->RTL) (default runtime_layout=False), leaving covered non-RTL shapes unchanged…
8 NEEDS-INFO While the C++ twin was hand-reviewed in the diff (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d_ctx_prologue.cpp:108) and parity cases were added…
9 N/A This PR does not contain a shared op-handler fix or a lowerer-only…
10 NEEDS-INFO Since the previous FAIL, the author updated dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention.json (+192 -0 lines) and…
11 N/A No new IR opcode is introduced in this PR.
12 N/A core/backend.py and CPP_UNPORTED_ARCHES are not touched in this PR.
13 N/A Emitted .ll formatting, stamps, and ROCKE_DEBUG_LOC gate settings are…
14 N/A No classic helper under platform/cpp/helpers/ is modified in this PR.
§B graded 6/6
# Verdict Evidence
1 PASS New rejections in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1190-1213…
2 PASS Preflight gates _supports_tiled_2d_spec (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:7407-7408) and…
3 PASS runtime_layout defaults to False and is excluded from rocKE dispatch (`dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:990-…
4 NEEDS-INFO Not resolved since previous review: runtime_layout: bool = False in public specs UnifiedAttention2DTiledSpec (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:528-537 and…
5 NEEDS-INFO The newly opened runtime_layout gate scope (fp16/bf16, head sizes 64/128/256, block sizes 16/32/64, GQA ratios 1..16, sinks, and sub-block tiles on gfx942) is covered in parity emitters…
6 NEEDS-INFO Resolved from previous review with diff lines: _runtime_layout_reject_reason and _supports_runtime_layout in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1190-1340 and…
§C graded 10/10
# Verdict Evidence
1 FAIL The gfx942 builder repeats the same runtime-layout window predicate directly at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:2678 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:5376 instead of using one named predicate as gfx950…
2 NEEDS-INFO The runtime-layout path is integrated into the existing tiled 2D builders via if spec.runtime_layout in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1650 and the existing gfx950 builder sets `RTL =…
3 NEEDS-INFO The kernel name only adds the algorithm/codegen token rtl for runtime layout at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1107 and…
4 NEEDS-INFO The arch-independent runtime-layout kernarg contract is forked in per-arch Python files as RUNTIME_LAYOUT_EXTRA_ARGS at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1156 and _RTL_PARAMS at…
5 PASS The visible fan-out tracks the inherent feature surface: pack registration at dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/IngestorP…
6 PASS The diff refactors shared launch-signature construction into appendKernelArgument / tiled2DKernelSignature at…
7 N/A This PR is primarily a new runtime-layout/AttentionPagedUnified feature…
8 PASS The diff reuses current shared helpers instead of reimplementing them, e.g. sdpa_graph::declaresAlignment in…
9 PASS Runtime-layout rejection is handled across sibling Python and C++ sites: gfx942 Python at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/att…
10 FAIL The runtime-layout kernarg widening is not consistently updated in prose: dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_helper_rocke.instances.gfx942.attention_tiled_2d.h:185 says the mode adds 15 appended kernel arguments, while…
§D graded 21/21
# Verdict Evidence
1 PASS The new runtime-layout feature is declared as a UnifiedAttention2DTiledSpec field rather than as an environment-derived knob…
2 N/A This PR does not add or modify Python dispatcher KernelCandidate or…
3 N/A The previous NEEDS-INFO is resolved as N/A because this diff does not add…
4 PASS The runtime-layout catalog helper takes arch as an explicit parameter and passes it through to the LDS-capacity lookup rather than probing a device…
5 PASS runtime_layout is included in the tuning spec identity only when true, is listed as a semantic field, and is kept out of default config keys while false…
6 N/A This PR does not add or modify Python dispatcher cohort routing.
7 N/A This PR does not add or modify hand-pinned dispatcher overrides or…
8 N/A This PR does not add or modify dispatcher cohort predicates shared by…
9 N/A The new executable path is registered as a hipDNN ingestor pack, not as a…
10 NEEDS-INFO The source wiring for runtime_layout is visible, but the dispatcher-level test hunks are omitted; need untruncated diff lines showing a request-level test exercises request/graph matching through spec selection and binding for this feature.
11 N/A This PR does not register or promote any Python dispatcher attention…
12 N/A This PR does not add or reprioritize a Python CandidateRegistry candidate…
13 N/A This PR does not add a CandidateRegistry pin wrapper or pinned dispatcher…
14 PASS The runtime_layout field is excluded from the tuning spec hash/key while false at dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.…
15 N/A This PR does not add a new Python dispatcher candidate or arch module in…
16 N/A This PR adds a kernel-ingestor engine and runtime-layout kernel mode, not a…
17 N/A No new CandidateRegistry candidate is added, so the CPU routing test…
18 N/A The diff does not modify dispatcher registry ratchets or CandidateRegistry…
19 PASS The diff explicitly states that dispatch does not set runtime_layout at dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:991…
20 PASS The Python kernel spec adds runtime_layout at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:526, and the C++ side…
21 N/A This is a new runtime-layout kernel mode and ingestor engine, not a…
§E graded 17/17
# Verdict Evidence
1 NEEDS-INFO Prior NEEDS-INFO is not resolved: platform test and platform tool hunks are omitted, so I cannot verify platform tests do not import library code or that the platform←library one-way dependency is preserved.
2 PASS The visible library dispatch change only filters semantic fields for spec hashing at dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/comm…
3 NEEDS-INFO Prior NEEDS-INFO is not resolved: the relevant test hunks are omitted, so torch use in tests and any ROCm-library-location logic cannot be verified as gated with pytest.importorskip or absent.
4 NEEDS-INFO Prior NEEDS-INFO is not resolved: library test hunks and complete Python import sections are omitted, so upward imports, function-body imports, and whether known violation lists only shrink cannot be verified.
5 NEEDS-INFO Prior NEEDS-INFO is not resolved: platform CMake/test hunks are omitted, so absolute or repo-escaping paths and helper-script hygiene under platform/ cannot be fully checked.
6 NEEDS-INFO Prior NEEDS-INFO is not resolved: emitter trees are modified, but no trusted diff artifact provides process evidence that no lint autofix removed intentional handles or that byte identity was rerun after lint/format changes.
7 NEEDS-INFO Prior FAIL appears partially addressed because the file list now includes SUPPORT_MATRIX.md and dsl_docs/instances/attention.md, but those hunks are omitted, so I cannot cite the required documentation lines for the runtime_layout feature…
8 NEEDS-INFO Prior NEEDS-INFO is not resolved: large documentation, descriptor, generated, and test hunks are omitted, so Jira IDs, PR numbers, reviewer handles, or tracker URLs in committed artifacts cannot be ruled out.
9 NEEDS-INFO Prior NEEDS-INFO is not resolved: large descriptor, generated, documentation, and test hunks are omitted, so I cannot determine whether every touched area is part of the same deliverable rather than unrelated fixes.
10 NEEDS-INFO Prior NEEDS-INFO is not resolved: no trusted diff artifact provides merge-simulation evidence listing overlapping open PRs, agreed land order, and rerun affected goldens/suites on develop plus those PRs.
11 NEEDS-INFO Prior NEEDS-INFO is not resolved: large documentation, descriptor, generated, and test hunks are omitted, so scratch notes or throwaway design files cannot be ruled out.
12 PASS Both target arches are covered by runtime_layout parity rows: gfx942 rows are added in dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/attenti…
13 N/A No dead-code removal or concept rename is evident from the provided diff.
14 NEEDS-INFO The review-response comparison is unavailable, so it is missing whether a rebase/cherry-pick/port occurred and whether removed upstream lines or KNOWN LIMITATION notes were audited.
15 NEEDS-INFO Prior NEEDS-INFO is not resolved: large new source and test hunks are omitted, so every newly accepted parameter cannot be checked for use in its function body.
16 N/A No reviewer-facing severity labels or review comments are part of this code…
17 PASS no absolute throughput, bandwidth, or latency figure in added lines
§F graded 11/11
# Verdict Evidence
1 PASS Branch users/dpeabody/hipdnn-cdna-paged-attn matches users/<user>/… (PR metadata; diff touches…
2 PASS title feat(hip-kernel-provider): paged/ragged SDPA engine for gfx942/gfx950 matches type(scope): summary and is 69 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/test_attention_tiled_2d_runtime_layout_gfx950.py is in the diff
5 PASS No keys, certs, or secrets appear in the changed file list; additions are source, config, and test files such as…
6 NEEDS-INFO Unresolved from previous review: local pre-commit execution status before push is not observable from the git diff.
7 NEEDS-INFO Unresolved from previous review: GitHub review-thread states, bot/Copilot findings, and deferred ticket IDs are not included in the diff context.
8 N/A The PR base branch is develop, so stacked-PR feature-branch policy does…
9 N/A Nothing in the PR metadata or diff indicates the change was artificially…
10 NEEDS-INFO The bodies of all added unit-test files (e.g. dnn-providers/hip-kernel-provider/rocke/library/tests/test_gfx942_attention_runtime_layout.py, `dnn-providers/hip-kernel-provider/src/tests/engines/kernel_ingestor_engine/packs/TestAttentionPaged…
11 NEEDS-INFO Unresolved from previous review: whether the author verified the PR title and body via gh api / gh pr view --json after editing cannot be determined from the diff.
§G graded 34/34
# Verdict Evidence
1 NEEDS-INFO Parity configs are added in dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.py:372 and dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.py:1194, while…
2 NEEDS-INFO The GPU test files (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, IntegrationGpuAttentionPagedUnified.cpp) are omitted from the truncated diff, and on-silicon execution results are not shown.
3 NEEDS-INFO Local non-GPU gate sweep output (tools/run_checks.py, check_byte_identity.py across llvm20/22/23, check_ir_validity.py, pytest) cannot be verified from the diff alone.
4 NEEDS-INFO All new unit and runtime-layout test files (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, TestAttentionPagedUnified*.cpp, TestGfx950AttentionDenseSchemaGate.cpp) are omitted from the truncated diff.
5 NEEDS-INFO The new test files (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, TestAttentionPagedUnified*.cpp, TestGfx950AttentionDenseSchemaGate.cpp) are omitted from the truncated diff, so test falsifiability…
6 FAIL Not resolved since the previous review: C++ comments state conflicting appended kernel argument counts (15 and 17) while the actual count is 16 in dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_attention_tiled_2d_rtl_internal.h:6 and :42 (ROCKE_ATTN2D_RTL_NUM_PARAMS = 16); see '15…
7 NEEDS-INFO The engine README (dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md) and device test files are omitted from the truncated diff.
8 NEEDS-INFO The test harnesses (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, IntegrationGpuAttentionPagedUnified.cpp) are omitted from the truncated diff.
9 NEEDS-INFO The PR body cites SGPR spill and instruction counts from dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md, which is omitted from the truncated diff.
10 N/A This PR introduces functional support for a new paged/ragged SDPA engine…
11 NEEDS-INFO dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/packs/AttentionPagedUnifiedNative.cpp is truncated at line 600 before the dispatch registration and handler implementation, and pack/integration tests are omitted from the…
12 PASS Both Python and C++ rocKE builders and bindings are updated in lockstep (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.p…
13 N/A No new LLVM intrinsic overload classes or mangled intrinsic symbols are…
14 N/A This PR does not make hardware bottleneck classification claims.
15 PASS The PR diff is inspected directly from head commit 3bed02a (e.g.…
16 PASS Symbols and shared helpers match existing codebase conventions across layers (e.g., check_head_counts / rocke_attention_check_head_counts in…
17 NEEDS-INFO Emitter scripts are updated in dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.py:337 and dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.py:1163, but…
18 NEEDS-INFO The guard test files (dnn-providers/hip-kernel-provider/src/tests/engines/kernel_ingestor_engine/packs/TestGfx950AttentionDenseSchemaGate.cpp, TestAttentionPagedUnifiedMatchers.cpp, and test_gfx950_attention_dense_schema_gate.py) are omitted…
19 NEEDS-INFO The test fixture files (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py, IntegrationGpuAttentionPagedUnified.cpp, TestAttentionPagedUnifiedMatchers.cpp) are omitted from the diff.
20 N/A This PR does not add or modify a claim-ban detector or linter.
21 N/A This PR does not add benchmarking or performance measurement code.
22 NEEDS-INFO The new test files (test_gfx942_attention_runtime_layout.py, test_attention_tiled_2d_runtime_layout_gfx950.py), golden contract test (TestIngestorPlanLaunchContractGoldens.cpp), and golden JSON files (attention.json, llvm_gfx_all.json) are…
23 N/A This PR adds a new paged/ragged SDPA engine and runtime_layout kernel mode…
24 N/A This PR does not re-bless existing IR goldens; runtime_layout=False leaves…
25 N/A This PR is a feature addition (runtime_layout mode and…
26 N/A This PR does not include performance benchmarks or speedup claims.
27 N/A This PR does not report a performance A/B comparison.
28 NEEDS-INFO dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:528 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:544 default runtime_layout=False so existing kernels remain…
29 N/A This PR does not swap guarded math primitives for fast approximations.
30 NEEDS-INFO The PR description cites SGPR-to-VGPR lane spill and zero scratch usage from dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md, which is omitted from the diff.
31 N/A This PR adds a new engine and kernel mode rather than claiming a fix or…
32 NEEDS-INFO The static register-pressure and instruction-share measurements cited in the PR body live in dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md, which is omitted from the…
33 PASS dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:995, dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/tuning_…
34 N/A runtime_layout is a layout/execution mode rather than an opt-in extra…
§H graded 5/5
# Verdict Evidence
1 PASS Host metadata checks bound the K/V per-block span before dispatch by computing kvSpanBytes from the K/V tensor layout and refusing spans over…
2 FAIL The runtime-layout K/V path reads phys from the block table and uses it directly in a 64-bit page base before creating only a per-page buffer resource, so an out-of-range physical id is not contained by the original allocation bound…
3 PASS Runtime-layout specs restrict the relevant compile-time constants to power-of-two production values, including head_size 64/128/256, block_size 16/32/64…
4 PASS The runtime seq_len drives max_seq_prefix_len and num_tiles, padding-only Q blocks force num_tiles to zero, and the softmax mask still requires col_abs <…
5 FAIL The new engine states that actual lengths and block ids are only read on device and execute performs no host synchronization, while block_tables feed device-side page bounds and physical page bases, so the required host validation is absent for at least the block-table value input…
§I graded 6/6
# Verdict Evidence
1 N/A This PR adds the runtime_layout kernel mode and the…
2 PASS Unsupported spec combinations, builders, and graph attributes are loudly rejected with explicit reasons in…
3 NEEDS-INFO The diffs for dnn-providers/hip-kernel-provider/rocke/library/tests/test_gfx942_attention_runtime_layout.py, dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_tiled_2d_runtime_layout_gfx950.py, and the C++ pack tests remain…
4 N/A This PR does not add or update performance benchmark or routing comparison…
5 NEEDS-INFO The diffs for dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md and projects/hipdnn/tools/IngestorGenerator/configs/attention_paged_unified.yaml remain omitted, so the…
6 N/A This PR does not introduce any new DSL ops or hardware-instruction-named…
§J graded 7/7
# Verdict Evidence
1 NEEDS-INFO Previous NEEDS-INFO is not resolved: the current diff still omits the Markdown hunks for dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/rocKE/attention_paged_unified/README.md and…
2 NEEDS-INFO Previous NEEDS-INFO is not resolved: the current diff still omits the Markdown hunks, so notation tables, symbol definitions, domains, bridges such as decay/log-decay, and log-base statements cannot be verified.
3 NEEDS-INFO Previous NEEDS-INFO is not resolved: the current diff still omits the new attention_paged_unified/README.md content, so it cannot be checked for restating an algorithm owned by another document rather than linking and stating only the delta.
4 NEEDS-INFO Previous NEEDS-INFO is not resolved: added Markdown hunks are still omitted, and origin/develop contents are not provided here, so any added/new/shared novelty claims cannot be enumerated or checked against develop.
5 NEEDS-INFO Previous NEEDS-INFO is not resolved: the Markdown hunks are still omitted, so scaling or complexity claims in the docs, and whether they state the aggregate the reader tunes at every site, cannot be checked.
6 NEEDS-INFO Previous NEEDS-INFO is not resolved: the Markdown hunks are still omitted, so constants, catalog/model tables, tuned band edges, out-of-scope rows, and stated assumptions cannot be checked for provenance.
7 NEEDS-INFO Previous NEEDS-INFO is not resolved: the new attention_paged_unified/README.md hunk is omitted, so its sibling filename convention, declared entry point, and ownership split cannot be verified.
§K graded 13/13
# Verdict Evidence
1 N/A This PR does not add or modify any tuned dispatch candidates.
2 NEEDS-INFO Resolved with exact line citations: the new runtime_layout spec field (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:528…
3 PASS runtime_layout is omitted from spec_hash items when False (dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:996) and…
4 PASS KnobSpace.defaults() in dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/unified_rules.py:401 excludes runtime_layout alongside…
5 N/A This PR does not add or modify any tuning knobs, axes, or drop/refuse rules.
6 NEEDS-INFO Resolved with exact line citations: dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/tuning_specs.py:185 classifies runtime_layout as a semantic field without invoking production heuristics, and spec legality is gated in…
7 N/A This PR does not modify KnobSpace.base(req) or request normalization.
8 N/A This PR does not touch knob normalization or axis type coercion.
9 N/A This PR does not touch benchmark recording, replay, or PinRefused handling.
10 N/A This PR does not modify KnobSpace.find or tuning_id resolution.
11 N/A This PR does not touch sweep entry points or core.pin_to_spec.
12 N/A This PR does not reorder production stacks, change accept_default, or…
13 N/A This PR does not modify any benchmark or sweep harness.
§L graded 10/10
# Verdict Evidence
1 N/A This PR does not add a new gfx target or modify the architecture/ISA target…
2 N/A This PR does not add a new MMA op_id or MMA fragment layout row.
3 N/A This PR does not change MMA atom selection, MMA catalog metadata, intrinsic…
4 FAIL Previous NEEDS-INFO is not resolved: new runtime-layout lane math hard-codes wave lanes as threads = spec.num_warps * 64 instead of deriving the wave size from the exact target at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1133.
5 NEEDS-INFO Previous NEEDS-INFO is resolved: gfx942 runtime-layout explicitly rejects gfx950-only recipe flags such as use_mfma_32x32/transpose-read paths in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1138…
6 N/A This PR does not add pipeline vocabulary to ArchTarget/core arch data…
7 PASS Previous NEEDS-INFO is resolved: runtime-layout rejects FP8 KV storage via kv_storage_dtype in…
8 N/A This PR does not add or modify scaled-MMA operations, scaled op ids, or…
9 N/A This PR does not add a new gfx architecture target.
10 N/A This PR does not touch intrinsic arch-domain artifacts or LLVM…
§M graded 10/10
# Verdict Evidence
1 NEEDS-INFO Host-time Python if statements test only spec booleans (if RTL:), while runtime SSA values use b.cmp_* with b.select, b.scf_if, and b.scf_for without escaping across regions…
2 N/A The diff does not touch legalize_for_target or any target legalization…
3 N/A The diff does not modify print_ir or ck.dsl.ir/v1 IR serialization…
4 N/A The diff does not modify the portable IR schema, importer, recipe…
5 NEEDS-INFO The runtime_layout async K/V loads keep the existing wave-uniform LDS base (K_wave_base/V_wave_base) and only supply the per-page buffer resource and masked voff…
6 NEEDS-INFO The grid-stride loop and rtl_rows_if/rtl_live_block conditions are CTA-uniform, b.sync() is placed after exiting the conditional before the next loop iteration…
7 PASS Buffer resources for runtime_layout are constructed via b.buffer_rsrc / rocke_b_buffer_rsrc without hard-coding word 3 or using illegal temporal…
8 N/A The diff does not add any new LLVM intrinsic declarations or modify backend…
9 PASS Existing public headers under cpp/include/rocke/ are not renamed, ROCKE_ABI_VERSION and BINARY_ABI are bumped in lockstep from 12 to 13…
10 N/A The diff does not touch !dbg emission, LLVM module debug flags, or DWARF…
§N graded 18/18
# Verdict Evidence
1 N/A The PR does not add or modify any module under `platform/python/rocke/helper…
2 PASS The runtime-layout paths reuse shared helpers binary_search_seq_idx (`dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py…
3 N/A The PR does not use or modify SoftwarePipeline.
4 N/A The PR does not add or modify fused GEMM/epilogue kernels or fusion helpers.
5 N/A The PR does not add or modify any rocke.helpers.tiling API.
6 N/A The PR does not touch front-door/toolbox tiling helpers or LayoutStyle.
7 N/A The PR does not use or modify helpers/tiling/ layout overrides or MMA…
8 N/A The PR does not use or modify helpers/tiling/ LDS or global memory verbs.
9 N/A The PR does not use or modify helpers/tiling/ axis roles, batch…
10 N/A The PR does not change helpers/tiling/ emit, LDS staging, or MMA issue.
11 N/A The PR does not use or modify helpers/tiling/ traits, TileMma…
12 N/A The PR makes no LDS bank-conflict or global coalescing claims under…
13 N/A The PR does not touch WarpDistributionEncoding, make_tile_desc, or…
14 N/A The PR does not use helpers/tiling/ data-movement verbs or add a tiling…
15 N/A The PR does not modify helpers/tiling/ allocation or bounds checks on…
16 N/A The PR does not use or modify optimize_layout or evaluate_transform.
17 N/A The PR does not touch the tiling recorder, round-trip analysis, or tiling…
18 N/A The PR does not touch InterleavedStyle or demo kernels under…
§O graded 16/16
# Verdict Evidence
1 N/A The diff does not modify core/verify.py, cpp/core/verify.cpp, or introduce…
2 PASS In dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1943 and…
3 NEEDS-INFO Resolved from previous NEEDS-INFO: in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1889 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1554, parameters and IR…
4 PASS In dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1892 and 2015, module singletons I32 and I64 are used directly and…
5 N/A The diff does not touch core/passes.py, cpp/core/passes.cpp, or op purity…
6 N/A The diff does not touch constant folding in core/passes.py or…
7 N/A The diff does not modify kernel attributes, COMGR options, or cluster_dims…
8 N/A The diff does not touch CodegenPolicy, SchedulerStrategy, or compile_kernel…
9 N/A The diff does not touch waitcnt encoding, named/split barriers, or…
10 PASS In dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1892 and…
11 PASS In dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1661 and…
12 N/A The diff does not touch core target catalogs or lowerer vector-load /…
13 PASS Adding runtime_layout to rocke_attention_tiled_2d_spec_t in dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_helper_rocke.instances…
14 PASS New header dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_attention_tiled_2d_rtl_internal.h:9 uses #ifndef include guards (no…
15 N/A The diff does not modify engine fallback logic, LLVM flavor resolution, or…
16 PASS In dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d.cpp:162 and 242, C++ spec validation faithfully mirrors Python…
§P graded 14/14
# Verdict Evidence
1 PASS dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:544 adds runtime_layout: bool = False as a defaulted field on…
2 PASS dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_attention_tiled_2d_rtl_internal.h:47 defines rocke_attn2d_rtl_declare_params…
3 N/A This PR touches attention runtime-layout kernels and a hipDNN ingestor…
4 PASS Resolved since previous review: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d.cpp:346 implements…
5 N/A The diff does not touch WMMA/RDNA GEMM, GemmMultiD, or GemmMultiAbd…
6 N/A The diff does not add or modify Stream-K, MoE sorting, fused MoE, or…
7 N/A The diff does not change Stream-K, persistent Stream-K…
8 N/A The diff does not change quantization helpers, MX GEMM, quantized clamps…
9 N/A The diff does not introduce numeric emulation or quantized decode-path…
10 N/A The diff adds hipDNN ingestor descriptors and packs rather than a new rocKE…
11 NEEDS-INFO The dispatch/launch tail of dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/packs/AttentionPagedUnifiedNative.cpp is truncated after line 535, and the new Python test files (test_gfx942_attention_runtime_layout.py and…
12 N/A The diff does not modify COMGR or HIP library loading logic or fix a HIP…
13 NEEDS-INFO The diff modifies attention dispatch/tuning cache key and defaults (dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:992, dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/unified_rules.py:401)…
14 PASS New runtime-layout kernel logic is placed inside the attention kernel and C++ instance modules (`dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx…
§Q graded 7/7
# Verdict Evidence
1 NEEDS-INFO The diffs for dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt, dnn-providers/hip-kernel-provider/ROCKE_ENGINE_test_categories_external.yaml, and the new test files…
2 PASS The new runtime_layout kernel variant adds 19 parity configs per architecture to the C/Python parity emitters and binding config mirror…
3 NEEDS-INFO The PR does not weaken any gate criteria or claim numerics from IR parity; it extends the parity emitters (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.c:353…
4 NEEDS-INFO The PR modifies dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention.json (+192 -0 lines) and `dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_al…
5 NEEDS-INFO The contents of the new test files (dnn-providers/hip-kernel-provider/rocke/library/tests/test_gfx942_attention_runtime_layout.py, dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_tiled_2d_runtime_layout_gfx950.py, and…
6 N/A This PR does not refactor any test harness or verification gate.
7 NEEDS-INFO The PR adds new spec rejection rules in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1190 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1068 (and C++ twins…
§R graded 11/11
# Verdict Evidence
1 N/A runtime_layout is a new semantic mode for hipDNN's paged/ragged SDPA…
2 NEEDS-INFO Resolved with exact current-head citations: runtime_layout geometry and LDS rules are separated per architecture in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:1216-1253 (LDS budget via…
3 PASS Resolved with exact current-head citations: LDS sizing accounts for non-aliasable live allocations across the grid-stride loop and bank-conflict row…
4 NEEDS-INFO The PR adds runtime_layout global/buffer load and barrier paths (dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:3245-3304…
5 PASS Global Q/O loads and stores retain global_load_vN / global_store_vN with width 8 (align=16) and unit dim stride…
6 N/A The diff does not change any LDS or MMA layout; runtime_layout reuses the…
7 N/A The PR makes no bottleneck classification claims (memory-bound…
8 PASS Resolved with exact current-head citations: no waits or barriers are removed, and the single added b.sync() at the end of each grid-stride loop iteration…
9 N/A This PR is not a performance regression fix and does not change any…
10 N/A The PR makes no benchmark timing claims, occupancy-from-wall-clock…
11 N/A The PR introduces no new optimization failure mode, heuristic, case study…
§S graded 8/8
# Verdict Evidence
1 NEEDS-INFO Still NEEDS-INFO: The PR adds public symbols (UnifiedAttention2DTiledSpec.runtime_layout at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:528 and…
2 NEEDS-INFO Still NEEDS-INFO: Existing APIs (UnifiedAttention2DTiledSpec, build_unified_attention_2d_tiled, build_gfx942_4warp_gqa, supports_tiled_2d_spec, ROCKE_ABI_VERSION / BINARY_ABI) are updated in…
3 N/A The diff does not remove or rename any public rocke platform or library API.
4 N/A The diff does not modify __all__ in any module, and new internal helpers…
5 NEEDS-INFO Still NEEDS-INFO: New fail-fast ValueError and NotImplementedError raises are added in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:537…
6 NEEDS-INFO Still NEEDS-INFO: Parity emitter configs are added in dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_tiled_2d_emit.py:372 and…
7 NEEDS-INFO Still NEEDS-INFO: The C++ twin and binding are updated in dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp:827…
8 FAIL Not resolved since previous review: The PR touches dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h:116 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/portable_ir/src/abi.py:84, but the PR description still does not state that `PYTHONDONTWRITEBYTECODE=1…

lblk_tok = b.mul(lblk, b.const_i32(BS))
blk_ok = b.cmp_lt(lblk_tok, seq_len)
bt_voff = b.select(
blk_ok, b.shl(b.add(seq_base, lblk), b.const_i32(2)), rtl_kv_oob

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

[BLOCKER] §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 —…

The runtime-layout K/V path reads phys from the block table and uses it directly in a 64-bit page base before creating only a per-page buffer resource, so an out-of-range physical id is not contained by the original allocation bound: dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:3254-3257.

Fix: Add a default-on host contents validation (with an explicit opt-out for graph-captured hot paths) that checks every K and V page-table physical id used by the launch is in range of the K/V container block count before launch.

Comment on lines 185 to +187
bool emit_lse; /* False (Python kw_only) */
/* Runtime-layout mode (hipDNN paged/ragged SDPA): the Q/O/KV strides and a
* dense-or-ragged Q become 15 appended kernel arguments (after the runtime

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

2 checklist findings on this snippet

[HIGH] §C.10 ★ (line 185) ★ WIDENING a shared structure is a CLASS too (the dual of the rule above) — adding a member to a gate SET, putting a change INSIDE a branch every member takes, or introducing a NEW source of truth read by ONE field of a multi-field object reaches every member / sibling field / consumer, in…

The runtime-layout kernarg widening is not consistently updated in prose: dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_helper_rocke.instances.gfx942.attention_tiled_2d.h:185 says the mode adds 15 appended kernel arguments, while dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_attention_tiled_2d_rtl_internal.h:26 starts an enum that runs through ROCKE_ATTN2D_RTL_NUM_PARAMS after 16 named arguments.

Fix: Correct all runtime-layout kernarg-count comments to the actual 16 appended arguments and grep the repo for stale 15 appended or 17 appended wording.


[HIGH] §G.6 ★ (line 187) ★ Every load-bearing PROSE claim maps to a check that would FAIL if it were false — numeric AND structural. Numbers: identical across code / docstring / README / commit and matching the measurement; a "numerically identical / unchanged / bit-identical" claim needs an A/B that can actually…

Not resolved since the previous review: C++ comments state conflicting appended kernel argument counts (15 and 17) while the actual count is 16 in dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_attention_tiled_2d_rtl_internal.h:6 and :42 (ROCKE_ATTN2D_RTL_NUM_PARAMS = 16); see '15 appended kernel arguments' in dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_helper_rocke.instances.gfx942.attention_tiled_2d.h:187 and 'the 17 appended kernel arguments' / 'the 17 appended arguments' in dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_tiled_2d_internal.h:188, dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d_ctx_prologue.cpp:570, dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx950_attention_tiled_2d_internal.h:187, and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_public_entry_glue.cpp:855.

Fix: Update the five C++ comments in helper_helper_rocke.instances.gfx942.attention_tiled_2d.h:187, instance_gfx942_attention_tiled_2d_internal.h:188, attention_tiled_2d_ctx_prologue.cpp:570, instance_gfx950_attention_tiled_2d_internal.h:187, and attention_tiled_2d_public_entry_glue.cpp:855 to say '16 appended' to match ROCKE_ATTN2D_RTL_NUM_PARAMS and RUNTIME_LAYOUT_EXTRA_ARGS.

* @file AttentionPagedUnifiedNative.cpp
* @brief The hipkernel:AttentionPagedUnified engine's native half: matching and dispatch.
*
* The kernels are rocKE's tiled 2D unified-attention builders

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

[HIGH] §H.5 Every launch input that feeds a device-side bound/mask has BOTH a host-validation and a NEGATIVE test; an asymmetry (one input guarded, a sibling not) is the tell.

The new engine states that actual lengths and block ids are only read on device and execute performs no host synchronization, while block_tables feed device-side page bounds and physical page bases, so the required host validation is absent for at least the block-table value input: dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/packs/AttentionPagedUnifiedNative.cpp:57-61.

Fix: Add host validation for every launch input that feeds a device bound or mask, including page-table physical ids, sequence lengths, ragged offsets, strides and counts, and add negative tests that exercise each rejected input.

# A q-block of padding rows only reads no key: its rows store zeros.
num_tiles = b.select(
b.cmp_lt(qb_start_pos, cur_batch_q_len), num_tiles, b.const_i32(0)
)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

[MEDIUM] §C.1 Repeated predicate/logic → one named helper (missed-site = silent bug).

The gfx942 builder repeats the same runtime-layout window predicate directly at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:2678 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:5376 instead of using one named predicate as gfx950 does.

Fix: Define a named SW_ACTIVE = SLIDING_WINDOW > 0 or RTL in the gfx942 builder and replace all repeated SLIDING_WINDOW > 0 or RTL sites with it.

if spec.tile_size is not None and spec.tile_size != spec.block_size:
return (
"runtime_layout requires tile_size == block_size "
"(one paged block per KV tile)"

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

[MEDIUM] §L.4 Wave size is a property of the exact target (ArchTarget.wave_size), never a default or a flag flipped across targets — gfx1250 is CDNA at wave32. Lane math derives from the target.

Previous NEEDS-INFO is not resolved: new runtime-layout lane math hard-codes wave lanes as threads = spec.num_warps * 64 instead of deriving the wave size from the exact target at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx950/attention_tiled_2d.py:1133.

Fix: Replace hard-coded wave-lane constants in the new runtime-layout geometry and launch logic with the exact target wave size, and mirror that in the C++ runtime-layout/pack code.

#: signatures, and calling a mismatched library with them is memory-unsafe rather
#: than merely wrong.
BINARY_ABI = 12
BINARY_ABI = 13

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

[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 previous review: The PR touches dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h:116 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/portable_ir/src/abi.py:84, but the PR description still does not state that PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py were run 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 the results in the PR description.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

7 participants