MIOpen error tolerance

projects/miopen/driver/bn_driver.hpp

@@ -2072,9 +2072,8 @@ int BatchNormDriver<TInput, Tref, TAcc, TScaleBias, TOut>::VerifyBackward()
     if(!back)
         return miopenStatusSuccess;
 
-    const Tref maxrms =
-        static_cast<Tref>(((sizeof(TInput) == 4) ? RMSTOL_FP32 : RMSTOL_FP16) * 1000);
-    bool anError = false;
+    const Tref maxrms = static_cast<Tref>((sizeof(TInput) == 4) ? RMSTOL_FP32 : RMSTOL_FP16);
+    bool anError      = false;
 
     RunBackwardCPU();
 

projects/miopen/test/verify.hpp

@@ -223,23 +223,25 @@ std::size_t mismatch_diff(R1&& r1, R2&& r2, T diff)
             float_equal, diff, std::bind(abs_diff, std::placeholders::_1, std::placeholders::_2)));
 }
 
-template <class R1, class R2>
-double rms_range(R1&& r1, R2&& r2)
+// Reference values are 'ref', observed (measured) values are 'obs'.
+template <class RefType, class ObsType>
+double rms_range(RefType&& ref, ObsType&& obs)
 {
-    std::size_t n = range_distance(r1);
-    if(n == range_distance(r2))
+    std::size_t n = range_distance(ref);
+    if(n == range_distance(obs))
     {
         if(n == 0)
             return 0;
-        double square_difference = range_product(r1, r2, 0.0, sum_fn{}, square_diff);
-        double mag1 = static_cast<double>(*std::max_element(r1.begin(), r1.end(), compare_mag));
-        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 square_difference = range_product(ref, obs, 0.0, sum_fn{}, square_diff);
+        double mag1 = static_cast<double>(*std::max_element(ref.begin(), ref.end(), compare_mag));
+        // If reference is all 0, avoid dividing by 0 when normalizing
+        mag1              = (mag1 == 0.0) ? 1.0 : std::fabs(mag1);
+        double normalizer = std::max(mag1, std::numeric_limits<double>::min());
+
+        return std::sqrt(square_difference) / (std::sqrt(n) * normalizer);
     }
     else
-        return double(std::numeric_limits<range_value<R1>>::max());
+        return double(std::numeric_limits<range_value<RefType>>::max());
 }
 } // namespace miopen
 #endif