feat(rocke): improve conv sweep benchmarks by bartekxk · Pull Request #14379 · ROCm/rocm-libraries

bartekxk · 2026-10-26T14:03:48Z

Motivation

Improve sweep time by new fast-rank mechanism. Change rocke conv measurements method to be the same as miopen.

Technical Details

-probe by less timing runs and the measure winners

applicability keys
load modules before a run
change timing to synchronize in each iteration

Test Plan

test_conv_*

Test Result

passed

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

JIRA ID: AICK-2563

therock-pr-bot · 2026-10-26T14:10:08Z

✅ All Checks Passed — Ready for Review

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

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

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

🙋 Wish to Override Policy?

therock-pr-bot · 2026-10-26T14:10:11Z

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

…arkocot/benchmark-impr

The conv sweeps timed N launches back to back between one event pair and averaged. hipDNN's dnn-benchmark (and MIOpen, where it records events around a launch) times each call on its own, so it pays the fixed device cost of an event pair per call that a back-to-back loop amortises away. On small kernels that made the sweep's numbers look faster than the same kernels measure through hipDNN, which skewed every rocKE-vs-MIOpen comparison built on them. - rocke.runtime.StallGate / time_launches_staged: the ctypes twin of hipdnn_data_sdk's StallGate (hipStreamWaitValue32 on signal memory, released by the host, with a watchdog) and of dnn-benchmark's staged loop: one stalled start/stop event pair per call, host submission outside it. - --timing {staged,block} on both conv sweeps and benchmark_conv_compare.py, default staged: every reported time is the median of stalled single-call samples. Early stopping and the --fast-rank probe still rank with cheap block launches; only the reported number changes protocol. - Per-case rows record their timing (rows without the field are block); --skip-measured / --resume skip only cases measured the same way, and a merge takes one protocol's rows instead of letting the cheaper protocol win the faster-row rule. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

Copilot

Copilot was unable to review this pull request because the user who requested the review has insufficient remaining quota.

therock-pr-bot · 2026-10-27T15:04:10Z

Pre-commit check failed

⛔ pre-commit failed

Please run locally:

python -m pip install pre-commit
pre-commit install
pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

CongMa13

I found two actionable issues in the AOT sweep support. The inline comments describe the triggers and impact.

- case_results.py: ResultsWriter no longer truncates --results in place when it names the same file as --skip-measured; measured_keys is read first and the writer appends instead of overwriting. - module_cache.py: an over-budget launcher is now released on the next miss (or at clear()) instead of leaking a HIP module per miss for the rest of the run. - benchmark_conv_compare.py: _run_script returns the child's exit status; main() now fails the shard when a sweep dies instead of always exiting 0. - CMakeLists.txt: wire the new CPU-only library tests into the conv pytest list. - batched_timing.py: time_batch_staged checks hipStreamWaitValue32 support once up front and falls back to block timing with a warning, instead of silently reporting every finalist as unranked. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

aledudek

LGTM

…arkocot/benchmark-impr # Conflicts: # dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_direct_conv.py

plan_for() now calls cache.entry_meta() (a meta dict) instead of cache.get() (a (bytes, meta) tuple), and _streamk_launch() now returns an extras dict plus rt.memset_async() instead of a free() closure and rt.memset(). The test fixtures in test_grouped_conv_dispatch.py and test_streamk_sweep_launch.py still matched the old contracts, which is what failed rocke_library_conv_pytest on CI.

The merge commit only ran pre-commit on conflicted files, so these two black violations from develop slipped through. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

cgmillette · 2026-10-29T15:24:35Z

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

PR #14379 — improve conv sweep benchmarks (bartekxk, AICK-2563) — 🟠 First-round: closer look — the module-cache leak fix unloads launches the fast-rank probe still holds

Prior item	Status
CongMa13 (10-28), aledudek +1: `--results` and `--skip-measured` on one file truncate its rows before the skip set is read	🟢 FIXED `5a1adb6f` — both sweeps now read `measured_keys` first, then open the writer in append mode when `same_results_file()` matches (`benchmark_direct_conv.py:2550`, `benchmark_implicit_gemm_conv.py:4683`). No test covers the same-file path. Thread resolved, no reply (read)
CongMa13 (10-28), aledudek +1: `module_cache.py:159`: an over-budget launcher is never unloaded, so modules leak	🟠 Leak FIXED `5a1adb6f`, but the fix causes N1: the next miss now unloads the previous over-budget launcher (`module_cache.py:145`), even when the caller has bound it and not yet launched it. Thread resolved, no reply (ran: probe)
aledudek (10-28): three `benchmark_conv_compare` tests fail with "too many values to unpack", and a sweep that dies still exits 0	🟢 FIXED `5a1adb6f` — `_run_script` now returns the child's exit status (`benchmark_conv_compare.py:162`), and `main()` fails the run on a non-zero status (`:759`, `:769`). The tests were right and are unchanged. All three pass on both gfx94X shards and in our run. Thread resolved (ran: CI logs)
aledudek (10-28): none of the new test files is wired into CTest	🟢 FIXED `5a1adb6f` — the five CPU files are in the conv pytest list (`platform/CMakeLists.txt:617`–`621`), and `test_stall_gate.py` runs in the platform pytest lane. At head, all six pass in both gfx94X shards under `ctest -L '^standard$'`: batched timing 14, capability keys 15, module cache 8, shard runner 65, launch zeroing 6 and stall gate 5, including both real-gate GPU tests. `test_early_stop.py`, which gains the two staged-timing tests, is still in no CTest list, as on develop. Thread resolved (ran: CI logs; read: CMake)
aledudek (10-28): `batched_timing.py:335`: with staged timing as the default, a device without `hipStreamWaitValue32` leaves every finalist unranked	🟢 FIXED `5a1adb6f` — `time_batch_staged` checks the gate once and falls back to block timing with a warning (`batched_timing.py:315`). 🟡 Still open: • The rows still record `timing: staged`, because the writer takes it from `--timing`. A fallback run therefore files block times as staged, the mix the module itself says must not be compared. • The serial path (`EarlyStop._timed`) has no up-front check, so it skips every candidate with the gate error. • No test covers the fallback. Thread resolved (read)
aledudek (10-29): APPROVED at `5a1adb6f`	🟢 Covers the code — given after the fix commit. The later commits are the develop merge `431375d4`, the test-double updates for the `entry_meta` API (`259d0da7`) and black formatting (`45484de0`) (read)
Bots / CI	🟢 • Policy bot pass. pre-commit passes at head, so the bot's 10-27 "pre-commit failed" comment is stale. Copilot could not review (quota). • At head `45484de0`, the one red check (`Multi-Arch CI Summary`) and all 8 cancelled checks belong to the first Multi-Arch run. That run was cancelled about 6 minutes in, when the removal of the "Not ready to Review" label started a second run. Its summary failed only because its build jobs were cancelled, so the red check is not this PR's (read: job log). • In the second run, both Linux gfx94X hipkernelprovider shards pass, including `rocke_library_conv_pytest` and `rocke_conv_gpu_gfx942_pytest`. The hipDNN superbuild passes on Linux and Windows. The ASAN, clang-tidy and Component CI runs skipped all their jobs for these paths. At 15:17Z the Windows gfx110X test shards had not started, and no gfx950 or gfx1250 test job ran (read: CI logs)
Our checks	🟢 • Merge: clean with develop `ad7f6dc5` (3 behind, 10 ahead). Develop has not changed any of the PR's paths since the base. • Build and tests: native build OK, ctest 33/33 (1 skip). Byte identity GREEN at llvm20/22/23 for the gemm families; the PR changes no kernel code. Library parity ×3, representative golden PASS ×3, differential golden OK (0 changed, 0 missing). PR tests: 241 passed, 4 skipped, none failed. ROCm clang 24 compile: no failures, no spills. Rules 1–3 clean; no links or tracker IDs added (ran). • Shipped data: none changes. The diff covers benchmark and runtime Python, tests and one CMake list. It touches no AOT tables, tuning records or goldens (ran: git). • Timing (read): a staged sample arms a `hipStreamWaitValue32` stall, records the start event, launches, records the stop event, releases the stall and syncs. The reported time is the median of these samples. This is the protocol the module documents as hipDNN's; we did not check the hipDNN/MIOpen equivalence. Early stop and the probe still rank with block launches. The default cut keeps 64 finalists, and every reported time comes from round 2. Modules load at bind time, outside the timed events. • Keys (read): the identity key covers every field `identity_supports_problem` reads except `wave_size`, which the split-K atomic rule reads (wave32 vs wave64). This is harmless while a cache holds one arch, and the coverage test's +1 perturbation cannot catch it. The direct problem key covers every attribute the predicate reads for direct kernels; the implicit key is the whole problem

#	Sev	file:line	Finding
N1	🟠	`benchmarks/common/module_cache.py:145`, `batched_timing.py:428`, `benchmark_direct_conv.py:2228`	Problem: once a run exceeds `--module-cache-bytes` (512 MiB by default; the module's docstring expects the widest caches to exceed it), each miss unloads the previous over-budget launcher. That is safe only if every launch has drained before the next lookup. `--fast-rank` binds a chunk of up to 128 candidates before timing it and keeps finalists for round 2. A direct-conv plan looks up all its steps before it binds any. Consequence: `bind()` keeps the raw function handle, so the probe and round 2 launch kernels whose module is already unloaded, which is undefined behaviour in HIP. Two-stage and stream-K candidates raise instead and silently drop out of the ranking. A multi-step or batch-split direct plan raises "was unloaded", which aborts the sweep. Evidence: we ran a probe with the real `ModuleCache`, `KernelLauncher` and `rank()`, HIP stubbed, 200 binaries and room for 50. 482 of 680 launches hit an unloaded module, and 10 of the 16 finalists were over budget. With the pre-fix cache (`aa5d590a`), no launch hit an unloaded module (ran: probe). Fix: unload over-budget launchers only after the case's final drain, and add a fast-rank test with a small budget

Utility: 🟢 Positive — hipDNN-comparable staged timing and resumable sharded sweeps; module-unload fix needs rework.

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

First-round review: a local qwen3.6 model made the independent first pass, and the orchestrator (Claude Opus 5.5) verified it. On a clean merge with develop, the orchestrator ran the native build and ctest, byte identity, library parity, the representative and differential goldens, the PR's test files and a ROCm clang 24 compile. It read both gfx94X hipkernelprovider logs and the failed summary job at head. It traced the module cache through the fast-rank and direct-conv bind paths and probed it on CPU under the real ranking, with HIP stubbed and the pre-fix cache as a control. It also re-ran the shard-runner tests the way the worker ran them. The orchestrator owns this merge-worthiness call. This review flags whether a closer look is needed; it is not an approval or a rejection.

cgmillette · 2026-10-29T15:30:43Z

Automated merge-queue status (rocKE merge monitor)

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

merged with today's develop branch, the build or the tests fail (its new test test_early_stop.py::test_staged_timing_reports_the_median_of_stalled_samples fails on the merge without a device (the stall gate is not stubbed))

What to do: merge the current develop branch into this PR and fix the failing build or tests. After that, turn auto-merge back on; the monitor re-checks every push.

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

Copilot

🟡 Changes recommended

Unresolved critical defects allow incorrect kernels to win, break JIT sweeps, unload modules still in use, and discard aggregate shard results.

11 open issues

_{Review effort: Max}

- module_cache: keep over-budget launchers alive until end_case() confirms the case's launches have drained, instead of unloading on the next miss (--fast-rank and multi-step direct plans could race an in-flight launch) - benchmark_direct_conv / case_results: exclude verification-failed rows from winner ranking and merge, so a fast-but-wrong kernel cannot become Best - batched_timing / early_stop: raise on unsupported staged-timing gate instead of silently falling back to block timing under a "staged" label - benchmark_implicit_gemm_conv: convert the remaining JIT zeroing closures to memset_async (matching the AOT path) and fix the stream-K caller that still unpacked _streamk_launch's old 3-tuple return - run_conv_shards: handle --script compare's already-merged shard rows in the final merge and in --resume's --skip-measured handling - benchmark_conv_compare: persist per-engine shard rows beside --results instead of an ephemeral tempdir, so a killed run leaves them discoverable Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

cgmillette · 2026-10-29T23:51:06Z

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

PR #14379 — improve conv sweep benchmarks (bartekxk, AICK-2563) — 🟡 Follow-up: no closer look needed — our module-cache 🟠 is fixed, and so are 10 of Copilot's 11 threads (compare-mode resume only partly); the fix commit breaks one of the PR's own CPU tests, which CI never runs, CongMa13's approval predates it, and CI at head is still building

Prior item	Status
CongMa13 (10-28), aledudek +1: `--results` and `--skip-measured` on one file truncate its rows before the skip set is read	🟢 FIXED `5a1adb6f`, unchanged at head — both sweeps read `measured_keys` first and append when `same_results_file()` matches (`benchmark_direct_conv.py:2559`, `benchmark_implicit_gemm_conv.py:4694`). Still no test of the same-file path (ran: grep). Thread resolved
CongMa13 (10-28), aledudek +1 `module_cache.py:157`: over-budget launchers are never unloaded, so modules leak	🟢 FIXED `5a1adb6f`, made race-free in `2bb985a3` (see our N1 row) — after each case, only the held modules stay loaded (ran: probe). 🟡 Two paths still hand out launchers that nothing unloads: `--module-cache-bytes 0` ("loads per use") and an unhashable signature. In our probe, `end_case()` and `clear()` unloaded 0 of 5 (ran). That matches develop's per-use loads, but it is the same class (Rule 9) → queue them in `_unadmitted` too. Thread resolved
aledudek (10-28) `benchmark_conv_compare.py:744`: three compare tests fail, and a sweep that dies still exits 0	🟢 FIXED `5a1adb6f`, unchanged — `_run_script` returns the child's status (`benchmark_conv_compare.py:162`), and `main()` fails the run on a non-zero one (`:775`, `:785`). The three tests pass on the merge (ran). Thread resolved
aledudek (10-28): none of the new test files is in CTest	🟢 FIXED `5a1adb6f`, unchanged — the five CPU files are in the conv pytest list (`platform/CMakeLists.txt:618`–`622`), and `test_stall_gate.py` runs in the platform lane. `test_early_stop.py`, which holds the PR's two staged-timing tests, is still in no CTest list, as on develop, so CI cannot see N1 (read). Thread resolved
aledudek (10-28) `batched_timing.py:336`; ours (10-29) 🟡; Copilot (10-29) `batched_timing.py:323`: without `hipStreamWaitValue32`, finalists went unranked; the `5a1adb6f` fallback then filed block times as `timing: staged`, and the serial path had no check	🟢 FIXED `2bb985a3` — both paths now check the gate once and raise "rerun with --timing block" instead of falling back (`batched_timing.py:321`, `early_stop.py:273`), so no block time is recorded as staged. With no device, `rank(timing="staged")` runs the block probe round, then raises (ran: probe). 🟡 No test covers the raise, and the new serial check breaks the PR's staged early-stop test (N1). Threads resolved, no reply
aledudek (10-29) APPROVED at `5a1adb6f`	🟡 Stale — predates `259d0da7`, `45484de0`, `2bb985a3` and two develop merges
Ours (10-29) N1 🟠; Copilot (10-29) `module_cache.py:147`: the leak fix unloads modules that bound fast-rank chunks, round-2 finalists and multi-step plans still use	🟢 FIXED `2bb985a3` — over-budget launchers wait in `_unadmitted` (`module_cache.py:156`) until `end_case()` (`:159`). Both sweeps call it only after the case's last drain (`benchmark_direct_conv.py:2690`, `benchmark_implicit_gemm_conv.py:5142`). We re-ran the first round's probe on the merge: real cache, launcher and `rank()`, HIP stubbed, 200 binaries with room for 50, two cases. Per case, 0 of 680 launches hit an unloaded module (482 of 680 at `45484de0`), and a two-step plan past the budget now binds (ran: probe). The two new cache tests pass (`test_module_cache.py:152`, `:171`) (ran). Thread resolved, no reply
Ours (10-29) merge monitor: auto-merge off until N1 is fixed	🟢 Resolved — N1 is fixed. bartekxk re-armed squash auto-merge at 22:34Z, with the push of `2bb985a3` and the develop merge. Nothing at head is 🟠 now (read: gh)
Copilot (10-29) ×8: a failed finalist can become the direct sweep's `Best`; a failed row can win the merge; the JIT `_launch_spk`, `_launch_two_stage` and `_launch_atomic` clear synchronously under the stall gate (two threads, one as Rule 9); the JIT stream-K caller unpacks the new extras dict; the runner's merge drops every compare row; absolute latencies in two docstrings (Rule 2)	🟢 FIXED `2bb985a3` • Failed finalists: `row()` keeps a FAIL plan out of the ranking (`benchmark_direct_conv.py:2291`), as the implicit fast-rank path already did (`benchmark_implicit_gemm_conv.py:5458`) (read). • Merge: `merge()` drops `passed: false` rows (`case_results.py:503`). A faster failed row no longer beats a slower passing one; at `45484de0` it did (ran: probe). • JIT clears: all three closures queue `memset_async` on the launch's resolved stream (`benchmark_implicit_gemm_conv.py:3770`, `:4007`, `:4511`). The remaining synchronous clears run at setup or around verification, outside the timed calls (ran: grep). • Stream-K: the JIT caller frees `sk_extras["ws_ptr"]` (`:4179`), the one buffer `_streamk_launch` allocates (`:5775`) (read). • Compare rows: the runner now keys compare rows directly (`run_conv_shards.py:570`). Its tests stub merged-format shards and pass (ran). • Rule 2: both docstrings are qualitative now (`module_cache.py:11`, `run_conv_shards.py:30`). The squash message is the PR body, with no numbers (read: gh). All 8 threads resolved, none answered
Copilot (10-29) `benchmark_conv_compare.py:740`: an interrupted comparison leaves no durable rows, so the runner can neither show progress nor resume them	🟡 Partly `2bb985a3` — with `--results`, the per-engine rows now persist beside it (`benchmark_conv_compare.py:742`). But `--resume` still re-measures a killed shard's finished cases. The runner carries those rows (its `shard-*.jsonl` glob also matches the new names, `run_conv_shards.py:226`), yet `expand_merged_for_skip` (`case_results.py:200`) expands only merged rows. In our probe, neither sweep would skip a case that a killed shard had finished (ran: probe). Progress still counts only finished shards → pass rows that already name an `engine` through unchanged. Thread resolved, no reply
CongMa13 (10-29) APPROVED at `45484de0`	🟡 Stale, still counts — given before `2bb985a3` (the fixes for our N1 and Copilot's threads) and the develop merge `4446a879`. GitHub reports APPROVED, so with auto-merge armed, this approval and green CI land the PR (read: gh)
Bots / CI	🟡 • At head `4446a879` (23:41Z): nothing has failed. Pre-commit, clang-tidy, Component CI, the ASAN summary, base freshness, both hipDNN superbuilds (Linux, Windows) and the policy bot pass. TheRock Multi-Arch CI is still building: Windows gfx110X math-libs since 23:02Z, and Linux math-libs for gfx94X, gfx950 and gfx125X queued since 23:33Z. No hipkernelprovider test shard has run at head. • At `45484de0`: both gfx94X shards passed `rocke_library_conv_pytest`. Since then, `2bb985a3` changed two of that lane's files; both pass on our merge (below). • GitHub: MERGEABLE, BLOCKED until the checks finish; review decision APPROVED (read: gh)
Our checks	🟡 • Range: one new commit, `2bb985a3` (benchmark Python in 8 files, `test_module_cache.py`, `test_run_conv_shards.py`). `4446a879` is a plain develop merge: git's own result, with no hand edits. It brings develop's edits to two PR files, `benchmark_direct_conv.py` and `platform/CMakeLists.txt` (#13889, #13965, #14385) (ran: git). • Merge `5f8f6502` (head + develop `3866fb55`; 3 behind, 12 ahead, 0 conflicts): native build OK and ctest 33/33 (1 skip); byte identity GREEN ×3 for the gemm families; library parity 2 passed ×3; representative golden PASS ×3; differential golden OK (ll at llvm20–24 and canonical; 0 changed, 0 missing); ROCm clang 24 gemm compile with 0 failures, 0 new rejections and 0 spills; the PR's tests: 242 passed, 4 skipped, 1 failed (N1); Rules 1–3 clean (ran). The newest develop `d730d162` still merges cleanly and changes no PR file (ran: merge-tree). • Shipped data: `2bb985a3` touches no kernel code, AOT tables, tuning records or goldens (ran: git). • Stack: #14628 sits on this branch; its rebase onto `2bb985a3` is tracked there. This PR lands first

#	Sev	file:line	Finding
N1	🟡	`library/tests/test_early_stop.py:209`, `benchmarks/common/early_stop.py:273`	Problem: `2bb985a3` makes `EarlyStop._timed` build the stall gate before its first staged sample. The PR's own `test_staged_timing_reports_the_median_of_stalled_samples` stubs `time_launches_staged` but not the gate. Consequence: wherever no gate can be built (a host without a GPU, or a device without `hipStreamWaitValue32`), the test raises `HipError`: on the merge, 1 failed and 12 passed (ran). It passed at `45484de0`, and develop does not have it. CI stays green only because no CTest list runs `test_early_stop.py`. Fix: stub `stall_gate.default_gate` in the test (the test then passes: ran, probe), and add the file to the conv pytest list (`platform/CMakeLists.txt:618`)

Utility: 🟢 Positive — hipDNN-comparable staged timing, sharded resumable sweeps, module reuse; unload race fixed.

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

Follow-up review: a local qwen3.6 model made the independent first pass, and the orchestrator (Claude Opus 5.5) verified it. It reviewed the one new commit since our review, 2bb985a3, and checked that 4446a879 is git's own develop merge. On a clean merge with develop, it ran the native build and ctest, byte identity, library parity, both golden gates, a ROCm clang 24 compile and the PR's tests. On CPU it re-ran the first round's module-cache probe on the merge, with 45484de0's cache as a control. It also probed the staged-timing raise, the failed-row merge, compare-mode resume and the failing test with the gate stubbed. It read every thread, the approvals, the auto-merge state and CI at head. The orchestrator owns this merge-worthiness call. This review flags whether a closer look is needed; it is not an approval or a rejection.

cgmillette

Automated merge-queue hold (rocKE merge monitor)

Hi @bartekxk, this PR is held in the rocKE merge queue: one of its own new tests fails on a merge with today's develop (587f9942).

What fails: library/tests/test_early_stop.py::test_staged_timing_reports_the_median_of_stalled_samples (1 failed; develop has no such failure). EarlyStop._timed now builds the stall gate before the first staged sample (benchmarks/common/early_stop.py:273), and the test does not stub it. On a host without a ROCm device it raises HipError: hipGetDevice: hipError(100) no ROCm-capable device is detected; this device cannot be timed with --timing staged, rerun with --timing block. Our 10-29 follow-up review noted this as N1.

Fix: stub the gate in the test, next to the time_launches_staged patch (test_early_stop.py:221), for example:

monkeypatch.setattr("rocke.runtime.stall_gate.default_gate", lambda: None)

With that stub, the test passes on the same merge.

Once it's fixed, push and re-request a review from @cgmillette. A held PR needs that explicit re-approval before it can merge.

fixed

Brings in #14393 (gfx1250 split-K forward conv) and #14379 (conv sweep benchmarks). Only the representative IR golden conflicted, on its summary counters: develop's split-K cases and this branch's five depthwise dgrad cases per flavor do not overlap. Merged case by case against the merge base, with the counters recomputed from the merged cases. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The direct predicate reads the column padding as PAD_W, which pad_w overrides, but DIRECT_PROBLEM_CAPABILITY_ATTRS from #14379 left pad_w out. Two problems differing only in pad_w then shared a capability key, so a cache filter could answer a per-axis padding with the verdict for an equal one. test_direct_problem_key_covers_every_extent_the_predicate_reads caught it after the merge; a new test pins the separation directly. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

## Motivation On wave32 targets, the row-streaming depthwise dgrad (`DirectDepthwiseDgradStreamSpec`, AOT variant `direct_depthwise_dgrad`) runs one channel per lane. It uses 16-bit loads and stores and scalar FMAs. Depthwise dgrad also has no direct kernel for filters past 11x11 or for non-square decomposed large kernels (51x5 / 5x51), so those layers fall back to implicit GEMM. This PR does three things: - it adds channel packing to the row-streaming depthwise dgrad; - it adds a rotated-filter dgrad mode to the output-stationary depthwise kernel, with per-axis padding for non-square filters; - it mirrors both in the C++ engine with byte-identical IR. The C++ engine had no port of the row-streaming depthwise dgrad, so this PR also adds its scalar form. Like `conv_direct_dgrad_wmma`, the new binaries are served by the direct-conv AOT cache. There is no dispatcher wiring. JIRA ID: AICK-2337 ## Technical Details **Channel-packed row-streaming dgrad (`DirectDepthwiseDgradStreamSpec.cpl`, `library/kernels/common/conv_direct_grouped.py`)** - With `cpl` 2, 4 or 8, every lane owns `cpl` adjacent channels: - dY rows load and dX rows store as one `cpl`-wide vector per lane; - every tap runs as `cpl / 2` packed `<2 x f32>` FMAs over channel pairs. - The schedule does not change: dY rows stream into `stride * KH` circular accumulator slots with a few rows prefetched, and a tail loop flushes the rows past the last dY row. - `cpl == 1` emits the unchanged kernel. - Validator rules: - `cpl` in {1, 2, 4, 8} and `groups % cpl == 0`; - `cpl * KH * KW <= 200` for the packed weights; - `cpl * block_w * KH <= 192` for the accumulator slots. - `block_ch` grows by `cpl`, and the kernel name gains `cpl{N}`. **Rotated-filter dgrad (`DirectDepthwiseTiledSpec.dgrad`)** - A stride-1 depthwise dgrad is the forward correlation of dY with the filter rotated by 180 degrees. The padding is `KH-1-PAD` rows and `KW-1-PAD_W` columns. - The output-stationary forward kernel computes it through the dgrad kernarg block: - the extents and the A-side strides swap roles; - weight row `r` is read from row `KH-1-r`, reversed; - the epilogue addresses dX through its kernarg strides; - the grid tiles `H x W`. - Only one filter row is register-resident, so this mode takes filters up to 64 on a side, including non-square ones. The forward mode keeps its 32 limit. - The kernel name gains the `dg` flag. `dgrad=False` emits the unchanged kernel. - The mode packs one channel per lane: both engines refuse `dgrad` with `cpl > 1` (the forward mode's channel-packed kernel, from #14392, is forward only). **Per-axis padding (`DirectConvProblem.pad_w`)** - `pad_w = -1` means `PAD`. `PAD_W` is the effective column padding, and `Wo` uses it. - Only the rotated-filter dgrad takes a problem with `PAD_W != PAD`: - the forward tiled mode refuses one; - the WMMA convolutions (`direct_fwd_wmma`, `fwd_smallc_wmma`, `fwd_smallk_wmma` through `wmma_fwd_gate`, `direct_dgrad_wmma`, and #14387's stem dgrad kernels `dgrad_patch_wmma` and `dgrad_smallc_wmma`) refuse one in their validators, in both engines. They offset columns by `PAD`. The check is `one_padding_reason` / `rocke_direct_conv_one_padding_reason`, shared with the row-streaming dgrad; - `direct_spec_for_problem` refuses one for every other variant; - `direct_job_for_identity` builds non-square / per-axis identities only for the new variant. - `kernel_identity` matches `pad_w` against `PAD_W`, and `conv_args` passes `pW = PAD_W`. - The direct problem's capability key (`DIRECT_PROBLEM_CAPABILITY_ATTRS`, from #14379) includes `pad_w`, so problems that differ only in their column padding never share a cached support verdict. - The benchmark accepts `-p != -q` only for a stride-1 depthwise dgrad. **AOT (`library/kernels/conv/direct_variants.py`, `direct_aot_jobs.py`)** - `direct_depthwise_dgrad` gains the `DW_DGRAD_CPL_GRID` knobs. The `cpl == 1` knobs keep their identities. - New variant `direct_depthwise_tiled_dgrad`, in `DIRECT_WAVE_PARAMETRIC`: - filters 13x13 to 31x31, plus 51x5 and 5x51, all at "same" padding; - tiles `DW_TILED_DGRAD_BLOCK_W` x `DW_TILED_DGRAD_BLOCK_H`. - `DirectCaps` gains `kw` / `pad_w` for those rows. **C++ engine** - New `rocke/instance_conv_direct_grouped_dw_dgrad.h` and `conv_direct_grouped_build_depthwise_dgrad_stream.cpp` port `DirectDepthwiseDgradStreamSpec`. This covers the defaults, `block_ch`, the kernel name, `validate` / `is_valid` (same order and reason text as Python), and the scalar and `cpl > 1` builders. They also add `_new` and `lower_to_llvm` entries. - `rocke_direct_depthwise_tiled_spec_t.dgrad` covers the builder mode, the `dg` name flag, the dgrad filter limit and the padding checks of `validate()`. - `rocke_direct_conv_problem_t.pad_w` (default -1), `rocke_direct_conv_problem_pad_w` and `rocke_direct_conv_one_padding_reason`. Every C++ mirror of `DirectConvProblem.Wo` in the direct-conv family uses it, as the Python property does. - The depthwise kinds have no `rocke_engine` binding, here or on `develop`. The parity pair is the cross-engine gate. **Docs:** `dsl_docs/instances/convolution.md` (the dgrad mode in the tiled section, and a new row-streaming depthwise dgrad section), `dsl_docs/instances/index.md`, `SUPPORT_MATRIX.md`. ## Test Plan - `check_byte_identity.py`, all families, at llvm20, llvm22 and llvm23. - `conv_direct_grouped_emit.{py,c}` parity configs 179-200 (after develop's wave32 forward configs 166-178): - row-streaming dgrad, scalar and `cpl` 2 / 4 / 8, in fp16 and bf16, at stride 1 and 2, on wave32 (gfx1250) and wave64 (gfx950); - rotated-filter dgrad on 13x13, 31x31, 51x5 (`pad_w`), 5x51 (`pad_w`, unrolled rows), a 7x7 with a padding that differs per axis and is not "same", and a 5x5 with `pad_w = 0` (the kernel name carries the per-axis padding); - nine specs both engines reject: - per-axis padding on the forward tiled mode; - the dgrad mode at stride 2; - a 65-row filter; - `groups % cpl != 0`; - `cpl = 3`; - the packed-accumulator bound; - `pad_w < -1` on the rotated-filter dgrad; - per-axis padding on the row-streaming kernel; - `pad_w < -1` on the row-streaming kernel. - Differential golden (`tests/instances/differential/golden/llvm_gfx_all.json`): the new `conv_direct_grouped` configs that both engines emit are blessed for `ll` (llvm20 / llvm22 / llvm23) and `ir_canonical`. - Golden IR: five new gfx1250 cases (row-streaming scalar, `cpl` 2, `cpl` 8 stride 2 bf16; rotated-filter 31x31 and 51x5) at every flavor. - `library/tests/test_direct_depthwise_dgrad_v2.py`, now in CTest: validators, AOT grid entries (including the packed-entry spill gate), the padding gate, and numerics against a numpy reference on any arch rocke models (each case built at the arch's wave size; cases the arch's validator refuses are skipped). - `rocke_conv_direct_dw_dgrad_host` ctest: the stream name-overflow cleanup, `pad_w < -1`, the per-axis padding name, the `dgrad` / `cpl > 1` refusal, and the per-axis padding refusal of the six WMMA kernels (including the two stem dgrad kernels). It is listed in the `quick` and `ffm-quick` categories of `ROCKE_ENGINE_test_categories_external.yaml`, so CI runs it. - Targeted CPU pytest, with and without the `rocke_engine` binding on the path. - black 25.12.0 / clang-format 18.1.4 on the changed files. ## Test Result - Byte identity GREEN at llvm20, llvm22 and llvm23. `conv_direct_grouped` has 201 configs: the 13 new compiled configs are IDENTICAL, and the 9 new reject configs are BOTH_REJECTED. - Golden `--check` passes at every flavor. Every existing hash is reproduced: `cpl == 1` and `dgrad=False` emit the old IR. The golden diff adds case keys and updates the summary counters. - Differential golden `run_diff.py --python-only --check-golden` is GOLDEN OK for `ll` at every flavor and for `ir_canonical`, with 0 unblessed changes and 0 missing. The bless adds 13 `conv_direct_grouped` configs (179-190, 197) to each section; no existing hash changes or drops. - `test_direct_depthwise_dgrad_v2.py` on gfx1250 (before the latest develop merge): 11 passed, 10 subtests (the GPU cases). - After merging develop at `311cb32338e` (#14387, wave32 WMMA stem dgrad), CPU only: - byte identity GREEN at llvm20, llvm22 and llvm23 for `--only conv,gfx1250`, including `conv_direct_grouped` (201 configs), `conv_dgrad_patch_wmma` (9) and `conv_dgrad_smallc_wmma` (9); - golden `--check` passes (llvm20 / 22 / 23 / 24); the differential golden is GOLDEN OK for `ll` at every flavor and for `ir_canonical`, with 0 unblessed changes and 0 missing; - native build and ctest: 34/34 passed, including `rocke_conv_direct_dw_dgrad_host`; - `test_direct_depthwise_dgrad_v2.py`, `test_direct_kernel_cache.py`, `test_direct_dgrad_wmma.py`, `test_direct_fwd_wmma.py`, `test_fwd_small_wmma.py`, `test_dgrad_stem_wmma.py`, `test_kernel_label.py`: 213 passed, 8 skipped with the `rocke_engine` binding; 203 passed, 18 skipped without it. The skips need a GPU or the binding; - `test_direct_depthwise_fwd_cpl.py`, `test_conv_direct_grouped_backend_parity.py`, `test_direct_depthwise_col_validate.py`, `test_conv_dgrad_depthwise_windowed.py`: 76 passed, 8 skipped with the binding; 70 passed, 14 skipped without it. - The merge's only conflict was the llvm23 summary counters of the representative golden: develop adds 4 gfx1250 stem-dgrad cases and this PR 5 gfx1250 depthwise-dgrad cases. Cases and expected failures are the union of both sides, and the counters are recomputed from them: `ok_cases` is 314 / 314 / 342 (llvm20 / 22 / 23). Every case hash equals develop's or this branch's. - Develop adds no `conv_direct_grouped` parity configs and changes no `conv_direct_grouped` entry of the differential golden (#14387 adds its own `conv_dgrad_patch_wmma` / `conv_dgrad_smallc_wmma` sections), so this PR's configs keep 179-200. - After that merge, #14387's two stem dgrad kernels also refuse a per-axis padding (they offset columns by `PAD`); their emitted IR is unchanged. - After merging develop at `32aed8707aa` (#14393 split-K forward conv, #14379 conv sweep benchmarks), CPU only: - the only conflict was again the representative golden's summary counters; merged case by case against the merge base, `ok_cases` is 318 / 318 / 347. The golden differs from develop only by this PR's 5 cases per flavor, and every hash equals develop's or this branch's; - #14379's `test_capability_keys.py` caught that its direct capability key left out `pad_w`; fixed in `23b898c8d2c` with a regression test; - byte identity GREEN at llvm20, llvm22 and llvm23 for `--only conv,gfx1250` (`conv_direct_grouped` 201 configs); golden `--check` passes (llvm20 / 22 / 23 / 24); the differential golden is GOLDEN OK for `ll` and `ir_canonical`; - native build and ctest: 34/34 passed; - the seven test files above: 213 passed, 8 skipped with the binding; the four above: 76 passed, 8 skipped with the binding; - develop's new sweep tests (`test_capability_keys`, `test_run_conv_shards`, `test_module_cache`, `test_aot_launch_zeroing`, `test_batched_timing`, `test_early_stop`, `test_streamk_sweep_launch`) plus the grouped-conv dispatch, split-K forward and this PR's tests: 445 passed, 10 skipped, 1 failed. The failure, `test_early_stop.py::test_staged_timing_reports_the_median_of_stalled_samples`, fails identically on `develop`. - After merging develop at `b33ff848a7b` (#14494 split the rocKE CMake test registration into `tests/cmake/*.cmake`), the `rocke_conv_direct_dw_dgrad_host` registration moves verbatim into `platform/tests/cmake/InstanceTests.cmake`; `tests/CMakeLists.txt` is develop's unchanged. CPU only: byte identity GREEN at llvm20 / 22 / 23 (`conv_direct_grouped` 201 configs), golden `--check` passes (llvm20 / 22 / 23 / 24), the differential golden is GOLDEN OK, native ctest 34/34 (`rocke_conv_direct_dw_dgrad_host` passes), `test_cmake_test_registration.py` 5 passed, and the test counts above are unchanged. - After merging develop at `64c2954a8e` (#14679 split the representative-IR golden into per-family shards), this PR's 5 cases per flavor move unchanged into the `conv_direct` and `conv_direct_depthwise_tiled` shards under `platform/tests/golden/rocke_representative_ir_sha256/`; the shards differ from develop's by those additions only. CPU only: golden `--check` against the shard directory passes (llvm20 / 22 / 23 / 24), byte identity GREEN at llvm20 / 22 / 23, the differential golden is GOLDEN OK, native ctest 34/34, and `test_ir_golden.py` / `test_rocke_ci_static.py` / `test_cmake_test_registration.py` pass; the test counts above are unchanged. - The full CPU pytest suite, run before the latest develop merges, showed no new failures against `develop`. The remaining failures are pre-existing and identical on `develop`: `test_conv_abi.py::test_cache_applies_runtime_launch_limits`, six cases in `test_conv_batch_and_miopen_parse.py`, `test_conv_dgrad_patch.py::TestPatchIdentity::test_off_keeps_the_historical_hash`, and `test_engine_linkage.py::test_installed_pytest_gets_distinct_artifact_environment_variables`. - No GPU run on gfx950 / gfx942 (no hardware here). Those targets are pinned by the parity and golden gates. - No performance results are included in this PR. ## Submission Checklist - [x] Both engines emit byte-identical IR for every new code path; the parity pair and golden cases are added. - [x] Defaults keep every existing configuration and its IR unchanged. - [x] New C++ code is in new files where the change allowed it. - [x] CTest entry, docs and support-matrix rows. - [ ] Dispatcher wiring: not part of this PR; the new binaries are served by the direct-conv AOT cache. --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

assistant-librarian · 2026-11-03T11:07:09Z

TheRock Submodule Bump Activity

Newest first

2026-11-03 — Included in TheRock via ROCm/TheRock#9698.

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

therock-pr-bot Bot added the Not ready to Review label Oct 26, 2026

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

bartekxk added 4 commits October 26, 2026 15:04

bench

2f92cd5

impr

2a28cb8

impr

7aa1999

impr

f3f8279

bartekxk force-pushed the users/barkocot/benchmark-impr branch from 9f0b99a to f3f8279 Compare October 26, 2026 15:04

bartekxk and others added 2 commits October 26, 2026 15:04

Merge branch 'develop' of github.com:ROCm/rocm-libraries into users/b…

b908e6d

…arkocot/benchmark-impr

bartekxk marked this pull request as ready for review October 27, 2026 14:56

bartekxk requested review from a team as code owners October 27, 2026 14:56

bartekxk changed the title ~~Users/barkocot/benchmark impr~~ feat(rocke): improve conv sweep benchmarks Oct 27, 2026

bartekxk requested a lite review from Copilot October 27, 2026 14:56

Copilot AI reviewed Oct 27, 2026

CongMa13 reviewed Oct 28, 2026

View reviewed changes

Comment thread dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/case_results.py Resolved

Comment thread dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/module_cache.py Resolved

aledudek reviewed Oct 28, 2026

View reviewed changes

aledudek approved these changes Oct 29, 2026

View reviewed changes

bartekxk and others added 3 commits October 29, 2026 08:05

Merge branch 'develop' of github.com:ROCm/rocm-libraries into users/b…

431375d

…arkocot/benchmark-impr # Conflicts: # dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_direct_conv.py

style(rocke): black-format files brought in by the develop merge

45484de

The merge commit only ran pre-commit on conflicted files, so these two black violations from develop slipped through. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

bartekxk enabled auto-merge (squash) October 29, 2026 13:23

therock-pr-bot Bot removed the Not ready to Review label Oct 29, 2026

cgmillette disabled auto-merge October 29, 2026 15:26

tenpercent requested a max review from Copilot October 29, 2026 17:33

Copilot started reviewing on behalf of tenpercent October 29, 2026 17:34

View session

Copilot AI reviewed Oct 29, 2026

View reviewed changes

CongMa13 approved these changes Oct 29, 2026

View reviewed changes

bartekxk and others added 2 commits October 29, 2026 17:29

Merge branch 'develop' into users/barkocot/benchmark-impr

4446a87

bartekxk enabled auto-merge (squash) October 29, 2026 22:34

cgmillette mentioned this pull request Oct 29, 2026

feat(rocke): faster conv fwd for depthwise, grouped, patch stems #14628

Merged

cgmillette mentioned this pull request Oct 30, 2026

fix(rocke): skip GDN/KDA tests on Windows CI #14710

Merged

cgmillette previously requested changes Oct 30, 2026

View reviewed changes

cgmillette disabled auto-merge October 30, 2026 02:28

cgmillette mentioned this pull request Oct 30, 2026

"[rocke] TestIrParityCoverage.test_ir_cases_match_golden_sha256 times out intermittently on Linux gfx94X (60 s per-test limit)" #14733

Closed

bartekxk enabled auto-merge (squash) October 30, 2026 10:19

bartekxk disabled auto-merge October 30, 2026 10:20

bartekxk merged commit 80cc53d into develop Oct 30, 2026
165 of 168 checks passed

bartekxk deleted the users/barkocot/benchmark-impr branch October 30, 2026 10:20

jakpiase mentioned this pull request Oct 30, 2026

feat(rocke): channel-packed and large-filter depthwise dgrad #14386

Merged

4 of 5 tasks

cgmillette mentioned this pull request Oct 30, 2026

feat(rocke): paged K/V for dense prefill attention on gfx942 and gfx950 #14552

Open

1 task done

Conversation

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

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

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

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

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

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

therock-pr-bot Bot commented Oct 27, 2026

Pre-commit check failed

Uh oh!

CongMa13 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

aledudek left a comment

Choose a reason for hiding this comment

Uh oh!

cgmillette commented Oct 29, 2026

PR #14379 — improve conv sweep benchmarks (bartekxk, AICK-2563) — 🟠 First-round: closer look — the module-cache leak fix unloads launches the fast-rank probe still holds

Uh oh!

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

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

cgmillette commented Oct 29, 2026

Uh oh!

cgmillette left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Nov 3, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

bartekxk commented Oct 26, 2026 •

edited

Loading

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

edited

Loading

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

edited

Loading

cgmillette commented Oct 29, 2026 •

edited

Loading