Skip to content

fix(miopen): pick the FP8 exponent bias from the device at run time - #14627

Merged
yassinsolim merged 8 commits into
developfrom
users/ysoliman/miopen-fp8-runtime-bias
Nov 4, 2026
Merged

yassinsolim merged 8 commits into
developfrom
users/ysoliman/miopen-fp8-runtime-bias

Conversation

@yassinsolim

Copy link
Copy Markdown
Member

Proposed changes

Closes #14375.

MIOpen picked MIOPEN_FP8_IEEE_EXPONENT_BIAS once at configure time from GPU_TARGETS. A host build for several targets that includes gfx950 or gfx12 got the IEEE bias everywhere, so on gfx942 the hipBLASLt FP8 and BF8 GEMM tests read FNUZ data with the wrong bias and failed (errors around 0.1 to 0.17). This PR picks the bias at run time from the device instead.

  • UsesIeeeExponentBias(device_name) in target_properties.hpp returns true for gfx950 and gfx12, the same targets as MIOPEN_CK_USE_OCP_FP8_TARGETS.
  • GetDataTypeKBP and GetDataTypeKernelParams take the device name and pass the bias for that device to kernel compiles. The one-argument versions are gone, and every caller passes its handle or context device name.
  • check_numerics, the naive conv solver and the MHA solvers use the helper in place of the config.h macro.
  • The CMake seed in MiopenCkCompileDefs.cmake, the MIOPEN_FP8_IEEE_EXPONENT_BIAS option and its config.h entry are removed. config.h is installed, so the define leaves the public headers.
  • Host FP8 types in hip_float8.hpp read the bias from a run-time setting (FNUZ by default), and kernel builds still get it as a -D flag. The conv solver test fixtures set the host bias from the device through a new SetHostFp8ExponentBias<T>() in conv_test_base.hpp. MIOpenDriver sets it too, but ParseBaseArg doesn't accept convfp8 or convbfp8 yet, so that path can't run.
  • The first commit is newling's 79a9d61 cherry-picked unchanged (thanks James). It makes the hipBLASLt GEMM test use explicit FNUZ host types and adds CPU tests for that encoding.

What could break:

  • Builds with CK off never got the seed, so gfx950 and gfx12 used bias 0 there. They now use 1, same as CK-on builds.
  • HIP's FP8 header treats gfx1170 and gfx1171 as OCP, but this rule gives them bias 0, same as MIOPEN_CK_USE_OCP_FP8_TARGETS today. I haven't tested MIOpen FP8 on those targets.
  • Out-of-tree code that read MIOPEN_FP8_IEEE_EXPONENT_BIAS from config.h, or called the one-argument GetDataTypeKBP, won't build.

Testing, one CK-on build for gfx942 and gfx950, run on MI325X and MI355X, with develop built the same way:

  • On gfx942 develop fails all 34 hipBLASLt FP8 and BF8 GEMM cases. They pass with this PR, along with the two new CPU encoding tests.
  • On gfx950 the GEMM tests skip by design (gfx94X only).
  • GPU_Fp8Conversion_FP8.DefaultAndSoftwareFallback (gfx950) and GPU_Bfp8Conversion_BFP8.DefaultAndSoftwareFallback (both) fail the same way on develop, on the NaN element of the ieee format, so I left them alone.
  • The FP8 conv solver tests are disabled upstream. Forced on, they fail on develop and here before checking any values, since the naive FP8 kernels never get the names the solver looks up (fp8_naive_conv.cpp uses FWD_KERNEL_NAME and friends, and nothing defines them) and FP8 GEMM conv needs rocBLAS gemm_ex3, which MIOpen only uses below rocBLAS 5. Their kernel build options do show the new bias, 0 on gfx942 and 1 on gfx950 (develop passes 1 on both).
  • A CK-off build compiles.

Checklist

  • I have added automated tests relevant to the introduced functionality
  • I have sufficient test coverage for the changes, and code coverage hasn't decreased as a result of my PR
  • I have ran the tests, and they are all passing locally
  • I have added relevant documentation for the changes
  • I have removed the stale documentation which is no longer relevant after this pull request
  • I have ran make format & make check_format to ensure the changes have been formatted

newling and others added 3 commits October 28, 2026 15:31
MIOpen chose MIOPEN_FP8_IEEE_EXPONENT_BIAS once at configure time from
GPU_TARGETS. A host build for several targets that included gfx950 or
gfx12 therefore used the IEEE bias on gfx942 as well, which broke the
FP8 and BF8 hipBLASLt GEMM tests there (issue #14375).

- Add miopen::UsesIeeeExponentBias(device_name): true for gfx950 and
  gfx12, matching MIOPEN_CK_USE_OCP_FP8_TARGETS.
- GetDataTypeKBP and GetDataTypeKernelParams take the device name and
  emit the bias for that device. Update every caller.
- check_numerics, the naive conv solver and the MHA solvers use the
  device helper instead of the config.h macro.
- Remove the CMake seed, the MIOPEN_FP8_IEEE_EXPONENT_BIAS option and
  its config.h entry.
- Host FP8 types in hip_float8.hpp read the bias from a run-time
  setting. MIOpenDriver and the FP8 solver test fixture set it from the
  device.
The FP8 conv solver tests generate FP8 data on the host and check the
kernel output against a reference. Now that kernels get the exponent
bias of the device they run on, the host values need the same bias, or
gfx950 kernels read the host bits as values twice as large.

Add SetHostFp8ExponentBias<T>() to conv_test_base.hpp and call it from
the forward, backward data and backward weights fixtures. Drop the call
from ConvFwdSolverTestF8 in solver_f8.hpp, which no test instantiates.
Copilot AI balanced review requested due to automatic review settings October 29, 2026 00:17
@yassinsolim
yassinsolim requested review from a team as code owners October 29, 2026 00:17
@therock-pr-bot

therock-pr-bot Bot commented Oct 29, 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 29, 2026 •

Copy link
Copy Markdown

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

Copilot AI 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.

🟡 Changes recommended

The added tests leave device-dependent bias selection untested across gfx942, gfx950, and gfx1200, including the same-process 0→1→0 sequence needed to catch multi-target regressions.

1 open issue
What changed in this PR

This PR addresses #14375 by selecting MIOpen’s FP8 exponent bias from the runtime device rather than the build’s target list, correcting mixed-target builds running on gfx942.

Changes:

  • Pass device-specific bias settings through kernel compilation and solver applicability checks.
  • Replace the build-wide bias option with runtime host configuration.
  • Use explicit FNUZ types in hipBLASLt GEMM tests and add CPU encoding checks.
File Description
projects/​miopen/​test/​gtest/​solver_wrw.hpp Initialize host FP8 bias for weight-gradient tests.
projects/​miopen/​test/​gtest/​solver_bwd.hpp Initialize host FP8 bias for backward tests.
projects/​miopen/​test/​gtest/​hipblaslt_gemm.cpp Use fixed FNUZ types and add encoding tests.
projects/​miopen/​test/​gtest/​conv_test_base.hpp Add device-aware host bias initialization.
projects/​miopen/​src/​tensor.cpp Pass device names to tensor kernel options.
projects/​miopen/​src/​solver/​pooling/​forwardNd.cpp Use device-aware forward pooling options.
projects/​miopen/​src/​solver/​pooling/​forwardNaive.cpp Use device-aware naive pooling options.
projects/​miopen/​src/​solver/​pooling/​forward2d.cpp Use device-aware 2D forward pooling options.
projects/​miopen/​src/​solver/​pooling/​backwardNd.cpp Use device-aware backward pooling options.
projects/​miopen/​src/​solver/​pooling/​backward2d.cpp Use device-aware 2D backward pooling options.
projects/​miopen/​src/​solver/​mha/​mha_solver_forward.cpp Gate FP8 applicability using device bias.
projects/​miopen/​src/​solver/​mha/​mha_solver_backward.cpp Gate FP8 applicability using device bias.
projects/​miopen/​src/​solver/​conv/​conv_direct_naive_conv.cpp Select naive convolution bias per device.
projects/​miopen/​src/​solver/​adam/​transformers_adam_w.cpp Pass device name to datatype options.
projects/​miopen/​src/​solver/​adam/​adam.cpp Pass device name to datatype options.
projects/​miopen/​src/​kernels/​hip_float8.hpp Support runtime host bias while retaining kernel flags.
projects/​miopen/​src/​include/​miopen/​utility/​transposing_solver.hpp Use device-aware transpose kernel options.
projects/​miopen/​src/​include/​miopen/​target_properties.hpp Add the device-to-bias helper.
projects/​miopen/​src/​include/​miopen/​datatype.hpp Require device names when generating datatype options.
projects/​miopen/​src/​hip/​utilocl.cpp Propagate device names through utility kernel options.
projects/​miopen/​src/​gpu_conv_reference.cpp Compile reference kernels with device-specific bias.
projects/​miopen/​src/​conv/​problem_description.cpp Use device-aware convolution datatype options.
projects/​miopen/​src/​CMakeLists.txt Remove the build-wide bias option.
projects/​miopen/​src/​check_numerics.cpp Select numerics-check bias per device.
projects/​miopen/​include/​miopen/​config.h.in Remove the installed bias configuration define.
projects/​miopen/​driver/​driver.hpp Initialize host FP8 bias from the driver handle.
projects/​miopen/​cmake/​MiopenCkCompileDefs.cmake Remove target-list-based bias seeding.
projects/​miopen/​CHANGELOG.md Document the runtime fix and removed option.

Review effort: Balanced

@yassinsolim
yassinsolim marked this pull request as draft October 29, 2026 15:16
Check that GetDataTypeKBP and GetDataTypeKernelParams pick bias 0 for gfx942
and 1 for gfx950 and gfx1200, and that SetHostFp8ExponentBias switches the
host bias 0 to 1 to 0 in one process. SetHostFp8ExponentBias gets an overload
that takes the device name so the test does not need a handle.
aalabsi-amd added a commit that referenced this pull request Oct 30, 2026
…#13107)

JIRA: ALMIOPEN-2602
## Summary

Jenkins runs every `miopen_gtest` case whose name contains
`FP32`/`FP16`/`BFP16`/`TF32`/`NONE`. The GitHub CI tiers in
`test_categories.yaml` are a whitelist, so about 21,000 cases in 191
suites were never selected by any tier. This PR adds the
missing families and tiers them by measured time. It also quarantines
the failures that the new coverage exposed: each
exclusion is scoped to the arch/OS where it fails and carries a ticket.

## Changes (`projects/miopen/test/gtest/test_categories.yaml`)

### Coverage added
| Where | Change |
|---|---|
| `positive_patterns` | **+30** prefixed families: `GPU_Adam`,
`GPU_AmpAdam`, `GPU_AddLayerNorm`, `GPU_LayerNormBwdTest`,
`GPU_T5LayerNormBwdTest`, `GPU_BNInputValidation`,
`GPU_BN_Fwd_Spatial_Welford`, `GPU_BN_Spatial`, `GPU_Bn3DPeract`,
`GPU_Bn3DSpatial`, `GPU_BnDerive`, `GPU_BnPeract`, `GPU_CbnaInference`,
`GPU_na_*`, `GPU_Allocator`, `GPU_CandidateSelection`, `GPU_Find2Conv`,
`GPU_FindDb`, `GPU_FindModeTrustVerify`, `GPU_GRU`, `GPU_LSTM`,
`GPU_LstmMSRnn`, `GPU_RNNSeqApi`, `GPU_RNNVanilla`, `GPU_Pooling2d`,
`GPU_Pooling3d`, `GPU_TensorReorder`, `GPU_TensorTransform`,
`GPU_binaryTensorOps`, `GPU_reference_kernel` |
| `positive_patterns` (unprefixed `TEST`/`TEST_F`) | **+16**:
`GPU_BN_Spatial_FP32`, `GPU_CKValidatorIntegration_FP32`,
`GPU_CkImplLoader_FP16`, `GPU_Conv{2D,3D}KernelTuning{AI,Solver}_FP32`,
`GPU_ConvAIModelLoad_FP32`, `GPU_ConvNDKernelTuningUtils_FP32`,
`GPU_ConvStridedTensors_FP32`,
`GPU_HipBLASLtGEMMTest_{FP32,FP16,BFP16}`,
`GPU_PerfConfig_HipImplicitGemm3DGroup`,
`GPU_RunAIHeuristics_Integration_FP32`, `GPU_SoftmaxFind20` |
| `positive_patterns` | **−2** dead patterns `*/GPU_GRUExtra_*` and
`*/GPU_lstm_extra*` (they matched nothing; now covered by `GPU_GRU_*` /
`GPU_LSTM_*`) |
| `standard.exclude` | **+21** `Full/` excludes, so standard runs only
the `Smoke/` and unprefixed instantiations of the new families |
| `comprehensive.exclude` | **+12** full-only suites (> 200 s serial or
untimed): BN find2/fusion tuning (`GPU_na_train_find2_*`,
`GPU_na_fusion_test_bn_*`), `GPU_BN_Spatial_FP32`,
`GPU_BN_Fwd_Spatial_Welford`, `GPU_LSTM_fused_infer`, `I8`/`I32`
variants |
| `comprehensive.exclude`, `full.exclude` |
`Full/GPU_LstmMSRnn_FP32.TestFloat/1095` (ALMIOPEN-1926), the same index
Jenkins filters on gfx94X |

### Quarantines (scoped to arch/OS)
| Key | Pattern | Cases removed | Reason | Ticket |
|---|---|---:|---|---|
| `exclude_gpu_gfx110X` | `Full/GPU_LstmMSRnn_FP32*` | 1,535
(comprehensive and full) | Intermittent CPU/GPU mismatch at
`rnn_seq_api.hpp:272`, bias-enabled cases only, a different index set on
every attempt. 88 of 4,316 executions failed (2.0%) at 76 distinct
indices. FP16 (0 / 3,312) and Linux gfx94X (0 / 2,302) are clean. |
ALMIOPEN-1926 |
| `exclude_gpu_gfx115X` | `Full/GPU_LstmMSRnn_FP32*` | 1,535 | Same
signature on Windows gfx1151: 18 of 2,316 executions failed (run
39358709264) | ALMIOPEN-1926 |
| `exclude_gpu_gfx115X` (Linux + Windows) |
`Smoke/CPU_DBSync_NONE.StaticFDBSync/9` | 1 (every tier) | `Db file does
not exist … gfx115128.HIP.fdb.txt` / `gfx115120.HIP.fdb.txt` on hosts
whose CU count has no shipped db. Windows: 5 of 5 executions on 5 hosts,
including develop nightly 38527180317, so it predates this PR. Linux: 12
of 12 on the 20-CU strixhalo hosts in run 39740910967 (passes on
CS-RORDMZ hosts). Only the gfx115128 (28-CU) db ships: #13680 fixed its
Windows install, but 20-CU hosts still have no gfx115120 db. | #13416 |
| `exclude_gpu_gfx115X_windows` |
`Full/GPU_GRU_Base_FP16.TestFloat16/batchSize_17_seqLength_23_inVecLen_13_hiddenSize_67_*biasMode_1_*`
| 8 | 3 of these 8 bias-enabled cases failed verification at
`gru_test.cpp:1985/2224/2473` (same bias-only signature); they passed on
the rerun and in the next run | ALMIOPEN-1926 |
| `exclude_gpu_gfx110X`, `exclude_gpu_gfx115X` (Linux + Windows) |
`GPU_ConvAIModelLoad_FP32*` | 1 (every tier) | `no fdeep .tn.model files
found`: AI heuristic models ship only for gfx908/gfx90a/gfx942/gfx950,
and each install carries only the runner target's db files. Linux
gfx1151 failed first (run 39358709264, reproduced locally with the CI
binaries). Since #13680 (merged 10-29) Windows installs per-target db
files too, so Windows gfx110X fails as well (run 40983714656, 2 of 2
attempts on 2 runners); gfx94X/gfx90a/gfx950 keep the test. |
**ALMIOPEN-2966** |
| `exclude_gpu_gfx115X_linux` | `CPU_PerfDb_NONE.*`,
`GPU_BN_Spatial_FP32.MIOpen3900Regression` | 3 + 1 (every tier) | Killed
by the OOM killer (SIGKILL / exit 137) on the 7.5 GB runners
`strixhalo-12` and `strixhalo-10` (`MIOpen3900Regression` uses ~2.1 GB
per tensor on a shared-memory APU); both pass on the 31–125 GB gfx1151
hosts. Only standard-tier cases are excluded (TheRock release tests run
standard on these hosts). Out-of-memory kills in comprehensive/full are
not excluded case by case, see open items. | **ALMIOPEN-2964** |

Jenkins doesn't run Windows gfx110X/gfx115X or Linux gfx1151, so these
quarantines need no Jenkins
`add_gtest_negative_filter`. gfx94X keeps full `LstmMSRnn` coverage
apart from `/1095`, which matches Jenkins's filter, and gets
re-evaluated if a GitHub CI run sees it fail.

## CI runs on this PR

| Run | Date | Commit | Tier | Archs tested | Result |
|---|---|---|---|---|---|
|
[37362130564](https://github.com/ROCm/rocm-libraries/actions/runs/37362130564)
| 10-05 | `7b72efe` | standard | gfx94X, gfx950, Win gfx110X | ✅ |
|
[37844419899](https://github.com/ROCm/rocm-libraries/actions/runs/37844419899)
| 10-08 | `7b72efe` | standard¹ | gfx94X | ✅ |
|
[37960128662](https://github.com/ROCm/rocm-libraries/actions/runs/37960128662)
| 10-09 | `7b72efe` | standard¹ | gfx94X | ✅ |
|
[38770867047](https://github.com/ROCm/rocm-libraries/actions/runs/38770867047)
| 10-15 | `7b72efe` | standard | gfx94X, Win gfx110X | ✅ |
|
[38815363944](https://github.com/ROCm/rocm-libraries/actions/runs/38815363944)
| 10-15 | `1d54535` | comprehensive | gfx94X, Win gfx110X | ❌ Win
gfx110X `LstmMSRnn_FP32` flaky (both attempts); one runner lost
(TheRock#3845) |
|
[38983909810](https://github.com/ROCm/rocm-libraries/actions/runs/38983909810)
| 10-16 | `1d54535` | full | gfx94X, Win gfx110X | ❌ Win gfx110X
`LstmMSRnn_FP32` flaky (both attempts) |
|
[39338174289](https://github.com/ROCm/rocm-libraries/actions/runs/39338174289)
| 10-19 | `d9701b2` | full | Win gfx110X (completed before cancel) | ✅
gfx110X green after its quarantine |
|
[39358709264](https://github.com/ROCm/rocm-libraries/actions/runs/39358709264)
| 10-20 | `e964061` | full | gfx94X, **gfx90a**, **gfx1151** Linux +
Windows, Win gfx110X | ❌ gfx1151 only: `LstmMSRnn_FP32` (Win),
`ConvAIModelLoad` (Linux), `StaticFDBSync/9` (Win), GRU FP16 flake
(Win), OOM on strixhalo-12, hang on CS-RORDMZ-DT245, GPU `unspecified
launch failure` and runner lost (Win) |
|
[39480689944](https://github.com/ROCm/rocm-libraries/actions/runs/39480689944)
| 10-20 | `48eccae` | full | same as above | ❌ gfx1151 only:
`StaticFDBSync/9` (Win), OOM on strixhalo-12 (Linux), hang on
CS-RORDMZ-DT245 (Linux). gfx94X, gfx90a, gfx110X green. |
|
[39529845480](https://github.com/ROCm/rocm-libraries/actions/runs/39529845480)
| 10-21 | `455cc52` | full | same as above | ❌ Linux gfx1151 only: OOM
on strixhalo-10 in `GPU_BN_Spatial_FP32.MIOpen3900Regression` (now
excluded under ALMIOPEN-2964) and `…/n8c192h256w512`, plus one
container-setup failure. **No gtest failure in any lane.** gfx94X,
gfx90a, gfx110X and Windows gfx1151 green. |
|
[39740910967](https://github.com/ROCm/rocm-libraries/actions/runs/39740910967)
| 10-22 | `7012ede` | full | same as above | ❌ Linux gfx1151 only: OOM
on strixhalo-10/-12 in **new** large cases
(`na_fusion_test_bn_spatial_FP32 nchw_8_192_256_512`,
`na_fusion_inference_test_FP32 nchw_1_16_4096_4096`) and
`StaticFDBSync/9` on Linux (now excluded on all gfx115X); infra:
container setup (halobox-b13), `BAD_COMMAND` on Windows host
wshctr-halobox-b48-05 (passed on rerun). gfx94X, gfx90a, gfx110X and
Windows gfx1151 green. |
|
[40053357276](https://github.com/ROCm/rocm-libraries/actions/runs/40053357276)
| 10-24 | `ffcb4d4` | full | gfx94X, Win gfx110X | ✅ all MIOpen shards |
|
[40746486850](https://github.com/ROCm/rocm-libraries/actions/runs/40746486850)
| 10-28 | `94c1c2b` | standard (opt-in labels removed) | gfx94X, Win
gfx110X | ✅ all MIOpen shards |
|
[40901558044](https://github.com/ROCm/rocm-libraries/actions/runs/40901558044)
| 10-29 | `98c7817` | standard (TEMP workflow change removed) | gfx94X,
Win gfx110X | ❌ Linux gfx94X: all 17 FP8 + 17 BFP8 `RunHipBLASLtGEMM`
cases (`hipblaslt_gemm.cpp:235`, error 0.12–0.17 vs 0.07). Not this PR:
the attempt-2 rerun rebuilt only gfx950, whose `*_generic` MIOpen
artifacts (built with the IEEE FP8 bias) replaced gfx94X's; tracked in
#14375, fix in #14627. Windows green. |
|
[40983714656](https://github.com/ROCm/rocm-libraries/actions/runs/40983714656)
| 10-30 | `7d71fd8` (develop merge) | standard | gfx94X, Win gfx110X | ❌
Windows gfx110X shard 3 only:
`GPU_ConvAIModelLoad_FP32.BundledFdeepModelsLoad` (2 of 2 attempts)
after #13680 moved Windows dbs to per-target packages; now excluded on
gfx110X/gfx115X (ALMIOPEN-2966). Linux gfx94X green, including FP8/BFP8.
|

¹ The `test_type:comprehensive` / `test_type:full` labels used then are
not wired into the multi-arch workflow, so these runs ran standard.
Labels used later: `test_filter:comprehensive`, `test_filter:full`,
`ci:gfx90a`, `ci:gfx1151`.

## Timing

All numbers are **per shard** (4 shards, serial ctest), as ctest wall
minutes; the TheRock Test step is capped at **120 min**.

### Measured
| OS / arch | Tier | develop | This PR | Δ | Runs (PR / develop) |
|---|---|---:|---:|---:|---|
| Linux gfx94X | standard | 26.8 (25.5–28.7) | 27.7 (26.8–29.2) |
**+0.9** | 37362130564, 37844419899, 37960128662, 38770867047 /
38541202878, 38945603498, 38998133138 |
| Linux gfx94X | comprehensive | 41.2 (38.8–42.7) | 57.4 (54.0–62.7) |
**+16.2** | 38815363944 / nightly 38697600236 |
| Linux gfx94X | full | not run² | 68.6 (63.9–72.5) | ≈ +27 |
38983909810, 39358709264, 39480689944 |
| Linux gfx90a | full | not run | 83.6 (78.3–86.6) | – | 39358709264,
39480689944 |
| Linux gfx950 | standard | not run | 15.4 (15.0–16.0) | – | 37362130564
|
| Linux gfx1151 | full | not run | 28 (27.1–29.0)³ | – | 39358709264,
39480689944 |
| Windows gfx110X | standard | 17.2 (10.9–19.5) | 17.1 (8.8–20.6) |
within runner-host noise⁴ | 37362130564, 38770867047 / same develop PR
runs |
| Windows gfx110X | comprehensive | not run⁵ | 46.3 (42.0–50.3) | – |
38815363944 |
| Windows gfx110X | full | not run | 51.9 (48.0–56.7) | – | 38983909810,
39338174289, 39358709264, 39480689944 |
| Windows gfx1151 | comprehensive | 28.3 (20.0–42.8) | – | – | – /
nightly 38697600236 |
| Windows gfx1151 | full | not run | 53.2 (35.3–73.2)⁶ | – | 39480689944
|

The slowest shard measured is gfx90a full at 86.6 min (72% of the cap).
Every other lane is at or below 73 min.

² Develop's full tier selects the same cases as develop comprehensive
(within 25 cases), so the delta is measured against the develop
comprehensive time.
³ Healthy shards only; the other shards were OOM-killed (ALMIOPEN-2964)
or hung at startup (ALMIOPEN-2965).
⁴ Windows gfx110X shards range from 9 to 21 min depending on the runner
host (CS-RORDMZ vs azure scale-set).
⁵ The develop nightly runs Windows gfx1151, not gfx110X.
⁶ Large spread between hosts: one gfx1151 runner is about 2× slower than
the others.

### Modelled delta vs develop (final yaml, cases added and per-shard
minutes)
`yaml_select.py diff --old git:<merge-base 4c5c20a> --new git:HEAD --gpu
<arch> --os <os>`, using this PR's per-case timings measured on each
arch.

| OS / arch | standard | comprehensive | full |
|---|---|---|---|
| Linux gfx94X | +11,284 cases · **+1.5 min** | +25,714 · **+15.6 min**
| +29,887 · **+23.8 min** |
| Linux gfx90a | +11,284 · **+1.9 min** | +25,714 · **+18.3 min** |
+29,887 · **+27.4 min** |
| Windows gfx110X | +11,281 · **+1.4 min** | +24,068 · **+15.8 min** |
+28,241 · **+21.8 min** |
| Windows gfx1151 | +11,281 · **+1.5 min** | +24,060 · **+16.9 min** |
+28,233 · **+23.2 min** |

The model agrees with the measurements where a develop baseline exists:
gfx94X comprehensive +15.6 modelled vs +16.2 measured;
standard +1.5 vs +0.9.

<details><summary>Largest additions per tier (Linux gfx94X, serial
minutes)</summary>

| Tier | Suites |
|---|---|
| standard | `Smoke/GPU_Find2Conv` 0.8, `Smoke/GPU_TensorReorder` 0.4,
`Smoke/GPU_reference_kernel_{bwd,bww}_2d_*` 0.3–0.4 each |
| comprehensive | `Full/GPU_Pooling3d` 9.0, `Full/GPU_LayerNormBwdTest`
7.8, `Full/GPU_Pooling2d` 7.0, `Full/GPU_BnPeract` 4.3,
`Full/GPU_TensorTransform` 3.7 |
| full only | `Full/GPU_na_fusion_test_bn_spatial` 10.8,
`Full/GPU_na_train_find2_spatial` 10.2, `Full/GPU_na_*_peract` 5.8 each
|

</details>

### Where each tier runs after merge
| Tier | Lanes |
|---|---|
| standard | rocm-libraries PR gate and develop postsubmit (Linux
gfx94X, Windows gfx110X); TheRock postsubmit (gfx94X, gfx90a, gfx950,
Windows gfx110X); TheRock release tests (Linux gfx1151, Windows gfx110X)
|
| comprehensive | rocm-libraries nightly (Linux gfx94X, Windows
gfx1151); ASAN nightly (gfx94X) |
| full | only on request, with a `test_filter:full` label |

> [!NOTE]
> **Standard-tier hot spot (pre-existing, not changed by this PR).** In
this PR's standard runs on Linux gfx94X, `Full/GPU_Softmax_*`
> is the most expensive group: **36.5 of 109.6 serial minutes (33%)**,
about 9 min per shard. Breakdown per run: `Full/GPU_Softmax_FP32`
> 18.5 min (1,008 cases), `_FP16` 7.1, `_BFP16` 7.0,
`_NonContiguous_FP32` 3.9. It has been in standard since #3513. On
gfx11/gfx12
> the FP32/FP16/BFP16 variants are already excluded as slow
(TheRock#3362); `Full/GPU_Softmax_NonContiguous_FP32` still costs 5.1
min
> (8%) on Windows gfx110X.
>
> Moving it to comprehensive (adding `Full/GPU_Softmax_*` to
`standard.exclude`) is modelled at **−9.1 min per shard on gfx942**
> (27.7 → about 18.6) and **−1.2 min per shard on Windows gfx110X**.
Comprehensive and full keep every case, so the nightly still covers it.
> Proposed as a follow-up rather than part of this PR. Next largest in
standard: `Standard/GPU_Reduce_{FP32,FP16}` 3.9 + 3.1 min,
> `Full/GPU_GroupConv3D_*` 1.4–2.2 min per dtype.

## ASAN validation of the added suites

Every gtest suite this PR adds was built and run under full ASAN, both
host-side and device-side, on gfx942 with XNACK on (`gfx942:xnack+`).
Apart from one pair of test files that leak memory, they are clean.

- **Setup.** This branch was built with `-fsanitize=address` against the
rockrel *Multi-Arch Release ASAN* nightly #139 artifacts (run
[40425965688](https://github.com/ROCm/rockrel/actions/runs/40425965688),
TheRock `e73fd19`), using the same test environment as TheRock CI
(`configure_asan_env.py`, `ROCM_PATH`, `HSA_XNACK=1`). It ran on MI300X:
28,354 selected cases in 4 shards, plus a canary run of the `Smoke/` and
unprefixed suites first.
- **Result.** **25,438 tests passed, with 0 gtest failures and 0 ASAN
memory errors** (no overflow or use-after-free in MIOpen).
- **One MIOpen problem, in the tests themselves:** `na_train_find2.cpp`
and `na_inference_find2.cpp` call `miopenFindSolutions` but never
`miopenDestroySolution`, so LeakSanitizer fails the process at exit.
They are excluded from ASAN builds only, using `SKIP_TESTS` in
`test/gtest/CMakeLists.txt` under `THEROCK_SANITIZER=ASAN`, next to the
existing `find_2_conv.cpp` leak skip. This removes the 12
`*/GPU_na_{train,inference}_find2_*` suites (528 cases) from ASAN
builds; normal CI keeps them. Tracked in **ALMIOPEN-3097**.
- **Runtime issues, not MIOpen:**
- **HIP start-up under ASAN with XNACK on (AILIROCR-587).** When two
ASAN processes run on the same node, the HIP runtime hangs while
initializing (`KernelBlitManager::initHeap`) or ASAN fires inside ROCr
(`AqlQueue::HandleInsufficientScratch`). A single process started 5 of 5
times. The run therefore used one ASAN process per node.
- **Bogus device-side leak report (ALMIOPEN-3098).** At exit, the device
ASAN runtime lists about 1 TB of "leaks" on a 192 GB GPU, and the stack
addresses are float data. It's informational and doesn't change the exit
code.
- **ROCr loader leak (AILIROCR-588).** A one-off 51-byte `strdup` leak
in ROCr's `LoadCodeObject`, seen in the canary only.

I didn't add an ASAN exclusion to `test_categories.yaml`. Its parser has
no sanitizer-specific key, so any yaml key would also drop these suites
from normal CI.

## Tickets

| Ticket | Status | Scope |
|---|---|---|
| ALMIOPEN-1926 | referenced | `LstmMSRnn_FP32` (and GRU FP16) bias-only
CPU/GPU mismatch: gfx110X, gfx115X, and the Jenkins gfx94X `/1026`,
`/1095` |
| **ALMIOPEN-2964** | **created for this PR** | 7.5 GB Linux gfx1151
runners `strixhalo-12` and `strixhalo-10` OOM-kill `miopen_gtest`;
matches the `exclude_gpu_gfx115X_linux` OOM exclusions. The fix is to
remove those hosts from the runner pool. |
| **ALMIOPEN-2965** | **created for this PR** | `miopen_gtest` hangs
before any output on Linux gfx1151 runners CS-RORDMZ-DT245 (3 of 3 jobs)
and ctr-halobox-b13; no test starts, so there is no yaml exclusion |
| **ALMIOPEN-2966** | **created for this PR** |
`GPU_ConvAIModelLoad_FP32.BundledFdeepModelsLoad` fails instead of
skipping when the system db has no `*.tn.model` (Linux gfx1151; Windows
gfx110X/gfx1151 since #13680); repro steps on any GPU via
`MIOPEN_SYSTEM_DB_PATH`; matches the `exclude_gpu_gfx110X` /
`exclude_gpu_gfx115X` entries |
| **ALMIOPEN-3097** | **created for this PR** | ASAN: `na_train_find2` /
`na_inference_find2` tests leak `miopenSolution_t` (no
`miopenDestroySolution`); matches the ASAN `SKIP_TESTS` entries |
| **AILIROCR-587** | **created for this PR** | ASAN: HIP init on gfx942
hangs in `KernelBlitManager::initHeap` or trips ASAN in ROCr
`AqlQueue::HandleInsufficientScratch` when two ASAN processes share a
node |
| **ALMIOPEN-3098** | **created for this PR** | ASAN: device-side leak
report at exit lists ~1 TB of impossible allocations (stacks are float
data) |
| **AILIROCR-588** | **created for this PR** | ASAN: 51 B `strdup` leak
in ROCr `ExecutableImpl::LoadCodeObject` (seen once) |
| #13416 / #13680 | referenced | `StaticFDBSync` "Db
file does not exist" on hosts whose CU count has no shipped db / Windows
db install fix |
| ROCm/TheRock#3845 | referenced | lost azure gfx1101 runner
(38815363944 attempt 1) |

## Test plan
- `python3 shared/ctest/parse_test_categories.py test_categories.yaml …`
returns rc 0 (21 suites on Linux).
- `yaml_select.py diff` after every quarantine: each change removes only
the intended cases on the intended arch/OS; Linux gfx94X
  and gfx90a selections are unchanged by the RDNA quarantines.
- CI: the runs above, read across all attempts. Run 39529845480 confirms
the quarantines take effect (`LstmMSRnn_FP32` ran 0
times on gfx110X/gfx115X, gfx94X/gfx90a ran it 3,454 times each with no
failures) and has no gtest failure in any lane.
- ASAN: local build of this branch against rockrel ASAN nightly #139;
all added suites run on MI300X gfx942:xnack+ (see *ASAN validation*).
25,438 passed, 0 failures, 0 ASAN errors.

## Open items
- [x] Drop the TEMP workflow commit `4b8890f` before merge: removed in
`98c7817` (workflow identical to develop's at the merge-base).
- [ ] Linux gfx1151 out-of-memory kills: **no more per-case
exclusions.** Each run kills a different large case (4 shapes across 3
runs; 82 cases of ≥ 32M elements; no size threshold separates killed
cases from survivors). Fixed by ALMIOPEN-2964 (remove the 7.5 GB hosts).
Linux gfx1151 comprehensive/full runs only when requested by label,
never by default.
- [ ] Final run: TEMP workflow change removed (`98c7817`) and labels
`ci:gfx1151`, `ci:gfx90a`, `test_filter:full` removed; the PR gate runs
standard on Linux gfx94X and Windows gfx110X. Run 40746486850
(`94c1c2b`) passed. 40901558044 hit the FP8 artifact overwrite (#14375)
and 40983714656 hit `ConvAIModelLoad` on Windows gfx110X after #13680;
that test is now excluded on gfx110X/gfx115X. Needs one green run with
that change.
- [ ] Re-enable the `na_*find2` suites under ASAN once ALMIOPEN-3097 is
fixed (remove the two `SKIP_TESTS` entries).
- [ ] After merge, check the first nightly (Linux gfx94X and Windows
gfx1151 comprehensive) against the modelled +15.6 / +16.9 min per shard.
@yassinsolim
yassinsolim marked this pull request as ready for review October 30, 2026 08:43
@rocking5566
rocking5566 force-pushed the users/ysoliman/miopen-fp8-runtime-bias branch from bad225e to 3a3c414 Compare October 30, 2026 12:12
…pen-fp8-runtime-bias

# Conflicts:
#	projects/miopen/CHANGELOG.md
@yassinsolim
yassinsolim merged commit 0341584 into develop Nov 4, 2026
103 checks passed
@yassinsolim
yassinsolim deleted the users/ysoliman/miopen-fp8-runtime-bias branch November 4, 2026 15:48
cycheng pushed a commit that referenced this pull request Nov 6, 2026
…14627)

## Proposed changes

Closes #14375.

MIOpen picked `MIOPEN_FP8_IEEE_EXPONENT_BIAS` once at configure time
from `GPU_TARGETS`. A host build for several targets that includes
gfx950 or gfx12 got the IEEE bias everywhere, so on gfx942 the hipBLASLt
FP8 and BF8 GEMM tests read FNUZ data with the wrong bias and failed
(errors around 0.1 to 0.17). This PR picks the bias at run time from the
device instead.

- `UsesIeeeExponentBias(device_name)` in target_properties.hpp returns
true for gfx950 and gfx12, the same targets as
`MIOPEN_CK_USE_OCP_FP8_TARGETS`.
- `GetDataTypeKBP` and `GetDataTypeKernelParams` take the device name
and pass the bias for that device to kernel compiles. The one-argument
versions are gone, and every caller passes its handle or context device
name.
- check_numerics, the naive conv solver and the MHA solvers use the
helper in place of the config.h macro.
- The CMake seed in MiopenCkCompileDefs.cmake, the
`MIOPEN_FP8_IEEE_EXPONENT_BIAS` option and its config.h entry are
removed. config.h is installed, so the define leaves the public headers.
- Host FP8 types in hip_float8.hpp read the bias from a run-time setting
(FNUZ by default), and kernel builds still get it as a `-D` flag. The
conv solver test fixtures set the host bias from the device through a
new `SetHostFp8ExponentBias<T>()` in conv_test_base.hpp. MIOpenDriver
sets it too, but `ParseBaseArg` doesn't accept convfp8 or convbfp8 yet,
so that path can't run.
- The first commit is newling's 79a9d61 cherry-picked unchanged (thanks
James). It makes the hipBLASLt GEMM test use explicit FNUZ host types
and adds CPU tests for that encoding.

What could break:
- Builds with CK off never got the seed, so gfx950 and gfx12 used bias 0
there. They now use 1, same as CK-on builds.
- HIP's FP8 header treats gfx1170 and gfx1171 as OCP, but this rule
gives them bias 0, same as `MIOPEN_CK_USE_OCP_FP8_TARGETS` today. I
haven't tested MIOpen FP8 on those targets.
- Out-of-tree code that read `MIOPEN_FP8_IEEE_EXPONENT_BIAS` from
config.h, or called the one-argument `GetDataTypeKBP`, won't build.

Testing, one CK-on build for gfx942 and gfx950, run on MI325X and
MI355X, with develop built the same way:
- On gfx942 develop fails all 34 hipBLASLt FP8 and BF8 GEMM cases. They
pass with this PR, along with the two new CPU encoding tests.
- On gfx950 the GEMM tests skip by design (gfx94X only).
- `GPU_Fp8Conversion_FP8.DefaultAndSoftwareFallback` (gfx950) and
`GPU_Bfp8Conversion_BFP8.DefaultAndSoftwareFallback` (both) fail the
same way on develop, on the NaN element of the ieee format, so I left
them alone.
- The FP8 conv solver tests are disabled upstream. Forced on, they fail
on develop and here before checking any values, since the naive FP8
kernels never get the names the solver looks up (fp8_naive_conv.cpp uses
`FWD_KERNEL_NAME` and friends, and nothing defines them) and FP8 GEMM
conv needs rocBLAS gemm_ex3, which MIOpen only uses below rocBLAS 5.
Their kernel build options do show the new bias, 0 on gfx942 and 1 on
gfx950 (develop passes 1 on both).
- A CK-off build compiles.

## Checklist

- [x] I have added automated tests relevant to the introduced
functionality
- [ ] I have sufficient test coverage for the changes, and code coverage
hasn't decreased as a result of my PR
- [ ] I have ran the tests, and they are all passing locally
- [x] I have added relevant documentation for the changes
- [x] I have removed the stale documentation which is no longer relevant
after this pull request
- [x] I have ran `make format` & `make check_format` to ensure the
changes have been formatted

---------

Co-authored-by: James Newling <james.newling@gmail.com>
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

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.

[Issue] Linux gfx94X-dcgpu: MIOpen HipBLASLt FP8/BFP8 GEMM tests fail with numerical error

4 participants