Skip to content

fix(miopen): reject reduce tensors that do not fit the int32 legacy kernels - #12874

Merged
averinevg merged 13 commits into
developfrom
users/evaverin/ALMIOPEN-2718/check-reduce-dims
Oct 9, 2026
Merged

averinevg merged 13 commits into
developfrom
users/evaverin/ALMIOPEN-2718/check-reduce-dims

Conversation

@averinevg

@averinevg averinevg commented Oct 1, 2026 •

Copy link
Copy Markdown
Contributor

Summary

The legacy CK reduction kernels take sizes as int and use 32-bit offsets, but miopenReduceTensor
narrowed size_t values with static_cast<int>, silently producing wrong results for large
tensors. Such tensors are now rejected with miopenStatusBadParm before any kernel is launched.

JIRA ID : ALMIOPEN-2718

Risk Assessment

Risk level 3. Host-side validation in one file, but a public API contract change: oversized tensors
now fail with miopenStatusBadParm, the sizing queries reject rank > 6, and reduction errors return
miopenStatusBadParm instead of miopenStatusUnknownError. Kernel sources are unchanged.

ASIC Coverage

Architecture-independent host-side change; passing PR CI is sufficient.

Testing Summary

  • New Smoke/GPU_ReduceShapeLimits_FP32 / _FP16 suites: rejection of rank > 6, lengths, strides
    and byte spans above INT32_MAX, and a single fp16 reduction whose padded length overflows; sizing
    queries stay permissive; a length-1 dimension with a stride above INT32_MAX still reduces correctly.

Testing Checklist

  • Reduce shape limit tests - ./bin/test_reduce_shape_limits - Devices: gfx90a - Status: Passed
  • PR CI - GitHub PR checks - Status: Pending

Technical Changes

  • ReduceTensor checks that lengths, strides (except of length-1 dims), byte spans, element counts,
    grid sizes and the multi-block padded length fit into int.
  • Shared descriptor validation for GetWorkspaceSize, GetIndicesSize and ReduceTensor.
  • LogCmdRedux logs lengths as size_t; CHANGELOG updated.

🤖 Generated with Claude Code

averinevg and others added 9 commits September 30, 2026 18:35
The lambda that stringifies the input lengths for the logged
MIOpenDriver command took an int, so lengths above INT32_MAX were
printed truncated (e.g. 2^31 as -2147483648) and the logged command
did not reproduce the failing call.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…(ALMIOPEN-2718)

The legacy reduction kernels take lengths, strides, element counts and
grid sizes as int and address memory through 32-bit byte offsets, but
ReduceTensor narrowed the size_t values with static_cast<int>, silently
producing wrong results for large tensors.

- ReduceTensor now rejects with miopenStatusBadParm any input or output
  tensor whose lengths or strides exceed INT32_MAX, or whose byte span
  does, plus element counts and grid sizes above INT32_MAX. Strides of
  length-1 dimensions are exempt: they never contribute to an offset,
  and the kernel is passed stride 1 for them.
- GetWorkspaceSize does not do the int32 checks, since the RNN uses it
  to size reductions that may never run.
- Descriptor validation (rank match, rank <= 6, output lengths) is
  shared by GetWorkspaceSize, GetIndicesSize and ReduceTensor, so the
  sizing queries now also reject rank > 6.
- All errors thrown from reducetensor.cpp use miopenStatusBadParm
  instead of miopenStatusUnknownError.
- Fix int overflow in get_padding_need (warpwise branch).
- Remove the unused GetIndicesTypeSize().

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
New Smoke/GPU_ReduceShapeLimits_FP32 suite in its own file. It checks
that ReduceTensor rejects with miopenStatusBadParm tensors that do not
fit the int32 legacy kernels (rank > 6, input/output lengths or strides
above INT32_MAX, byte span above INT32_MAX), that GetWorkspaceSize and
GetIndicesSize reject rank > 6 but stay permissive otherwise (RNN
sizing path), and that a length-1 dimension with a stride above
INT32_MAX still reduces correctly on the GPU.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…-2718)

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ALMIOPEN-2718)

For a multi-block reduction with a single output the number of blocks
is not capped, so the padded reduced length the kernel computes as int
(reduceSizePerBlock * BlkGroupSize, about twice the reduced length) can
exceed INT32_MAX even when the tensor itself fits: e.g. a packed fp16
tensor of 2^30 - 1 elements reduced to one value. The kernel then
reduces nothing and silently returns a wrong result.

get_padding_need now rejects such reductions with miopenStatusBadParm.
ToIntOrThrow moves to an unnamed namespace ahead of detailDynamic so
get_padding_need can use it.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…PEN-2718)

Mention single reductions of about 2^30 or more fp16, bf16 or int8
elements under Resolved Issues, and move the rank > 6 rejection in the
sizing queries and the UnknownError -> BadParm status change to a new
Changed section, since they are behavior changes rather than fixes.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ests (ALMIOPEN-2718)

Add a Smoke/GPU_ReduceShapeLimits_FP16 suite with a packed fp16 tensor
of INT32_MAX / 2 elements reduced to one value: the tensor fits into
INT32_MAX bytes, but the padded length of the multi-block kernel does
not. The test body is shared between the FP32 and FP16 suites.

Also replace the magic numbers with named limits (the fp32 byte span
case now uses the smallest rejected row length) and clarify the
comments on the sizing cases and on the null workspace that keeps a
regression from launching kernels.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
@therock-pr-bot

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

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

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

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

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

🙋 Wish to Override Policy?

@therock-pr-bot

therock-pr-bot Bot commented Oct 1, 2026

Copy link
Copy Markdown

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

@averinevg
averinevg enabled auto-merge (squash) October 8, 2026 12:39
@averinevg
averinevg disabled auto-merge October 8, 2026 12:40
@averinevg
averinevg enabled auto-merge (squash) October 8, 2026 12:40

@iq136boy iq136boy 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

@averinevg
averinevg merged commit 681bc9e into develop Oct 9, 2026
98 checks passed
@averinevg
averinevg deleted the users/evaverin/ALMIOPEN-2718/check-reduce-dims branch October 9, 2026 10:39
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

mattwill-amd pushed a commit that referenced this pull request Oct 14, 2026
JIRA ID: ALMIOPEN-2913

## Motivation

Only 5 of the 30 commits that landed against `projects/miopen` since the
3.6.2 version bump (`2d20a8f775e`, #12211) were reflected in
`CHANGELOG.md`. The remaining 25 were reviewed and the 15
externally-visible ones documented, so the ROCm 10.2.0 section
accurately reflects what shipped.

14 entries added (two of the fixes share one entry), plus a
duplicate-heading fix.

### Added
* [Conv] NHWC (channel-last) support in the GEMM convolution solvers,
with new im2col/col2im kernels (#8693)
* [Conv] Cross-architecture AI heuristics for immediate mode
(`MIOPEN_FIND_MODE=2`), extending data-driven selection to Navi (#9680)
* [Conv] Missing 3x3 stride-1 pad-1 tile-28 CK depthwise backward-data
instance — ALMIOPEN-2423 (#11271)

### Changed
* [Conv] `ConvHipConv` held out of the AI perf-config picker until its
model is retrained — ALMIOPEN-2494 (#13337)
* [Conv] CK byte-span gate lifted for forward and backward-data, now
that CK fixes the 32-bit offset overflow itself — ALMIOPEN-2645 (#12396)

### Optimized
* [Conv] `ConvDirectNaiveConv`: per-direction kernel files, float
accumulators, spatial tiling, bfloat16 path — ALMIOPEN-2485/2486 (#7529)
* Kernel cache reuses cached SQLite connections instead of reopening per
load/store — ALMIOPEN-2598 (#11933)

### Resolved Issues
* [Conv] Incorrect backward-weights results from the gfx950 direct
kernel at `waves_q == 1` (missing cross-lane ordering on LDS staging
reuse) — ROCM-31508 (#12406)
* [Conv] Up to 8x slower hipconv `direct`/`direct_l1`/`direct_wgrad`
under ROCm 10.1's full AMDGCN LTO default; fixed with `-fno-offload-lto`
— ROCM-31334 (#12328)
* [Conv] gfx942 multi-node training throughput regression from backward
GEMM on small-batch 1x1 NHWC shapes — ROCM-31155 (#13232)
* [Conv] Invalid CK perf-config entries in the gfx1201 and gfx1100
system databases — ROCM-29026 (#12088, #12117)
* [BatchNorm] JIT compilation failures on non-GFX9 targets; Welford DPP
assembly restricted to GFX9 — ALMIOPEN-2838 (#12889)
* Out-of-bounds write in `miopenLoadSolution` when a serialized solution
carries more than three work dimensions — ROCM-31082 (#12242)
* `uint64_t` typedef conflict between `miopen_cstdint.hpp` and system
`<stdint.h>` on LP64 Linux (#12623)

## Technical Details

Single-file change to `projects/miopen/CHANGELOG.md`.

The 10.2.0 section had **two** `### Resolved Issues` headings. The two
entries under the trailing one (`ConvHipConv` gfx1250 crash/LDS race,
and the perf-db reload config loss) were merged into the first, and the
trailing heading removed — this accounts for the 2 removed bullets in
the diff. Section order is now `Resolved Issues → Added → Changed →
Optimized → Removed`, matching the 3.6.1 section. `Optimized` is new to
this release section.

Already documented, so untouched: #12874 (reduce int32 rejection),
#12000 / #13016 / #13213 (hipconv v0.3.1 / v0.3.2 / v0.4.0), #11761
(gfx803 removal).

## Notes

Deliberately excluded as not externally visible (10 commits):

* **hipDNN convolution forwarding (#12276)** —
`MIOPEN_ENABLE_HIPDNN_WRAPPER` defaults to `OFF`
(`projects/miopen/CMakeLists.txt:162`), so no shipped build compiles the
forwarding code and the runtime `MIOPEN_HIPDNN_FORWARDING` flag has
nothing to enable.
* **Test-only**: #13289 (gfx1250 depthwise test applicability,
ALMIOPEN-2841), #11610 (hipDNN parity harness), #12660 (Jenkins status
post).
* **Internal kernel-math cleanup, generated ISA unchanged**: #13312
(drop direct `__ocml` calls), #13190 (remove broken single-argument
`__half` overloads).
* **Add/revert pairs, net no change**: #12466 + #13066 (gfx1101 Navi
system DBs), #12496 + #12800 (gfx1250 in the dev image arch lists).

## Risk Assessment

Documentation-only change to a single Markdown file; no code, build, or
runtime impact.

**Risk Level: 🟢 Low**

## Test Plan

None required — Markdown-only. Verified that the entry topics are not
already documented elsewhere in the file, that no content was lost when
merging the duplicate heading, and that every new entry maps to a commit
in `2d20a8f775e..origin/develop -- projects/miopen`.

## Test Result

16 bullets added, 2 relocated — 14 net-new entries covering 15 commits.
Duplicate heading removed; section ordering matches 3.6.1.

Co-authored-by: Claude <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