Skip to content

[MIOpen] Fix softmax kernel NaN failure when beta=0 (ROCM-26026) - #8426

Merged
SreecharanGundaboluAMD merged 3 commits into
developfrom
users/SreecharanGundaboluAMD/fix_softmax_beta
Jun 17, 2026
Merged

SreecharanGundaboluAMD merged 3 commits into
developfrom
users/SreecharanGundaboluAMD/fix_softmax_beta

Conversation

@SreecharanGundaboluAMD

@SreecharanGundaboluAMD SreecharanGundaboluAMD commented Jun 15, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

When beta=0 is passed to the MIOpen softmax kernel, the output buffer (y forward, dx backward) may contain uninitialized memory or NaN. The kernel was unconditionally reading and scaling these buffers by beta, which produces NaN under IEEE 754 semantics (NaN * 0.0 == NaN), poisoning the result even though beta=0 means the buffer should be ignored entirely. This fixes ROCM-26026 and adds a regression gtest to catch any future reintroduction.

Technical Details

Kernel fix (projects/miopen/src/kernels/MIOpenSoftmax.cpp)

Four accumulation sites shared the same pattern:

value = computed * alpha + CVT_FLOAT2ACCUM(output_buf[idx]) * CVT_FP32_2ACCUM(beta);

Each is replaced with a runtime guard:

value = computed * alpha +
        (beta != 0.0f ? CVT_FLOAT2ACCUM(output_buf[idx]) * CVT_FP32_2ACCUM(beta)
                      : FLOAT_ACCUM{0});

Sites fixed:

  • softmaxfwd scalar path (line ~261)
  • softmaxfwd vectorized path (line ~390)
  • softmaxbwd scalar path (line ~460)
  • softmaxbwd vectorized path (line ~529)

This matches the standard BLAS contract: when beta=0, the output buffer must not be read regardless of its contents. The branch is cheap relative to the memory load it avoids.

Regression test (projects/miopen/test/gtest/soft_max.cpp)

Added GPU_Softmax_BetaZeroNaN_FP32, a parameterized fixture with two test methods:

  • ForwardTest: pre-poisons the y output buffer with quiet_NaN, calls SoftmaxForward with alpha=1, beta=0, asserts all output elements are finite.

  • BackwardTest: pre-poisons the dx output buffer with quiet_NaN, calls SoftmaxBackward with alpha=1, beta=0, asserts all output elements are finite.

  • Smoke: all 3 algos (FAST, ACCURATE, LOG) × 2 modes (INSTANCE, CHANNEL) = 6 forward + 6 backward cases

Risk Assessment

Risk Level: 🟢 Low

Level Indicators
🟢 Low Single kernel file changed; fix is additive (guard added, no logic removed); only affects the beta=0 code path which was previously broken; no architectural changes

Impacted Components

  • projects/miopen/src/kernels/MIOpenSoftmax.cpp — forward and backward softmax GPU kernels (scalar and vectorized paths)
  • projects/miopen/test/gtest/soft_max.cpp — softmax gtest suite

Potential Side Effects

  • No behavioral change when beta != 0; all four sites are untouched for that case.
  • When beta == 0, the output buffer is now skipped — this is the correct BLAS contract. Previously broken inputs now produce correct results.

Mitigation Steps

  • Change is isolated to one kernel source file and four symmetric sites.
  • Regression test directly reproduces the failure mode and will catch any future reintroduction.
  • No new dependencies or API surface.

Test Plan

Added GPU_Softmax_BetaZeroNaN_FP32 gtest fixture that pre-poisons the output buffer with NaN before each kernel call. Without the fix, both ForwardTest and BackwardTest fail on every element. With the fix, both pass.

Run smoke tests:

./test_soft_max --gtest_filter="Smoke/GPU_Softmax_BetaZeroNaN_FP32*"

Test Result

  • Both ForwardTest and BackwardTest pass with the kernel fix applied across all algo/mode combinations. Confirmed that without the fix the tests fail on every output element with NaN. No regressions in existing SoftmaxCommon (FP16, BF16) tests.
  • Passes CI.

Submission Checklist

MIOpen's HIP softmax kernel (softmaxfwd and softmaxbwd) unconditionally reads
the output buffer (y[] / dx[]) in the alpha/beta blend step, even when the
caller passes beta == 0. Per IEEE-754, 0.0 * NaN = NaN. PyTorch's CTC loss
path allocates uninitialized workspace and calls SoftmaxForward(alpha=1,
beta=0) into it, causing stale NaN/Inf bytes to poison the log-probability
output.

The original OpenCL kernel guarded this read with `#if USE_BETA`; the OpenCL→HIP
port (PR #1722) removed it. Fix by adding a runtime `beta != 0.0f` ternary guard
at all four alpha/beta blend sites (forward and backward, CSR-Vector and
CSR-Stream paths) to skip reading the uninitialized output buffer when beta=0.

The GPU compiler converts the uniform warp-wide ternary to a conditional move
with no divergence cost.

Fixes ROCM-26026.

Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>

@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 with a few additions:

  • can we add a changelog item for this fix?
  • let's add a Risk Assessment category to this PR
  • can we add a test or tests to cover this scenario?

Thanks!

@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, thanks!

@SreecharanGundaboluAMD
SreecharanGundaboluAMD enabled auto-merge (squash) June 16, 2026 22:57
@SreecharanGundaboluAMD
SreecharanGundaboluAMD merged commit fb8157c into develop Jun 17, 2026
55 checks passed
@SreecharanGundaboluAMD
SreecharanGundaboluAMD deleted the users/SreecharanGundaboluAMD/fix_softmax_beta branch June 17, 2026 03:14
JonathanLichtnerAMD added a commit that referenced this pull request Jul 15, 2026
JIRA ID: ALMIOPEN-2312

## Motivation

`projects/miopen/CHANGELOG.md`'s "MIOpen 3.5.2 for ROCm 7.14.0" section
only had 4 entries, but 68 commits touched `projects/miopen` between the
`therock-7.13` tag and `release/therock-7.14`. Auditing that commit
range (via `git log --oneline
therock-7.13..remotes/origin/release/therock-7.14 -- projects/miopen`)
found 18 additional user-facing changes — new gfx950/gfx1250
architecture support, SystemDB refreshes, AI-heuristic improvements,
several correctness/crash bug fixes, and OpenCL solver removals — that
were never documented. This PR backfills those entries so the 7.14.0
changelog accurately reflects what shipped.

## Technical Details

Added 24 lines to the existing "MIOpen 3.5.2 for ROCm 7.14.0" section of
`projects/miopen/CHANGELOG.md`, merging new bullets into the existing
`Added`/`Changed`/`Resolved Issues` sub-sections (no new entries were
duplicated or reworded) and adding a new `Removed` sub-section for
solver removals. No existing content was altered.

Each new bullet was traced to a specific commit + PR:

| Entry | Commit | PR |
|---|---|---|
| Added gfx950 (MI350X/MI355X) SystemDB | `60d8e1e5175` | #8081 |
| Added initial gfx1250 architecture support | `d3f057b2f22` | #7587 |
| Added JSON-formatted performance logs (`MIOPEN_PERFORMANCE_LOGS`) |
`ffbde871260` | #5736 |
| Refreshed gfx942 SystemDB | `faf754a9ef6` | #8174 |
| Refined 2D AI heuristics (Tunanet/KTN) for gfx942 & gfx950 |
`c24ad415884` | #7109 |
| Fixed thread-safety issue in AI-heuristic model caches | `cd957402309`
| #9149 (cherry-pick of #9131) |
| Fixed CK grouped-conv solvers not always registered in host |
`c971d54d7ae` | #9105 |
| Fixed legacy CK buffer resource descriptor (gfx1250) | `1dee011b8b0` |
#8758 |
| Removed stale gfx942 SystemDB naive-solver entries | `f42668f6316` |
#7028 |
| Fixed softmax NaN when `beta=0` | `fb8157c8512` | #8426 |
| Fixed BN backward spatial variant 1 vectorized memory access bug |
`1351a7e8b10` | #6765 |
| Fixed use-after-free of `kernel_id` in `InitInvokerFactoryNCHW` |
`394ebffb50d` | #7922 |
| Fixed SoftMarginLoss forward reduce-kernel kernarg corruption |
`1e65a31570f` | #8065, #8135 |
| Fixed build with `MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF` | `e6288c3ca61`
| #6800 |
| Fixed install RPATH for MIOpenDriver / CK backend libs | `ef8a4cf2f93`
| #7485 |
| Fixed dangling lambda causing log-level 6 segfault | `190f9791737` |
#6953 |
| Enabled Windows long-path support for `addkernels` | `1f9d69724cf` |
#7387 |
| Removed OCL solver `ConvOclDirectFwd1x1` | `126c720cadd` | #6491 |
| Removed OCL solver `conv_ocl_dir2Dfwdgen` | `2d499ddf709` | #6633 |
| Removed OCL solvers `ConvOclBwdWrW1x1`, `ConvOclBwdWrW2` |
`b5d43a33dc6` | #6859 |
| Removed OCL solvers `OpTensorLeadingOnes`, `Op5dTensorGeneric` |
`8642c83b7de` | #7347 |

The remaining ~40 commits in the 68-commit range were deliberately
excluded (tests/CI/build/docs/internal refactors/1:1 OCL→HIP swaps with
no behavior change, per each commit's own stated scope) and are not
represented here.

## Risk Assessment

**Risk Level: 🟢 Low**

This is a documentation-only change to a single Markdown file. No code,
build configuration, tests, or public API surface is touched.

### Impacted Components

- `projects/miopen/CHANGELOG.md` only.

### Potential Side Effects

None expected. This does not change any runtime behavior; it only
documents changes that already shipped on `release/therock-7.14`.

### Mitigation Steps

None required beyond normal review — a reviewer familiar with the 7.14
release can spot-check a few entries against the linked commit hashes/PR
numbers for accuracy.

## Test Plan

Not applicable (no code changes). Verification was done by
cross-referencing every added bullet against a specific commit hash and
PR number pulled from `git log`, and by diffing the edited section
against `git show
remotes/origin/release/therock-7.14:projects/miopen/CHANGELOG.md` to
confirm no existing entry was duplicated or altered.

## Test Result

All 21 new bullets map 1:1 to a real commit/PR in the
`therock-7.13..release/therock-7.14` range for `projects/miopen`. No
pre-existing changelog text was modified — `git diff` shows a pure
addition (24 insertions, 0 deletions). Two borderline commits
(`a67834413e8` "Standalone build", `7c7cd5774a9` "enable warnings about
deprecated API/features") were deliberately left out pending author
judgment on whether they warrant an entry.
pschang-phy pushed a commit that referenced this pull request Jul 17, 2026
JIRA ID: ALMIOPEN-2312

## Motivation

`projects/miopen/CHANGELOG.md`'s "MIOpen 3.5.2 for ROCm 7.14.0" section
only had 4 entries, but 68 commits touched `projects/miopen` between the
`therock-7.13` tag and `release/therock-7.14`. Auditing that commit
range (via `git log --oneline
therock-7.13..remotes/origin/release/therock-7.14 -- projects/miopen`)
found 18 additional user-facing changes — new gfx950/gfx1250
architecture support, SystemDB refreshes, AI-heuristic improvements,
several correctness/crash bug fixes, and OpenCL solver removals — that
were never documented. This PR backfills those entries so the 7.14.0
changelog accurately reflects what shipped.

## Technical Details

Added 24 lines to the existing "MIOpen 3.5.2 for ROCm 7.14.0" section of
`projects/miopen/CHANGELOG.md`, merging new bullets into the existing
`Added`/`Changed`/`Resolved Issues` sub-sections (no new entries were
duplicated or reworded) and adding a new `Removed` sub-section for
solver removals. No existing content was altered.

Each new bullet was traced to a specific commit + PR:

| Entry | Commit | PR |
|---|---|---|
| Added gfx950 (MI350X/MI355X) SystemDB | `60d8e1e5175` | #8081 |
| Added initial gfx1250 architecture support | `d3f057b2f22` | #7587 |
| Added JSON-formatted performance logs (`MIOPEN_PERFORMANCE_LOGS`) |
`ffbde871260` | #5736 |
| Refreshed gfx942 SystemDB | `faf754a9ef6` | #8174 |
| Refined 2D AI heuristics (Tunanet/KTN) for gfx942 & gfx950 |
`c24ad415884` | #7109 |
| Fixed thread-safety issue in AI-heuristic model caches | `cd957402309`
| #9149 (cherry-pick of #9131) |
| Fixed CK grouped-conv solvers not always registered in host |
`c971d54d7ae` | #9105 |
| Fixed legacy CK buffer resource descriptor (gfx1250) | `1dee011b8b0` |
#8758 |
| Removed stale gfx942 SystemDB naive-solver entries | `f42668f6316` |
#7028 |
| Fixed softmax NaN when `beta=0` | `fb8157c8512` | #8426 |
| Fixed BN backward spatial variant 1 vectorized memory access bug |
`1351a7e8b10` | #6765 |
| Fixed use-after-free of `kernel_id` in `InitInvokerFactoryNCHW` |
`394ebffb50d` | #7922 |
| Fixed SoftMarginLoss forward reduce-kernel kernarg corruption |
`1e65a31570f` | #8065, #8135 |
| Fixed build with `MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF` | `e6288c3ca61`
| #6800 |
| Fixed install RPATH for MIOpenDriver / CK backend libs | `ef8a4cf2f93`
| #7485 |
| Fixed dangling lambda causing log-level 6 segfault | `190f9791737` |
#6953 |
| Enabled Windows long-path support for `addkernels` | `1f9d69724cf` |
#7387 |
| Removed OCL solver `ConvOclDirectFwd1x1` | `126c720cadd` | #6491 |
| Removed OCL solver `conv_ocl_dir2Dfwdgen` | `2d499ddf709` | #6633 |
| Removed OCL solvers `ConvOclBwdWrW1x1`, `ConvOclBwdWrW2` |
`b5d43a33dc6` | #6859 |
| Removed OCL solvers `OpTensorLeadingOnes`, `Op5dTensorGeneric` |
`8642c83b7de` | #7347 |

The remaining ~40 commits in the 68-commit range were deliberately
excluded (tests/CI/build/docs/internal refactors/1:1 OCL→HIP swaps with
no behavior change, per each commit's own stated scope) and are not
represented here.

## Risk Assessment

**Risk Level: 🟢 Low**

This is a documentation-only change to a single Markdown file. No code,
build configuration, tests, or public API surface is touched.

### Impacted Components

- `projects/miopen/CHANGELOG.md` only.

### Potential Side Effects

None expected. This does not change any runtime behavior; it only
documents changes that already shipped on `release/therock-7.14`.

### Mitigation Steps

None required beyond normal review — a reviewer familiar with the 7.14
release can spot-check a few entries against the linked commit hashes/PR
numbers for accuracy.

## Test Plan

Not applicable (no code changes). Verification was done by
cross-referencing every added bullet against a specific commit hash and
PR number pulled from `git log`, and by diffing the edited section
against `git show
remotes/origin/release/therock-7.14:projects/miopen/CHANGELOG.md` to
confirm no existing entry was duplicated or altered.

## Test Result

All 21 new bullets map 1:1 to a real commit/PR in the
`therock-7.13..release/therock-7.14` range for `projects/miopen`. No
pre-existing changelog text was modified — `git diff` shows a pure
addition (24 insertions, 0 deletions). Two borderline commits
(`a67834413e8` "Standalone build", `7c7cd5774a9` "enable warnings about
deprecated API/features") were deliberately left out pending author
judgment on whether they warrant an entry.
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