From 6737d66a66206f1ad90ab9da5871ada2eeb02875 Mon Sep 17 00:00:00 2001 From: atamazov Date: Fri, 6 Oct 2023 21:36:24 +0300 Subject: [PATCH] rocblas-remove-int8x4-uses(02) Removed the rest of INT8x4 support (except the driver) --- driver/conv_driver.hpp | 4 ++++ include/miopen/miopen.h | 1 - src/check_numerics.cpp | 1 - src/driver_arguments.cpp | 4 +--- src/gemm_v2.cpp | 16 ++++++++-------- src/hip/general_tensor_reorder_sol.cpp | 2 -- src/include/miopen/conv/problem_description.hpp | 1 - src/include/miopen/datatype.hpp | 5 ----- src/include/miopen/tensor.hpp | 1 - src/include/miopen/visit_float.hpp | 1 - src/ocl/tensorocl.cpp | 6 ------ src/pooling_api.cpp | 1 - src/reducetensor.cpp | 3 --- src/solver/batchnorm/forward_inference_ck.cpp | 2 -- .../conv_ck_igemm_fwd_bias_activ_fused.cpp | 4 ---- ...v_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp | 4 ---- ...v_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp | 4 ---- ...v_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp | 4 ---- .../conv_hip_implicit_gemm_bwd_data_xdlops.cpp | 4 ---- src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp | 4 ---- ...conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp | 4 ---- src/solver/gemm.cpp | 2 -- src/solver/mlir_common.cpp | 1 - src/tensor.cpp | 1 - test/driver.hpp | 2 -- test/gtest/conv_embed_db.cpp | 1 - test/gtest/conv_hip_igemm_xdlops.cpp | 1 - test/gtest/conv_igemm_dynamic.cpp | 1 - test/gtest/conv_igemm_mlir.cpp | 1 - test/gtest/conv_igemm_mlir_xdlops.cpp | 1 - test/gtest/conv_trans.cpp | 1 - test/gtest/db_sync.cpp | 2 -- 32 files changed, 13 insertions(+), 77 deletions(-) diff --git a/driver/conv_driver.hpp b/driver/conv_driver.hpp index 5f67b83588..3c28859776 100644 --- a/driver/conv_driver.hpp +++ b/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(4)) + #if MIOPEN_BACKEND_OPENCL #define STATUS_SUCCESS CL_SUCCESS typedef cl_int status_t; diff --git a/include/miopen/miopen.h b/include/miopen/miopen.h index 7a85ec5875..fd978c6200 100644 --- a/include/miopen/miopen.h +++ b/include/miopen/miopen.h @@ -350,7 +350,6 @@ typedef enum 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) */ miopenBFloat16 = 5, /*!< 16-bit binary floating point (8-bit exponent, 7-bit fraction) (Partially supported) */ miopenDouble = 6, /*!< 64-bit floating point (Partially supported) */ diff --git a/src/check_numerics.cpp b/src/check_numerics.cpp index 50cb0af4af..12210069df 100644 --- a/src/check_numerics.cpp +++ b/src/check_numerics.cpp @@ -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 {""}; } diff --git a/src/driver_arguments.cpp b/src/driver_arguments.cpp index 53aa6b41ab..d9e82d270a 100644 --- a/src/driver_arguments.cpp +++ b/src/driver_arguments.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(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; diff --git a/src/gemm_v2.cpp b/src/gemm_v2.cpp index 35add906a1..a44daf962d 100644 --- a/src/gemm_v2.cpp +++ b/src/gemm_v2.cpp @@ -610,9 +610,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; } @@ -876,10 +876,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; } @@ -1140,10 +1140,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; } diff --git a/src/hip/general_tensor_reorder_sol.cpp b/src/hip/general_tensor_reorder_sol.cpp index f78aa358f8..ae85048f63 100644 --- a/src/hip/general_tensor_reorder_sol.cpp +++ b/src/hip/general_tensor_reorder_sol.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); } diff --git a/src/include/miopen/conv/problem_description.hpp b/src/include/miopen/conv/problem_description.hpp index 438e1d306a..12d6f08811 100644 --- a/src/include/miopen/conv/problem_description.hpp +++ b/src/include/miopen/conv/problem_description.hpp @@ -51,7 +51,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"; diff --git a/src/include/miopen/datatype.hpp b/src/include/miopen/datatype.hpp index 29057e3de9..2260cb5b4f 100644 --- a/src/include/miopen/datatype.hpp +++ b/src/include/miopen/datatype.hpp @@ -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,7 +153,6 @@ 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; diff --git a/src/include/miopen/tensor.hpp b/src/include/miopen/tensor.hpp index a133ab4adc..1bf2bcbaae 100644 --- a/src/include/miopen/tensor.hpp +++ b/src/include/miopen/tensor.hpp @@ -101,7 +101,6 @@ inline std::size_t GetTypeSize(miopenDataType_t d) case miopenFloat: return 4; case miopenHalf: case miopenBFloat16: return 2; - case miopenInt8x4: break; case miopenInt8: case miopenFloat8: case miopenBFloat8: return 1; diff --git a/src/include/miopen/visit_float.hpp b/src/include/miopen/visit_float.hpp index 35e1ae6ae7..ae84b7f0a9 100644 --- a/src/include/miopen/visit_float.hpp +++ b/src/include/miopen/visit_float.hpp @@ -91,7 +91,6 @@ void visit_float(miopenDataType_t t, F f) f(as_float{}); break; } - case miopenInt8x4: MIOPEN_THROW("miopenInt8x4: Support discontinued."); } } diff --git a/src/ocl/tensorocl.cpp b/src/ocl/tensorocl.cpp index 1e1f8b1df9..6329d33f69 100644 --- a/src/ocl/tensorocl.cpp +++ b/src/ocl/tensorocl.cpp @@ -1944,7 +1944,6 @@ std::string GetCastTensorBuildOptionFromType(const std::string& buildOption, mio case miopenDouble: // TODO MIOPEN_THROW(miopenStatusBadParm, "miopenDouble data type not supported in cast tensor."); - case miopenInt8x4: // fallthrough default: MIOPEN_THROW(miopenStatusBadParm, "Invalid data type in cast tensor desc."); } } @@ -1968,11 +1967,6 @@ void CastTensor(const Handle& handle, MIOPEN_THROW(miopenStatusBadParm, "Tensor dimension lengths do not match."); } - if(srcDesc.GetType() == miopenInt8x4 || dstDesc.GetType() == miopenInt8x4) - { - MIOPEN_THROW(miopenStatusBadParm, "Tensor cast operation is not supported for int8x4."); - } - auto flat_descriptors = GetConsistentFlattenedTensorDescriptors(srcDesc, dstDesc); const TensorDescriptor& srcDesc_flat = std::get<0>(flat_descriptors); const TensorDescriptor& dstDesc_flat = std::get<1>(flat_descriptors); diff --git a/src/pooling_api.cpp b/src/pooling_api.cpp index bf318f7b78..a2c792f8bb 100644 --- a/src/pooling_api.cpp +++ b/src/pooling_api.cpp @@ -50,7 +50,6 @@ inline void Pooling_logging_cmd(const miopenPoolingDescriptor_t poolDesc, case miopenFloat: ss << "pool"; break; case miopenInt32: case miopenInt8: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: case miopenFloat8: diff --git a/src/reducetensor.cpp b/src/reducetensor.cpp index 9cabbc03f6..6ca1067fb8 100644 --- a/src/reducetensor.cpp +++ b/src/reducetensor.cpp @@ -211,7 +211,6 @@ inline int GetDataTypeSize(miopenDataType_t t) case miopenFloat8: case miopenBFloat8: case miopenInt8: return (1); - case miopenInt8x4: return (4); // Support discontinued. case miopenBFloat16: return (2); case miopenInt32: return (4); default: MIOPEN_THROW("Only float, half, double, bfloat16, int8 data types are supported."); @@ -268,7 +267,6 @@ inline int GetDataTypeId(miopenDataType_t t) case miopenBFloat16: return (static_cast('B')); case miopenDouble: return (static_cast('D')); case miopenInt8: - case miopenInt8x4: // Support discontinued. case miopenFloat8: case miopenBFloat8: case miopenInt32: return (static_cast('O')); @@ -308,7 +306,6 @@ static ck::DataTypeEnum_t mapDataTypeId(miopenDataType_t t) case miopenBFloat16: return DataTypeEnum_t::BFloat16; case miopenDouble: return DataTypeEnum_t::Double; case miopenInt8: return DataTypeEnum_t::Int8; - case miopenInt8x4: return DataTypeEnum_t::Int8x4; case miopenInt32: return DataTypeEnum_t::Int32; case miopenFloat8: case miopenBFloat8: diff --git a/src/solver/batchnorm/forward_inference_ck.cpp b/src/solver/batchnorm/forward_inference_ck.cpp index ff17432a62..d1d67ddea6 100644 --- a/src/solver/batchnorm/forward_inference_ck.cpp +++ b/src/solver/batchnorm/forward_inference_ck.cpp @@ -200,7 +200,6 @@ bool BnCKFwdInference::IsApplicable(const ExecutionContext& context, return (CheckCKApplicability(bn_problem) != -1); case miopenInt32: case miopenInt8: - case miopenInt8x4: // Support discontinued. case miopenFloat8: case miopenBFloat8: default: MIOPEN_THROW("Unsupported datatype"); @@ -244,7 +243,6 @@ BnCKFwdInference::GetSolution(const ExecutionContext& context, break; case miopenInt8: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenFloat8: case miopenBFloat8: default: MIOPEN_THROW("Unsupported datatype"); diff --git a/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp b/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp index e2df6f8097..d9c0410d98 100644 --- a/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp +++ b/src/solver/conv_ck_igemm_fwd_bias_activ_fused.cpp @@ -287,7 +287,6 @@ void PerformanceConfigConvCKIgemmFwdBiasActivFused::HeuristicInit( case miopenInt8: case miopenFloat: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: default: MIOPEN_THROW("Unsupported datatype"); @@ -342,7 +341,6 @@ bool PerformanceConfigConvCKIgemmFwdBiasActivFused::IsValid( case miopenInt8: case miopenFloat: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: default: MIOPEN_THROW("Unsupported datatype"); @@ -435,7 +433,6 @@ bool ConvCKIgemmFwdBiasActivFused::IsApplicable(const FusionContext& ctx, case miopenInt8: case miopenFloat: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: default: MIOPEN_THROW("Unsupported datatype"); @@ -469,7 +466,6 @@ ConvSolution ConvCKIgemmFwdBiasActivFused::GetSolution( case miopenInt8: case miopenFloat: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: default: MIOPEN_THROW("Unsupported datatype"); diff --git a/src/solver/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp index 94983c7f0e..cd95c85cf6 100644 --- a/src/solver/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp @@ -210,7 +210,6 @@ void PerformanceConfigHipImplicitGemm3DGroupBwdXdlops::HeuristicInit( case miopenFloat: Init(problem); break; case miopenInt8: Init(problem); break; case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenFloat8: case miopenBFloat8: @@ -253,7 +252,6 @@ bool PerformanceConfigHipImplicitGemm3DGroupBwdXdlops::IsValid( case miopenFloat: return CheckIsSupportCKArgs(problem); case miopenInt8: return CheckIsSupportCKArgs(problem); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenFloat8: case miopenBFloat8: @@ -324,7 +322,6 @@ bool ConvHipImplicitGemm3DGroupBwdXdlops::IsApplicable( case miopenFloat: return CheckCKApplicability(problem); case miopenInt8: return CheckCKApplicability(problem); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenFloat8: case miopenBFloat8: @@ -352,7 +349,6 @@ ConvSolution ConvHipImplicitGemm3DGroupBwdXdlops::GetSolution( return InitInvokerFactory, CKArgs, conv::DataInvokeParams>( problem, config.kernel_id); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: case miopenFloat8: diff --git a/src/solver/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp index a21c9ba300..4a5ecc8a3f 100644 --- a/src/solver/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp @@ -210,7 +210,6 @@ void PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::HeuristicInit( case miopenFloat: Init(problem); break; case miopenInt8: Init(problem); break; case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenFloat8: case miopenBFloat8: case miopenBFloat16: @@ -253,7 +252,6 @@ bool PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::IsValid( case miopenFloat: return CheckIsSupportCKArgs(problem); case miopenInt8: return CheckIsSupportCKArgs(problem); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenFloat8: case miopenBFloat8: case miopenBFloat16: @@ -322,7 +320,6 @@ bool ConvHipImplicitGemm3DGroupFwdXdlops::IsApplicable( case miopenFloat: return CheckCKApplicability(problem); case miopenInt8: return CheckCKApplicability(problem); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenFloat8: case miopenBFloat8: case miopenBFloat16: @@ -350,7 +347,6 @@ ConvSolution ConvHipImplicitGemm3DGroupFwdXdlops::GetSolution( return InitInvokerFactory, CKArgs, conv::DataInvokeParams>( problem, config.kernel_id); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: case miopenFloat8: diff --git a/src/solver/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp index 5c92f57142..f09fca29fd 100644 --- a/src/solver/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp @@ -206,7 +206,6 @@ void PerformanceConfigHipImplicitGemm3DGroupWrwXdlops::HeuristicInit( case miopenFloat: Init(problem); break; case miopenInt8: Init(problem); break; case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenFloat8: case miopenBFloat8: @@ -249,7 +248,6 @@ bool PerformanceConfigHipImplicitGemm3DGroupWrwXdlops::IsValid( case miopenFloat: return CheckIsSupportCKArgs(problem); case miopenInt8: return CheckIsSupportCKArgs(problem); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenFloat8: case miopenBFloat8: @@ -320,7 +318,6 @@ bool ConvHipImplicitGemm3DGroupWrwXdlops::IsApplicable( case miopenFloat: return CheckCKApplicability(problem); case miopenInt8: return CheckCKApplicability(problem); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenFloat8: case miopenBFloat8: @@ -348,7 +345,6 @@ ConvSolution ConvHipImplicitGemm3DGroupWrwXdlops::GetSolution( return InitInvokerFactory, CKArgs, conv::WrWInvokeParams>( problem, config.kernel_id); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenFloat8: case miopenBFloat8: diff --git a/src/solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp index 5aa17e75bf..d55ded78de 100644 --- a/src/solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp @@ -172,7 +172,6 @@ void PerformanceConfigHipImplicitGemmBwdXdlops::HeuristicInit( case miopenBFloat8: case miopenInt8: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: break; } @@ -214,7 +213,6 @@ bool PerformanceConfigHipImplicitGemmBwdXdlops::IsValid( case miopenBFloat8: case miopenInt8: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: break; } @@ -294,7 +292,6 @@ bool ConvHipImplicitGemmBwdXdlops::IsApplicable( case miopenBFloat8: case miopenInt8: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: break; } @@ -318,7 +315,6 @@ ConvSolution ConvHipImplicitGemmBwdXdlops::GetSolution( problem, config.kernel_id); case miopenInt8: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: case miopenFloat8: diff --git a/src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp index 7a9c322bc7..697ab33f70 100644 --- a/src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp @@ -172,7 +172,6 @@ void PerformanceConfigHipImplicitGemmFwdXdlops::HeuristicInit( case miopenFloat8: case miopenBFloat8: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: break; } @@ -215,7 +214,6 @@ bool PerformanceConfigHipImplicitGemmFwdXdlops::IsValid( case miopenFloat8: case miopenBFloat8: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: break; } @@ -294,7 +292,6 @@ bool ConvHipImplicitGemmFwdXdlops::IsApplicable( case miopenFloat8: case miopenBFloat8: case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: break; } @@ -320,7 +317,6 @@ ConvSolution ConvHipImplicitGemmFwdXdlops::GetSolution( return InitInvokerFactory, CKArgs, conv::DataInvokeParams>( problem, config.kernel_id); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: case miopenFloat8: diff --git a/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp index 0a8b6eeb03..0c5a91b627 100644 --- a/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp @@ -198,7 +198,6 @@ void PerformanceConfigHipImplicitGemmGroupFwdXdlops::HeuristicInit( case miopenFloat: Init(problem); break; case miopenInt8: Init(problem); break; case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenFloat8: case miopenBFloat8: @@ -240,7 +239,6 @@ bool PerformanceConfigHipImplicitGemmGroupFwdXdlops::IsValid( case miopenFloat: return CheckIsSupportCKArgs(problem); case miopenInt8: return CheckIsSupportCKArgs(problem); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenFloat8: case miopenBFloat8: @@ -311,7 +309,6 @@ bool ConvHipImplicitGemmGroupFwdXdlops::IsApplicable( case miopenFloat: return CheckCKApplicability(problem); case miopenInt8: return CheckCKApplicability(problem); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenFloat8: case miopenBFloat8: @@ -339,7 +336,6 @@ ConvSolution ConvHipImplicitGemmGroupFwdXdlops::GetSolution( return InitInvokerFactory, CKArgs, conv::DataInvokeParams>( problem, config.kernel_id); case miopenInt32: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: case miopenFloat8: diff --git a/src/solver/gemm.cpp b/src/solver/gemm.cpp index 6ae4f5dde7..2cfbd257dc 100644 --- a/src/solver/gemm.cpp +++ b/src/solver/gemm.cpp @@ -89,8 +89,6 @@ bool GemmFwdBase::IsApplicable(const ExecutionContext& ctx, const auto& yDesc = problem.GetOut(); // rocBlas needs the output to be 32-bit always - if(xDesc.GetType() == miopenInt8x4) - return false; if(xDesc.GetType() == miopenInt8 // && (yDesc.GetType() != miopenFloat // && yDesc.GetType() != miopenInt32)) diff --git a/src/solver/mlir_common.cpp b/src/solver/mlir_common.cpp index 4e41b8944a..e54da1ac89 100644 --- a/src/solver/mlir_common.cpp +++ b/src/solver/mlir_common.cpp @@ -57,7 +57,6 @@ static const char* DTypeName(miopenDataType_t ty) case miopenBFloat16: return "bf16"; case miopenInt32: return "i32"; case miopenInt8: return "i8"; - case miopenInt8x4: return "i8x4"; // Support discontinued. case miopenFloat8: return "fp8"; case miopenBFloat8: return "bfp8"; } diff --git a/src/tensor.cpp b/src/tensor.cpp index df0d7c2819..ba1b2ba106 100644 --- a/src/tensor.cpp +++ b/src/tensor.cpp @@ -50,7 +50,6 @@ bool IsDataTypeSupported(miopenDataType_t t) case miopenFloat8: case miopenBFloat8: case miopenInt8: - case miopenInt8x4: // Support discontinued. case miopenBFloat16: case miopenDouble: return true; } diff --git a/test/driver.hpp b/test/driver.hpp index fd83dd1fc5..bcbfe7d676 100644 --- a/test/driver.hpp +++ b/test/driver.hpp @@ -274,7 +274,6 @@ struct test_driver { case miopenHalf: ss << "--half "; break; case miopenBFloat16: ss << "--bfloat16 "; break; - case miopenInt8x4: ss << "--UNSUPPORED_TYPE "; break; case miopenInt8: ss << "--int8 "; break; case miopenInt32: ss << "--int32 "; break; case miopenFloat: ss << "--float "; break; @@ -303,7 +302,6 @@ struct test_driver { case miopenHalf: ret.emplace_back("--half"); break; case miopenBFloat16: ret.emplace_back("--bf16"); break; - case miopenInt8x4: ret.emplace_back("--UNSUPPORTED_TYPE"); break; case miopenInt8: ret.emplace_back("--int8"); break; case miopenInt32: ret.emplace_back("--int32"); break; case miopenFloat: ret.emplace_back("--float"); break; diff --git a/test/gtest/conv_embed_db.cpp b/test/gtest/conv_embed_db.cpp index 4672bb0404..36d1e28123 100644 --- a/test/gtest/conv_embed_db.cpp +++ b/test/gtest/conv_embed_db.cpp @@ -73,7 +73,6 @@ void Run2dDriver(miopenDataType_t prec) case miopenHalf: params = ConfigWithHalf::GetParam(); break; case miopenInt8: params = ConfigWithInt8::GetParam(); break; case miopenBFloat16: params = ConfigWithBFloat16::GetParam(); break; - case miopenInt8x4: // Support discontinued. case miopenInt32: case miopenFloat8: case miopenBFloat8: diff --git a/test/gtest/conv_hip_igemm_xdlops.cpp b/test/gtest/conv_hip_igemm_xdlops.cpp index 3407446557..14e1738a35 100644 --- a/test/gtest/conv_hip_igemm_xdlops.cpp +++ b/test/gtest/conv_hip_igemm_xdlops.cpp @@ -65,7 +65,6 @@ void Run2dDriver(miopenDataType_t prec) case miopenHalf: case miopenBFloat16: case miopenFloat: - case miopenInt8x4: // Support discontinued. case miopenInt32: case miopenDouble: FAIL() << "miopenHalf, miopenBFloat16, miopenFloat, miopenInt32, " diff --git a/test/gtest/conv_igemm_dynamic.cpp b/test/gtest/conv_igemm_dynamic.cpp index 59341ac8a0..6e2863eb8f 100644 --- a/test/gtest/conv_igemm_dynamic.cpp +++ b/test/gtest/conv_igemm_dynamic.cpp @@ -68,7 +68,6 @@ void Run2dDriver(miopenDataType_t prec) case miopenHalf: case miopenInt8: case miopenBFloat16: - case miopenInt8x4: // Support discontinued. case miopenInt32: case miopenDouble: case miopenFloat8: diff --git a/test/gtest/conv_igemm_mlir.cpp b/test/gtest/conv_igemm_mlir.cpp index 29d240645c..c99d0a03c4 100644 --- a/test/gtest/conv_igemm_mlir.cpp +++ b/test/gtest/conv_igemm_mlir.cpp @@ -81,7 +81,6 @@ void Run2dDriver(miopenDataType_t prec) case miopenInt8: params = ConfigWithInt8::GetParam(); break; case miopenFloat: params = ConfigWithFloat::GetParam(); break; case miopenBFloat16: - case miopenInt8x4: // Support discontinued. case miopenInt32: case miopenFloat8: case miopenBFloat8: diff --git a/test/gtest/conv_igemm_mlir_xdlops.cpp b/test/gtest/conv_igemm_mlir_xdlops.cpp index e2c5a3ed8f..ea14e3b68e 100644 --- a/test/gtest/conv_igemm_mlir_xdlops.cpp +++ b/test/gtest/conv_igemm_mlir_xdlops.cpp @@ -53,7 +53,6 @@ void Run2dDriver(miopenDataType_t prec) case miopenInt8: params = ConfigWithInt8::GetParam(); break; case miopenBFloat16: case miopenFloat: - case miopenInt8x4: // Support discontinued. case miopenInt32: case miopenDouble: case miopenFloat8: diff --git a/test/gtest/conv_trans.cpp b/test/gtest/conv_trans.cpp index 185a459947..a5266d9f9c 100644 --- a/test/gtest/conv_trans.cpp +++ b/test/gtest/conv_trans.cpp @@ -55,7 +55,6 @@ void Run2dDriver(miopenDataType_t prec) case miopenBFloat8: case miopenInt8: case miopenBFloat16: - case miopenInt8x4: // Support discontinued. case miopenInt32: case miopenDouble: FAIL() << "miopenHalf, miopenInt8, miopenBFloat16, miopenInt32, miopenDouble " diff --git a/test/gtest/db_sync.cpp b/test/gtest/db_sync.cpp index 13b647ceac..13f0297724 100644 --- a/test/gtest/db_sync.cpp +++ b/test/gtest/db_sync.cpp @@ -72,8 +72,6 @@ miopenDataType_t GetDataTypeFromString(const std::string& data_type) return miopenHalf; else if(data_type == "INT8") return miopenInt8; - else if(data_type == "INT8x4") - return miopenInt8x4; // Support discontinued. Maintain compatibility with old databases. else if(data_type == "INT32") return miopenInt32; else if(data_type == "BF16")