Skip to content

[MIOpen] fix bn fwd training edge case#2193

Closed
bghimireamd wants to merge 9 commits into
developfrom
users/bghimire/fix_bn_variance_issue
Closed

[MIOpen] fix bn fwd training edge case#2193
bghimireamd wants to merge 9 commits into
developfrom
users/bghimire/fix_bn_variance_issue

Conversation

@bghimireamd
Copy link
Copy Markdown
Contributor

@bghimireamd bghimireamd commented Oct 21, 2025

Motivation

./bin/MIOpenDriver bnorm -n 1 -c 256 -H 1 -W 1 -m 1 --forw 1 -b 0 -r 1 -s 1 --layout NCHW

This edge case was causing issue ( hardcode variance to zero for single samples). This PR is to fix this edge case.

Technical Details

Test Plan

Test Result

Submission Checklist

@bghimireamd bghimireamd marked this pull request as ready for review October 21, 2025 01:05
@bghimireamd bghimireamd requested a review from a team as a code owner October 21, 2025 01:05
Copy link
Copy Markdown
Contributor

@randyspauldingamd randyspauldingamd left a comment

Choose a reason for hiding this comment

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

LGTM

Copy link
Copy Markdown
Contributor

@JonathanLichtnerAMD JonathanLichtnerAMD left a comment

Choose a reason for hiding this comment

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

LGTM but could we add a test for this?

Comment thread projects/miopen/src/kernels/MIOpenBatchNormFwdTrainSpatialHIP.cpp Outdated
@EnricoDeg
Copy link
Copy Markdown
Contributor

LGTM

@bghimireamd
Copy link
Copy Markdown
Contributor Author

Closing: Test case N=1,H=1,W=1 is invalid for training mode (PyTorch rejects with "Expected more than 1 value per channel when training"). Proper fix is API-level input validation, not kernel changes. Will address in separate PR.

bghimireamd added a commit that referenced this pull request Dec 5, 2025
## Motivation

Adds validation in both hipDNN frontend and MIOpen to reject batch norm
forward training and backward shapes where N × spatial_dimensions ≤ 1
(e.g., {1, 256, 1, 1}).

## Technical Details

This supersedes the [closed kernel-level
approach](#2193) .

## Test Plan

This prevents undefined behavior, aligns with PyTorch, and includes 22
hipDNN frontend tests. Invalid test cases removed from backend fixtures.

## Test Result

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
assistant-librarian Bot pushed a commit to ROCm/MIOpen that referenced this pull request Dec 5, 2025
[MIOpen] [HipDNN] [ALMIOPEN-493] Handle Batch norm invalid
 cases (#3030)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

## Motivation

Adds validation in both hipDNN frontend and MIOpen to reject batch norm
forward training and backward shapes where N × spatial_dimensions ≤ 1
(e.g., {1, 256, 1, 1}).

## Technical Details

This supersedes the [closed kernel-level
approach](ROCm/rocm-libraries#2193) .

## Test Plan

This prevents undefined behavior, aligns with PyTorch, and includes 22
hipDNN frontend tests. Invalid test cases removed from backend fixtures.

## Test Result

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
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.

4 participants