Fix the f8 reference kernel issue that failed CI (ROCm#2586)

atamazov · Dec 7, 2023 · 8ee1ad7 · 8ee1ad7
1 parent 762ae81
commit 8ee1ad7
Show file tree

Hide file tree

Showing 4 changed files with 28 additions and 8 deletions.
diff --git a/src/kernels/hip_f8_impl.hpp b/src/kernels/hip_f8_impl.hpp
@@ -202,12 +202,13 @@ MIOPEN_HIP_HOST_DEVICE uint8_t cast_to_f8(T _x, bool stoch, uint32_t rng)
         }
         mantissa += (1 << mfmt); // Add the implicit 1 into mantissa
     }
-    const long tmp = (mfmt - wm + exponent_diff);
-    if(tmp == 33)
-        printf("Gotcha");
 
-    bool midpoint = (mantissa & ((static_cast<uint32_t>(1) << (mfmt - wm + exponent_diff)) - 1)) ==
-                    (static_cast<uint32_t>(1) << (mfmt - wm + exponent_diff - 1));
+    bool midpoint;
+    if(exponent_diff <= wm)
+        midpoint = (mantissa & ((1 << (mfmt - wm + exponent_diff)) - 1)) ==
+                   (1 << (mfmt - wm + exponent_diff - 1));
+    else
+        midpoint = false;
     /* This part is a bit tricky. The judgment of whether it is a tie needs to be done before we
        shift right as shift right could rip off some residual part and make something not midpoint
        look like midpoint. For example, the fp16 number 0x1002 (0 00100 0000000010), it is larger

diff --git a/test/gtest/conv_f8_bwd.cpp b/test/gtest/conv_f8_bwd.cpp
@@ -34,7 +34,14 @@
 
 std::vector<Conv3DTestCase> ConvTestConfigs()
 { // g    n   c   d    h   w   k   z  y  x pad_x pad_y pad_z stri_x stri_y stri_z dia_x dia_y dia_z
-    return {{1, 16, 16, 1, 14, 14, 16, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution}};
+    return {{1, 16, 16, 1, 14, 14, 16, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 128, 64, 1, 28, 28, 64, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 64, 32, 1, 28, 28, 32, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {32, 128, 32, 1, 28, 28, 32, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {16, 128, 16, 1, 28, 28, 16, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {8, 128, 8, 1, 28, 28, 8, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {4, 128, 4, 1, 28, 28, 4, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {2, 128, 2, 1, 28, 28, 2, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution}};
 }
 
 template <typename T = float>

diff --git a/test/gtest/conv_f8_fwd.cpp b/test/gtest/conv_f8_fwd.cpp
@@ -34,7 +34,12 @@
 
 std::vector<Conv3DTestCase> ConvTestConfigs()
 { // g    n   c   d    h   w   k   z  y  x pad_x pad_y pad_z stri_x stri_y stri_z dia_x dia_y dia_z
-    return {{1, 16, 16, 1, 14, 14, 16, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution}};
+    return {{1, 16, 16, 1, 14, 14, 16, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 64, 64, 1, 14, 14, 64, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 64, 32, 1, 28, 28, 32, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {2, 128, 32, 1, 28, 28, 32, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {32, 128, 32, 1, 28, 28, 32, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {5, 120, 60, 1, 28, 28, 60, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution}};
 }
 
 template <typename T = float>

diff --git a/test/gtest/conv_f8_wrw.cpp b/test/gtest/conv_f8_wrw.cpp
@@ -35,7 +35,14 @@
 
 std::vector<Conv3DTestCase> ConvTestConfigs()
 { // g   n   c   d    h   w   k   z  y  x pad_x pad_y pad_z stri_x stri_y stri_z dia_x dia_y dia_z
-    return {{1, 16, 16, 1, 14, 14, 16, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution}};
+    return {{1, 16, 16, 1, 14, 14, 16, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 64, 128, 1, 28, 3, 128, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 64, 64, 1, 28, 3, 64, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 32, 64, 1, 14, 14, 64, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 32, 32, 1, 14, 14, 32, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 64, 32, 1, 14, 14, 32, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 128, 64, 1, 7, 7, 64, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {1, 128, 32, 1, 7, 7, 32, 1, 3, 3, 1, 1, 0, 1, 1, 1, 1, 1, 1, miopenConvolution}};
 }
 
 template <typename T = float>