Revert "MIOpen:feature:tf32:demonstrate tf32 in conv3d on MI30X platform"

projects/miopen/Jenkinsfile

@@ -156,10 +156,6 @@ pipeline {
             name: "DATATYPE_FP32",
             defaultValue: true,
             description: "")
-        booleanParam(
-            name: "DATATYPE_TF32",
-            defaultValue: true,
-            description: "")
         booleanParam(
             name: "DATATYPE_FP16",
             defaultValue: true,
@@ -203,7 +199,6 @@ pipeline {
         Bf16_flags      = " -DMIOPEN_TEST_BFLOAT16=On"
         Int8_flags      = " -DMIOPEN_TEST_INT8=On"
         Full_test       = " -DMIOPEN_TEST_ALL=On"
-        Tf32_flags      = " -DMIOPEN_TEST_TF32=On"
 
         gfx908_flags    = " -DMIOPEN_INSTALL_GPU_DATABASES=gfx908"
         gfx90a_flags    = " -DMIOPEN_INSTALL_GPU_DATABASES=gfx90a"
@@ -531,24 +526,6 @@ pipeline {
                         runBuildAndSingleGtestJob(Full_test + gfx1101_flags, Build_timeout_minutes)
                     }
                 }
-                stage('TF32 Hip All gfx942') {
-                    when {
-                        beforeAgent true
-                        expression { params.TARGET_GFX942 && params.DATATYPE_TF32 }
-                    }
-                    options {
-                        retry(2)
-                    }
-                    agent{ label rocmnode("gfx942") }
-                    steps{
-                        runBuildAndSingleGtestJob(Full_test + Tf32_flags + gfx942_flags, Build_timeout_minutes)
-                    }
-                    post {
-                        always {
-                            cleanWs()
-                        }
-                    }
-                }
             }
         }
         stage("Nightly Tests") {

projects/miopen/driver/conv_driver.hpp

@@ -456,12 +456,6 @@ class ConvDriver : public Driver
         constexpr bool is_bfp8 = std::is_same<Tgpu, bfloat8_fnuz>::value;
         if(is_bfp8 || is_fp8 || TensorsCasted())
             tolerance *= 37.0;
-
-        { // tf32 has same mantissa length as fp16
-            auto math_type_ = inflags.GetValueInt("math_type");
-            if(std::is_same_v<Tgpu, float> && (miopen::EnvEnableTF32() || math_type_))
-                tolerance = 8.2e-3;
-        }
         return tolerance;
     }
 
@@ -874,8 +868,6 @@ int ConvDriver<Tgpu, Tref>::GetandSetData()
             warmupConvDesc,
             static_cast<int>(miopenConvolutionFindModeNormal)); // Repeat via hidden API.
         miopenSetConvolutionGroupCount(warmupConvDesc, group_count);
-        miopenSetConvolutionAttribute(
-            warmupConvDesc, MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE, inflags.GetValueInt("math_type"));
 
         int warmup_out_len_size = miopen::deref(warmupInputTensor).GetNumDims();
         std::vector<int> warmup_out_len(warmup_out_len_size);
@@ -1030,8 +1022,6 @@ int ConvDriver<Tgpu, Tref>::AddCmdLineArgs()
                          "0",
                          "MIOpen tuning policy (Default=0, or no tuning policy set)",
                          "int");
-    // TODO:(LYM) change back to 0
-    inflags.AddInputFlag("math_type", 'M', "1", "math type of compute (Default=1)", "int");
 
     return 0;
 }
@@ -1236,14 +1226,6 @@ int ConvDriver<Tgpu, Tref>::SetConvDescriptorFromCmdLineArgs()
         miopenSetTransposeConvNdOutputPadding(convDesc, spatial_dim, trans_output_pads.data());
     }
 
-    auto math_type_ = inflags.GetValueInt("math_type");
-    if(math_type_ < miopenMathDefault || math_type_ > miopenMathPedantic)
-    {
-        std::cout << "Invalid math_type value: " << math_type_ << std::endl;
-        exit(0); // NOLINT (concurrency-mt-unsafe)
-    }
-    miopenSetConvolutionAttribute(convDesc, MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE, math_type_);
-
     return miopenStatusSuccess;
 }
 

projects/miopen/include/miopen/miopen.h

@@ -120,14 +120,6 @@ typedef enum
     miopenStatusVersionMismatch = 10, /*!< Version mismatch of the supplied binary data argment. */
 } miopenStatus_t;
 
-typedef enum
-{
-    // TODO:(LYM) temporary use Pedantic as default until TF32 is fully supported
-    miopenMathDefault = 0, /*!< Use TF32 if possible */
-    miopenMathPedantic =
-        1, /*!< Default MathType. Strict IEEE compliance. Don't allow datatype down conversion. */
-} miopenMathType_t;
-
 #ifdef MIOPEN_BETA_API
 typedef enum
 {
@@ -647,9 +639,6 @@ typedef enum
 #else
 // miopenReserved1 = 2,
 #endif
-    // TODO:(LYM) temporarily use Pedantic as default until TF32 is fully supported
-    MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE =
-        3, /*!< refer to miopenMathType_t,default is miopenMathPedantic >*/
 } miopenConvolutionAttrib_t;
 
 /*! @ingroup convolutions

projects/miopen/src/conv/problem_description.cpp

@@ -186,14 +186,9 @@ void ProblemDescription::MakeNetworkConfig(std::string& conf_key) const
         ss << 'x' << GetWeightsLayout();
         ss << 'x' << GetOutLayout();
     }
-    const auto data_type =
-        EncodeDataTypesForKey(GetInDataType(), GetWeightsDataType(), GetOutDataType());
-    ss << 'x' << data_type;
+    ss << 'x' << EncodeDataTypesForKey(GetInDataType(), GetWeightsDataType(), GetOutDataType());
 
     std::ostringstream optional;
-    if(data_type == "FP32" && UseTF32())
-        optional << "TF32" << 'x';
-
     if(const auto ct = GetInCastType())
         optional << "ci" << GetDataTypeName(*ct);
     if(const auto ct = GetWeightsCastType())
@@ -244,12 +239,10 @@ void ProblemDescription::Serialize(std::ostream& stream) const
         stream << sep << GetWeightsLayout();
         stream << sep << GetOutLayout();
     }
-    // clang-format on
-    const auto data_type =
-        EncodeDataTypesForKey(GetInDataType(), GetWeightsDataType(), GetOutDataType());
-    stream << sep << data_type;
+    stream << sep << EncodeDataTypesForKey(GetInDataType(), GetWeightsDataType(), GetOutDataType());
     stream << sep << GetDirectionStr();
 
+    // clang-format on
     // New performance config entries shall come into variable/optional part of db key.
     // This is to support backward compatibility with previous versions of databases.
     std::ostringstream optional;
@@ -264,10 +257,6 @@ void ProblemDescription::Serialize(std::ostream& stream) const
             optional << "_cw" << GetDataTypeName(*ct);
         if(const auto ct = GetOutCastType())
             optional << "_co" << GetDataTypeName(*ct);
-
-        // cx indicates compute datatype
-        if(data_type == "FP32" && UseTF32())
-            optional << "_cxTF32";
     }
     if(!optional.str().empty())
     {
@@ -327,14 +316,6 @@ void ProblemDescription::SetupFloats(ExecutionContext& ctx) const
                  << "x" << GetDataTypeName(GetOutDataType()));
 }
 
-void ProblemDescription::SetupComputeType(const ExecutionContext& ctx) const
-{
-    if(ctx.GetStream().GetDeviceName() == "gfx942" && conv.EnableTF32())
-    {
-        use_tf32 = true;
-    }
-}
-
 std::string ProblemDescription::ComputeLayout(const TensorDescriptor& td) const
 {
     return td.GetLayout_str();

projects/miopen/src/conv/solver_finders.cpp

@@ -302,8 +302,7 @@ std::vector<Solution> EvaluateInvokers(const Handle& handle,
 
             MIOPEN_THROW_IF(elapsed <= 0, "Invalid elapsed time detected in EvaluateInvokers");
 
-            MIOPEN_LOG_I("solution(current vs best):" << sol << ": " << elapsed
-                                                      << (elapsed < best ? " < " : " >= ") << best);
+            MIOPEN_LOG_I(sol << ": " << elapsed << (elapsed < best ? " < " : " >= ") << best);
             if(elapsed < best)
             {
                 best         = elapsed;

projects/miopen/src/convolution.cpp

@@ -468,18 +468,6 @@ std::size_t ConvolutionDescriptor::GetWorkSpaceSize(ExecutionContext ctx,
     return workspace_size;
 }
 
-bool ConvolutionDescriptor::EnableTF32() const
-{
-    /* true only when both EnvEnableTF32() and (MathType==Default) are true. */
-    // temporarily disable TF32 until tf32 feature are fully complete validated with database.
-    // TODO:(LYM) change back to &&
-    if((miopen::EnvEnableTF32() ||
-        (static_cast<miopenMathType_t>(attribute.Get(MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE)) ==
-         miopenMathDefault)))
-        return true;
-    return false;
-}
-
 std::ostream& operator<<(std::ostream& stream, const ConvolutionDescriptor& c)
 {
     stream << "conv" << c.spatialDim << "d, ";
@@ -552,18 +540,6 @@ void ConvolutionAttribute::Set(miopenConvolutionAttrib_t attr, int value)
         }
         fp8rounding_mode.rounding_mode = rounding_mode;
     }
-    else if(attr == MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE)
-    {
-        const auto math_type_ = static_cast<miopenMathType_t>(value);
-        if(math_type_ != miopenMathDefault && math_type_ != miopenMathPedantic)
-        {
-            MIOPEN_THROW(miopenStatusBadParm,
-                         "[Set conv attribute] Error: Attempt to set invalid value for "
-                         "MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE: " +
-                             std::to_string(value));
-        }
-        math_type.value = math_type_;
-    }
     else
     {
         MIOPEN_THROW(miopenStatusBadParm,
@@ -580,8 +556,6 @@ int ConvolutionAttribute::Get(miopenConvolutionAttrib_t attr) const
         return static_cast<int>(fp8rounding_mode.rounding_mode);
     else if(attr == MIOPEN_CONVOLUTION_ATTRIB_DETERMINISTIC)
         return deterministic.value;
-    else if(attr == MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE)
-        return math_type.value;
     MIOPEN_THROW(miopenStatusBadParm,
                  "[Get conv attribute] Error: Attribute [" +
                      std::to_string(static_cast<int>(attr)) + "] does not exist.");

projects/miopen/src/convolution_api.cpp

@@ -67,7 +67,6 @@ static inline auto MakeFwdCtxAndProblem(miopenHandle_t handle,
 
     auto ctx = ExecutionContext{&miopen::deref(handle)};
     problem.SetupFloats(ctx);
-    problem.SetupComputeType(ctx);
     return std::make_tuple(std::move(ctx), std::move(problem));
 }
 
@@ -87,7 +86,6 @@ static inline auto MakeBwdCtxAndProblem(miopenHandle_t handle,
 
     auto ctx = ExecutionContext{&miopen::deref(handle)};
     problem.SetupFloats(ctx);
-    problem.SetupComputeType(ctx);
     return std::make_tuple(std::move(ctx), std::move(problem));
 }
 
@@ -113,7 +111,6 @@ static inline auto MakeWrWCtxAndProblem(miopenHandle_t handle,
 
     auto ctx = ExecutionContext{&miopen::deref(handle)};
     problem.SetupFloats(ctx);
-    problem.SetupComputeType(ctx);
     return std::make_tuple(std::move(ctx), std::move(problem));
 }
 

projects/miopen/src/include/miopen/conv/problem_description.hpp

@@ -296,8 +296,6 @@ struct MIOPEN_INTERNALS_EXPORT ProblemDescription : ProblemDescriptionBase
         return GetInCastType() || GetWeightsCastType() || GetOutCastType();
     }
 
-    bool UseTF32() const { return use_tf32; }
-
     // To be used in Solvers that do not implement ALT FP16 kernels.
     // Those Solvers must be non-applicable for gfx90a when this function returns true.
     bool IsGfx90aFp16altRequired() const
@@ -398,8 +396,6 @@ struct MIOPEN_INTERNALS_EXPORT ProblemDescription : ProblemDescriptionBase
             self.GetInDataType(), self.GetWeightsDataType(), self.GetOutDataType());
         f(data_type, "data_type");
         f(self.GetDirectionStr(), "direction");
-        if(data_type == "FP32" && self.UseTF32())
-            f("TF32", "compute_datatype");
     }
 
     template <class Self, class Visitor>
@@ -411,7 +407,6 @@ struct MIOPEN_INTERNALS_EXPORT ProblemDescription : ProblemDescriptionBase
     }
 
     void SetupFloats(ExecutionContext& ctx) const;
-    void SetupComputeType(const ExecutionContext& ctx) const;
 
 private:
     std::string ComputeLayout(const TensorDescriptor& td) const;
@@ -431,7 +426,6 @@ struct MIOPEN_INTERNALS_EXPORT ProblemDescription : ProblemDescriptionBase
     Scalar alpha                          = Scalar(1.0);
     Scalar beta                           = Scalar(0.0);
     miopenAlphaBetaCase_t alpha_beta_case = DEFAULT;
-    mutable bool use_tf32                 = false;
 };
 
 } // namespace conv

projects/miopen/src/include/miopen/conv/solvers.hpp

@@ -4524,15 +4524,13 @@ struct PerformanceConfigHipImplicitGemm3DGroupFwdXdlops
     MIOPEN_INTERNALS_EXPORT bool IsValid(const miopen::conv::ProblemDescription&) const;
     MIOPEN_INTERNALS_EXPORT bool
     operator==(const PerformanceConfigHipImplicitGemm3DGroupFwdXdlops& other) const;
-    bool UseTF32() const { return use_tf32; }
 
 private:
-    template <typename DataType, typename ComputeType = DataType>
-    bool Init(const miopen::conv::ProblemDescription&);
-    template <typename DataType, typename ComputeType = DataType>
+    template <typename DataType>
+    void Init(const miopen::conv::ProblemDescription&);
+    template <typename DataType>
     bool CheckIsSupportCKArgs(const miopen::conv::ProblemDescription&) const;
     void InitValidKernels(const miopen::conv::ProblemDescription& problem);
-    mutable bool use_tf32 = false;
 };
 
 struct ConvHipImplicitGemm3DGroupFwdXdlops final
@@ -4569,7 +4567,7 @@ struct ConvHipImplicitGemm3DGroupFwdXdlops final
     bool MayNeedWorkspace() const override { return true; }
 
 private:
-    template <typename DataType, typename ComputeType = DataType>
+    template <typename DataType>
     bool CheckCKApplicability(const miopen::conv::ProblemDescription&) const;
 };
 

projects/miopen/src/include/miopen/convolution.hpp

@@ -53,15 +53,8 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC)
 MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP8_ROUNDING_MODE)
 MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP8_ROUNDING_SEED)
 
-// disable TF32 by default temporarily until we fully complete this feature.
-// TODO:(LYM) change back
-MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TF32_OVERRIDE, 0);
-MIOPEN_DECLARE_ENV_VAR_BOOL(NVIDIA_TF32_OVERRIDE, 0);
-
 namespace miopen {
 
-MIOPEN_INTERNALS_EXPORT bool EnvEnableTF32();
-
 namespace conv {
 struct ProblemDescription;
 } // namespace conv
@@ -139,17 +132,6 @@ struct MIOPEN_INTERNALS_EXPORT ConvolutionAttribute
         }
     } deterministic;
 
-    class MathType
-    {
-        // temporary set default to pedantic until we fully complete this feature.
-        // TODO:(LYM) change back
-        miopenMathType_t value = miopenMathPedantic;
-        friend struct ConvolutionAttribute;
-
-    public:
-        inline int Get() const { return value; }
-    } math_type;
-
     /// Tri-state attribute values:
     /// * -1: Default (attribute-specific).
     /// * 0: Disabled/Yes.
@@ -368,7 +350,6 @@ struct MIOPEN_INTERNALS_EXPORT ConvolutionDescriptor : miopenConvolutionDescript
                                     Data_t dw,
                                     Data_t workSpace,
                                     std::size_t workSpaceSize) const;
-    miopenMathType_t GetMathType() const;
 
     std::size_t spatialDim;
     miopenConvolutionMode_t mode;
@@ -392,8 +373,6 @@ struct MIOPEN_INTERNALS_EXPORT ConvolutionDescriptor : miopenConvolutionDescript
     std::size_t GetSolutionCountFallback(const ExecutionContext& ctx,
                                          const conv::ProblemDescription& problem) const;
 
-    bool EnableTF32() const;
-
     friend void to_json(nlohmann::json& json, const ConvolutionDescriptor& conv);
     friend void from_json(const nlohmann::json& json, ConvolutionDescriptor& conv);
 

projects/miopen/src/include/miopen/fusion/context.hpp

@@ -42,7 +42,6 @@ struct FusionContext : ExecutionContext
     {
         auto ctx = ExecutionContext{*this};
         conv_problem.SetupFloats(ctx);
-        conv_problem.SetupComputeType(ctx);
         return ctx;
     }
 };