MIOpen error tolerance by Beanavil · Pull Request #3638 · ROCm/MIOpen

Beanavil · 2025-03-17T20:02:29Z

Scope

While testing the variant 0 for bnorm backward spatial single, we observed that the error tolerance in MIOpenDriver is sub-optimal. For instance, the following results were observed

Backwards prop batch norm verification passed on dx (0.151171)
Backwards prop batch norm verification passed on dscale (0.277024)
Backwards prop batch norm verification passed on dbias (0.217736)
Backwards Prop Batch Norm Verifies on CPU and GPU.

Notice that the execution is successful, according to the driver checks. However, the error reported is high so this should actually be failing.

Notes for the reviewer

Fixed the maxrms definition because it was too big compared to the rms values being computed (due to the rms normalization done). The VerifyForward method seems to also be using the same value for maxrms as the new implementation in VerifyBackward.
Fixed the normalization of the rms. Previously we were finding the maximum absolute value from both the reference values and the results, but this may not be correct because it reduces the significance of the rms: if the results differ a lot from the reference values (e.g. ref values are order of 10 and results are order 1000) then the normalization will divide the rms obtained (which will be order of 100) by a number order of 1000 and then the normalized rms will be 100/1000 = 0.1 which is way lower than what it should be.
- The new implementation only takes into account the reference results for computing the normalization factor

BrianHarrisonAMD · 2025-03-19T14:32:14Z


+// Reference values are r1, measured values are r2
 template <class R1, class R2>
 double rms_range(R1&& r1, R2&& r2)


I think we should rename these arguments to make it clear that r1 is intended to be the reference.

BrianHarrisonAMD · 2025-03-19T14:32:52Z

-        double mag2 = static_cast<double>(*std::max_element(r2.begin(), r2.end(), compare_mag));
-        double mag =
-            std::max({std::fabs(mag1), std::fabs(mag2), std::numeric_limits<double>::min()});
-        return std::sqrt(square_difference) / (std::sqrt(n) * mag);
+        double normalizer = std::max({std::fabs(mag1), std::numeric_limits<double>::min()});
+        return std::sqrt(square_difference) / (std::sqrt(n) * normalizer);


Hehe this will be a spicy one.
Interesting find though!

BrianHarrisonAMD · 2025-03-19T14:37:19Z

Starting CI!

averinevg

LGTM

BrianHarrisonAMD · 2025-03-28T19:36:33Z

Restarting CI again.

BrianHarrisonAMD · 2025-03-30T17:43:01Z

Failing on formatting.
Please run the formatting script.

Beanavil · 2025-03-31T08:23:01Z

@BrianHarrisonAMD opsie, missed to run this before. Should be ok now!

Beanavil · 2025-03-31T12:43:42Z

Btw I see that some smoke tests are failing for gfx90a (MI200), but I have double-checked on our side and they pass on our MI200s, how should we proceed for those?

BrianHarrisonAMD · 2025-03-31T13:26:44Z

Btw I see that some smoke tests are failing for gfx90a (MI200), but I have double-checked on our side and they pass on our MI200s, how should we proceed for those?

Looks like it was stopped / aborted before fully finishing.
Ill run the internal CI to see if it passes.

Beanavil · 2025-05-14T09:03:21Z

@BrianHarrisonAMD any updates on this?

BrianHarrisonAMD · 2025-05-14T14:03:59Z

I can start the build again, and let @BradPepersAMD know about it.

BrianHarrisonAMD · 2025-05-14T16:44:00Z

CI running now.

BradPepersAMD · 2025-07-14T06:30:35Z

Did this pass CI and are we good to merge this now?

BrianHarrisonAMD · 2025-07-24T16:50:10Z

We need to update this with latest, and re-run CI on it.
Since it impacts all tests, and driver commands we need to do a large sweep on these changes before they can be safely merged.

amd-hsivasun · 2025-07-28T20:29:14Z

Imported to ROCm/rocm-libraries

@Beanavil

## Scope While testing the variant 0 for bnorm backward spatial single, we observed that the error tolerance in MIOpenDriver is sub-optimal. For instance, the following results were observed Notice that the execution is successful, according to the driver checks. However, the error reported is high so this should actually be failing. ## Notes for the reviewer - Fixed the definition because it was too big compared to the rms values being computed (due to the rms normalization done). The method seems to also be using the same value for as the new implementation in . - Fixed the normalization of the rms. Previously we were finding the maximum absolute value from both the reference values and the results, but this may not be correct because it reduces the significance of the rms: if the results differ a lot from the reference values (e.g. ref values are order of 10 and results are order 1000) then the normalization will divide the rms obtained (which will be order of 100) by a number order of 1000 and then the normalized rms will be 100/1000 = 0.1 which is way lower than what it should be. - The new implementation only takes into account the reference results for computing the normalization factor --- 🔁 Imported from [ROCm/MIOpen#3638](ROCm/MIOpen#3638) 🧑‍💻 Originally authored by @Beanavil --------- Co-authored-by: Beatriz Navidad Vilches <beatriz@streamhpc.com> Co-authored-by: assistant-librarian[bot] <assistant-librarian[bot]@users.noreply.github.com> Co-authored-by: JonathanLichtnerAMD <195780826+JonathanLichtnerAMD@users.noreply.github.com> Co-authored-by: Milo Lurati <70884255+MiloLurati@users.noreply.github.com> Co-authored-by: Bálint Siklósi <25982046+sikba@users.noreply.github.com> Co-authored-by: Balint Siklosi <balint.siklosi@streamhpc.com> Co-authored-by: SreecharanGundaboluAMD <sgundabo@amd.com> Co-authored-by: Ewan Crawford <ewan.crawford@streamhpc.com>

Beanavil requested review from BradPepersAMD, BrianHarrisonAMD, JonathanLichtnerAMD and adickin-amd as code owners March 17, 2025 20:02

Beanavil changed the title ~~MiOpen error tolerance~~ MIOpen error tolerance Mar 17, 2025

BrianHarrisonAMD reviewed Mar 19, 2025

View reviewed changes

averinevg previously approved these changes Mar 19, 2025

View reviewed changes

Comment thread test/verify.hpp Outdated

Beanavil force-pushed the miopen-error-tolerance branch from dd06398 to 389e109 Compare March 21, 2025 15:59

cderb previously approved these changes Mar 21, 2025

View reviewed changes

Beanavil force-pushed the miopen-error-tolerance branch from 389e109 to fea5f7d Compare March 31, 2025 08:21

Fix maxrms computation

ed9e2d9

Beanavil force-pushed the miopen-error-tolerance branch from fea5f7d to 83c6070 Compare May 13, 2025 07:30

Do not use results to normalize the residuals' stddev

0ecf7af

Beanavil force-pushed the miopen-error-tolerance branch from 83c6070 to 0ecf7af Compare May 21, 2025 10:47

Beanavil dismissed stale reviews from cderb and averinevg via 0ecf7af July 28, 2025 18:54

assistant-librarian Bot mentioned this pull request Jul 28, 2025

MIOpen error tolerance ROCm/rocm-libraries#912

Merged

amd-hsivasun closed this Jul 28, 2025

Conversation

Beanavil commented Mar 17, 2025

Scope

Notes for the reviewer

Uh oh!

BrianHarrisonAMD Mar 19, 2025

Choose a reason for hiding this comment

Uh oh!

Beanavil Mar 21, 2025

Choose a reason for hiding this comment

Uh oh!

BrianHarrisonAMD Mar 19, 2025

Choose a reason for hiding this comment

Uh oh!

BrianHarrisonAMD commented Mar 19, 2025

Uh oh!

averinevg left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

BrianHarrisonAMD commented Mar 28, 2025

Uh oh!

BrianHarrisonAMD commented Mar 30, 2025

Uh oh!

Beanavil commented Mar 31, 2025

Uh oh!

Beanavil commented Mar 31, 2025

Uh oh!

BrianHarrisonAMD commented Mar 31, 2025

Uh oh!

Beanavil commented May 14, 2025

Uh oh!

BrianHarrisonAMD commented May 14, 2025

Uh oh!

BrianHarrisonAMD commented May 14, 2025

Uh oh!

BradPepersAMD commented Jul 14, 2025

Uh oh!

BrianHarrisonAMD commented Jul 24, 2025

Uh oh!

amd-hsivasun commented Jul 28, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants