feat(hip-kernel-provider): FlyDSL SDPA forward parity with AOTriton, plus vLLM forms by BradPepersAMD · Pull Request #15066 · ROCm/rocm-libraries

BradPepersAMD · 2026-11-04T14:00:50Z

JIRA ID : ALMIOPEN-3144

Summary

This PR extends the FlyDSL SDPA forward engines from #14111 so that they match all of AOTriton's forward features on gfx11, gfx12, gfx90a, gfx942 and gfx950. Every feature was verified by running AOTriton's own forward test grid through hipDNN with the AOTriton test harness in dnn-benchmarking; that grid ran before the bf16 P rounding and packed dropout, which are covered by their own device tests and A/B timing. It is one commit on develop and is forward only; backward is separate work.

Features

Variable length: packed THD sequences, per-batch padding lengths and seqused_k (ragged Q with seq_len_kv alone) are served on every family.
Ragged tensors at per-batch origins (feat(hipdnn): ragged tensors with per-batch origins, lengths and strides (ALMIOPEN-3151) #15080): packed sequences follow hipDNN's ragged rules, which are cuDNN's: with per-batch lengths each sequence is its length's tokens from its own table entry, in any order and with gaps (Q/K/V on RDNA prefill and CDNA, and K under seqused_k on aiw). INT64 tables, other batch strides, element-unit offsets and head-major layouts are declined.
Dropout: Philox dropout is bit-compatible with AOTriton at every head-dim width, including with bias, padding and packed sequences (device-tested in 16-bit and f32), plus a mask-writer op for the RNG dump.
fp32 I/O: exact fp32 attention is served on CDNA and on the RDNA prefill engine.
Bias: additive bias is served at every head-dim tier, in f32 or the query dtype, vectorized or strided.
Attention sinks: cuDNN's FLOAT sink is served on every family.
Softmax scale: zero, negative and subnormal scales are served.
Head dims: head dims up to 512, Q/K ≠ V head dims, MLA 192/128 and odd head dims (through a padded-copy helper op) are served.
Windows: left-only windows on CDNA and the FlashAttention/vLLM window mapping (SEM-027) are served.
gfx950 head dim 256: it runs on a new dual-wave kernel, ported from AOTriton's gfx950 forward.
Integration bundle: a dense attention-sink quick-tier bundle and its tests are folded in from feat(hipdnn): validate SDPA forward attention sinks and support them in the CPU reference #14696 (Daming Feng, ALMIOPEN-3094), whose validation and CPU-reference sink support were already on develop via feat(hipdnn): SDPA CPU and GPU references cover the full SDPA graph API #14358; this PR supersedes feat(hipdnn): validate SDPA forward attention sinks and support them in the CPU reference #14696.

Defects fixed

Large bf16 scores: they no longer produce NaN, and the head-dim 512, scale 1.2 case now meets AOTriton's error bound, through an exact softmax exponent at the max key (both reported against feat(hip-kernel-provider): FlyDSL kernels: RMSNorm fwd/bwd + SDPA forward (RDNA, CDNA) #14111).
bf16 probabilities: every bf16 engine now rounds P to nearest before P·V instead of truncating it, which had biased O by about 2^-8.5 on one-signed V (AOTriton PR 266, In gfx1201's FlyDSL kernel, replace bf16_trunc_pack_v8 with bf16_rne_pack_v8 aotriton#266).
Hidden instruction hazard: a new checker, tools/check_trans_hazards.py, flags a transcendental result read by the next VALU instruction in the gfx942/gfx950 objects, the hazard an inline-asm convert hid from the compiler on gfx950.
gfx950 16-bit decode at huge logits: it no longer returns O = 0 when every score is very large, because it now rounds the row max down like the other CDNA kernels.
Packed lengths past a table slot: the RDNA prefill and CDNA generic kernels no longer clamp a per-batch length to offset[b + 1] - offset[b] or write a slot's rows past its length, which read out-of-order origins wrong and wrote tokens the tensor does not address.
Window-removed keys: FlyDSL kernels no longer read V at keys a window or seq_len_kv removes, so NaN/Inf in stale or nulled KV no longer reaches O.
Reference behavior: the CPU and GPU references now treat probability-0 dropout without an RNG dump as no dropout, and never read V at keys a window removes, matching the engines.
Frontend descriptor unpacking: it no longer drops a tensor's ragged offset, ragged multiplier or alignment, which had made packed graphs loaded from JSON dense.
GPU reference Q padding: the GPU reference no longer reads Q padding on padded query rows.
GPU reference fp8 dequantize: the fp8 dequantize pass no longer reads Q/K/V padding.
Tests that missed their target: three reference tests now fail when the code they name breaks (follow-ups from Adam Dickin's feat(hipdnn): SDPA CPU and GPU references cover the full SDPA graph API #14358 review).
rocKE golden-graph collector: it now skips sweep templates in any bundle directory, so rocke_library_sdpa_pytest no longer reads the new SinkToken sweep template as a graph.
ASM ragged path: the per-engine schema guard now lets develop's ASM ragged path (feat(hip-kernel-provider): Add wiring for ragged tensors to the ASM SDPA Engine #14202) through, and a new test checks its output against the CPU reference.

Validation

AOTriton test harness, through dnn-benchmarking on all five arches (run before the bf16 P rounding and packed dropout, which are covered by their own device tests and A/B timing):
- Sets: T0–T2, PERF, a 170-graph model-shape tier, and a stratified 105,656-graph sample of AOTriton's 462k-case forward grid per arch. The full grid also ran on gfx950.
- References: every result was checked against PyTorch/AOTriton and hipDNN's GPU reference, with the CPU reference on a subset.
- Outcome: every graph passes except two gfx942 out-of-memory errors caused by another job on the node, which pass on a dedicated GPU, and one known AOTriton 0.14.1 LSE deviation on gfx950.
Functional suites, on this commit: green on all five arches (build with 0 warnings; FlyDSL suite 0 failures). The only failures are 11 existing TestGfx950AttentionDense kpack tests, which fail the same way on develop.
Integration quick tier: every FlyDSL SDPA lane passes, including develop's padded-sequence bundles, which these engines now serve. Their padded rows have LSE = -inf on both sides, so the lanes compare LSE with allclose_matching_infinities on the host.
Test tiers: the broader RDNA object-form and sweep cases this PR adds run in the nightly comprehensive tier (Comprehensive*), so TheRock's Windows gfx110X test shard stays under its 30-minute limit; every case still runs nightly.
Determinism: every FlyDSL SDPA forward gives the same bits on every launch (no atomics; split partials merge in a fixed order), and new repeat-launch tests check it on every family.
Speed: model-shape geomean speedup vs AOTriton 0.14.1 is 1.58–2.61× by arch.

Risk and mitigation

Risk	Mitigation
A wrong answer from a changed kernel	AOTriton's grid ran on every arch, checked against PyTorch and both hipDNN references, before the bf16 P rounding and packed dropout, which device tests and the hazard checker cover. The new tests are mutation-checked.
A slowdown for graphs that don't use a new feature	Unchanged objects are byte-identical. Changed objects were A/B-timed against the previous ones, except gfx90a's changed bf16 objects, which were not A/B-timed. Four groups run slower and are deferred items in COVERAGE.md §7.5: gfx950 dual-wave, one gfx942 head-dim-256 tile, one gfx12 MLA 192/128 causal kernel, and item 27, gfx942 bf16 P rounding (+0.9% geomean, up to +7% on some objects, partly recovered with a per-object scheduler option).
Effect on code outside the FlyDSL engines	The backend, the flatbuffers, data and plugin SDKs, the ASM engine and the Python bindings are byte-identical to develop. SDPA cache keys and the plugin API are unchanged. The only frontend change is the unpack bug fix. The provider-wide schema guard (`ApplicabilityManager`, `GraphSchema.hpp`) declines a graph for every engine whose declared schema (built-ins 1.4.0, ingestor engines their UED `sdk_version`) is below the graph's floor. For ingestor engines this is the same decline develop's `GenericPlanBuilder::understandsGraph` already makes on the graphs hipDNN stamps. The gfx950 rocKE conv dgrad engine (#14822), whose matcher reads alignment, now declares 1.3.0 (UED and `make_conv_dgrad_aot_config.py`), so it serves 32/64-byte aligned graphs; a gfx950 test covers that.
Missing AOT objects in CI	The 1,912 SDPA objects (57 MiB; 64.8 MiB with the other FlyDSL ops) are on the DVC `ingestor` remote.
gfx1100 node instability	Three gfx1100 nodes failed during the sweep. A serialized replay of the failing window did not reproduce it, and the same graphs pass on another gfx1100 card type. It stays tracked as an open item.

Related PRs and follow-ups

feat(hipdnn): complete TensorAttributes metadata Python bindings #14332 (Nikhar Asthana): merged; it provides the Python ragged-tensor accessors, so this PR has none of its own, and its packed-vs-padded Python reference tests (test_sdpa_ragged.py) use them.
feat(hipdnn,hip-kernel-provider): add cu_seq_len_q/kv to SDPA forward (ALMIOPEN-2935) #14560 (Sergei Grigoriev): adds cuDNN 9.25's cu_seq_len_q/kv. Follow-up after it lands: FlyDSL engines and both references accept cu_seq_len_q as cuDNN's spelling of the ragged-Q + seq_len_kv form they already serve (SEM-028), so seqused_k moves to the cuDNN API.
feat(hip-kernel-provider): trained UHD heuristics for the FlyDSL engines #14831 follow-up: trained UHD heuristics for the FlyDSL SDPA engines, retrained on this catalog. feat(hip-kernel-provider): trained UHD heuristics for the FlyDSL engines #14831's SDPA models were trained on the previous catalog, so this PR removes them; until the retrained models land, each FlyDSL SDPA engine ranks its kernels by descriptor priority, and autotune measures the order. The RMSNorm fwd/bwd models stay (catalog unchanged).
Other follow-ups, each in its own PR: paged KV in the engines, softcap end to end, and 16-bit attention sinks.

🤖 Generated with Claude Code

therock-pr-bot · 2026-11-04T14:06:39Z

✅ All Checks Passed — Ready for Review

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

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

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

🙋 Wish to Override Policy?

therock-pr-bot · 2026-11-04T14:06:42Z

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

…plus vLLM forms Extends the FlyDSL SDPA forward engines of #14111 to match all of AOTriton's forward features, verified on AOTriton's forward test suite (forward only; backward is separate work), and serves the forms vLLM needs. Families: RDNA3 (gfx11-generic), RDNA4 (gfx12-generic) and CDNA (gfx90a, gfx942, gfx950). Sits directly on develop, which carries #14111 (FlyDSL core), #14358 (SDPA references) and #14332 (TensorAttributes Python bindings); this commit adds only the parity work on top of them. Features (all families unless noted): - Variable length: padding by per-batch lengths (padded rows O = 0, LSE = -inf); packed THD via ragged offsets; packed + both lengths (PyTorch seqused_k); ragged Q with seq_len_kv alone (SEM-028), which develop's frontend and references already accept (cuDNN 9.25 spells it cu_seq_len_q + seq_len_kv; hipDNN maps that once #14560 lands; no API change here). - Dropout: Philox4x32-10 bit-compatible with AOTriton (0.14.2b/0.15 word order), the rng_dump mask writer op, dropout with bias and under padding, at every AOTriton head-dim width on RDNA and CDNA. - Packed + Philox dropout (PyTorch's varlen flash path) on CDNA, gfx11 and flydsl_sdpa_aiw, compact or with per-batch lengths where the family serves them: each sequence draws AOTriton's varlen counters (its index, positions within it, the graph's declared extents), so dense, padded and packed storage of one problem draw one mask, and the rng_dump stays the dense FLOAT [B, Hq, Sq, Skv]. Matchers only, no object changes; packed with a bias still declines. SdpaSemanticsRegister SEM-015 records the mask layout. - fp32 I/O: exact fp32 on CDNA (with bias, dropout, any scale, head dims 136-248) and exact IEEE fp32 on the RDNA prefill engine (VALU path). - Bias at every head-dim tier, f32 or query dtype, vectorized and strided (any layout incl. key-broadcast), padded bias on the vectorized objects. - Attention sinks in FLOAT, cuDNN's dtype (SEM-022): every FlyDSL family now reads the FLOAT sink the frontend accepts (it read only a query-dtype sink before); any other sink dtype is declined. - Zero, negative and subnormal softmax scales. - Head dims 264-512, Q/K != V head-dim pairs, MLA 192/128, Hk != Hv on CDNA. - Head dims that are not a multiple of 8 (any row pitch) through padded copies (new sdpa_head_pad helper op), which like the engines declines K off Q's head dim and O off V's; multiple-of-8 graphs are unchanged. - Windows/bands: left-only windows on CDNA; FA/vLLM window mapping (SEM-027). No FlyDSL kernel reads V at keys no row of its block sees (window, band, past seq_len_kv), so NaN/Inf in nulled or stale KV no longer reaches O. - Exact softmax exponent at the max key (fixes two #14111 findings: NaN at large bf16 scores, and a 2^-8 output bias at e.g. head_dim 512, scale 1.2): CDNA generic/decode round the max down and clamp P in v_exp_f32 with DX10_CLAMP off (NaN rows stay NaN); CDNA dual-wave and RDNA take s - m first; flydsl_sdpa_aiw no longer contracts the scale into the exponent. - bf16 P rounded to nearest before the second product in every bf16 FlyDSL SDPA forward (truncation biased O about 2^-8.5 low wherever V does not have zero mean, AOTriton's lesson 16): CDNA generic and decode, bias and dropout builds included (gfx950 v_cvt_pk_bf16_f32; gfx942/gfx90a AOTriton's bf16_round_nonan_pack_v8), and flydsl_sdpa_aiw; the dual-wave and RDNA packs already rounded. The gfx942 bf16 generic objects cost +0.9% geomean (COVERAGE 7.5, deferred item 27). - gfx950 16-bit decode takes the rounded-down max like every other plain build: rounded to nearest, c * m could round up where |c * m| passes 2^31 (Q = K = V = 132,096), every exponent of the row fell below -126, and O was 0. - tools/check_trans_hazards.py (REGEN.md step 4) flags a transcendental result read by the next VALU instruction in the gfx942/gfx950 objects (the hazard an inline-asm convert hid from the backend's hazard recognizer). - gfx950 decode: the 16-bit decode builds pack bf16 P rounded to nearest, so the exact exponent no longer serializes their K loads (Llama-style long-cache decode was 14-22% slower; now within 0.1% geomean of before, at most +1.7%), and the gfx950 DMA V path reads a causal block's V window through a per-workgroup buffer descriptor instead of clamping every row (decode head_dim 128 causal +2.3-3.3% -> +0.2-1.4%). - Performance: CDNA head_dim 256 tiles with K chunking on gfx942/gfx90a, and gfx950 head_dim 256 on the dual-wave kernel (4 waves, ported from AOTriton's gfx950 FlyDSL forward). - Determinism: the FlyDSL SDPA forward kernels are deterministic by construction (no atomics; fixed-order split-KV combine; per-wave LDS partials); the split count follows the device's CU count. Documented in COVERAGE 5.7; no determinism tests yet. Frontend and references: - p = 0 dropout without an RNG dump (and with fp8 scaling) is no dropout in the references. - The CPU and GPU references never read V at keys a window or block mask removes (NaN/Inf in nulled or stale KV blocks no longer leaks into O). - GPU reference: a padded query row exits before the softmax max, and the fp8 dequantize pass skips Q/K/V padding, so padding is never read (SEM-008). Tests now exercise the native ragged plans (CPU BF16 densify comparison, GPU executor routing) and a paged fp8 K/V case with scalar descales (follow-ups from Adam Dickin's #14358 review). - Fix: frontend descriptor unpacking keeps ragged offsets, the ragged multiplier and byte alignment (graphs loaded from JSON and re-lowered were made dense). It asks for the multiplier only on a ragged tensor, and reads an extension a backend of the same major version does not serve (NOT_SUPPORTED, or BAD_PARAM from a therock-10.0 backend) as unset, so dense graphs still deserialize there. - Python: test_sdpa_ragged.py builds a packed THD SDPA graph through #14332's ragged-offset accessors, checks what validation and JSON/binary serialization keep, and compares it on the CPU reference with its padded equivalent and a float64 oracle (that case skips without HIPDNN_REF_CAPI_LIB). - Folds in #14696 (Daming Feng, ALMIOPEN-3094): dense attention-sink quick-tier bundle (quick/SdpaFwd/SinkToken: FLOAT [1, H_q, 1, 1] sink, GQA 8:2, causal, stats) and tests (an untyped sink rejected through Graph::validate, sink lowering, dense CPU plan execution with a sink and a dangling sink uid, CPU-reference sink rescaling per query head, -inf sink, wrong-shape sink); its validation and CPU-reference sink support were already on develop via #14358. Provider: - A per-engine graph-schema guard under engine plugin API 1.4.0: each engine sees only graphs up to the schema it declares. The built-in engines (HIP_MLOPS, HIP_FLASH2, ASM_SDPA) declare 1.4.0, which #14202 reviewed them against, so ASM's ragged (group-mode) path still reaches it; the FlyDSL engines declare 1.4.0 through their UEDs; every other engine stays at the provider's 1.2.0 baseline. The FlyDSL engines decline override-shape graphs. The guard holds engines back one by one inside the threaded applicability path, so the engines it admits keep their pool offload. - ASM_SDPA forward: its group-mode kernels are now checked against the CPU reference (THD, head dims 128 and 192/128, no mask and bottom-right causal per sequence); test only, the engine is unchanged. - FlyDSL SDPA integration lanes (test configuration only): the generated lane TOMLs grade the LSE/stats output with allclose_matching_infinities, since a padded row's LSE is -inf in the engine and the reference alike, and the lanes compare on the host (HIPDNN_TEST_VALIDATOR=cpu), which that validator needs; O keeps the default validator. Support claims for the FlyDSL engines on the SinkToken and padneq quick bundles (gfx942, gfx1201). Validation: AOTriton's forward test grid run through hipDNN with dnn-benchmarking, checked against PyTorch/AOTriton and hipDNN's GPU and CPU references, on gfx90a, gfx942, gfx950, gfx1201 and gfx1100 (details and results in the PR). Every existing #14111 object either changed only through a listed kernel modification or stayed byte-identical; see flydsl/COVERAGE.md and the vendored-kernel modification lists in each kernel README. AOT objects (DVC remote ingestor): SDPA objects 688 -> 1,912 across the five targets (incl. the dropout-mask and head-pad helpers); FlyDSL object dirs 25.6 -> 64.8 MiB. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

codecov · 2026-11-05T07:42:20Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #15066      +/-   ##
===========================================
+ Coverage    78.08%   78.10%   +0.02%     
===========================================
  Files         2977     2978       +1     
  Lines       508809   509371     +562     
  Branches     79009    79102      +93     
===========================================
+ Hits        397280   397813     +533     
- Misses       84394    84411      +17     
- Partials     27135    27147      +12

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`53.89% <ø> (+0.01%)`	⬆️	Carriedforward from b707fcf
TensileLite-Unit	`85.08% <ø> (+<0.01%)`	⬆️	Carriedforward from b707fcf
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from b707fcf
hipBLASLt	`70.07% <ø> (+0.02%)`	⬆️	Carriedforward from b707fcf
hipCUB	`82.68% <ø> (ø)`		Carriedforward from b707fcf
hipDNN	`87.40% <ø> (+0.06%)`	⬆️
hipFFT	`42.65% <ø> (ø)`		Carriedforward from b707fcf
hipRAND	`76.12% <ø> (ø)`		Carriedforward from b707fcf
hipSOLVER	`69.67% <ø> (ø)`		Carriedforward from b707fcf
hipSPARSE	`87.03% <ø> (ø)`		Carriedforward from b707fcf
rocALUTION	`86.47% <ø> (ø)`		Carriedforward from b707fcf
rocBLAS	`69.20% <ø> (ø)`		Carriedforward from b707fcf
rocFFT	`46.49% <ø> (ø)`		Carriedforward from b707fcf
rocRAND	`57.32% <ø> (ø)`		Carriedforward from b707fcf
rocSOLVER	`76.66% <ø> (ø)`		Carriedforward from b707fcf
rocSPARSE	`83.61% <ø> (ø)`		Carriedforward from b707fcf
rocThrust	`93.15% <ø> (ø)`		Carriedforward from b707fcf

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines	Coverage Δ
...hipdnn_frontend/detail/DescriptorUnpackHelpers.hpp	`86.05% <ø> (-3.87%)`	⬇️
...e/hipdnn_test_sdk/utilities/CpuFpReferenceSdpa.hpp	`85.32% <ø> (+0.83%)`	⬆️
...nn_test_sdk/utilities/CpuFpReferenceSdpaRagged.hpp	`94.42% <ø> (+0.10%)`	⬆️
...test_sdk/utilities/RaggedSdpaDensifyTestGraphs.hpp	`94.78% <ø> (+0.36%)`	⬆️
...de/hipdnn_test_sdk/utilities/SdpaFwdTestGraphs.hpp	`88.57% <ø> (+1.61%)`	⬆️
...pdnn_test_sdk/utilities/SdpaVllmFormTestGraphs.hpp	`96.90% <ø> (ø)`
.../utilities/cpu_graph_executor/detail/PlanUtils.hpp	`82.43% <ø> (+0.79%)`	⬆️

... and 13 files with indirect coverage changes

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

- Build the frontend lane's arch list from a const char* so clang-tidy's readability-redundant-string-init passes when no FlyDSL content is built. - Skip *.template.json in the rocKE SDPA golden-graph collector (pytest and CMake), so the SinkToken sweep template is not read as a graph. - Device-test packed THD sequences with f32 dropout on CDNA and RDNA, and accept them in the f32 host tables; the RDNA skip messages now say 16-bit. - COVERAGE.md: describe how each family rounds bf16 P, and drop the rows that still listed packed dropout as declined; remove leftover plan labels. - aiw README: note that the now-unused pack helpers stay for upstream parity. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

… every launch The FlyDSL SDPA forwards are deterministic by construction: no shipped object (gfx90a, gfx942, gfx950, gfx11-generic, gfx12-generic) has an atomic instruction, and a split decode's merge reads each split's partial from its own workspace slot in split order. Nothing tested it. These tests do, and they catch the two ways a forward stops being deterministic: a race, and a read of memory the launch did not write. - The CDNA and RDNA device harnesses take a launch count. Each further launch of the same object runs on the same inputs, with its O, its LSE and its workspace filled first with another byte (NaN, zero, then a large finite value), and must give the first launch's O and LSE bit for bit. A packed launch refills only the token rows, so the guard rows still hold the NaN the packed check reads. - TestGpuFlydslSdpaCdnaDispatch.GivesTheSameBitsOnEveryLaunch: every object the graph admits, at the model shape (dual-wave objects where they ship), f16 at head_dim 256 under a causal mask with the LSE, a single-token decode over 20,000 keys (split, merged from the workspace), and packed sequences. - TestGpuFlydslSdpaRdnaDispatch.GivesTheSameBitsOnEveryLaunch, TestGpuFlydslSdpaRdnaDispatch.DecodesTheSameBitsOnEveryLaunch and TestGpuFlydslSdpaAiwDispatch.GivesTheSameBitsOnEveryLaunch: bf16 at the default shape, f16 causal with the LSE, a split decode over 2,100 keys, and packed sequences on the aiw objects. Mutation-checked on gfx950, gfx1100 and gfx1201: a merge launched on a stream of its own, unordered after the split launch, and a merge that reads its partials from workspace no launch writes, each fail the new tests. The race passes a single launch against the reference; the repeated launches catch it. The intermittent ServesHeadDim256OnTheDualwaveKernel failure on one MI300X GPU is that GPU, not the kernels. On it, a few percent of launches of some f16 causal objects at head_dim 256 change O and the LSE in rows 12 and 28 (mod 32) of query heads 4 and 20 only, by up to 0.023 (the f16 bound is 0.004). The same binaries and inputs give identical bits on every launch on two other MI300X GPUs, on gfx950 and on gfx90a. On a healthy GPU those objects' worst error over 20 input draws is 0.00039 to 0.00041, a tenth of the bound. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…rehensive tier The Windows gfx110X hipkernelprovider standard shards run at 28-30 min against a 30-min step limit on develop, and this branch adds ~1-1.5 min per shard, mostly gfx11 time in hip_kernel_provider_tests and hip_kernel_provider_integration_tests. - HIP_MLOPS_ENGINE_test_categories.yaml: a Comprehensive* GTest instantiation prefix, excluded from quick, standard and ffm-quick and run by comprehensive (nightly: Linux gfx94X, Windows gfx1151) and full, as integration-tests' test_categories.yaml does. - The frontend lane's DualwaveHeadDim256 cases (gfx950's dual-wave kernel, ~30 s on a gfx11 runner) run in every tier on gfx950, the architecture they are pinned to, and on other GPUs from ComprehensiveDualwaveHeadDim256. - TestGpuFlydslSdpaRdnaDispatch.DropsThroughABias keeps head_dim 128 and DropsAtTheGenericTiersWidths 48 and 224; the other widths run in Comprehensive/TestGpuFlydslSdpaRdnaComprehensiveTier. TestGpuFlydslSdpaObjects still launches every dropout object form in every tier. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…comprehensive tier The first trim leaves this branch ~60 s over develop on the Windows gfx110X hipkernelprovider standard shard 1, which develop alone runs at 28-29 min of its 30-min limit. Moved to Comprehensive* instantiations (comprehensive and full tiers), every case still run there: - TestGpuFlydslSdpaObjects: OnePerForm leaves the RDNA prefill and decode forms isNightlyForm names (a bias with a sink, a generic tier or a V head dim of its own, or on prefill a split V; exact f32; V-mask pairs; the generic decode tiers split four ways) to Comprehensive/...OnePerNightlyForm. OnePerForm keeps the other 44 forms per engine, develop's set. - TestGpuFlydslSdpaRdnaDispatch.DecodesZeroAndNegativeScalesAcrossSplits, DecodesWithinAotritonsBoundAtScale1Point2 and AppliesALeftOnlyWindowAsTheCpuReferenceDoes run as Comprehensive/TestGpuFlydslSdpaRdnaComprehensiveTier; their prefill and non-causal siblings stay in standard. - The frontend lane's AotritonScale1Point2HeadDim512 cases run in every tier on CDNA (the decode objects they pin) and in the comprehensive tier elsewhere (ComprehensiveAotriton). COVERAGE.md lists the cases the comprehensive tier runs and standard does not. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

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

The trained UHD models for the FlyDSL engines were fit to the previous FlyDSL catalog. This catalog has different objects and kernel fields, so the models are removed and every FlyDSL UED binds only its default slots: the native ranker orders each engine's kernels by descriptor priority, the engine estimate stays the disabled default slot, and autotune measures the order. Models retrained on this catalog are a follow-up; gen_descriptors.py binds them per arch once installed under heuristics/. The generator's model binding, its tests and the UHD infrastructure are unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…igins hipDNN's ragged tensors now follow cuDNN's per-batch rules (RFC 0014, SEM-026): with per-batch lengths, sequence b is seq_len[b] tokens from offset[b], the tables may list the origins in any order and leave gaps, and their last entry need not bound the last sequence. The RDNA prefill kernel (vendored modification 37) and the CDNA generic kernel (modification 47) clamped each length to offset[b + 1] - offset[b] and filled a slot's rows past its length. They now read each sequence at its origin for its own length and write no token outside it. A new launch flag keeps the query counts from Q's table for seq_len_kv alone (PyTorch's seqused_k). The aiw kernel already read K at per-batch origins for its used lengths. Tests: device cases on every packed family with origins out of order, gaps and a last table entry that bounds nothing, run against the per-sequence oracle and, through the frontend lane, hipDNN's CPU reference graph executor; the harnesses check every token no sequence owns stays unwritten. Decline cases for lengths and tables at a batch stride of 2 and head-major packing with lengths, beside the existing INT64, element-unit and head-major declines. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…r-batch origins The RDNA prefill (gfx11-generic, gfx12-generic) and CDNA generic prefill (gfx90a, gfx942, gfx950) objects rebuilt from the kernels that read packed sequences at per-batch origins for their own lengths. The RDNA decode, CDNA decode and dual-wave objects are byte-identical; only their manifests' input digests move. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…dsl-aotriton-parity SUPPORT_MATRIX_OVERVIEW.md takes develop's copy: since #15293 the overview is refreshed in a PR of its own, not with claim changes. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

cgmillette · 2026-11-07T00:14:39Z

Automated overlap notice (rocKE merge monitor)

Hi @BradPepersAMD,

This PR and #14908 (AnaghaRaoAMD) add the same two pieces. The recorded owner of both is 14908: it lands first in the merge plan, and for the helper, the shared integration-test header IngestorIntegrationSupport.hpp is the natural home.

Piece	This PR at `c3502de66b`	14908 at `3b5cd4de6c` (owner)
SDPA tensor helper	`bshdStrides` at `FlydslSdpaFrontendGraphs.hpp:236`, `bshdTensor` at `:242`	`sdpaStrides` at `SdpaLayoutTestSupport.hpp:40`, `makeSdpaTensor` at `IngestorIntegrationSupport.hpp:539`
`Comprehensive*` tier in `HIP_MLOPS_ENGINE_test_categories.yaml`	`58b054a`: tier comment `:16-27`, comprehensive description `:62`, excludes `:41`, `:54`, `:89`	`00846da`: tier comment `:16-26`, comprehensive description `:61`, excludes `:40`, `:53`, `:88`

Paths are under dnn-providers/hip-kernel-provider/: the two integration headers in src/integration_tests/kernel_ingestor_engine/, and SdpaLayoutTestSupport.hpp (new in 14908) in src/tests/engines/kernel_ingestor_engine/packs/.

Helper: bshdTensor is makeSdpaTensor with its last parameter fixed to SdpaLayout::BSHD; the other seven parameters and the body match. bshdStrides is the BSHD case of sdpaStrides (SdpaLayoutTestSupport.hpp:60-61).
Tier: the excludes and the quick and standard descriptions are the same lines; the tier comment and the comprehensive description are worded differently, so whichever PR lands second conflicts there. Our review of 14908 asked it to take your comment wording (the integration-tests precedent, no suite list) in place of its "matching miopen-provider's convention" line.

Resolution: once 14908 lands, this PR rebases, drops its copies, and reuses makeSdpaTensor / sdpaStrides and 14908's tier change.

Landing order: the merge plan has 14908 first (wave 2, a cgmillette priority); this PR is not ready to merge yet. If this PR lands first instead, it needs nothing more: 14908 keeps the shared-header helper and the tier comment, and replaces this PR's copies in its own rebase.

Rebase steps, once 14908 is in develop:

Rebase on develop.
HIP_MLOPS_ENGINE_test_categories.yaml: keep develop's version (14908's tier change) and drop your hunk.
FlydslSdpaFrontendGraphs.hpp: delete bshdStrides and bshdTensor (:235-257), add #include "IngestorIntegrationSupport.hpp", and call detail::makeSdpaTensor(…, SdpaLayout::BSHD) at the three bshdTensor calls (:466, :468, :471) and sdpaStrides(SdpaLayout::BSHD, …) at the bshdStrides call (:572).
Build hip_kernel_provider_integration_tests and rerun the tests that include the header: IntegrationGpuFlydslSdpaFrontend.cpp, IntegrationGpuFlydslSdpaFrontendFeatures.cpp and asm_sdpa_engine/IntegrationGpuSdpaForwardRagged.cpp.

Checked by reading both heads; nothing was built. If you think this PR should own either piece instead, reply here and on 14908, and the reviews will follow the new decision.

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

…dsl-aotriton-parity Conflicts: - HIP_MLOPS_ENGINE_test_categories.yaml: took develop's (#14908 owns the Comprehensive* tier); quick/standard/ffm-quick still exclude Comprehensive*. - src/integration_tests/CMakeLists.txt: kept develop's attention knob-suite table (it absorbs the Unified2D/3D and gfx942 unified 3D blocks) plus the FlyDSL frontend-lane sources and HIPKERNELPROVIDER_FLYDSL_SDPA_ARCHS. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

FlydslSdpaFrontendGraphs.hpp drops its bshdStrides/bshdTensor copies and builds Q, K, V with detail::makeSdpaTensor(..., SdpaLayout::BSHD) and O's strides with sdpaStrides(SdpaLayout::BSHD, ...) from IngestorIntegrationSupport.hpp, which #14908 owns. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

cgmillette · 2026-11-07T06:49:54Z

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

PR #15066 — FlyDSL SDPA forward parity with AOTriton, plus vLLM forms (BradPepersAMD, ALMIOPEN-3144) — 🟠 First-round: conflicts, copies, dropped models

Prior item	Status
cgmillette (11-07) overlap notice: reuse 14908's SDPA helpers	🟠 OPEN • `bshdTensor`, `bshdStrides` still in this PR (ran) • 14908 has landed; see N2
Bots / CI	🟡 Running • running: Linux gfx94X hipkernelprovider shards 1/2, 2/2; miopenprovider passed • Windows gfx110X hipkernelprovider, hipdnn, integration ✓; 50 ✓ • merge: conflicting; CI's merge ref predates 14908
Our checks	🟠 Conflicts • merge vs `1702865a2f`: 2 files conflict (ran) • resolved merge: rocke SDPA-from-hipDNN tests 194 ✓ / 1 skip (ran) • C++ build and ctest not run here; no ABI change (read) • rules 1–3: nothing in `rocke/` (read)
Code quality	🟠 Should-fix • keeps its own copy of 14908's tensor helper: N2 • two new copies of develop helpers: N4, N5 • helper-op engines, Philox plumbing repeated per variant (minor)

#	Sev	file:line	Finding
N1	🟠	`dnn-providers/hip-kernel-provider/src/integration_tests/CMakeLists.txt:123`	Conflicts with develop in 2 files (also `HIP_MLOPS_ENGINE_test_categories.yaml:16`) • impact: no merge; 14908 turned the knob suites into `_knob_suites` • fix: merge develop; keep its table, add the FlyDSL block after it
N2	🟠	`dnn-providers/hip-kernel-provider/src/integration_tests/kernel_ingestor_engine/FlydslSdpaFrontendGraphs.hpp:248`	Second home for the SDPA tensor builder (`bshdTensor`, `bshdStrides` at `:242`) • impact: same body as develop's `makeSdpaTensor` (`IngestorIntegrationSupport.hpp:539`) • owner: 14908, as recorded (landed first; shared header) • fix: include that header; `makeSdpaTensor(..., SdpaLayout::BSHD)`; drop both copies
N3	🟠	`dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/FlyDSL/rmsnorm/flydsl_rmsnorm.ued.json:13`	Deletes RMSNorm fwd/bwd trained UHD models too (also `rmsnorm_bwd`; `db915abec2`) • impact: RMSNorm falls back to priority order on 6 arches; its catalog is unchanged • fix: keep the RMSNorm models and slots; drop only the SDPA ones
N4	🟠	`dnn-providers/hip-kernel-provider/flydsl/generators/gen_sdpa_head_pad.py:85`	New `_verify` re-implements develop's `verify_layout` (also `gen_sdpa_dropout_mask.py:73`) • impact: the same three layout checks in three places • fix: call `verify_layout`; move it beside `verify_arch` in `_codeobject.py`
N5	🟠	`dnn-providers/hip-kernel-provider/flydsl/kernels_src/kernels/dropout_mask/sdpa_dropout_mask.py:44`	Two new copies of the null-safe Philox u64 load (also `flash_attn_func_rdna.py:457`) • impact: three homes beside develop's `fmha_common_gfx1201.py` one • fix: one copy in `attention_rdna4/philox.py`, imported by all three
N6	🟡	`dnn-providers/hip-kernel-provider/src/core/ApplicabilityManager.cpp:133`	New schema guard drops develop's gfx950 dgrad on aligned graphs • impact: dgrad (UED 1.0.0) serves 32/64-byte alignment; floor is 1.3.0 • fix: dgrad UED `sdk_version` 1.3.0; name the guard in the Risk row

Utility: 🟢 Positive — broad, well-tested SDPA forward coverage; conflicts and copies remain

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

First-round review: first pass by qwen3.6 (local), verified by the orchestrator (Claude Opus 5.5). Checks: merge with develop, tests on a resolved merge, CI shard logs, overlap, rules 1–3, saved probes, code quality. Flags whether a closer look is needed; not an approval or rejection.

The UHD removal took the RMSNorm fwd/bwd trained models and their UED slots along with the SDPA ones. Only the SDPA catalog changed, so restore the RMSNorm heuristics (heuristic.uhd.json + model.bin.dvc per arch) and UED bindings exactly as develop has them; gen_descriptors regenerates them unchanged. COVERAGE §6 now says only the SDPA engines lost their models. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

… the newer fields The gfx950 rocKE conv dgrad engine (#14822) gates tensor alignment (multiples of 16 bytes), declines ragged, pass-by-value and override-shape graphs, but its UED claimed graph schema 1.0.0. Both the provider's schema guard and the SDK plan builder's understandsGraph therefore declined every graph declaring 32- or 64-byte alignment (floor 1.3.0), which the matcher serves. Declare 1.3.0 at the source (make_conv_dgrad_aot_config.py) and in the shipped UED, update the generated census expectation, and add the engine to TestContainer's opt-in list. Tests: the graph matcher accepts 32/64-byte alignment; on gfx950 an aligned dgrad graph built through the frontend gets a plan on this engine and matches the CPU reference. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…enerators gen_sdpa_head_pad.py and gen_sdpa_dropout_mask.py each carried a _verify that repeated develop's verify_layout (signature, kernarg segment, no VGPR spill). Call verify_arch and the shared verify_layout from _sdpa_rdna_build.py instead. It stays there, not in _codeobject.py: every op's generators import _codeobject.py, so moving it would restamp every FlyDSL op's inputs_sha256 and need a full regeneration. The two generators are inputs of their ops, so sdpa_head_pad and sdpa_dropout_mask were regenerated for all five targets: the ten objects are byte-identical, only each manifest's inputs_sha256 (and its SOURCE.md and .dvc pointer) changed. test_inputs_digest now lists the two helper ops among _sdpa_rdna_build.py's readers. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

cgmillette · 2026-11-07T08:16:03Z

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

PR #15066 — FlyDSL SDPA forward parity with AOTriton, plus vLLM forms (BradPepersAMD, ALMIOPEN-3144) — 🟠 Follow-up: CI update, shard timeout

Prior item	Status
Our 11-07 first round, N1–N6	🟠 OPEN • same head `b707fcf487`: nothing changed
Bots / CI	🟠 Red, from this PR • Linux gfx94X hipkernelprovider shard 1/2 hit the step time limit in both attempts, before its last tests ran (read: logs) • shard 2/2, miopenprovider, Windows gfx110X ✓; 52 ✓ • the other shared tests match develop's times (read: develop's shard log)
Code quality	🟠 Should-fix • unchanged: N2, N4, N5

#	Sev	file:line	Finding
N7	🟠	`hip_kernel_provider_tests_standard_suite`	The new FlyDSL GPU cases push gfx94X shard 1/2 past CI's step time limit • impact: `TestGpuFlydslSdpaCdnaDispatch` and `TestGpuFlydslSdpaObjects` make the standard suite run several minutes longer than on develop, on a shard develop already runs near the limit (read: logs) • impact: `ServesHeadDim256OnTheDualwaveKernel` is the slowest new case • fix: move the slow cases to the comprehensive tier, or shrink them

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

Follow-up review (CI update): no new commits since the first-round review; the orchestrator (Claude Opus 5.5) compared this head's gfx94X shard logs with develop's. Flags whether a closer look is needed; not an approval or rejection.

…ehensive tier TheRock's Linux gfx94X hipkernelprovider shard 1/2 runs the standard tier under a 30-minute step limit. Develop alone takes 1676-1737 s there, and this branch adds about 220 s to hip_kernel_provider_tests_standard_suite on that shard, so the shard times out. Every case below still runs nightly, in Comprehensive* (the gfx94X nightly runs the comprehensive tier). The standard gtest list keeps the same names, so the gtest shard assignment does not change. - ServesHeadDim256OnTheDualwaveKernel runs pre-merge only where the build ships the dual-wave kernel for the device (gfx950). Elsewhere it only exercises the generic objects (about 115 s on gfx942), and ServesHeadDim256WhereShipped already covers head_dim 256 there pre-merge. The comprehensive copy runs it on the other archs. - TestGpuFlydslSdpaObjects OnePerForm, CDNA generic engine: isNightlyCdnaForm sends the f16 and causal builds of the V-mask, bias-with-sink, bucketed-bias, strided-f32-bias and split-V forms to OnePerNightlyForm. It also sends split V beside a bias, bucket or V mask, and exact f32 beyond its plain non-causal form. That leaves 31 of 88 forms per target pre-merge: develop's 22 plus the bf16 non-causal build of each feature. Checked offline against the gfx942/gfx950/gfx90a kdp metadata. - ComputesF32WithABiasDropoutAndAZeroScale runs pre-merge at head_dim 192 (the column-split bucket). 64, 128 and 96 move to the comprehensive tier. - HoldsAotritonsBoundAtScale1Point2 runs pre-merge at 64 and 256; 512 moves to the comprehensive tier. The frontend lane's AotritonScale1Point2HeadDim512 still runs on CDNA in every tier. - KeepsValuesAtKeysNoRowSeesOutOfO runs pre-merge with NaN plus the f32 cases; infinity and -infinity move to the comprehensive tier. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

cgmillette · 2026-11-07T12:19:33Z

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

PR #15066 — FlyDSL SDPA forward parity with AOTriton, plus vLLM forms (BradPepersAMD, ALMIOPEN-3144) — 🟠 Follow-up: fixes landed, one copy left

Prior item	Status
Our N1 (11-07) 2 files conflict with develop	🟢 FIXED `c88866208a` • merge vs `1702865a2f`: 0 conflicts (ran) • keeps develop's knob-suite table, adds the FlyDSL block (read)
Our N2 + overlap notice (11-07) reuse 14908's SDPA tensor helper	🟢 FIXED `7a72bae8fa` • `bshdTensor`, `bshdStrides` gone; calls `makeSdpaTensor(..., SdpaLayout::BSHD)` (ran: probe)
Our N3 (11-07) RMSNorm trained models deleted	🟢 FIXED `04ba10ea7c` • RMSNorm fwd/bwd descriptors now match develop's file for file (ran)
Our N4 (11-07) generators' `_verify` repeats `verify_layout`	🟢 FIXED `3b7748c5e9` • both helper-op generators call `verify_arch` + `verify_layout` (ran: probe)
Our N5 (11-07) third copy of `_load_u64_or_zero`	🟠 OPEN • unchanged: `sdpa_dropout_mask.py:44`, `flash_attn_func_rdna.py:457` beside develop's `fmha_common_gfx1201.py:1348` (ran: probe) • fix: one copy, imported by all three kernels
Our N6 (11-07) schema guard drops gfx950 conv dgrad	🟢 FIXED `b04ff8a6ff` • dgrad UED and its config generator declare 1.3.0 (read) • new matcher test serves 32/64-byte alignment; body's Risk row names the guard (read)
Our N7 (11-07) gfx94X shard 1/2 past the step limit	🟢 FIXED `ce8eb50b69` • slowest CDNA cases now `Comprehensive/`, out of the standard tier (read) • correction: shard time depends on the runner pool (15334): about 15 minutes on `linux-gfx942-1gpu-ccs-ossci-rocm`, 28–30 on `linux-gfx942-1gpu-ccs-csp-ossci-rocm`, where develop also times out; both our timeouts ran there, so the PR's few added minutes needed that pool too (read) • this head's gfx94X shards have not run yet
Bots / CI	🟡 Running • running: math-libs gfx94X, gfx950, gfx125X; Windows compiler-runtime; both hipDNN superbuilds • gfx94X hipkernelprovider shards wait on math-libs • 0 failed so far; the rest ✓ or skipped
Our checks	🟢 Clean • merge vs `1702865a2f`: clean, nothing reverted (ran) • build + ctest 60/60 (1 skip); differential golden unchanged (ran) • rocke companion tests only, none of the PR's own (ran) • `ROCKE_ABI_VERSION`: develop 12, no bump (read)
Code quality	🟠 Should-fix • fixes reuse the owners' helpers, test the dgrad claim: N2, N4, N6 • open: N5's null-safe load copies • minor: `verify_layout` stays in the RDNA build module; per-variant copies unchanged

Utility: 🟢 Positive — six of seven open rows fixed; one helper copy remains

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: merge with develop, build and ctest, differential golden, saved probes, CI runner pools, code quality. Flags whether a closer look is needed; not an approval or rejection.

cgmillette · 2026-11-07T16:43:42Z

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

PR #15066 — FlyDSL SDPA forward parity with AOTriton, plus vLLM forms (BradPepersAMD, ALMIOPEN-3144) — 🟠 Follow-up: CI update, shards at the limit

Prior item	Status
Our 11-07 N7: gfx94X shard past the step limit	🟡 PARTLY FIXED `ce8eb50b69` • the moved cases took out most of the extra time; the dispatch cases left still put each slower-pool shard over the limit (ran: per-case times vs develop's same shard and pool) • fix: fewer `TestGpuFlydslSdpaCdnaDispatch` / `TestGpuFlydslSdpaRdnaDispatch` cases in the standard tier, the rest in `Comprehensive/`
Our 11-07 N5: third copy of `_load_u64_or_zero`	🟠 OPEN • same head `ce8eb50b69`: nothing changed
Bots / CI	🟡 Red: step limit on the slower pools • gfx94X shards 1/2 and 2/2 (`ccs-csp-ossci` pool) and Windows gfx110X shard 1/2 stopped at the Test step limit; no test failed (read: logs) • develop's own shards finish just under that limit on these pools (15334); the suites this PR does not touch keep develop's pace (ran) • the faster `ccs-ossci` pool has room; the other checks ✓
Code quality	🟠 Should-fix • unchanged: N5

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

Follow-up review (CI update): no new commits since the last follow-up; the orchestrator (Claude Opus 5.5) compared this head's shard logs with develop's same shards on the same runner pools, case by case. Flags whether a closer look is needed; not an approval or rejection.

github-actions Bot added documentation project: hipdnn project: hip-kernel-provider project: hipdnn-integration-tests labels Nov 4, 2026

assistant-librarian Bot added the organization: ROCm label Nov 4, 2026

BradPepersAMD force-pushed the users/brpepers/flydsl-aotriton-parity branch from 71ba557 to ae4e94c Compare November 4, 2026 16:41

BradPepersAMD force-pushed the users/brpepers/flydsl-aotriton-parity branch from ae4e94c to dbf110a Compare November 5, 2026 02:15

github-actions Bot added the rocke: platform Touches rocKE platform label Nov 5, 2026

BradPepersAMD and others added 3 commits November 5, 2026 09:29

BradPepersAMD marked this pull request as ready for review November 6, 2026 20:30

BradPepersAMD requested review from a team as code owners November 6, 2026 20:30

BradPepersAMD and others added 5 commits November 6, 2026 15:00

Merge branch 'develop' into users/brpepers/flydsl-aotriton-parity

348597a

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

Merge remote-tracking branch 'origin/develop' into users/brpepers/fly…

b707fcf

…dsl-aotriton-parity SUPPORT_MATRIX_OVERVIEW.md takes develop's copy: since #15293 the overview is refreshed in a PR of its own, not with claim changes. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

BradPepersAMD and others added 2 commits November 6, 2026 20:03

BradPepersAMD and others added 2 commits November 7, 2026 00:29

cgmillette mentioned this pull request Nov 7, 2026

[rocke] Linux gfx94X Test hipkernelprovider shards run at the 30-minute Test step limit on develop #15334

Open

Navigation Menu

Repository navigation

feat(hip-kernel-provider): FlyDSL SDPA forward parity with AOTriton, plus vLLM forms - #15066

feat(hip-kernel-provider): FlyDSL SDPA forward parity with AOTriton, plus vLLM forms#15066
BradPepersAMD wants to merge 16 commits into
developfrom
users/brpepers/flydsl-aotriton-parity

BradPepersAMD commented Nov 4, 2026 •

edited

Loading

Uh oh!

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

edited

Loading

Uh oh!

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

edited

Loading

Uh oh!

codecov Bot commented Nov 5, 2026 •

edited

Loading

Uh oh!

cgmillette commented Nov 7, 2026

Uh oh!

cgmillette commented Nov 7, 2026

Uh oh!

cgmillette commented Nov 7, 2026

Uh oh!

cgmillette commented Nov 7, 2026

Uh oh!

cgmillette commented Nov 7, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

BradPepersAMD commented Nov 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Features

Defects fixed

Validation

Risk and mitigation

Related PRs and follow-ups

Uh oh!

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

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

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

Uh oh!

Uh oh!

codecov Bot commented Nov 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

cgmillette commented Nov 7, 2026

Uh oh!

cgmillette commented Nov 7, 2026

PR #15066 — FlyDSL SDPA forward parity with AOTriton, plus vLLM forms (BradPepersAMD, ALMIOPEN-3144) — 🟠 First-round: conflicts, copies, dropped models

Uh oh!

cgmillette commented Nov 7, 2026

PR #15066 — FlyDSL SDPA forward parity with AOTriton, plus vLLM forms (BradPepersAMD, ALMIOPEN-3144) — 🟠 Follow-up: CI update, shard timeout

Uh oh!

cgmillette commented Nov 7, 2026

PR #15066 — FlyDSL SDPA forward parity with AOTriton, plus vLLM forms (BradPepersAMD, ALMIOPEN-3144) — 🟠 Follow-up: fixes landed, one copy left

Uh oh!

cgmillette commented Nov 7, 2026

PR #15066 — FlyDSL SDPA forward parity with AOTriton, plus vLLM forms (BradPepersAMD, ALMIOPEN-3144) — 🟠 Follow-up: CI update, shards at the limit

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

BradPepersAMD commented Nov 4, 2026 •

edited

Loading

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

edited

Loading

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

edited

Loading

codecov Bot commented Nov 5, 2026 •

edited

Loading