Skip to content

fix(miopen): fix batch-split miscalculation for large WrW tensors (ROCM-26011) - #11188

Merged
SreecharanGundaboluAMD merged 1 commit into
developfrom
users/SreecharanGundaboluAMD/fixASM_kernels_conv_1x1
Aug 31, 2026
Merged

SreecharanGundaboluAMD merged 1 commit into
developfrom
users/SreecharanGundaboluAMD/fixASM_kernels_conv_1x1

Conversation

@SreecharanGundaboluAMD

Copy link
Copy Markdown
Contributor

JIRA ID : ROCM-26011

Motivation

nn.Conv2d(96, 64, kernel_size=1) on MI350X with a [64, 96, 512, 512] fp32 input silently
returns weight.grad == 0.0 (every element exactly zero) after backward(), while
bias.grad and a mathematically-equivalent F.linear on the reshaped tensor both produce
correct, non-zero gradients. No warning or error is raised — the model trains, but the conv1x1
layer never receives gradients (ROCM-26011).

The root cause is in igemm_split_batch_size(), which decides how many dispatches to split a
convolution's batch across so no single dispatch addresses more than the ~4GB tensor limit.
For batch sizes in (max_n, 2*max_n), the split-count search returned a degenerate divisor of
1 ("no split needed") even though the tensor genuinely exceeded the 4GB limit. This let
ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC launch a single dispatch that addressed more data
than the hardware limit allows, silently producing near-zero weight gradients instead of
failing loudly.

Technical Details

In projects/miopen/src/include/miopen/conv/asm_implicit_gemm.hpp, the split-multiplier
search used:

size_t m = std::ceil(n / max_n);

n and max_n are both integer types, so n / max_n truncates via integer division.

The fix replaces the float-ceil idiom with an integer-safe ceiling division:

size_t m = (static_cast<size_t>(n) + max_n - 1) / max_n;

Risk Assessment

Risk Level: 🟡 Medium

The diff is a single line, but it sits in a core MIOpen dispatch path (igemm_split_batch_size)
that gates numerical correctness — not just performance — for
ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC across every batch size that falls in the affected
(max_n, 2*max_n) range.

Potential Side Effects

  • Convolutions whose batch size previously fell in the buggy (max_n, 2*max_n) range will now
    be correctly split into multiple dispatches. This changes solver behavior (more dispatches,
    slightly different perf characteristics) for those specific shapes, but corrects silently
    wrong output to numerically correct output.

Mitigation Steps

  • Previously failing shape passes.
  • Change is minimal and localized to the split-count computation, with no change to the
    surrounding search algorithm's structure, keeping regression surface small.

Test Plan

  • Reproduced the reported bug with a standalone repro matching the ticket's nn.Conv2d(96, 64, kernel_size=1) / [64, 96, 512, 512] fp32 shape, and additionally swept batch sizes spanning
    n=43..192 with multiple channel configurations to cover the full (max_n, 2*max_n) boundary
    range where the split-count search previously degenerated to m=1. Compared MIOpen
    ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC weight-gradient output against a reference
    (F.linear-equivalent / naive) computation before and after the fix.

Test Result

  • Before the fix, all swept shapes in the affected range produced weight.grad that was zero or
    otherwise numerically incorrect. After the fix, all 11+ verified shapes (n=43..192, multiple
    channel configs) match the reference computation to ~1e-6/1e-7 precision, including the exact
    shape from the reported ticket. No regressions observed in shapes outside the affected range,
    where the split-count computation was already correct.

Submission Checklist

igemm_split_batch_size() decides how many dispatches to split a
convolution's batch across so that no single dispatch needs to address
more than ~4GB of tensor data. The split-count search computed
`std::ceil(n / max_n)` where n and max_n are both integers, so the
division truncated before ceil ever saw it (e.g. n=43, max_n=42 gives
1, not the required 2). Whenever the batch size n fell in
(max_n, 2*max_n), this produced a degenerate starting divisor of 1,
which the subsequent search accepted immediately (everything is
divisible by 1) and returned "no split needed" even though the tensor
exceeded the 4GB limit that had just been computed.

This let ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC launch with a
single >4GB-addressed dispatch, producing silently wrong (near-zero)
weight gradients instead of failing loudly. Verified against 11+
shapes spanning the affected range (n=43..192, multiple channel
configs) that all previously produced incorrect results and now
verify correctly to ~1e-6/1e-7 precision.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
@therock-pr-bot

therock-pr-bot Bot commented Aug 24, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ⚠️ Warning Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/miopen/src/include/miopen/conv/asm_implicit_gemm.hpp; no test file found
🔎 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

Copy link
Copy Markdown

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

@JonathanLichtnerAMD JonathanLichtnerAMD left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

LGTM

@SreecharanGundaboluAMD
SreecharanGundaboluAMD merged commit 34833f8 into develop Aug 31, 2026
193 of 203 checks passed
@SreecharanGundaboluAMD
SreecharanGundaboluAMD deleted the users/SreecharanGundaboluAMD/fixASM_kernels_conv_1x1 branch August 31, 2026 12:39
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

JonathanLichtnerAMD added a commit that referenced this pull request Sep 2, 2026
## Summary

Backfills the **MIOpen 3.6.1 / ROCm 10.1.0** changelog section with
externally-visible changes that landed since the `release/therock-10.0`
branch but weren't yet documented. All 43 MIOpen commits in that range
were reviewed; CI/build-infra, test-only, docs, version-bump, and
already-documented changes were excluded.

21 entries added:

### Added
* [RNN] Fused single-layer LSTM forward-inference kernel, gated behind
`MIOPEN_DEBUG_RNN_FUSED_INFERENCE` (#8915)
* [Conv] Naive reference kernel enabling 3D INT8 forward convolution
(#10051)
* [Conv] `ConvHipConv` solver backed by vendored hipconv kernels (#9746)
* [Conv] float8 vector type support + deduplicated math kernels (#10130)
* [Conv] Validated gfx1151 and gfx1200 (Navi) SystemDBs (#10050)
* [Conv] `MIOPEN_DEBUG_DISABLE_SYSTEM_DB` /
`MIOPEN_DEBUG_DISABLE_USER_DB` env vars (#10728)
* [Conv] CK depthwise convolution on RDNA wave32 GPUs (#9373)

### Changed
* [Conv] Winograd Rage RxS f2x3 on gfx950 + gfx942 v4_6_1/v4_9_1 kernels
(#9014)
* [Conv] Updated 3D AI heuristics for gfx942 & gfx950 (#8954)

### Optimized
* [Conv] Single GEMM for point (1x1) output convolution, all directions,
2D/3D (#10912)
* [Conv] Naive backward-data kernel work group widened to 1024 (#10684)

### Resolved Issues
* [Conv] Batch-split miscalculation on large WrW tensors — ROCM-26011
(#11188)
* [Conv] Conv failure on XNACK-unsupported gfx9 APUs (gfx90c) (#11211)
* [BatchNorm] Device heap-overflow in fwd spatial `FinalMeanVariance` —
ALMIOPEN-2424 (#11093)
* [Conv] Deterministic `SetNextValue` polarity for grouped Bwd/Wrw CK
solvers (#10240)
* [LRN][LayerNorm] Grid alignment + OOB threads on gfx1250 (#10251)
* [Conv] Stale gfx950 grouped-wrw SystemDB entries (#11130)
* [Conv][Windows] System-DB path resolution from the MIOpen module
(#10838)
* [Conv] System-DB fallback so MI308X (gfx942) resolves a usable DB
(#8802)
* JSON perf logs (`MIOPEN_PERFORMANCE_LOGS`) dropping the last solver
during Find (#11166)
* Build failure against GCC 15 libstdc++ under C++20 (#9642)

## Notes

Deliberately excluded: internal/default-off changes (#10219, #10088,
#10649, #10231), a near-duplicate of an existing 10.0.0 entry (#10188),
and CI/build-infra/test-only/docs commits.

Docs-only change; no code impact.

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

---------

Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants