Repository navigation
fix(hip-kernel-provider): ASM SDPA declines non-positive attention scales - #14688
Open
brockhargreaves-amd wants to merge 3 commits into
Open
brockhargreaves-amd wants to merge 3 commits into
brockhargreaves-amd wants to merge 3 commits into
Conversation
1 task done
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
This was referenced Oct 29, 2026
brockhargreaves-amd
changed the title
brockhargreaves-amd
force-pushed
the
users/brockhargreaves-amd/asm-sdpa-decline-nonpositive-scale
branch
from
October 30, 2026 18:57
11cd3c2 to
6013fba
Compare
brockhargreaves-amd
marked this pull request as ready for review
October 30, 2026 19:26
CMiservaAMD
approved these changes
Nov 2, 2026
CMiservaAMD
reviewed
Nov 2, 2026
CMiservaAMD
left a comment
CMiservaAMD
left a comment
Contributor
There was a problem hiding this comment.
Found by AI. Severity: none blocking.
Thanks, this looks right to me.
- The check
FLT_MIN <= scale <= FLT_MAXrefuses +0, -0, negatives, subnormals, NaN of either sign and +-inf. Because comparisons with NaN are false, NaN fails both sides. - It applies to the baked scale in both
isApplicablepaths, and to a runtime pass-by-value scale before the first launch (BAD_PARAM).
Built and ran it on gfx942 (MI300A):
- the 10 new tests pass;
- the ASM SDPA plan-builder suites pass (155 passed, 1 gfx950-only skip);
- the tests properly caught two deliberate breaks of
isSupportedAttnScale: letting +0 and -0 through, and letting NaN through. Each break failed 5 of the 10 new tests.
It merges cleanly onto the current develop. No existing bundle or test uses a scale of 0 or less.
Our forward test-plan runs and an AOTriton-suite run also show that +0 with Skv not a multiple of 64 gives an all-NaN O. I added a comment on ALMIOPEN-3043; this check covers it.
A few suggestions follow inline.
10 of 13 tasks
…ales SdpaFwdPlanBuilder::isApplicable and SdpaBwdPlanBuilder::isApplicable accepted any attention scale, and execute passed it to the kernels unchecked. cuDNN and hipDNN accept any scale value. On gfx942 (MI300X) a scale of 0, -0.0, a negative value or a subnormal gave NaN or Inf outputs and gradients instead of a decline. The mechanism comes from the .co disassembly. The forward and backward kernels write -inf into every masked score (causal and sliding-window masks, and the padding past S_kv in the last K tile), and only then multiply the scores by scale * log2(e). A scale of 0 turns those scores into NaN, and a negative scale turns them into +inf. The bf16 forward kernels also start the running row max at -FLT_MAX / (scale * log2(e)), which is positive for a negative scale or -0.0, so every exp2 overflows even with no mask. 32-bit denormals are flushed, so a subnormal scale acts as 0. An infinite scale gives NaN everywhere. The gfx942 runs used the plugin from a CI build of #14312, which accepts all of these scales, and compared results with the CPU reference: - Backward: 0, -0.0, 1e-40 and -1.2 gave NaN or Inf dQ, dK and dV with top-left or bottom-right causal masks and sliding windows, in bf16 and fp16, hd128 and hd192. With no mask and S_kv = 128 (the hd128 K tile is 192), all of dQ was NaN. - Forward: -1.2, -0.088 and -0.0 gave NaN O and Inf LSE even with no mask. 0 gave NaN O with a bottom-right causal mask or S_kv = 200. - A runtime pass-by-value scale tensor gave the same results as attn_scale_value. Both isApplicable functions now decline a scale that is not a finite, positive, normal float (FLT_MIN <= scale <= FLT_MAX; NaN fails). The check covers attn_scale_value, and the value of a compile-time constant or runtime-with-default pass-by-value scale tensor, narrowed to the f32 kernarg as execute narrows it. They also decline a scale tensor whose value the plugin SDK cannot read, on which buildPlan would throw. The value of a pure runtime pass-by-value scale tensor arrives only at execute, so both execute functions now throw BAD_PARAM for such a scale before the first launch, as the gfx1151 WMMA attention pack does for its runtime scale. The check is shared in SdpaPlanUtils.hpp. An unset scale (1.0) and every finite, positive, normal scale stay accepted; 1e-37 and 1/sqrt(128) were correct on gfx942. The rule also declines some cases that happened to compute correctly: no mask with S_kv a multiple of the K tile, for any of these scales in the backward and for +0.0 or a subnormal scale in the forward. A rule keyed on the mask and the tile would break once group mode makes the K tail per-sequence, so the decline does not depend on them. Scales near FLT_MAX (1e37, 1e38) still give NaN in the forward; that is a separate issue. New tests: - TestSdpaPlanUtils (host only): IsSupportedAttnScaleRequiresFinitePositiveNormalFloat and AttnScaleBeforeExecute. - TestSdpaFwdPlanBuilder and TestSdpaBwdPlanBuilder (gfx942, and gfx950 for the forward, like the other isApplicable tests): IsApplicableRejectsUnsupportedAttnScale declines 0, -0.0, -1.2, -1/sqrt(128), -FLT_MIN, three subnormals, NaN and +-inf, each from attn_scale_value, a constant scale tensor and a runtime scale tensor default, plus an FP8_E4M3 scale tensor. IsApplicableAcceptsSupportedAttnScale accepts 1e-37, 1/sqrt(128) and 1 from each source, the unset default and a pure runtime scale tensor. - TestSdpaFwdPlanExecute and TestSdpaBwdPlanExecute: execute returns BAD_PARAM for each unsupported runtime scale (host only, on a plan with no loaded kernels), and passes supported ones on to the launch. Checked on gfx942 (MI300X). This commit's plugin replaced the one in a CI build of #14312, and results were compared with the CPU reference on the 82 scale cases behind the results above: - The 61 cases whose unsupported scale comes from attn_scale_value or a constant scale tensor are now declined. With the old plugin 48 of them failed, 12 passed (the no-mask cases named above) and one was already declined (fp16 forward). - The 4 cases with an unsupported runtime scale tensor (0 and -1.2, forward and backward) now fail at execute with BAD_PARAM, and the message gives the scale. The old plugin gave NaN or Inf in 3 of them. - The 15 cases with a supported scale pass with both plugins: 1/sqrt(128) with each mask, dtype and scale source, and 1e-37. The forward with 1e37 or 1e38 fails with both, as noted above. - On the develop CI build, develop's quick and standard SdpaFwd, SdpaFwdRuntimeScale and SdpaBackward bundles give the same result for each case as with develop's plugin: 42 pass and 39 are declined, for the same reasons, none of them the scale. - The unit tests pass. Each of 14 mutants that removes, widens or narrows one arm of these checks makes one of the new tests fail. Rebased onto develop 98b504d, after #14075, whose range gate now runs first in isApplicable. The GPU results above are from develop 61c2675. On the new base, the four ALMIOPEN-3035 ASM SDPA fix branches (backward layouts, attention scale, a16 accumulator, forward layouts) were built and run together: - hip_kernel_provider_tests: 912 pass and 7 skip on gfx942 (MI325X), 886 pass and 33 skip on gfx950 (MI355X); none fail. - hip_kernel_provider_integration_tests: none fail on either. - develop's SDPA bundles on ASM_SDPA_ENGINE, run like the CI lane (claims enforced): 62 pass and 56 skip on gfx942, 15 pass and 103 skip on gfx950, and no support claim fails. A host-side run of the real isApplicable over these bundles finds no outcome change from develop. - #14622's SdpaBackward bundles (138218c) with its temporary [[test_skips]] removed: 275 pass and 225 skip on gfx942, none fail, and no support claim fails. This change declines def06a, def06b and def06c (scale 0 and -1.2), the cases behind the DEF-06 skip, and no other case (host-checked for this branch alone). The execute-path scale tests (PassesSupportedRuntimeAttnScaleToLaunch) pass on both GPUs. - clang-tidy 20 and clang-format 18 are clean on the changed files. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
brockhargreaves-amd
force-pushed
the
users/brockhargreaves-amd/asm-sdpa-decline-nonpositive-scale
branch
from
November 3, 2026 19:59
6013fba to
8bef35d
Compare
…e-nonpositive-scale
15 of 16 tasks
brockhargreaves-amd
added a commit
that referenced
this pull request
Nov 5, 2026
…rated bundles (#14622) ## Summary This PR adds the SDPA backward test plan required by RFC 0022 (Per-Op Test Plans, merged in #13937). It also adds a generator that writes the SdpaBackward bundles from the plan, and the bundles it writes. The 16 inherited SdpaBackward cases (all B = H = 1, packed BHSD, fixed scale) are replaced by 503 generated cases in four levels. Each case carries a purpose and traces to the plan items it covers. > Written with AI assistance (Claude Code). The plan follows the merged RFC 0022; it records one remaining deviation (DV-4) for reviewers to approve. JIRA ID : ALMIOPEN-2894 ## Risk Assessment Risk 2 of 5. The PR changes test data, test tooling, test configuration and documentation only; no library or provider code changes. The changes that reach beyond the SdpaBackward bundles: - `support_matrix_units.collect_sweep_units()` leaves out cases whose `metadata.purpose` is `probe`, as RFC 0022 §9 requires. Today only SdpaBackward cases carry a `purpose`, so other ops' matrices are unchanged. The committed `SUPPORT_MATRIX_OVERVIEW.md` is re-rendered. - `ASM_SDPA_ENGINE.toml`: two dtype-only SdpaBackward tolerance rows are replaced by 12 generated rows (seven tolerance classes), plus temporary skips for the open engine defects, each naming its ticket. - `ASM_SDPA_ENGINE_test_categories_integration.yaml`: six rejection probes that ASM_SDPA_ENGINE wrongly accepts are quarantined with their tickets (RFC 0022 §6). Both files affect the ASM lane only. - #14358's eight PyTorch golden cases move, unchanged, from `quick/SdpaBackward/{Default,AdditiveMask}/` to `quick/SdpaBackward/Golden/{Default,AdditiveMask}/`, so the generator owns the `Default` sweeps. They keep their fp16 tolerance row. `bhsd/` and the `*_fwd_pair` bundles are untouched. The main risk is CI noise from new cases. Every level has been run on gfx942 and gfx950 with no failure (see Testing Checklist). ## ASIC Coverage The bundles run in every provider lane whose tier YAML matches `*Sdpa*`. Only `ASM_SDPA_ENGINE` accepts SdpaBackward graphs today, and only on gfx942: its backward plan builder declines every other device ("Device string does not match gfx942"). Claims are harvested for gfx942 only. PR CI tests gfx942. On gfx950 every generated backward case is declined, so a gfx950 run only shows that nothing fails; one is recorded below. PR CI on gfx942 is sufficient. ## Testing Summary - **Generator tests** (`hipdnn_sdpa_bwd_bundle_generator_python_tests`, label `quick`, 19 tests). They fail when the checked-in sweeps or the TOML block drift from the generator. They also check legality, all-pairs coverage up to nightly, the per-case ceilings and cumulative budgets of each level, two-way traceability between plan items and cases, that every TOML skip still matches a written case, and that the tier-YAML quarantine holds only rejection probes. - **`generate_sdpa_bwd_bundles.py --dry-run --report`.** Reports counts, budgets, the all-pairs proof, legality, and each declared-but-gated case with its gate. - **Claims run** (2026-10-26, MI300X, gfx942). All generated bundles ran on `ASM_SDPA_ENGINE`, configured like CI. `harvest_support_claims.py` turned the passing cases into the 272 claims. Every claimed case passes its tolerance class with a margin of at least 2.4 (median 4.2). - **Measured run, every level on GPU** (2026-11-04, plan §4.5). `hipdnn_integration_tests` ran on `ASM_SDPA_ENGINE` configured like the CI lane: `--test-engine` (claims enforced), the PR's `ASM_SDPA_ENGINE.toml`, the tier YAML's quarantine as a GTest filter, device input fill and the GPU reference as oracle. It ran all four levels, on a fresh MI300X (gfx942) and an MI355X (gfx950) node, built from this branch. The bundles, TOML and YAML are identical to `6e014ad50e0`. - **PR CI** runs the quick and standard levels on gfx942. ## Testing Checklist - [x] Generator tests - CTest `hipdnn_sdpa_bwd_bundle_generator_python_tests` - 19 of 19 - Status: Passed - [x] Plan report - `python3 reference-data-scripts/generate_sdpa_bwd_bundles.py --dry-run --report` - 503 written, 77 gated, 0 legality violations, all-pairs 669 of 669 - Status: Passed - [x] Bundle and support-matrix checks - `verify_support_claims.py`, `render_support_matrix.py --output -`, `feature_matrix.py --check`; the bundle, claim, render and support-feature pytest suites - Status: Passed (`verify_golden_bundles.py` reports no SdpaBackward error; its 322 errors in other ops' bundles are on develop too) - [x] pre-commit on the changed files - Status: Passed - [x] Claims run - `hipdnn_integration_tests`, `ASM_SDPA_ENGINE`, MI300X (gfx942), 2026-10-26 - Status: Passed - [x] Measured run, gfx942 - `hipdnn_integration_tests --test-engine ASM_SDPA_ENGINE --test-config ASM_SDPA_ENGINE.toml '--gtest_filter=*SdpaBackward*-<the tier YAML quarantine>'` - MI300X, Slurm job 68275259 - Status: Passed. 553 tests in 43.3 s, 0 failed, no claim broken: - 289 executed and passed against the GPU reference, and 15 expected rejections rejected; - 226 declined by ASM, and 23 hit the DEF skips; - per level: quick 2.4 s, PR 2.8 s, nightly 16.5 s, full 21.4 s. - [x] Quarantine check, gfx942 - the six quarantined cases run alone - Status: all six fail as expected ("Engine ASM_SDPA_ENGINE accepted expected-rejection bundle"), which shows the quarantine is needed. - [x] Every level, gfx950 - same command - MI355X, Slurm job 68275260 - Status: Passed. 0 failed: ASM declines every backward graph on gfx950, and the rejection probes, quarantined ones included, are correctly rejected. - [x] The other SDPA bundles, same lane - `--gtest_filter=*Sdpa*:-*SdpaBackward*` - Status: Passed. gfx942: 47 passed, 84 skipped. gfx950: 16 passed, 115 skipped. 0 failed on either. - [ ] PR CI on gfx942 - TheRock multi-arch CI - Status: Pending on `6e014ad50e0`. The 2026-11-02 run failed one case, `prbr01`, which is now gated. Later runs failed only on the hipCCL configure break from #14976, which develop reverted in #15089 (merged into this branch). ## Technical Changes ### What this adds - **The plan:** `dnn-providers/integration-tests/docs/test-plans/sdpa-bwd.md`. It follows the RFC 0022 template: the operation, the feature inventory, the shape catalog, building the sweep (levels, budgets, probes, regressions, trade-off register), the golden reference, pass criteria, and exhaustive mode. - **The generator:** `reference-data-scripts/sdpa_bwd_plan.py` holds the declarations: axes, legality, the shape catalog, level rules and the cost model, one record per case. `generate_sdpa_bwd_bundles.py` writes the sweeps, the support-claim files and the engine-TOML tolerance block from them. Nobody edits the bundles by hand. RFC 0022 open question 1 cites this module as the per-op Python option. - **The bundles:** 503 written cases, spread across the levels: | Level | Directory | Cases | |---|---|---| | quick | `quick/` | 56 | | PR | `standard/` | 152 | | nightly | `comprehensive/` | 220 | | full | `full/` | 75 | - Each case carries `purpose` (`claim`, `probe` or `regression`) and its plan items. - 77 more cases are declared but not written yet, each waiting on a named gate: an oracle extension, a harness feature or golden data. - One of the 77 is PRB-R01 (`dropout_probability = 0`). Since #14358 the references decline any graph that sets `dropout_probability`, so it waits on ALMIOPEN-3149 or ALMIOPEN-3150. - **Levels:** the cost model assigns each case to the lowest level whose per-case limits it fits: | Level | Per-case time | Per-case bytes | Cumulative budget | |---|---|---|---| | quick | 1.5 s | 4 MiB | 20 s | | PR | 5 s | 16 MiB | 120 s | | nightly | 75 s | 64 MiB | 480 s | | full | 300 s | no limit | 1,800 s | Each cumulative budget covers that level and every level below it, and keeps 5 % in reserve. The budgets come from a CPU-reference cost model; with the GPU reference, the measured run above is far below them. The ASM lane YAML picks the cases up by GTest prefix (`quick_*Sdpa*`, `standard_*Sdpa*` and so on), and each level also runs every level below it. - **Rejection probes** (21) carry `metadata.expected_rejection` (#14346): the harness never executes them, and an engine that accepts one fails the lane. PRB-M02 and PRB-M02f (`right_bound = -2`) are not written, because the frontend rejects them since #14358; `TestSdpaBwdNode.PreValidateChecksDiagonalBandBounds` covers that rule. - **dBias:** the five dBias probes and the live GL-14 case are written, now that #14358's references model dBias. To keep PR and nightly inside their reserve, the plan's trim rule moved M07.pr to nightly and M26.n to full. - **Tolerance classes:** seven classes, written as 12 `[[tolerance_overrides]]` rows between generated BEGIN and END markers in `ASM_SDPA_ENGINE.toml`. - **Support claims:** 272 claims for `ASM_SDPA_ENGINE` on gfx942. - **Probes left out of the support matrices** (RFC 0022 §9). This is the `support_matrix_units` change above. - **Generator file I/O is UTF-8**, so its tests pass on Windows. ### Temporary ASM_SDPA_ENGINE exclusions These stay until each fix lands; the PR with the fix deletes its entries. **Open engine defects** (`[[test_skips]]` in `ASM_SDPA_ENGINE.toml`). These are valid graphs that ASM accepts and computes wrong. Other engines still run them. | Cases | What goes wrong | Ticket | Fix | |---|---|---|---| | DEF-01 | head-dim stride other than 1 | ALMIOPEN-3036 | #14684 | | DEF-02 | non-packed Stats | ALMIOPEN-3037 | #14684 | | DEF-04 | Hk = Hq > Hv | ALMIOPEN-3039 | #14684 | | DEF-05 and H61 | odd element strides | ALMIOPEN-3038 | #14684 | | DEF-06 | scale of 0 or below | ALMIOPEN-3043 | #14688 | | DEF-03 | fp16 dV precision in the AITER kernel | ALMIOPEN-3046 | none in hipDNN (kernel) | **Quarantined rejection probes** (`_common_excludes` in `ASM_SDPA_ENGINE_test_categories_integration.yaml`, as RFC 0022 §6 asks). These are undefined graphs ASM should decline but accepts. | Cases | What ASM accepts | Ticket | Fix | |---|---|---|---| | prbu01, prbu03 | O or dQ dims that differ from Q's | ALMIOPEN-3039 | #14684 | | prbu13 | a dQ sequence stride below D (overlapping rows) | ALMIOPEN-3038 | #14684 | | prbt04, prbt04n | an O dtype that differs from Q's | ALMIOPEN-3050 | none yet | | prbr05 | `dropout_scale_inv` without a dropout mask | ALMIOPEN-3051 | none yet | ### Dependencies (all merged) - #13937, RFC 0022 (2026-10-30). Its first-phase documents, the methodology guide and the `docs/test-plans/` index, are not on develop yet, so the plan links the RFC's sections directly. - #14312: O and Stats derived from the forward of the same inputs. - #14483: `left_bound` uses cuDNN's meaning. The bundles write `left_bound = W`. - #14358 (2026-11-02): SDPA references cover the full graph API, and the GPU reference is the first oracle in `auto` mode. It also adds frontend rule R14 (`right_bound >= -1`). - #14346 (2026-11-03): the `expected_rejection` and `no_*_reference` bundle flags. - #14709 (2026-11-02): zero strides in the data SDK `Tensor`. ### Interaction with other open PRs - **#14684, #14686, #14687, #14688** (ASM declines, ALMIOPEN-3035). Each deletes its own entries: TOML skips for the DEF cases, and tier-YAML lines for the quarantined probes it fixes (#14684: prbu01, prbu03, prbu13). Whichever of these and this PR lands second removes the entries. - **#14999** (SDPA forward plan bundles, ALMIOPEN-2893). Both PRs edit two files the same way: - `ASM_SDPA_ENGINE.toml`: keep this PR's generated block and put #14999's `SdpaFwdPlan` skips after the `# END SdpaBackward tolerance classes` line. The generator test only checks the BEGIN to END span. The filters do not overlap (`SdpaBackward` vs `SdpaFwdPlan`). - `SUPPORT_MATRIX_OVERVIEW.md`: re-render it; do not hand-merge. The two PRs also add lines near each other in `integration-tests/CMakeLists.txt` and `test_categories_external.yaml`; these auto-merge. ### Before review - [x] `prbr01` (`dropout_probability = 0`) gated; ALMIOPEN-3149 and ALMIOPEN-3150 track the decision. - [x] Rejection probes carry `expected_rejection`; the six ASM accepts are quarantined in the ASM tier YAML with their tickets. - [x] `prbm02` and `prbm02f` removed; the frontend unit test of R14 covers them. - [x] The five dBias probes and GL-14 written; REG-08 (rank-3 Stats) still waits for the derivation extension. - [x] The plan refreshed against the merged RFC: DV-1 and DV-2 now agree with it, DV-3 is an RFC §11 gap, and only DV-4 remains a deviation. - [x] The measured run (plan §4.5) recorded above. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
ASM_SDPA_ENGINEcan produce NaN or Inf for zero, negative, or subnormal attention scales. Forward and backward now decline known unsupported scales during applicability, and reject unsupported pure runtime scales withBAD_PARAMbefore launching a kernel. DOUBLE scale tensors are checked for float-range overflow before narrowing.JIRA ID : ALMIOPEN-3043
Very large finite float scales remain a separate numerical issue, tracked by ALMIOPEN-3137 under ALMIOPEN-3035. Accepting
FLT_MAXin the guard tests does not assert numerical correctness at that scale.Risk Assessment
Risk: 2/5. This changes ASM SDPA engine selection and execution validation for scales the kernels cannot handle. The scale policy and diagnostics are shared by forward and backward. Ordinary supported scales retain the same float rounding and kernel dispatch. No kernel, public API, or schema is changed.
ASIC Coverage
Forward is affected on gfx942 and gfx950; backward is supported on gfx942. Both architectures require device validation because applicability and launch behavior are involved. The build targets both architectures with SDPA enabled and the kernel ingestor disabled.
Testing Summary
Validation of commit
8bef35d9a40, rebased onto develop63939d78257, using thehip-kernel-providerpreset in Release mode,GPU_TARGETS=gfx942;gfx950, andHIPDNN_ENABLE_SDPA=ON.hip_kernel_provider_tests --gtest_filter='TestSdpa*:TestAsmSdpa*'hip_kernel_provider_integration_tests --gtest_filter='*Sdpa*-*DISABLED*'Host scale tests (
TestSdpaPlanUtils.*:TestSdpaFwdPlanExecute.*:TestSdpaBwdPlanExecute.*): 6 passed, 2 GPU-dependent launch tests skipped, 0 failed. Those launch tests passed in both GPU runs. GPU skips reflect existing architecture and reference restrictions. Both jobs ran with themiopenaccount in themiopenpartition. Cross-provider bundle tests were not rerun for this revision.The tests cover constant values, runtime defaults and pure runtime scale tensors; positive normal float endpoints; negative NaN and subnormal values; unsupported DOUBLE values including the first double above
FLT_MAX; and diagnostic precision at the normal/subnormal boundary. Endpoint acceptance tests check applicability and pre-launch validation, not large-scale numerics.Testing Checklist
Technical Changes
FLT_MINthat round up to it.max_digits10precision.Interaction with #14622: its DEF-06 skips for unsupported scales can be removed once this guard lands; whichever PR lands second should remove them.