perf(rocke/conv): add the pointwise fast path dgrad was missing by aledudek · Pull Request #11928 · ROCm/rocm-libraries

aledudek · 2026-09-10T15:41:06Z

ISSUE ID: AICK-1763

Motivation

Forward conv and wgrad both have a fast path for 1×1 convolutions. dgrad never got one.
For 1×1 with stride 1 and no padding, the address math is trivial, the offsets are simple multiply-adds. But dgrad was still running the general-case code: a runtime integer divide plus a bounds check that can never fail, and both of them inside the K-loop, so the cost repeats on every iteration. 1×1 shapes are common in real backward-data workloads, so this was worth fixing.

Technical Details

Four commits, branched off users/aledudek/transpose_load_bwd_weight_data_gfx1250 @ f02fad7.

The fast path. When the shape is 1×1, stride 1, unpadded and ungrouped, use the simple offset formula directly instead of the general one. Two early-returns in the dY and W descriptors, mirrored in the C++ engine in the same commit. The 1×1 kernel goes from 1295 to 991 instructions, and most of the integer divides and bounds selects disappear.
A vector_size_b bug. That field is meant as an upper limit, not a hard request, wgrad treats it that way. The K-outer path clamped it properly, the M-outer path didn't, so a too-wide value passed validation and then crashed the builder. Both engines now clamp. This changes no emitted kernel.
Two benchmark flags. --lds-k-outer {auto,on,off} so the two LDS layouts can be compared (the automatic choice is the same for every config of a given shape, so a normal sweep has nothing to compare against), and --csv-top N because the CSV writer was hardcoded to 5 rows and ignored --top. Both keep the old behaviour by default.
Docs and a driver. Both LDS case studies end their trace command with "single-config driver", that driver didn't exist, so it's added here. Also a short write-up of two other optimisations that were built, measured, and dropped because they didn't help.

Test Plan

tools/check_byte_identity.py at llvm20 and llvm22
tests/instances/differential/run_diff.py --only conv_implicit_gemm_dgrad
tests/instances/test_conv_dgrad_correctness.py
Numeric check of the fast path against dgrad_reference on three 1×1 shapes plus a 3×3 control
hsaco comparison of 3×3 stride-1 and stride-2 with and without the change
Emission-neutrality of commit 2 by hashing lowered IR across the full sweep grid, before and after

Test Result

Byte-identity GREEN, 69 configs, at both LLVM flavors. dgrad parity config 4 is pointwise, so the gate covers the new path.
Differential GREEN (69 families; dgrad 14 configs).
dgrad correctness: 22 passed.
3×3 stride-1 and stride-2 hsacos byte-identical before/after. only pointwise emission moves.
Fast path is numerically exact: relative error against dgrad_reference identical to the generic path on all four cases.
Commit 2: zero IR-hash changes across 4564 previously-building configs, zero regressions, 428 previously-unbuildable configs now build.
check_golden reports no drift.

Submission Checklist

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

therock-pr-bot · 2026-09-10T15:47:16Z

✅ 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-09-10T15:47:18Z

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

Copilot

🟡 Changes recommended

Two moderate driver correctness issues and additional review findings remain unresolved.

Get a fresh assessment by requesting another Copilot review.

Pull request overview

Adds a pointwise 1×1 dgrad fast path, vector-width clamping, benchmark controls, a single-config driver, and updated case-study documentation.

Changes:

Adds mirrored Python/C++ dgrad fast paths.
Clamps oversized vector widths.
Adds LDS-layout and CSV benchmark options.
Adds dgrad driver support and documentation updates.

File summaries

File	Review findings
`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/conv_implicit_gemm_dgrad.py`	Nit (2 votes): add numeric regression coverage for ungrouped 1×1 dgrad, including a partial M tile. Nit (1 vote): add regression coverage for oversized `vector_size_b`.
`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx950/conv_dgrad/run_one_dgrad.py`	Moderate (3 votes): resolve or reject `--split-k -1` before packing and launching. Nit (1 vote): qualify the per-combo LDS-layout explanation. Moderate (1 vote): reject non-dgrad `-F` values.
`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx950/conv_dgrad/dgrad_lds_layout_case_study.md`	Nit (3 votes): scope the single-config driver claim to dgrad or provide a wgrad driver. Nit (1 vote): remove or redact measured performance data prohibited by `platform/AGENTS.md`.
`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/benchmark/benchmark_implicit_gemm_conv.py`	Nit (3 votes): qualify the `auto` LDS-layout help text per configuration or scope it to gfx950. Nit (1 vote): make the worker comment reflect pipeline- and tile-dependent selection.
`dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_implicit_gemm_dgrad.cpp`	Reviewed matching C++ implementation for the dgrad fast path and vector-width clamping.

Review details

Suppressed comments (5)

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/benchmark/benchmark_implicit_gemm_conv.py:1653

The same shape-only assertion is repeated in this worker comment, but default_lds_k_outer also keys on warp_tile_n and pipeline. On gfx1250, a single auto sweep can therefore contain both K-outer and M-outer specs. Please qualify this comment to the gfx950 case study or state that the predicate is evaluated per configuration.

        # be compared on identical configs -- under "auto" the predicate is
        # constant across a sweep of one shape, so a sweep measures one layout
        # and silently has no baseline to compare against.

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx950/conv_dgrad/dgrad_lds_layout_case_study.md:183

platform/AGENTS.md explicitly prohibits committing software-achieved performance data, including benchmarks, latency/throughput, and profiling results. This new Finding 4 records measured stall/occupancy findings, instruction counts, and runtime outcomes throughout the section. Keep the qualitative methodology and algorithmic explanation, but remove or redact the measured results before merging.

## Finding 4: what an ATT trace says is expensive is not what is on the critical path

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx950/conv_dgrad/run_one_dgrad.py:77

This help repeats that the dispatch predicate has the same answer for every combo, but this script accepts arbitrary --arch values and the predicate is pipeline-dependent: on gfx1250 mem can select K-outer while valid wavelet selects M-outer. Please make the explanation per-combo or scope it to gfx950.

        help="force the K-outer B tile. The dispatch policy deduces this via "
        "DgradConvSpec.default_lds_k_outer, which answers the same way for "
        "every combo of a given shape -- so forcing it is the only way to A/B "
        "the two layouts.",

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx950/conv_dgrad/run_one_dgrad.py:115

parse_miopen_cmd returns the raw -F value, but _fwd is discarded here. A command using -F 1, -F 4, or -F 0 is therefore silently accepted even though this is a dgrad-only driver and the help documents -F 2; the script then reports a dgrad kernel for a command describing a different operation. Reject non-2 values (or make the direction an explicit driver option) before building.

    p, _dt, _fwd = parse_miopen_cmd(cmd)

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/conv_implicit_gemm_dgrad.py:1463

This fixes a builder failure for explicit oversized vector_size_b, but no committed test covers the new behavior. The existing dgrad correctness suite has no vector-size parameterized case, and byte-identity hashes cannot exercise a previously unbuildable configuration; add a regression test that constructs an oversized cap, verifies the build succeeds with a narrowed width, and preserves existing valid-width emission.

        load_vec_b = CoalescedTileLoader.choose_vec(
            tile_rows=block_n,
            tile_cols=block_k,
            block_size=threads,
            max_vec=min(_def_vec_b, spec.vector_size_b),
            vector_axis="row",

Files reviewed: 5/5 changed files
Comments generated: 4
Review effort level: Lite

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

…ying it vector_size_* is a cap, not a demand -- wgrad documents it that way and passes vector_size_c through as max_store_vec. The dgrad K-outer branch clamped an explicit vector_size_b through choose_vec; the M-outer branch took it verbatim, so an over-wide width passed is_valid_dgrad_spec and then raised from the coalesced tile loader deep in the builder. Route the M-outer branch through choose_vec as the K-outer branch already does, in both engines. Emission-neutral: choose_vec's accepted set is a strict subset of vecs_per_thread's, and the tile_n divisibility rule the validator already enforces makes the axis condition free, so it yields exactly the requested width wherever the verbatim path built. Verified by hashing the lowered IR across the whole sweep grid before and after -- zero changes on configs that already built, zero regressions, and 428 previously unbuildable M-outer configs now build. Byte-identity green at llvm20 and llvm22; differential green; dgrad correctness suite passes. This does not make is_valid_dgrad_spec sound -- other builder-side raises remain, and load_vec_a is still taken verbatim in both layouts. Those need their own change, since the A-side clamp is not emission-neutral.

--lds-k-outer {auto,on,off} forces the dgrad LDS layout. Under auto the deducer answers the same way for every combo of a given shape, so a sweep measures one layout and silently has no baseline to compare against; forcing it is what makes an A/B possible. auto stays the default, so normal runs are unchanged. --csv-top N replaces a hardcoded rocke_results[:5] that ignored --top, so --csv wrote 5 rows out of thousands. Defaults to 5, preserving the old behaviour.

…ejected fixes Both LDS-layout case studies end their capture command with "-- python3 <single-config driver>" and no such driver existed. Add it: one config, one dispatch, so --kernel-regex matches one thing. It does no numeric verification on purpose -- torch's HIP runtime and rocke's fight over the process HIP context and timings come out multiples wrong when a verify and a timing loop share a process. --print-name-only gets the kernel name without touching the GPU queue. Record Finding 4 in the dgrad study: an ATT capture classifies the kernel as latency-bound on global memory and attributes roughly three quarters of stall to two structures -- a non-pipelining tilde load phase and a dependent-scalar-load block search. Both were implemented in both engines, verified, measured and rejected: the ISA changed exactly as designed and runtime did not move. A large stall bucket means "this is where waves sit", not "this cost is removable". The consequence is the useful part: strided dgrad is bound by total memory time, so the scheduling family is a dead end and a real win must reduce bytes moved. Also document the two ATT traps (hit-weighted totals; inline_frames stacks live under "stacks" and are outermost-first), and correct two stale claims -- the sweep now has --lds-k-outer, and compv3/compv4 are rejected only on WMMA, not on gfx950.

Forward conv and wgrad both special-case pointwise convolution; dgrad never did. grep for is_pointwise hits conv_implicit_gemm.py and conv_implicit_gemm_wgrad.py in nine places each and the dgrad instance in none, in both engines. For Y=X=1, stride 1, pad 0, dilation 1, ungrouped, the tilde decomposition is the identity: h_tilde_slice == Ho, w_tilde_slice == Wo, y_dot_slice == x_dot_slice == 1 and gemm_k == kpg, so the dY offset reduces exactly to m_sub*K + k_sub and the W offset to k_sub*C + c_val. Instead dgrad emitted a runtime signed divide plus a tautological bounds predicate INSIDE the K-loop -- the descriptor closures run per element per K-tile from emit_load_phase, which sits inside the scf_for. Take the algebraic form directly on that shape class. The 1x1 kernel drops 1295 -> 991 instructions; v_mul_lo 54 -> 16, v_mul_hi 29 -> 16, v_cndmask 49 -> 16. This is worth real time because short-K pointwise dgrad is execution-bound, not latency-bound -- ATT wave state is EXEC ~55% / WAIT ~20%, the inverse of a long-K shape -- so deleting instructions converts to runtime. Two earlier attempts on latency-bound regimes changed the ISA exactly as designed and moved runtime by zero; this one is measured on hardware and does not. Non-pointwise emission is untouched: the 3x3 stride-1 and stride-2 hsacos are byte-identical with and without the change, so only pointwise dgrad moves. Numerically the fast path is exact, not approximate -- relative error against dgrad_reference is identical to the generic path on three 1x1 shapes and a 3x3 control. Byte-identity green at llvm20 and llvm22 (dgrad parity config 4 is pointwise, so the gate covers the new path); differential green; dgrad correctness suite passes; check_golden reports no drift. Follow-ups, deliberately not in this commit: - No representative-IR case is pointwise dgrad (all 11 are 3x3), so the new path has parity coverage but no golden. Worth adding one. - Grouped pointwise still takes the generic path. Extending it means folding k_out_group_base into the dY offset and using cpg as the W stride; every pointwise shape measured so far is ungrouped, so it would be untested code today. - make_dgrad_dx_descriptor composes unmerge_magic with a naive NHWC re-merge that is algebraically m*C + c, leaving dead magic-division in the non-strided epilogue for every shape, not just pointwise.

run_one_dgrad.py took --split-k -1 but never saw the degree it resolved to. The builder resolves auto onto a dataclasses.replace copy and hands back only a KernelDef, so the driver kept a spec at -1: sub-GEMM records packed with split_k=1 K-padding, needs_atomic False for a kernel that is in fact atomic (dX never re-zeroed between iterations), and gridDim.z going out as -1. Rejecting is the right answer rather than resolving it here -- the shipped dgrad candidate pins split_k=1 and declines the CK formula, so there is no auto dispatch for a single-config trace to reproduce. Fixed --split-k N is unaffected. While in here, three claims that were true only on gfx950 or only for dgrad: --lds-k-outer's help, its call-site comment and --kouter's help all said the deducer answers the same way for every combo of a shape, which stops holding on gfx1250 where wavelet builds; the case study said this driver serves the wgrad study too, when it builds DgradConvSpec and nothing else; and --miopen-cmd's help documented -F 4 as meaningful when the direction field is discarded outright. Also adds test_fp16_pointwise_1x1 to the dgrad correctness suite. The pointwise flat-offset branches only had emitter-vs-emitter parity, which a shared wrong offset sails straight through -- mutating the dY stride to K+1 now fails the suite. The docstring is explicit that the branch's validity predicate is *not* covered: on a flat offset the out-of-range lanes land past the end of the buffer and the descriptor clamp has already zeroed them, so deleting either clause leaves the test green no matter what dims you pick.

--lds-k-outer was removed in #11602/#11762 under "selection, not knobs"; re-adding it to the sweep driver contradicted convolution.md, the wgrad case study and lds-optimization-rocke.md, and let a sweep measure a kernel dispatch cannot select. The sweep driver now calls DgradConvSpec.default_lds_k_outer unconditionally, as wgrad does. run_one_dgrad.py gains --kouter auto (the new default), which asks the same predicate so a plain trace captures the kernel that ships; on/off remain for the in-process A/B the numeric tests use.

--csv-top was a bare slice bound, so 0 wrote a headers-only CSV and a negative value dropped that many worst-ranked rows -- both after a full sweep and both exiting 0. Reject < 1 before the sweep starts. The dgrad case study's K-outer selector collected eight tests while the prose promised six (-k LdsKOuter selects exactly the six), and its dispatch-policy command pointed at library/ as if it were under platform/ rather than beside it, so it failed at collection.

PR:11978 moved the conv kernels, tests and sweep driver out of platform/. run_one_dgrad.py still imported rocke.benchmark / rocke.instances.common and failed at startup; the case study's replay commands and two comment references pointed at the pre-move paths.

assistant-librarian · 2026-09-17T19:13:44Z

TheRock Submodule Bump Activity

Newest first

2026-09-17 — Included in TheRock via ROCm/TheRock#8289.

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

therock-pr-bot Bot added the Not ready to Review label Sep 10, 2026

assistant-librarian Bot added the organization: ROCm label Sep 10, 2026

aledudek force-pushed the users/aledudek/dgrad-wavescope-analysis branch from 4d45946 to d986356 Compare September 14, 2026 16:16

aledudek marked this pull request as ready for review September 14, 2026 19:04

aledudek requested review from a team as code owners September 14, 2026 19:04

aledudek force-pushed the users/aledudek/dgrad-wavescope-analysis branch from d986356 to 25cd02d Compare September 15, 2026 05:41

aledudek requested a lite review from Copilot September 15, 2026 05:42

Copilot started reviewing on behalf of aledudek September 15, 2026 05:43 View session

therock-pr-bot Bot removed the Not ready to Review label Sep 15, 2026

Copilot AI reviewed Sep 15, 2026

View reviewed changes

aledudek force-pushed the users/aledudek/dgrad-wavescope-analysis branch from f80028b to 8b8824d Compare September 15, 2026 11:09

bartekxk reviewed Sep 15, 2026

View reviewed changes

Comment thread ...rs/hip-kernel-provider/rocke/platform/python/rocke/benchmark/benchmark_implicit_gemm_conv.py Outdated

aledudek added 8 commits September 16, 2026 04:26

aledudek force-pushed the users/aledudek/dgrad-wavescope-analysis branch from 8df053f to 9938ac0 Compare September 16, 2026 09:39

bartekxk approved these changes Sep 16, 2026

View reviewed changes

aledudek enabled auto-merge (squash) September 16, 2026 11:39

Merge branch 'develop' into users/aledudek/dgrad-wavescope-analysis

f986787

aledudek merged commit 3fa0e12 into develop Sep 16, 2026
112 checks passed

aledudek deleted the users/aledudek/dgrad-wavescope-analysis branch September 16, 2026 14:22

Navigation Menu

perf(rocke/conv): add the pointwise fast path dgrad was missing - #11928

perf(rocke/conv): add the pointwise fast path dgrad was missing#11928
aledudek merged 9 commits into
developfrom
users/aledudek/dgrad-wavescope-analysis

aledudek commented Sep 10, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 10, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 10, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

assistant-librarian Bot commented Sep 17, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

aledudek commented Sep 10, 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 Sep 10, 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 Sep 10, 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

Suppressed comments (5)

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

assistant-librarian Bot commented Sep 17, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

aledudek commented Sep 10, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 10, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 10, 2026 •

edited

Loading