Finally remove INT8x4 support. (#2452)

Author: atamazov

docs/datatypes.md

@@ -12,6 +12,9 @@ typedef enum {
     miopenInt8     = 3,
     /* Value 4 is reserved. */
     miopenBFloat16 = 5,
+    miopenDouble   = 6,
+    miopenFloat8   = 7,
+    miopenBFloat8  = 8
 } miopenDataType_t;
 ```
 
@@ -23,6 +26,9 @@ Type descriptions:
  * `miopenInt32` - 32-bit integer, used primarily for `int8` convolution outputs
  * `miopenInt8` - 8-bit integer, currently only supported by `int8` convolution forward path, tensor set, tensor copy, tensor cast, tensor transform, tensor transpose, and im2col.
  * `miopenBFloat16` - brain float fp-16 (8-bit exponent, 7-bit fraction), currently only supported by convolutions, tensor set, and tensor copy.
+ * `miopenDouble` - 64-bit floating point, supported only by reduction, layerNorm, batchNorm.
+ * `miopenFloat8` - 8-bit floating point (layout 1.4.3, exponent bias 7), currently only supported by convolutions.
+ * `miopenBFloat8` - 8-bit floating point (layout 1.5.2, exponent bias 15), currently only supported by convolutions.
 
 
 Note: In addition to the standard datatypes above, pooling contains its own indexing datatypes:

driver/conv_driver.hpp

@@ -83,6 +83,10 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DRIVER_PAD_BUFFERS_2M)
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DRIVER_USE_GPU_REFERENCE)
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DRIVER_SUBNORM_PERCENTAGE)
 
+// Support in the library discontinued, but left in the driver
+// for reference in the future.
+#define miopenInt8x4 (static_cast<miopenDataType_t>(4))
+
 #if MIOPEN_BACKEND_OPENCL
 #define STATUS_SUCCESS CL_SUCCESS
 typedef cl_int status_t;

fin

@@ -348,11 +348,11 @@ MIOPEN_DECLARE_OBJECT(miopenReduceTensorDescriptor);
  */
 typedef enum
 {
-    miopenHalf     = 0, /*!< 16-bit floating point (Fully supported) */
-    miopenFloat    = 1, /*!< 32-bit floating point (Fully supported) */
-    miopenInt32    = 2, /*!< 32-bit int point (Partially supported) */
-    miopenInt8     = 3, /*!< 8-bit int point (Partially supported) */
-    miopenInt8x4   = 4, /*!< Pack of four Int8 in NCHW_VECT_C format (Support discontinued) */
+    miopenHalf  = 0, /*!< 16-bit floating point (Fully supported) */
+    miopenFloat = 1, /*!< 32-bit floating point (Fully supported) */
+    miopenInt32 = 2, /*!< 32-bit integer (Partially supported) */
+    miopenInt8  = 3, /*!< 8-bit integer (Partially supported) */
+    // miopenInt8x4   = 4, /*!< Pack of 4x Int8 in NCHW_VECT_C format (Support discontinued) */
     miopenBFloat16 = 5, /*!< 16-bit binary floating point (8-bit exponent, 7-bit fraction)
                            (Partially supported) */
     miopenDouble = 6,   /*!< 64-bit floating point (Partially supported) */

include/miopen/miopen.h

@@ -64,7 +64,6 @@ std::string GetKernelName(miopenDataType_t data_type)
     case miopenBFloat8: return {"check_numerics_bf8"};
     case miopenInt32:
     case miopenInt8:
-    case miopenInt8x4: // Support discontinued.
     case miopenDouble:
     default: return {""};
     }

src/check_numerics.cpp

@@ -50,7 +50,7 @@ void ConvDataType(std::stringstream& ss, const miopen::TensorDescriptor& desc)
     {
         ss << "convbfp16";
     }
-    else if(desc.GetType() == miopenInt8 || desc.GetType() == miopenInt8x4)
+    else if(desc.GetType() == miopenInt8)
     {
         ss << "convint8";
     }
@@ -190,8 +190,6 @@ std::string ConvArgsForMIOpenDriver(const miopen::TensorDescriptor& xDesc,
     ss << " -g " << convDesc.group_count;
     if(print_for_conv_driver)
         ss << " -F " << std::to_string(static_cast<int>(conv_dir)) << " -t 1"; // clang-format on
-    if(xDesc.GetType() == miopenInt8x4)
-        ss << " -Z 1";
     if(immediate_mode_solver_id.has_value())
     {
         ss << " -S " << *immediate_mode_solver_id;

src/driver_arguments.cpp

@@ -612,9 +612,9 @@ miopenStatus_t CallGemm(const Handle& handle,
         };
         break;
 
-        case miopenInt8x4:
         case miopenDouble: {
-            MIOPEN_THROW(miopenStatusBadParm, "Unknown or unsupported data type.");
+            MIOPEN_THROW(miopenStatusBadParm,
+                         "miopenDouble data type not supported by MIOpenGEMM.");
         };
         break;
         }
@@ -879,10 +879,10 @@ miopenStatus_t CallGemmStridedBatched(const Handle& handle,
             break;
         }
 
-        case miopenInt8x4:
         case miopenDouble: {
-            MIOPEN_THROW(miopenStatusBadParm, "Unknown or unsupported data type.");
-        };
+            MIOPEN_THROW(miopenStatusBadParm,
+                         "miopenDouble data type not supported by MIOpenGEMM.");
+        }
         break;
         }
 
@@ -1144,10 +1144,10 @@ miopenStatus_t CallGemmStridedBatchedSequential(const Handle& handle,
             break;
         }
 
-        case miopenInt8x4:
         case miopenDouble: {
-            MIOPEN_THROW(miopenStatusBadParm, "Unknown or unsupported data type.");
-        };
+            MIOPEN_THROW(miopenStatusBadParm,
+                         "miopenDouble data type not supported by MIOpenGEMM.");
+        }
         break;
         }
 

src/gemm_v2.cpp

@@ -137,8 +137,6 @@ GenericReorderSolutionImpl::GenericReorderSolutionImpl(miopenDataType_t data_typ
       order_2(order_2_),
       order_3(order_3_)
 {
-    if(data_type == miopenInt8x4)
-        MIOPEN_THROW("These data type are not supported");
     std::size_t data_size  = miopen::GetTypeSize(data_type);
     kernel_param_heuristic = tensor_reorder::HeuristicGet(data_size, dim_0, dim_1, dim_2, dim_3);
 }

src/hip/general_tensor_reorder_sol.cpp

@@ -53,10 +53,6 @@ inline std::string GetDataType(miopenDataType_t type)
         type_str = "bfloat16";
     }
     break;
-    case miopenInt8x4: {
-        type_str = "UNSUPPORTED_TYPE";
-    }
-    break;
     case miopenInt8: {
         type_str = "int8_t";
     }
@@ -157,10 +153,7 @@ inline KernelBuildParameters GetDataTypeKBP(miopenDataType_t type)
     case miopenDouble: use_fp64 = 1; break;
     case miopenFloat8: use_fp8 = 1; break;
     case miopenBFloat8: use_bfp8 = 1; break;
-    case miopenInt8x4: // fallthrough
-    default:
-        MIOPEN_THROW("Only float, half, bfloat16, int8, float8, bfloat8 data types are supported.");
-        break;
+    default: MIOPEN_THROW("Unsupported data type."); break;
     }
 
     auto kbp = KernelBuildParameters{
@@ -180,7 +173,7 @@ inline KernelBuildParameters GetDataTypeKBP(miopenDataType_t type)
     if(use_fp8 != 0)
         kbp.Define("MIOPEN_USE_FP8", use_fp8);
     if(use_bfp8 != 0)
-        kbp.Define("MIOPEN_USE_FP8", use_bfp8);
+        kbp.Define("MIOPEN_USE_BFP8", use_bfp8);
     return kbp;
 }
 

src/include/miopen/datatype.hpp

@@ -37,7 +37,6 @@ inline std::string GetDataTypeName(miopenDataType_t data_type)
     case miopenFloat: return "FP32";
     case miopenHalf: return "FP16";
     case miopenInt8: return "INT8";
-    case miopenInt8x4: return "INT8x4";
     case miopenInt32: return "INT32";
     case miopenBFloat16: return "BF16";
     case miopenDouble: return "FP64";