diff --git a/projects/miopen/Jenkinsfile b/projects/miopen/Jenkinsfile index c25457a59835..d8df7570bcca 100644 --- a/projects/miopen/Jenkinsfile +++ b/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") { diff --git a/projects/miopen/driver/conv_driver.hpp b/projects/miopen/driver/conv_driver.hpp index cb9ad68022a8..23523c6eb986 100644 --- a/projects/miopen/driver/conv_driver.hpp +++ b/projects/miopen/driver/conv_driver.hpp @@ -456,12 +456,6 @@ class ConvDriver : public Driver constexpr bool is_bfp8 = std::is_same::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 && (miopen::EnvEnableTF32() || math_type_)) - tolerance = 8.2e-3; - } return tolerance; } @@ -874,8 +868,6 @@ int ConvDriver::GetandSetData() warmupConvDesc, static_cast(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 warmup_out_len(warmup_out_len_size); @@ -1030,8 +1022,6 @@ int ConvDriver::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::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; } diff --git a/projects/miopen/include/miopen/miopen.h b/projects/miopen/include/miopen/miopen.h index 85e1f709ede5..b7b9214905e0 100644 --- a/projects/miopen/include/miopen/miopen.h +++ b/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 diff --git a/projects/miopen/src/conv/problem_description.cpp b/projects/miopen/src/conv/problem_description.cpp index ce31e3f9da6a..0d70ae9c9bd7 100644 --- a/projects/miopen/src/conv/problem_description.cpp +++ b/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(); diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index 0dd79302d186..1b7a39e0e980 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -302,8 +302,7 @@ std::vector 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; diff --git a/projects/miopen/src/convolution.cpp b/projects/miopen/src/convolution.cpp index 4b99f79ab85f..5bd6c1756865 100644 --- a/projects/miopen/src/convolution.cpp +++ b/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(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(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(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(attr)) + "] does not exist."); diff --git a/projects/miopen/src/convolution_api.cpp b/projects/miopen/src/convolution_api.cpp index d84774bc2772..1bcbe53b162c 100644 --- a/projects/miopen/src/convolution_api.cpp +++ b/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)); } diff --git a/projects/miopen/src/include/miopen/conv/problem_description.hpp b/projects/miopen/src/include/miopen/conv/problem_description.hpp index 8d96c2c55c5f..d91606b36abf 100644 --- a/projects/miopen/src/include/miopen/conv/problem_description.hpp +++ b/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 @@ -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 diff --git a/projects/miopen/src/include/miopen/conv/solvers.hpp b/projects/miopen/src/include/miopen/conv/solvers.hpp index a2ebe7ec8986..af40816e7db2 100644 --- a/projects/miopen/src/include/miopen/conv/solvers.hpp +++ b/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 - bool Init(const miopen::conv::ProblemDescription&); - template + template + void Init(const miopen::conv::ProblemDescription&); + template 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 + template bool CheckCKApplicability(const miopen::conv::ProblemDescription&) const; }; diff --git a/projects/miopen/src/include/miopen/convolution.hpp b/projects/miopen/src/include/miopen/convolution.hpp index 05c4ee670ea2..e613de98f680 100644 --- a/projects/miopen/src/include/miopen/convolution.hpp +++ b/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); diff --git a/projects/miopen/src/include/miopen/fusion/context.hpp b/projects/miopen/src/include/miopen/fusion/context.hpp index 9cc91f519340..435097a68a9a 100644 --- a/projects/miopen/src/include/miopen/fusion/context.hpp +++ b/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; } }; diff --git a/projects/miopen/src/include/miopen/solver/implicitgemm_ck_util.hpp b/projects/miopen/src/include/miopen/solver/implicitgemm_ck_util.hpp index 00d9406bc926..88f81c93f998 100644 --- a/projects/miopen/src/include/miopen/solver/implicitgemm_ck_util.hpp +++ b/projects/miopen/src/include/miopen/solver/implicitgemm_ck_util.hpp @@ -1467,8 +1467,7 @@ template ConvSolution MakeSolutionGroupConvImplicitGemmXdlops(const miopen::conv::ProblemDescription& problem, InvokerFactoryMakerNCHW&& invoker_factory_maker_ncdhw, - InvokerFactoryMakerNHWC&& invoker_factory_maker_ndhwc, - const bool use_tf32 = false) + InvokerFactoryMakerNHWC&& invoker_factory_maker_ndhwc) { #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL @@ -1476,14 +1475,10 @@ MakeSolutionGroupConvImplicitGemmXdlops(const miopen::conv::ProblemDescription& { switch(problem.GetInDataType()) { - case miopenInt8: return invoker_factory_maker_ncdhw(int8_t{}, int8_t{}); - case miopenHalf: return invoker_factory_maker_ncdhw(ck::half_t{}, ck::half_t{}); - case miopenFloat: - if(use_tf32) - return invoker_factory_maker_ncdhw(float{}, ck::tf32_t{}); - else - return invoker_factory_maker_ncdhw(float{}, float{}); - case miopenBFloat16: return invoker_factory_maker_ncdhw(ck::bhalf_t{}, ck::bhalf_t{}); + case miopenInt8: return invoker_factory_maker_ncdhw(int8_t{}); + case miopenHalf: return invoker_factory_maker_ncdhw(ck::half_t{}); + case miopenFloat: return invoker_factory_maker_ncdhw(float{}); + case miopenBFloat16: return invoker_factory_maker_ncdhw(ck::bhalf_t{}); case miopenInt64: case miopenInt32: case miopenDouble: @@ -1499,14 +1494,10 @@ MakeSolutionGroupConvImplicitGemmXdlops(const miopen::conv::ProblemDescription& { switch(problem.GetInDataType()) { - case miopenInt8: return invoker_factory_maker_ndhwc(int8_t{}, int8_t{}); - case miopenHalf: return invoker_factory_maker_ndhwc(ck::half_t{}, ck::half_t{}); - case miopenFloat: - if(use_tf32) - return invoker_factory_maker_ndhwc(float{}, ck::tf32_t{}); - else - return invoker_factory_maker_ndhwc(float{}, float{}); - case miopenBFloat16: return invoker_factory_maker_ndhwc(ck::bhalf_t{}, ck::bhalf_t{}); + case miopenInt8: return invoker_factory_maker_ndhwc(int8_t{}); + case miopenHalf: return invoker_factory_maker_ndhwc(ck::half_t{}); + case miopenFloat: return invoker_factory_maker_ndhwc(float{}); + case miopenBFloat16: return invoker_factory_maker_ndhwc(ck::bhalf_t{}); case miopenInt64: case miopenInt32: case miopenDouble: diff --git a/projects/miopen/src/kernels/gpu_reference_kernel/naive_conv.cpp b/projects/miopen/src/kernels/gpu_reference_kernel/naive_conv.cpp index fc2d54389f7e..ff09bbff4a19 100644 --- a/projects/miopen/src/kernels/gpu_reference_kernel/naive_conv.cpp +++ b/projects/miopen/src/kernels/gpu_reference_kernel/naive_conv.cpp @@ -71,7 +71,7 @@ inline __device__ __host__ ushort convert_fp32_to_bf16(float src_val) return target_val.ushortvec[1]; } -template +template inline __device__ __host__ dst_data_t cast_to(const src_data_t& val) { return static_cast(val); @@ -102,25 +102,6 @@ inline __device__ __host__ int8_t cast_to(const int32_t& val) return static_cast(val & 0xff); } -template <> -inline __device__ __host__ float cast_to(const float& val) -{ - union - { - float fp32; - uint32_t int32; - } u = {val}; - - u.int32 = u.int32 & 0xffffe000; - return u.fp32; -} - -template <> -inline __device__ __host__ double cast_to(const float& val) -{ - return static_cast(cast_to(val)); -} - inline __device__ __host__ bool IsZero(double val) { return val == 0.0; } inline __device__ __host__ bool IsOne(double val) { return val == 1.0; } @@ -160,11 +141,7 @@ inline __device__ void applyalphaBetaUpdate(dst_data_t* __restrict__ p_array, /// alpha and beta are double to ensure high precision. -template +template inline __device__ void naive_conv_fwd_nchw(const src_data_t* __restrict__ p_in, const src_data_t* __restrict__ p_wei, const double alpha, @@ -260,8 +237,8 @@ inline __device__ void naive_conv_fwd_nchw(const src_data_t* __restrict__ p_in, size_t f_idx = static_cast(ic) * fy * fx + static_cast(iy) * fx + static_cast(ix); - value += cast_to(p_in[i_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_wei[f_idx]); } else { @@ -273,8 +250,8 @@ inline __device__ void naive_conv_fwd_nchw(const src_data_t* __restrict__ p_in, static_cast(iy) * wei_strides[1] + static_cast(ix) * wei_strides[0]; - value += cast_to(p_in[i_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_wei[f_idx]); } } } @@ -294,11 +271,7 @@ inline __device__ void naive_conv_fwd_nchw(const src_data_t* __restrict__ p_in, } } -template +template inline __device__ void naive_conv_bwd_nchw(dst_data_t* __restrict__ p_in, const src_data_t* __restrict__ p_wei, const double alpha, @@ -399,8 +372,8 @@ inline __device__ void naive_conv_bwd_nchw(dst_data_t* __restrict__ p_in, size_t f_idx = static_cast(ik) * c_per_group * fy * fx + static_cast(iy) * fx + static_cast(ix); - value += cast_to(p_out[o_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_out[o_idx]) * + cast_to(p_wei[f_idx]); } else { @@ -412,8 +385,8 @@ inline __device__ void naive_conv_bwd_nchw(dst_data_t* __restrict__ p_in, static_cast(iy) * wei_strides[1] + static_cast(ix) * wei_strides[0]; - value += cast_to(p_out[o_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_out[o_idx]) * + cast_to(p_wei[f_idx]); } } } @@ -434,11 +407,7 @@ inline __device__ void naive_conv_bwd_nchw(dst_data_t* __restrict__ p_in, } } -template +template inline __device__ void naive_conv_wrw_nchw(const src_data_t* __restrict__ p_in, dst_data_t* __restrict__ p_wei, const double alpha, @@ -534,8 +503,8 @@ inline __device__ void naive_conv_wrw_nchw(const src_data_t* __restrict__ p_in, size_t o_idx = static_cast(in) * k * ho * wo + static_cast(iho) * wo + static_cast(iwo); - value += cast_to(p_in[i_idx]) * - cast_to(p_out[o_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_out[o_idx]); } else { @@ -548,8 +517,8 @@ inline __device__ void naive_conv_wrw_nchw(const src_data_t* __restrict__ p_in, static_cast(iho) * out_strides[1] + static_cast(iwo) * out_strides[0]; - value += cast_to(p_in[i_idx]) * - cast_to(p_out[o_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_out[o_idx]); } } } @@ -573,12 +542,7 @@ inline __device__ void naive_conv_wrw_nchw(const src_data_t* __restrict__ p_in, } // design block_size 256 - -template +template inline __device__ void naive_conv_fwd_ncdhw(const src_data_t* __restrict__ p_in, const src_data_t* __restrict__ p_wei, const double alpha, @@ -692,8 +656,8 @@ inline __device__ void naive_conv_fwd_ncdhw(const src_data_t* __restrict__ p_in, static_cast(iy) * fx + static_cast(ix); - value += cast_to(p_in[i_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_wei[f_idx]); } else { @@ -707,8 +671,8 @@ inline __device__ void naive_conv_fwd_ncdhw(const src_data_t* __restrict__ p_in, static_cast(iy) * wei_strides[1] + static_cast(ix) * wei_strides[0]; - value += cast_to(p_in[i_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_wei[f_idx]); } } } @@ -732,11 +696,7 @@ inline __device__ void naive_conv_fwd_ncdhw(const src_data_t* __restrict__ p_in, } } -template +template inline __device__ void naive_conv_bwd_ncdhw(dst_data_t* __restrict__ p_in, const src_data_t* __restrict__ p_wei, const double alpha, @@ -859,8 +819,8 @@ inline __device__ void naive_conv_bwd_ncdhw(dst_data_t* __restrict__ p_in, static_cast(iz) * fy * fx + static_cast(iy) * fx + static_cast(ix); - value += cast_to(p_out[o_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_out[o_idx]) * + cast_to(p_wei[f_idx]); } else { @@ -874,8 +834,8 @@ inline __device__ void naive_conv_bwd_ncdhw(dst_data_t* __restrict__ p_in, static_cast(iy) * wei_strides[1] + static_cast(ix) * wei_strides[0]; - value += cast_to(p_out[o_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_out[o_idx]) * + cast_to(p_wei[f_idx]); } } } @@ -899,11 +859,7 @@ inline __device__ void naive_conv_bwd_ncdhw(dst_data_t* __restrict__ p_in, } } -template +template inline __device__ void naive_conv_wrw_ncdhw(const src_data_t* __restrict__ p_in, dst_data_t* __restrict__ p_wei, const double alpha, @@ -1015,8 +971,8 @@ inline __device__ void naive_conv_wrw_ncdhw(const src_data_t* __restrict__ p_in, static_cast(iho) * wo + static_cast(iwo); - value += cast_to(p_in[i_idx]) * - cast_to(p_out[o_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_out[o_idx]); } else { @@ -1031,8 +987,8 @@ inline __device__ void naive_conv_wrw_ncdhw(const src_data_t* __restrict__ p_in, static_cast(iho) * out_strides[1] + static_cast(iwo) * out_strides[0]; - value += cast_to(p_in[i_idx]) * - cast_to(p_out[o_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_out[o_idx]); } } } @@ -1060,12 +1016,7 @@ inline __device__ void naive_conv_wrw_ncdhw(const src_data_t* __restrict__ p_in, /***************************** nhwc *****************************/ // design block_size 256 - -template +template inline __device__ void naive_conv_fwd_nhwc(const src_data_t* __restrict__ p_in, const src_data_t* __restrict__ p_wei, const double alpha, @@ -1163,8 +1114,8 @@ inline __device__ void naive_conv_fwd_nhwc(const src_data_t* __restrict__ p_in, static_cast(iy) * fx * c_per_group + static_cast(ix) * c_per_group + static_cast(ic); - value += cast_to(p_in[i_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_wei[f_idx]); } else { @@ -1179,8 +1130,8 @@ inline __device__ void naive_conv_fwd_nhwc(const src_data_t* __restrict__ p_in, static_cast(ix) * wei_strides[1] + static_cast(ic) * wei_strides[0]; - value += cast_to(p_in[i_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_wei[f_idx]); } } } @@ -1203,11 +1154,7 @@ inline __device__ void naive_conv_fwd_nhwc(const src_data_t* __restrict__ p_in, } } -template +template inline __device__ void naive_conv_bwd_nhwc(dst_data_t* __restrict__ p_in, const src_data_t* __restrict__ p_wei, const double alpha, @@ -1309,8 +1256,8 @@ inline __device__ void naive_conv_bwd_nhwc(dst_data_t* __restrict__ p_in, static_cast(iy) * fx * c_per_group + static_cast(ix) * c_per_group + static_cast(ic); - value += cast_to(p_out[o_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_out[o_idx]) * + cast_to(p_wei[f_idx]); } else { @@ -1325,8 +1272,8 @@ inline __device__ void naive_conv_bwd_nhwc(dst_data_t* __restrict__ p_in, static_cast(ix) * wei_strides[1] + static_cast(ic) * wei_strides[0]; - value += cast_to(p_out[o_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_out[o_idx]) * + cast_to(p_wei[f_idx]); } } } @@ -1349,11 +1296,7 @@ inline __device__ void naive_conv_bwd_nhwc(dst_data_t* __restrict__ p_in, } } -template +template inline __device__ void naive_conv_wrw_nhwc(const src_data_t* __restrict__ p_in, dst_data_t* __restrict__ p_wei, const double alpha, @@ -1449,8 +1392,8 @@ inline __device__ void naive_conv_wrw_nhwc(const src_data_t* __restrict__ p_in, static_cast(iho) * wo * k + static_cast(iwo) * k; - value += cast_to(p_in[i_idx]) * - cast_to(p_out[o_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_out[o_idx]); } else { @@ -1463,8 +1406,8 @@ inline __device__ void naive_conv_wrw_nhwc(const src_data_t* __restrict__ p_in, static_cast(iho) * out_strides[3] + static_cast(iwo) * out_strides[2]; - value += cast_to(p_in[i_idx]) * - cast_to(p_out[o_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_out[o_idx]); } } } @@ -1488,12 +1431,7 @@ inline __device__ void naive_conv_wrw_nhwc(const src_data_t* __restrict__ p_in, } // design block_size 256 - -template +template inline __device__ void naive_conv_fwd_ndhwc(const src_data_t* __restrict__ p_in, const src_data_t* __restrict__ p_wei, const double alpha, @@ -1606,8 +1544,8 @@ inline __device__ void naive_conv_fwd_ndhwc(const src_data_t* __restrict__ p_in, static_cast(iy) * fx * c_per_group + static_cast(ix) * c_per_group + static_cast(ic); - value += cast_to(p_in[i_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_wei[f_idx]); } else { @@ -1622,8 +1560,8 @@ inline __device__ void naive_conv_fwd_ndhwc(const src_data_t* __restrict__ p_in, static_cast(ix) * wei_strides[1] + static_cast(ic) * wei_strides[0]; - value += cast_to(p_in[i_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_wei[f_idx]); } } } @@ -1647,11 +1585,7 @@ inline __device__ void naive_conv_fwd_ndhwc(const src_data_t* __restrict__ p_in, } } -template +template inline __device__ void naive_conv_bwd_ndhwc(dst_data_t* __restrict__ p_in, const src_data_t* __restrict__ p_wei, const double alpha, @@ -1772,8 +1706,8 @@ inline __device__ void naive_conv_bwd_ndhwc(dst_data_t* __restrict__ p_in, static_cast(iy) * fx * c_per_group + static_cast(ix) * c_per_group + static_cast(ic); - value += cast_to(p_out[o_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_out[o_idx]) * + cast_to(p_wei[f_idx]); } else { @@ -1788,8 +1722,8 @@ inline __device__ void naive_conv_bwd_ndhwc(dst_data_t* __restrict__ p_in, static_cast(ix) * wei_strides[1] + static_cast(ic) * wei_strides[0]; - value += cast_to(p_out[o_idx]) * - cast_to(p_wei[f_idx]); + value += cast_to(p_out[o_idx]) * + cast_to(p_wei[f_idx]); } } } @@ -1813,11 +1747,7 @@ inline __device__ void naive_conv_bwd_ndhwc(dst_data_t* __restrict__ p_in, } } -template +template inline __device__ void naive_conv_wrw_ndhwc(const src_data_t* __restrict__ p_in, dst_data_t* __restrict__ p_wei, const double alpha, @@ -1929,8 +1859,8 @@ inline __device__ void naive_conv_wrw_ndhwc(const src_data_t* __restrict__ p_in, static_cast(iho) * wo * k + static_cast(iwo) * k; - value += cast_to(p_in[i_idx]) * - cast_to(p_out[o_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_out[o_idx]); } else { @@ -1946,8 +1876,8 @@ inline __device__ void naive_conv_wrw_ndhwc(const src_data_t* __restrict__ p_in, static_cast(iho) * out_strides[3] + static_cast(iwo) * out_strides[2]; - value += cast_to(p_in[i_idx]) * - cast_to(p_out[o_idx]); + value += cast_to(p_in[i_idx]) * + cast_to(p_out[o_idx]); } } } @@ -1973,323 +1903,297 @@ inline __device__ void naive_conv_wrw_ndhwc(const src_data_t* __restrict__ p_in, } } -#define DEFINE_2D_NAIVE_CONV_KERNEL( \ - direction, tensor_layout, src_data_t, acc_data_t, dst_data_t, use_tf32) \ - extern "C" __global__ void \ - naive_conv_ab_packed_##direction##_##tensor_layout##_##src_data_t##_##acc_data_t##_##dst_data_t##_##use_tf32( \ - src_data_t* __restrict__ p_in, \ - src_data_t* __restrict__ p_wei, \ - double alpha, \ - double beta, \ - dst_data_t* __restrict__ p_out, \ - Strides5D in_strides, \ - Strides5D wei_strides, \ - Strides5D out_strides, \ - int hi, \ - int wi, \ - int n, \ - int k_per_group, \ - int c_per_group, \ - int ho, \ - int wo, \ - int sy, \ - int sx, \ - int dy, \ - int dx, \ - int py, \ - int px, \ - int fy, \ - int fx, \ - int group) \ - { \ - naive_conv_##direction##_##tensor_layout(p_in, \ - p_wei, \ - alpha, \ - beta, \ - p_out, \ - in_strides, \ - wei_strides, \ - out_strides, \ - hi, \ - wi, \ - n, \ - k_per_group, \ - c_per_group, \ - ho, \ - wo, \ - sy, \ - sx, \ - dy, \ - dx, \ - py, \ - px, \ - fy, \ - fx, \ - group); \ - } \ - extern "C" __global__ void \ - naive_conv_ab_nonpacked_##direction##_##tensor_layout##_##src_data_t##_##acc_data_t##_##dst_data_t##_##use_tf32( \ - src_data_t* __restrict__ p_in, \ - src_data_t* __restrict__ p_wei, \ - double alpha, \ - double beta, \ - dst_data_t* __restrict__ p_out, \ - Strides5D in_strides, \ - Strides5D wei_strides, \ - Strides5D out_strides, \ - int hi, \ - int wi, \ - int n, \ - int k_per_group, \ - int c_per_group, \ - int ho, \ - int wo, \ - int sy, \ - int sx, \ - int dy, \ - int dx, \ - int py, \ - int px, \ - int fy, \ - int fx, \ - int group) \ - { \ - naive_conv_##direction##_##tensor_layout(p_in, \ - p_wei, \ - alpha, \ - beta, \ - p_out, \ - in_strides, \ - wei_strides, \ - out_strides, \ - hi, \ - wi, \ - n, \ - k_per_group, \ - c_per_group, \ - ho, \ - wo, \ - sy, \ - sx, \ - dy, \ - dx, \ - py, \ - px, \ - fy, \ - fx, \ - group); \ +#define DEFINE_2D_NAIVE_CONV_KERNEL(direction, tensor_layout, src_data_t, acc_data_t, dst_data_t) \ + extern "C" __global__ void \ + naive_conv_ab_packed_##direction##_##tensor_layout##_##src_data_t##_##acc_data_t##_##dst_data_t( \ + src_data_t* __restrict__ p_in, \ + src_data_t* __restrict__ p_wei, \ + double alpha, \ + double beta, \ + dst_data_t* __restrict__ p_out, \ + Strides5D in_strides, \ + Strides5D wei_strides, \ + Strides5D out_strides, \ + int hi, \ + int wi, \ + int n, \ + int k_per_group, \ + int c_per_group, \ + int ho, \ + int wo, \ + int sy, \ + int sx, \ + int dy, \ + int dx, \ + int py, \ + int px, \ + int fy, \ + int fx, \ + int group) \ + { \ + naive_conv_##direction##_##tensor_layout( \ + p_in, \ + p_wei, \ + alpha, \ + beta, \ + p_out, \ + in_strides, \ + wei_strides, \ + out_strides, \ + hi, \ + wi, \ + n, \ + k_per_group, \ + c_per_group, \ + ho, \ + wo, \ + sy, \ + sx, \ + dy, \ + dx, \ + py, \ + px, \ + fy, \ + fx, \ + group); \ + } \ + extern "C" __global__ void \ + naive_conv_ab_nonpacked_##direction##_##tensor_layout##_##src_data_t##_##acc_data_t##_##dst_data_t( \ + src_data_t* __restrict__ p_in, \ + src_data_t* __restrict__ p_wei, \ + double alpha, \ + double beta, \ + dst_data_t* __restrict__ p_out, \ + Strides5D in_strides, \ + Strides5D wei_strides, \ + Strides5D out_strides, \ + int hi, \ + int wi, \ + int n, \ + int k_per_group, \ + int c_per_group, \ + int ho, \ + int wo, \ + int sy, \ + int sx, \ + int dy, \ + int dx, \ + int py, \ + int px, \ + int fy, \ + int fx, \ + int group) \ + { \ + naive_conv_##direction##_##tensor_layout( \ + p_in, \ + p_wei, \ + alpha, \ + beta, \ + p_out, \ + in_strides, \ + wei_strides, \ + out_strides, \ + hi, \ + wi, \ + n, \ + k_per_group, \ + c_per_group, \ + ho, \ + wo, \ + sy, \ + sx, \ + dy, \ + dx, \ + py, \ + px, \ + fy, \ + fx, \ + group); \ } -#define DEFINE_3D_NAIVE_CONV_KERNEL( \ - direction, tensor_layout, src_data_t, acc_data_t, dst_data_t, use_tf32) \ - extern "C" __global__ void \ - naive_conv_ab_packed_##direction##_##tensor_layout##_##src_data_t##_##acc_data_t##_##dst_data_t##_##use_tf32( \ - src_data_t* __restrict__ p_in, \ - src_data_t* __restrict__ p_wei, \ - double alpha, \ - double beta, \ - dst_data_t* __restrict__ p_out, \ - Strides6D in_strides, \ - Strides6D wei_strides, \ - Strides6D out_strides, \ - int di, \ - int hi, \ - int wi, \ - int n, \ - int k_per_group, \ - int c_per_group, \ - int do_, \ - int ho, \ - int wo, \ - int sz, \ - int sy, \ - int sx, \ - int dz, \ - int dy, \ - int dx, \ - int pz, \ - int py, \ - int px, \ - int fz, \ - int fy, \ - int fx, \ - int group) \ - { \ - naive_conv_##direction##_##tensor_layout(p_in, \ - p_wei, \ - alpha, \ - beta, \ - p_out, \ - in_strides, \ - wei_strides, \ - out_strides, \ - di, \ - hi, \ - wi, \ - n, \ - k_per_group, \ - c_per_group, \ - do_, \ - ho, \ - wo, \ - sz, \ - sy, \ - sx, \ - dz, \ - dy, \ - dx, \ - pz, \ - py, \ - px, \ - fz, \ - fy, \ - fx, \ - group); \ - } \ - extern "C" __global__ void \ - naive_conv_ab_nonpacked_##direction##_##tensor_layout##_##src_data_t##_##acc_data_t##_##dst_data_t##_##use_tf32( \ - src_data_t* __restrict__ p_in, \ - src_data_t* __restrict__ p_wei, \ - double alpha, \ - double beta, \ - dst_data_t* __restrict__ p_out, \ - Strides6D in_strides, \ - Strides6D wei_strides, \ - Strides6D out_strides, \ - int di, \ - int hi, \ - int wi, \ - int n, \ - int k_per_group, \ - int c_per_group, \ - int do_, \ - int ho, \ - int wo, \ - int sz, \ - int sy, \ - int sx, \ - int dz, \ - int dy, \ - int dx, \ - int pz, \ - int py, \ - int px, \ - int fz, \ - int fy, \ - int fx, \ - int group) \ - { \ - naive_conv_##direction##_##tensor_layout(p_in, \ - p_wei, \ - alpha, \ - beta, \ - p_out, \ - in_strides, \ - wei_strides, \ - out_strides, \ - di, \ - hi, \ - wi, \ - n, \ - k_per_group, \ - c_per_group, \ - do_, \ - ho, \ - wo, \ - sz, \ - sy, \ - sx, \ - dz, \ - dy, \ - dx, \ - pz, \ - py, \ - px, \ - fz, \ - fy, \ - fx, \ - group); \ +#define DEFINE_3D_NAIVE_CONV_KERNEL(direction, tensor_layout, src_data_t, acc_data_t, dst_data_t) \ + extern "C" __global__ void \ + naive_conv_ab_packed_##direction##_##tensor_layout##_##src_data_t##_##acc_data_t##_##dst_data_t( \ + src_data_t* __restrict__ p_in, \ + src_data_t* __restrict__ p_wei, \ + double alpha, \ + double beta, \ + dst_data_t* __restrict__ p_out, \ + Strides6D in_strides, \ + Strides6D wei_strides, \ + Strides6D out_strides, \ + int di, \ + int hi, \ + int wi, \ + int n, \ + int k_per_group, \ + int c_per_group, \ + int do_, \ + int ho, \ + int wo, \ + int sz, \ + int sy, \ + int sx, \ + int dz, \ + int dy, \ + int dx, \ + int pz, \ + int py, \ + int px, \ + int fz, \ + int fy, \ + int fx, \ + int group) \ + { \ + naive_conv_##direction##_##tensor_layout( \ + p_in, \ + p_wei, \ + alpha, \ + beta, \ + p_out, \ + in_strides, \ + wei_strides, \ + out_strides, \ + di, \ + hi, \ + wi, \ + n, \ + k_per_group, \ + c_per_group, \ + do_, \ + ho, \ + wo, \ + sz, \ + sy, \ + sx, \ + dz, \ + dy, \ + dx, \ + pz, \ + py, \ + px, \ + fz, \ + fy, \ + fx, \ + group); \ + } \ + extern "C" __global__ void \ + naive_conv_ab_nonpacked_##direction##_##tensor_layout##_##src_data_t##_##acc_data_t##_##dst_data_t( \ + src_data_t* __restrict__ p_in, \ + src_data_t* __restrict__ p_wei, \ + double alpha, \ + double beta, \ + dst_data_t* __restrict__ p_out, \ + Strides6D in_strides, \ + Strides6D wei_strides, \ + Strides6D out_strides, \ + int di, \ + int hi, \ + int wi, \ + int n, \ + int k_per_group, \ + int c_per_group, \ + int do_, \ + int ho, \ + int wo, \ + int sz, \ + int sy, \ + int sx, \ + int dz, \ + int dy, \ + int dx, \ + int pz, \ + int py, \ + int px, \ + int fz, \ + int fy, \ + int fx, \ + int group) \ + { \ + naive_conv_##direction##_##tensor_layout( \ + p_in, \ + p_wei, \ + alpha, \ + beta, \ + p_out, \ + in_strides, \ + wei_strides, \ + out_strides, \ + di, \ + hi, \ + wi, \ + n, \ + k_per_group, \ + c_per_group, \ + do_, \ + ho, \ + wo, \ + sz, \ + sy, \ + sx, \ + dz, \ + dy, \ + dx, \ + pz, \ + py, \ + px, \ + fz, \ + fy, \ + fx, \ + group); \ } -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, float, double, float, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, float, double, float, 1) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, half, double, half, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, ushort, double, ushort, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, int8_t, int32_t, int8_t, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, int8_t, int32_t, int32_t, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, int8_t, int32_t, float, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, float, double, float, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, float, double, float, 1) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, half, double, half, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, ushort, double, ushort, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, int8_t, int32_t, int8_t, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, int8_t, int32_t, int32_t, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, int8_t, int32_t, float, 0) - -DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nchw, float, double, float, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nchw, float, double, float, 1) -DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nchw, half, double, half, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nchw, ushort, double, ushort, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nhwc, float, double, float, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nhwc, float, double, float, 1) -DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nhwc, half, double, half, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nhwc, ushort, double, ushort, 0) - -DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nchw, float, double, float, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nchw, float, double, float, 1) -DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nchw, half, double, half, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nchw, ushort, double, ushort, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nhwc, float, double, float, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nhwc, float, double, float, 1) -DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nhwc, half, double, half, 0) -DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nhwc, ushort, double, ushort, 0) - -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, float, double, float, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, float, double, float, 1) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, half, double, half, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, ushort, double, ushort, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, int8_t, int32_t, int32_t, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, int8_t, int32_t, float, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, float, double, float, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, float, double, float, 1) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, half, double, half, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, ushort, double, ushort, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, int8_t, int32_t, int32_t, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, int8_t, int32_t, float, 0) - -DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ncdhw, float, double, float, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ncdhw, float, double, float, 1) -DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ncdhw, half, double, half, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ncdhw, ushort, double, ushort, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ndhwc, float, double, float, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ndhwc, float, double, float, 1) -DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ndhwc, half, double, half, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ndhwc, ushort, double, ushort, 0) - -DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ncdhw, float, double, float, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ncdhw, float, double, float, 1) -DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ncdhw, half, double, half, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ncdhw, ushort, double, ushort, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ndhwc, float, double, float, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ndhwc, float, double, float, 1) -DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ndhwc, half, double, half, 0) -DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ndhwc, ushort, double, ushort, 0) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, float, double, float) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, half, double, half) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, ushort, double, ushort) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, int8_t, int32_t, int8_t) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, int8_t, int32_t, int32_t) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nchw, int8_t, int32_t, float) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, float, double, float) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, half, double, half) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, ushort, double, ushort) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, int8_t, int32_t, int8_t) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, int8_t, int32_t, int32_t) +DEFINE_2D_NAIVE_CONV_KERNEL(fwd, nhwc, int8_t, int32_t, float) + +DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nchw, float, double, float) +DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nchw, half, double, half) +DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nchw, ushort, double, ushort) +DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nhwc, float, double, float) +DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nhwc, half, double, half) +DEFINE_2D_NAIVE_CONV_KERNEL(bwd, nhwc, ushort, double, ushort) + +DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nchw, float, double, float) +DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nchw, half, double, half) +DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nchw, ushort, double, ushort) +DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nhwc, float, double, float) +DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nhwc, half, double, half) +DEFINE_2D_NAIVE_CONV_KERNEL(wrw, nhwc, ushort, double, ushort) + +DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, float, double, float) +DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, half, double, half) +DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, ushort, double, ushort) +DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, int8_t, int32_t, int32_t) +DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ncdhw, int8_t, int32_t, float) +DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, float, double, float) +DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, half, double, half) +DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, ushort, double, ushort) +DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, int8_t, int32_t, int32_t) +DEFINE_3D_NAIVE_CONV_KERNEL(fwd, ndhwc, int8_t, int32_t, float) + +DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ncdhw, float, double, float) +DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ncdhw, half, double, half) +DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ncdhw, ushort, double, ushort) +DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ndhwc, float, double, float) +DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ndhwc, half, double, half) +DEFINE_3D_NAIVE_CONV_KERNEL(bwd, ndhwc, ushort, double, ushort) + +DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ncdhw, float, double, float) +DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ncdhw, half, double, half) +DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ncdhw, ushort, double, ushort) +DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ndhwc, float, double, float) +DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ndhwc, half, double, half) +DEFINE_3D_NAIVE_CONV_KERNEL(wrw, ndhwc, ushort, double, ushort) /// \todo discuss whether we should split the kernels into separate files, or /// figure out a mechanism to compile each kernel separately to reduce hipRTC diff --git a/projects/miopen/src/ocl/convolutionocl.cpp b/projects/miopen/src/ocl/convolutionocl.cpp index aa208486a7cb..6b69c9f08cb4 100644 --- a/projects/miopen/src/ocl/convolutionocl.cpp +++ b/projects/miopen/src/ocl/convolutionocl.cpp @@ -198,7 +198,6 @@ static Invoker PrepareInvoker(ExecutionContext ctx, solver::Id solver_id) { problem.SetupFloats(ctx); - problem.SetupComputeType(ctx); ctx.do_search = false; ctx.disable_search_enforce = true; @@ -594,7 +593,6 @@ void ConvolutionDescriptor::FindConvFwdAlgorithm(const Handle& handle, const auto ctx = [&] { auto tmp = ExecutionContext{&handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); tmp.do_search = exhaustiveSearch; return tmp; }(); @@ -812,8 +810,6 @@ void ConvolutionDescriptor::ConvolutionForward(const Handle& handle, Scalar beta_val(beta, GetScalarDataType(yDesc)); const auto problem = conv::ProblemDescription{ xDesc, wDesc, yDesc, *this, conv::Direction::Forward, 0, alpha_val, beta_val}; - auto ctx = ExecutionContext{&handle}; - problem.SetupComputeType(ctx); ValidateAlphaBeta(problem); ConvForwardCheckNumerics(handle, tensors, [&]() { @@ -1136,7 +1132,6 @@ void ConvolutionDescriptor::FindConvBwdDataAlgorithm(const Handle& handle, const auto ctx = [&] { auto tmp = ExecutionContext{&handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); tmp.do_search = exhaustiveSearch; return tmp; }(); @@ -1344,7 +1339,6 @@ void ConvolutionDescriptor::FindConvBwdWeightsAlgorithm(const Handle& handle, const auto ctx = [&] { auto tmp = ExecutionContext{&handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); tmp.do_search = exhaustiveSearch; return tmp; }(); @@ -1510,11 +1504,6 @@ void ConvolutionDescriptor::ConvolutionWrwImmediate(const Handle& handle, }); } -miopenMathType_t ConvolutionDescriptor::GetMathType() const -{ - return static_cast(this->attribute.Get(MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE)); -} - void ConvolutionBackwardBias(const Handle& handle, const void* alpha, const TensorDescriptor& dyDesc, @@ -1610,18 +1599,4 @@ void ConvolutionBackwardBias(const Handle& handle, } } -bool EnvEnableTF32() -{ - // disable TF32 by default temporarily until we fully complete this feature. - // so either one is set to true, we enable TF32 - // TODO:(LYM) change back - bool bool_miopen = miopen::env::enabled(MIOPEN_TF32_OVERRIDE); - bool bool_nvidia = miopen::env::enabled(NVIDIA_TF32_OVERRIDE); - if(bool_miopen != bool_nvidia) - MIOPEN_LOG_I2("TF32_OVERRIDE is set to different values for MIOPEN_TF32_OVERRIDE (" - << bool_miopen << ") and NVIDIA_TF32_OVERRIDE (" << bool_nvidia - << "). TF32 will be treated as enabled."); // TODO:(LYM) back to disabled - return bool_miopen || bool_nvidia; // TODO:(LYM) back to disabled -} - } // namespace miopen diff --git a/projects/miopen/src/problem.cpp b/projects/miopen/src/problem.cpp index afb4daff9bb1..d3d4a198224d 100644 --- a/projects/miopen/src/problem.cpp +++ b/projects/miopen/src/problem.cpp @@ -511,7 +511,6 @@ std::vector Problem::FindSolutionsImpl(const Handle& handle, auto ctx = ExecutionContext{&handle}; conv_problem.SetupFloats(ctx); - conv_problem.SetupComputeType(ctx); ctx.do_search = options.exhaustive_search; const auto invoke_ctx = diff --git a/projects/miopen/src/solution.cpp b/projects/miopen/src/solution.cpp index 2b9105252dac..dc575643efbf 100644 --- a/projects/miopen/src/solution.cpp +++ b/projects/miopen/src/solution.cpp @@ -222,7 +222,6 @@ void Solution::RunImpl(const Handle& handle, { auto ctx = ExecutionContext{&handle}; conv_problem.SetupFloats(ctx); - conv_problem.SetupComputeType(ctx); const auto invoker_factory = GetSolver().GetSolver().GetInvokeFactory(ctx, conv_problem, perf_cfg.value_or("")); auto kernel_handles = std::vector{std::begin(kernels), std::end(kernels)}; @@ -246,7 +245,6 @@ void Solution::RunImpl(const Handle& handle, auto conv_ctx = ExecutionContext{&handle}; conv_problem.SetupFloats(conv_ctx); - conv_problem.SetupComputeType(conv_ctx); decltype(auto) db = MakeConvDbGetter(conv_ctx); const auto conv_solution = GetSolver().GetSolver().FindSolution( diff --git a/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp b/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp index 1779075a52a5..9b934ec041a5 100644 --- a/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp +++ b/projects/miopen/src/solver/conv/conv_MP_bidirectional_winograd.cpp @@ -745,7 +745,6 @@ ExecutionContext ConvMPBidirectWinograd_xdlops(use_tf32); - return kernel_name.str(); } diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp index 6ba73b42f456..931906ecf9d4 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp @@ -647,7 +647,7 @@ ConvSolution ConvHipImplicitGemm3DGroupBwdXdlops::GetSolution( #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL return MakeSolutionGroupConvImplicitGemmXdlops( problem, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { using T = decltype(data_type_val); switch(problem.GetAlphaBetaCase()) { @@ -674,7 +674,7 @@ ConvSolution ConvHipImplicitGemm3DGroupBwdXdlops::GetSolution( ctx, problem, config.kernel_id); } }, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { using T = decltype(data_type_val); switch(problem.GetAlphaBetaCase()) { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp index 57a1d0ca9d53..3cf19866e94a 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp @@ -64,7 +64,7 @@ using Bilinear = ck::tensor_operation::element_wise: using Scale = ck::tensor_operation::element_wise::Scale; static constexpr ck::index_t NumDimSpatial = 3; -template +template using DeviceOpGFwdBilinear = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD; + Bilinear>; -template +template using DeviceOpGFwdScale = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD; -template + Scale>; + +template using DeviceOpGFwdDefault = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD; + PassThrough>; -template +template using DeviceOpGFwdBilinearPtrs = ck::tensor_operation::device::instance::DeviceOperationInstanceFactory< - DeviceOpGFwdBilinear>; + DeviceOpGFwdBilinear>; -template +template using DeviceOpGFwdScalePtrs = ck::tensor_operation::device::instance::DeviceOperationInstanceFactory< - DeviceOpGFwdScale>; + DeviceOpGFwdScale>; -template +template using DeviceOpGFwdDefaultPtrs = ck::tensor_operation::device::instance::DeviceOperationInstanceFactory< - DeviceOpGFwdDefault>; + DeviceOpGFwdDefault>; namespace { -template +template struct CKArgs { CKArgs(const ::miopen::conv::ProblemDescription& problem) @@ -224,7 +219,7 @@ struct CKArgs { (void)alpha; (void)beta; - static_assert(std::is_same_v>, + static_assert(std::is_same_v>, "Default should be fwd pass through"); return MakeDefaultArgPtr(conv_ptr, in, w, out); } @@ -352,7 +347,7 @@ struct CKArgs miopenAlphaBetaCase_t alpha_beta_case; }; -template +template std::vector FillValidKernelsByAlphaBeta(const ::miopen::conv::ProblemDescription& problem) { @@ -365,25 +360,22 @@ FillValidKernelsByAlphaBeta(const ::miopen::conv::ProblemDescription& problem) return miopen::solver::FillValidKernelsIDs, CKArgs>(problem); default: - return miopen::solver::FillValidKernelsIDs, - CKArgs>(problem); + return miopen::solver::FillValidKernelsIDs, + CKArgs>(problem); } } } // namespace -template -bool PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::Init( - const miopen::conv::ProblemDescription& problem) +template +void PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::Init( + const ::miopen::conv::ProblemDescription& problem) { - valid_kernels = FillValidKernelsByAlphaBeta(problem); - if(valid_kernels.empty()) - return false; - index = 0; - kernel_id = valid_kernels[index]; - return true; + valid_kernels = FillValidKernelsByAlphaBeta(problem); + index = 0; + kernel_id = valid_kernels[index]; } -template +template bool PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::CheckIsSupportCKArgs( const ::miopen::conv::ProblemDescription& problem) const { @@ -396,12 +388,12 @@ bool PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::CheckIsSupportCKArgs( return IsCKArgsSupported, CKArgs>(problem, kernel_id); default: - return IsCKArgsSupported, - CKArgs>(problem, kernel_id); + return IsCKArgsSupported, CKArgs>(problem, + kernel_id); } } -template +template bool ConvHipImplicitGemm3DGroupFwdXdlops::CheckCKApplicability( const ::miopen::conv::ProblemDescription& problem) const { @@ -410,9 +402,7 @@ bool ConvHipImplicitGemm3DGroupFwdXdlops::CheckCKApplicability( case BILINEAR: return IsCKApplicable, CKArgs>(problem); case SCALE: return IsCKApplicable, CKArgs>(problem); - default: - return IsCKApplicable, - CKArgs>(problem); + default: return IsCKApplicable, CKArgs>(problem); } } @@ -422,17 +412,7 @@ void PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::InitValidKernels( switch(problem.GetInDataType()) { case miopenHalf: Init(problem); break; - case miopenFloat: - if(problem.UseTF32() && Init(problem)) - { - use_tf32 = true; - } - else - { - use_tf32 = false; - Init(problem); - } - break; + case miopenFloat: Init(problem); break; case miopenInt8: Init(problem); break; case miopenBFloat16: Init(problem); break; default: break; // Unsupported data types - valid_kernels remains empty @@ -586,12 +566,11 @@ void PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::HeuristicInit( bool ai_success = false; miopen::ai::tuning::candidate_selection::CandidateSelectionResult result; - auto run_ai_heuristics = [&](auto CKDataType, auto CKComputeType) { - using T = decltype(CKDataType); - using TCompute = decltype(CKComputeType); + auto run_ai_heuristics = [&](auto CKDataType) { + using T = decltype(CKDataType); auto fill_valid_kernels = [=](const ::miopen::conv::ProblemDescription& problem) -> std::vector { - return FillValidKernelsByAlphaBeta(problem); + return FillValidKernelsByAlphaBeta(problem); }; // Validation lambda for AI-predicted kernel + split_k combinations // Note: This solver currently doesn't use split_k (always 0), but validation @@ -620,27 +599,9 @@ void PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::HeuristicInit( }; switch(problem.GetInDataType()) { - case miopenHalf: - std::tie(ai_success, result) = run_ai_heuristics(ck::half_t{}, ck::half_t{}); - break; - case miopenFloat: - if(problem.UseTF32()) - { - std::tie(ai_success, result) = run_ai_heuristics(float{}, ck::tf32_t{}); - if(!ai_success || result.IsEmpty()) - { - MIOPEN_LOG_I2("Step 3: AI heuristics with TF32 failed, retrying with FP32"); - std::tie(ai_success, result) = run_ai_heuristics(float{}, float{}); - } - } - else - { - std::tie(ai_success, result) = run_ai_heuristics(float{}, float{}); - } - break; - case miopenBFloat16: - std::tie(ai_success, result) = run_ai_heuristics(ck::bhalf_t{}, ck::bhalf_t{}); - break; + case miopenHalf: std::tie(ai_success, result) = run_ai_heuristics(ck::half_t{}); break; + case miopenFloat: std::tie(ai_success, result) = run_ai_heuristics(float{}); break; + case miopenBFloat16: std::tie(ai_success, result) = run_ai_heuristics(ck::bhalf_t{}); break; default: break; } if(ai_success && !result.IsEmpty()) @@ -711,17 +672,7 @@ bool PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::IsValid( switch(problem.GetInDataType()) { case miopenHalf: return CheckIsSupportCKArgs(problem); - case miopenFloat: - if(problem.UseTF32() && CheckIsSupportCKArgs(problem)) - { - use_tf32 = true; - return true; - } - else - { - use_tf32 = false; - return CheckIsSupportCKArgs(problem); - } + case miopenFloat: return CheckIsSupportCKArgs(problem); case miopenInt8: return CheckIsSupportCKArgs(problem); case miopenBFloat16: return CheckIsSupportCKArgs(problem); case miopenInt64: @@ -798,15 +749,7 @@ bool ConvHipImplicitGemm3DGroupFwdXdlops::IsApplicable( switch(problem.GetInDataType()) { case miopenHalf: return CheckCKApplicability(problem); - case miopenFloat: - if(problem.UseTF32() && CheckCKApplicability(problem)) - { - return true; - } - else - { - return CheckCKApplicability(problem); - } + case miopenFloat: return CheckCKApplicability(problem); case miopenInt8: return CheckCKApplicability(problem); case miopenBFloat16: return CheckCKApplicability(problem); case miopenInt64: @@ -855,9 +798,8 @@ ConvSolution ConvHipImplicitGemm3DGroupFwdXdlops::GetSolution( #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL return MakeSolutionGroupConvImplicitGemmXdlops( problem, - [&](auto data_type_val, auto compute_type_val) { - using T = decltype(data_type_val); - using TCompute = decltype(compute_type_val); + [&](auto data_type_val) { + using T = decltype(data_type_val); switch(problem.GetAlphaBetaCase()) { case BILINEAR: @@ -877,15 +819,14 @@ ConvSolution ConvHipImplicitGemm3DGroupFwdXdlops::GetSolution( default: return InitInvokerFactoryFwdNCHW<3, false, - DeviceOpGFwdDefaultPtrs, - CKArgs, + DeviceOpGFwdDefaultPtrs, + CKArgs, miopen::conv::DataInvokeParams>( ctx, problem, config.kernel_id); } }, - [&](auto data_type_val, auto compute_type_val) { - using T = decltype(data_type_val); - using TCompute = decltype(compute_type_val); + [&](auto data_type_val) { + using T = decltype(data_type_val); switch(problem.GetAlphaBetaCase()) { case BILINEAR: @@ -902,13 +843,12 @@ ConvSolution ConvHipImplicitGemm3DGroupFwdXdlops::GetSolution( ctx, problem, config.kernel_id); default: return InitInvokerFactoryNHWC, - CKArgs, + DeviceOpGFwdDefaultPtrs, + CKArgs, miopen::conv::DataInvokeParams>( ctx, problem, config.kernel_id); } - }, - config.UseTF32()); + }); #else return {}; diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp index 04c07521f6ce..b7db5a901ec5 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp @@ -691,7 +691,7 @@ ConvSolution ConvHipImplicitGemm3DGroupWrwXdlops::GetSolution( #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL return MakeSolutionGroupConvImplicitGemmXdlops( problem, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { using T = decltype(data_type_val); switch(problem.GetAlphaBetaCase()) { @@ -718,7 +718,7 @@ ConvSolution ConvHipImplicitGemm3DGroupWrwXdlops::GetSolution( ctx, problem, config.kernel_id); } }, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { using T = decltype(data_type_val); switch(problem.GetAlphaBetaCase()) { diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp index b6ec1fd90a78..b60d062219eb 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp @@ -620,7 +620,7 @@ ConvSolution ConvHipImplicitGemmGroupBwdXdlops::GetSolution( #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL return MakeSolutionGroupConvImplicitGemmXdlops( problem, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { using T = decltype(data_type_val); return InitInvokerFactoryBwdNCHW<2, false, @@ -629,7 +629,7 @@ ConvSolution ConvHipImplicitGemmGroupBwdXdlops::GetSolution( miopen::conv::DataInvokeParams>( ctx, problem, config.kernel_id); }, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { using T = decltype(data_type_val); return InitInvokerFactoryNHWC, diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp index dd921e298ecb..51f49e1a3a97 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp @@ -555,7 +555,7 @@ ConvSolution ConvHipImplicitGemmGroupFwdXdlops::GetSolution( #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL return MakeSolutionGroupConvImplicitGemmXdlops( problem, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { using T = decltype(data_type_val); return InitInvokerFactoryFwdNCHW<2, false, @@ -564,7 +564,7 @@ ConvSolution ConvHipImplicitGemmGroupFwdXdlops::GetSolution( miopen::conv::DataInvokeParams>( ctx, problem, config.kernel_id); }, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { using T = decltype(data_type_val); return InitInvokerFactoryNHWC, diff --git a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp index ff8f8d3d8f11..60f364569933 100644 --- a/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp +++ b/projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp @@ -657,7 +657,7 @@ ConvSolution ConvHipImplicitGemmGroupWrwXdlops::GetSolution( #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL return MakeSolutionGroupConvImplicitGemmXdlops( problem, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { using T = decltype(data_type_val); return InitInvokerFactoryWrwNCHW<2, false, @@ -666,7 +666,7 @@ ConvSolution ConvHipImplicitGemmGroupWrwXdlops::GetSolution( miopen::conv::WrWInvokeParams>( ctx, problem, config.kernel_id); }, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { using T = decltype(data_type_val); return InitInvokerFactoryNHWC, diff --git a/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp b/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp index 3be3c0cb79e1..03c4f35874e4 100644 --- a/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp +++ b/projects/miopen/src/solver/conv_ck_igemm_grp_fwd_activ_fused.cpp @@ -661,7 +661,7 @@ GetSolutionForDimensionality(const FusionContext& ctx, using Layouts = LayoutsSelector; return MakeSolutionGroupConvImplicitGemmXdlops( conv_problem, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { (void)data_type_val; return InitInvokerFactoryFwdNCHW( ctx, conv_problem, config.kernel_id); }, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { (void)data_type_val; return InitInvokerFactoryNHWC; return MakeSolutionGroupConvImplicitGemmXdlops( conv_problem, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { (void)data_type_val; return InitInvokerFactoryFwdNCHW( ctx, conv_problem, config.kernel_id); }, - [&](auto data_type_val, [[maybe_unused]] auto compute_type_val) { + [&](auto data_type_val) { (void)data_type_val; return InitInvokerFactoryNHWC= 2) { key = opt[0]; @@ -294,12 +293,7 @@ void ParseProblemKey(const std::string& key_, conv::ProblemDescription& prob_des conv::ProblemDescription tmp{in, wei, out, conv, dir}; } conv.group_count = group_cnt; - if(precision == miopenFloat) - { - const auto math_type_ = use_tf32 ? miopenMathDefault : miopenMathPedantic; - conv.attribute.Set(MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE, static_cast(math_type_)); - } - prob_desc = conv::ProblemDescription{in, wei, out, conv, dir}; + prob_desc = conv::ProblemDescription{in, wei, out, conv, dir}; } struct FDBVal @@ -633,7 +627,6 @@ void CheckDynamicFDBEntry(size_t thread_index, miopen::conv::ProblemDescription problem; miopen::ParseProblemKey(kinder.first, problem); problem.SetupFloats(ctx); // TODO: Check if this is necessary - problem.SetupComputeType(ctx); std::stringstream ss; problem.Serialize(ss); ASSERT_TRUE(ss.str() == kinder.first) @@ -739,7 +732,6 @@ void CheckFDBEntry(size_t thread_index, miopen::conv::ProblemDescription problem; miopen::ParseProblemKey(kinder.first, problem); problem.SetupFloats(ctx); // TODO: Check if this is necessary - problem.SetupComputeType(ctx); std::stringstream ss; problem.Serialize(ss); // moment of truth diff --git a/projects/miopen/test/gtest/group_conv.hpp b/projects/miopen/test/gtest/group_conv.hpp index d9ab9e080898..27251a14f051 100644 --- a/projects/miopen/test/gtest/group_conv.hpp +++ b/projects/miopen/test/gtest/group_conv.hpp @@ -317,16 +317,7 @@ struct GroupConvTestFix double threshold = 80; if(CONV_DIR == Direction::Forward) { - if constexpr(std::is_same_v) - { - // float use tf32 compute which share same mantissa bits - threshold *= (compute_type == "TF32") ? std::numeric_limits::epsilon() - : std::numeric_limits::epsilon(); - } - else - { - threshold *= std::numeric_limits::epsilon(); - } + threshold *= std::numeric_limits::epsilon(); } else { @@ -362,9 +353,6 @@ struct GroupConvTestFix ctx.SetStream(&handle); - if(compute_type == "TF32") - problem.SetupComputeType(ctx); - if(!solv.IsApplicable(ctx, problem)) { test_skipped = true; @@ -475,11 +463,6 @@ struct GroupConvTestFix test_skipped = true; GTEST_SKIP() << "bf16 tests skipped on this hardware."; } - if(!IsTestSupportedByDevice(Gpu::gfx94X) && compute_type == "TF32") - { - test_skipped = true; - GTEST_SKIP() << "tf32 tests skipped on this hardware."; - } } float alpha_val; @@ -494,10 +477,6 @@ struct GroupConvTestFix weights = tensor{tensor_layout, conv_config.GetWeights()}; conv_desc = conv_config.GetConv(); - if(compute_type == "TF32") - conv_desc.attribute.Set(MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE, miopenMathDefault); - else - conv_desc.attribute.Set(MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE, miopenMathPedantic); miopen::TensorDescriptor output_desc = conv_desc.GetForwardOutputTensor(input.desc, weights.desc, miopen_type{}); @@ -582,8 +561,6 @@ struct GroupConvTestFix miopen::Scalar alpha{1.0}; miopen::Scalar beta{0.0}; - - std::string compute_type; }; template @@ -636,7 +613,6 @@ std::vector GetBetaValues() struct GPU_GroupConv##ndim##D_##dir##_##naming_type \ : GroupConvTestFix \ { \ - GPU_GroupConv##ndim##D_##dir##_##naming_type() { compute_type = #naming_type; } \ }; \ TEST_P(GPU_GroupConv##ndim##D_##dir##_##naming_type, GroupConv##ndim##D_##dir##_##type##_Test) \ { \ diff --git a/projects/miopen/test/gtest/group_conv2d_fwd.cpp b/projects/miopen/test/gtest/group_conv2d_fwd.cpp index 188daa298a32..bf1332f97042 100644 --- a/projects/miopen/test/gtest/group_conv2d_fwd.cpp +++ b/projects/miopen/test/gtest/group_conv2d_fwd.cpp @@ -30,7 +30,6 @@ using namespace group_conv; DEFINE_GROUP_CONV2D_TEST(float, FP32, Forward); -DEFINE_GROUP_CONV2D_TEST(float, TF32, Forward); DEFINE_GROUP_CONV2D_TEST(half, FP16, Forward); DEFINE_GROUP_CONV2D_TEST(bfloat16, BFP16, Forward); DEFINE_GROUP_CONV2D_TEST(int8_t, I8, Forward); diff --git a/projects/miopen/test/gtest/group_conv3d_fwd.cpp b/projects/miopen/test/gtest/group_conv3d_fwd.cpp index df3d61732ba7..7d661ffe123b 100644 --- a/projects/miopen/test/gtest/group_conv3d_fwd.cpp +++ b/projects/miopen/test/gtest/group_conv3d_fwd.cpp @@ -30,7 +30,6 @@ using namespace group_conv; DEFINE_GROUP_CONV3D_TEST(float, FP32, Forward); -DEFINE_GROUP_CONV3D_TEST(float, TF32, Forward); DEFINE_GROUP_CONV3D_TEST(half, FP16, Forward); DEFINE_GROUP_CONV3D_TEST(bfloat16, BFP16, Forward); /// \todo int8_t tests don't work. Need debugging diff --git a/projects/miopen/test/gtest/solver_bwd.hpp b/projects/miopen/test/gtest/solver_bwd.hpp index b27f7565c4e6..e9affc9f6aec 100644 --- a/projects/miopen/test/gtest/solver_bwd.hpp +++ b/projects/miopen/test/gtest/solver_bwd.hpp @@ -56,7 +56,6 @@ struct ConvBwdSolverTest const miopen::ExecutionContext ctx = [&] { auto tmp = miopen::ExecutionContext{&handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); return tmp; }(); diff --git a/projects/miopen/test/gtest/solver_fwd.hpp b/projects/miopen/test/gtest/solver_fwd.hpp index 17bfcd0a6aaf..6d74ab151c52 100644 --- a/projects/miopen/test/gtest/solver_fwd.hpp +++ b/projects/miopen/test/gtest/solver_fwd.hpp @@ -58,7 +58,6 @@ struct ConvFwdSolverTest const miopen::ExecutionContext ctx = [&] { auto tmp = miopen::ExecutionContext{&handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); return tmp; }(); diff --git a/projects/miopen/test/gtest/solver_wrw.hpp b/projects/miopen/test/gtest/solver_wrw.hpp index 4ef34f0faf8c..1f2285bc051c 100644 --- a/projects/miopen/test/gtest/solver_wrw.hpp +++ b/projects/miopen/test/gtest/solver_wrw.hpp @@ -56,7 +56,6 @@ struct ConvWrwSolverTest const miopen::ExecutionContext ctx = [&] { auto tmp = miopen::ExecutionContext{&handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); return tmp; }(); diff --git a/projects/miopen/test/gtest/unit_FinInterface.cpp b/projects/miopen/test/gtest/unit_FinInterface.cpp index 2e4744ce714a..3ddb91bc1c86 100644 --- a/projects/miopen/test/gtest/unit_FinInterface.cpp +++ b/projects/miopen/test/gtest/unit_FinInterface.cpp @@ -417,7 +417,6 @@ auto GetContext(miopen::Handle* handle, const miopen::conv::ProblemDescription& { auto tmp = miopen::ExecutionContext{handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); return tmp; } diff --git a/projects/miopen/test/gtest/unit_conv_ConvolutionDescriptor.hpp b/projects/miopen/test/gtest/unit_conv_ConvolutionDescriptor.hpp index 16f5796e655d..737d4193dfca 100644 --- a/projects/miopen/test/gtest/unit_conv_ConvolutionDescriptor.hpp +++ b/projects/miopen/test/gtest/unit_conv_ConvolutionDescriptor.hpp @@ -37,14 +37,12 @@ struct ConvolutionDescriptorParams std::vector&& strides_in, std::vector&& dilations_in, int group_count_in = 1, - bool deterministic_in = false, - bool tf32_compute_in = false) + bool deterministic_in = false) : pads(std::move(pads_in)), strides(std::move(strides_in)), dilations(std::move(dilations_in)), group_count(group_count_in), - deterministic(deterministic_in), - tf32_compute(tf32_compute_in) + deterministic(deterministic_in) { } @@ -62,10 +60,6 @@ struct ConvolutionDescriptorParams { desc.attribute.Set(MIOPEN_CONVOLUTION_ATTRIB_DETERMINISTIC, 1); } - // SET TF32 COMPUTE ATTRIBUTE - miopenMathType_t math_type = tf32_compute ? miopenMathDefault : miopenMathPedantic; - desc.attribute.Set(MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE, math_type); - return desc; } @@ -84,7 +78,6 @@ struct ConvolutionDescriptorParams std::vector dilations; int group_count; bool deterministic; - bool tf32_compute; }; } // namespace unit_tests diff --git a/projects/miopen/test/gtest/unit_conv_solver.cpp b/projects/miopen/test/gtest/unit_conv_solver.cpp index 5a30790b90d9..7fc295566cef 100644 --- a/projects/miopen/test/gtest/unit_conv_solver.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver.cpp @@ -336,18 +336,10 @@ miopen::solver::ConvSolution FindSolution(const miopen::solver::conv::ConvSolver template double GetThreshold(miopenConvAlgorithm_t algo, miopen::conv::Direction direction, - const Tolerances& tolerances, - const bool use_tf32_compute) + const Tolerances& tolerances) { double tolerance = tolerances.Get(GetDevGpuType(), miopen_type{}); double threshold = std::numeric_limits::epsilon() * tolerance; - if constexpr(std::is_same_v) - { - if(use_tf32_compute) - { - threshold = std::numeric_limits::epsilon() * tolerance; - } - } return threshold; } @@ -356,8 +348,7 @@ void VerifyData(const std::vector& data, const std::vector& ref_data, miopenConvAlgorithm_t algo, miopen::conv::Direction direction, - const Tolerances& tolerances, - bool use_tf32_compute = false) + const Tolerances& tolerances) { ASSERT_FALSE(miopen::range_zero(ref_data)) << "Reference data is all zeros"; if constexpr(!std::is_integral_v) @@ -384,7 +375,7 @@ void VerifyData(const std::vector& data, else { const auto error = miopen::rms_range(ref_data, data); - const double threshold = GetThreshold(algo, direction, tolerances, use_tf32_compute); + const double threshold = GetThreshold(algo, direction, tolerances); ASSERT_LT(error, threshold) << "Error beyond tolerance"; // std::cout << "error: " << error << " threshold: " << threshold << std::endl; } @@ -440,17 +431,9 @@ void RunSolverFwd(const miopen::solver::conv::ConvSolverInterface& solv, const auto ctx = [&] { auto tmp = miopen::ExecutionContext{&handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); return tmp; }(); - if(!(ctx.GetStream().GetDeviceName() == "gfx942") && - conv_config.GetXDataType() == miopenFloat && - conv_config.GetConv().GetMathType() == miopenMathDefault) - { - GTEST_SKIP() << "TF32 test is not supported on this device"; - } - if(!solv.IsApplicable(ctx, problem)) { // Do not put GTEST_SKIP here. @@ -498,12 +481,8 @@ void RunSolverFwd(const miopen::solver::conv::ConvSolverInterface& solv, output.data = handle.Read(out_dev, output.data.size()); - VerifyData(output.data, - ref_out.data, - algo, - miopen::conv::Direction::Forward, - params.tolerances, - problem.UseTF32()); + VerifyData( + output.data, ref_out.data, algo, miopen::conv::Direction::Forward, params.tolerances); } template @@ -565,7 +544,6 @@ void RunSolverBwd(const miopen::solver::conv::ConvSolverInterface& solv, const auto ctx = [&] { auto tmp = miopen::ExecutionContext{&handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); return tmp; }(); @@ -679,7 +657,6 @@ void RunSolverWrw(const miopen::solver::conv::ConvSolverInterface& solv, const auto ctx = [&] { auto tmp = miopen::ExecutionContext{&handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); return tmp; }(); @@ -864,7 +841,6 @@ void UnitTestConvSolverDevApplicabilityBase::RunTestImpl( const auto ctx = [&] { auto tmp = miopen::ExecutionContext{&handle}; problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); return tmp; }(); diff --git a/projects/miopen/test/gtest/unit_conv_solver.hpp b/projects/miopen/test/gtest/unit_conv_solver.hpp index aa50b6fb995a..dc95c6632818 100644 --- a/projects/miopen/test/gtest/unit_conv_solver.hpp +++ b/projects/miopen/test/gtest/unit_conv_solver.hpp @@ -227,8 +227,6 @@ using GPU_UnitTestConvSolverFwd_FP32 = miopen::unit_tests::UnitTestConvSolverFwd using GPU_UnitTestConvSolverBwd_FP32 = miopen::unit_tests::UnitTestConvSolverBwd; using GPU_UnitTestConvSolverWrw_FP32 = miopen::unit_tests::UnitTestConvSolverWrw; -using GPU_UnitTestConvSolverFwd_TF32 = miopen::unit_tests::UnitTestConvSolverFwd; - using GPU_UnitTestConvSolverFwd_I8 = miopen::unit_tests::UnitTestConvSolverFwd; using GPU_UnitTestConvSolverBwd_I8 = miopen::unit_tests::UnitTestConvSolverBwd; using GPU_UnitTestConvSolverWrw_I8 = miopen::unit_tests::UnitTestConvSolverWrw; diff --git a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupFwdXdlops.cpp b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupFwdXdlops.cpp deleted file mode 100644 index 76155951ede4..000000000000 --- a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupFwdXdlops.cpp +++ /dev/null @@ -1,218 +0,0 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2025 Advanced Micro Devices, Inc. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - * - *******************************************************************************/ - -#include "unit_conv_solver.hpp" - -namespace { - -auto GetConvSmokeTestCases(miopenDataType_t datatype) -{ - using TestCase = miopen::unit_tests::ConvTestCase; - - return std::vector{ - // clang-format off - TestCase{{datatype, miopenTensorNDHWC, {1, 64, 8, 8, 8}}, - {datatype, miopenTensorNDHWC, {96, 64, 1, 1, 1}}, - datatype, {{0, 0, 0}, {1, 1, 1}, {1, 1, 1}}}, - // clang-format on - }; -} -// for tf32 -auto GetConvSmokeTestCases(const std::string& compute_type_str) -{ - using TestCase = miopen::unit_tests::ConvTestCase; - if(compute_type_str != "TF32") - { - throw std::runtime_error("Unsupported compute_type in smoke tests"); - } - - return std::vector{ - // clang-format off - TestCase{{miopenFloat, miopenTensorNDHWC, {1, 64, 8, 8, 8}}, - {miopenFloat, miopenTensorNDHWC, {96, 64, 1, 1, 1}}, - miopenFloat, {{0, 0, 0}, {1, 1, 1}, {1, 1, 1}, 1, false, true}}, - // clang-format on - }; -} - -auto GetConvFullTestCases(miopenDataType_t datatype) -{ - using TestCase = miopen::unit_tests::ConvTestCase; - - return std::vector{ - // clang-format off - TestCase{{datatype, miopenTensorNDHWC, {1, 64, 8, 8, 8}}, - {datatype, miopenTensorNDHWC, {96, 64, 1, 1, 1}}, - datatype, {{1, 1, 1}, {1, 1, 1}, {1, 1, 1}}}, // non-zero padding - TestCase{{datatype, miopenTensorNDHWC, {1, 64, 8, 8, 8}}, - {datatype, miopenTensorNDHWC, {96, 64, 1, 1, 1}}, - datatype, {{0, 0, 0}, {2, 2, 2}, {1, 1, 1}}}, // stride > 1 - TestCase{{datatype, miopenTensorNDHWC, {1, 64, 8, 8, 8}}, - {datatype, miopenTensorNDHWC, {96, 64, 1, 1, 1}}, - datatype, {{0, 0, 0}, {1, 1, 1}, {2, 2, 2}}}, // dilation > 1 - TestCase{{datatype, miopenTensorNDHWC, {1, 64, 12, 24, 48}}, - {datatype, miopenTensorNDHWC, {384, 64, 1, 1, 1}}, - datatype, {{0, 0, 0}, {1, 1, 1}, {1, 1, 1}}}, // some different NCHW and k parameters - // clang-format on - }; -} -// for tf32 -auto GetConvFullTestCases(const std::string& compute_type_str) -{ - using TestCase = miopen::unit_tests::ConvTestCase; - // using datatype = miopenDataType_t::miopenFloat; - if(compute_type_str != "TF32") - { - throw std::runtime_error("Unsupported compute_type in smoke tests"); - } - - return std::vector{ - // clang-format off - TestCase{{miopenFloat, miopenTensorNDHWC, {1, 64, 8, 8, 8}}, - {miopenFloat, miopenTensorNDHWC, {96, 64, 1, 1, 1}}, - miopenFloat, {{1, 1, 1}, {1, 1, 1}, {1, 1, 1}, 1, false, true}}, // non-zero padding - TestCase{{miopenFloat, miopenTensorNDHWC, {1, 64, 8, 8, 8}}, - {miopenFloat, miopenTensorNDHWC, {96, 64, 1, 1, 1}}, - miopenFloat, {{0, 0, 0}, {2, 2, 2}, {1, 1, 1}, 1, false, true}}, // stride > 1 - TestCase{{miopenFloat, miopenTensorNDHWC, {1, 64, 8, 8, 8}}, - {miopenFloat, miopenTensorNDHWC, {96, 64, 1, 1, 1}}, - miopenFloat, {{0, 0, 0}, {1, 1, 1}, {2, 2, 2}, 1, false, true}}, // dilation > 1 - TestCase{{miopenFloat, miopenTensorNDHWC, {1, 64, 12, 24, 48}}, - {miopenFloat, miopenTensorNDHWC, {384, 64, 1, 1, 1}}, - miopenFloat, {{0, 0, 0}, {1, 1, 1}, {1, 1, 1}, 1, false, true}}, // some different NCHW and k parameters - // clang-format on - }; -} - -const auto& GetTestParams() -{ - static const auto params = [] { -// If MIOpen is built without CK these tests will fail, skip them to avoid failing -#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL - Gpu supportedDevices = Gpu::gfx908 | Gpu::gfx90A | Gpu::gfx94X | Gpu::gfx950; -#else - Gpu supportedDevices = Gpu::None; -#endif - auto p = miopen::unit_tests::UnitTestConvSolverParams(supportedDevices); - p.Tunable(5); - return p; - }(); - return params; -} - -} // namespace - -using GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_FP16 = GPU_UnitTestConvSolverFwd_FP16; -using GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_BFP16 = GPU_UnitTestConvSolverFwd_BFP16; -using GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_FP32 = GPU_UnitTestConvSolverFwd_FP32; -using GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_TF32 = GPU_UnitTestConvSolverFwd_TF32; -using CPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlopsDevApplicability_NONE = - CPU_UnitTestConvSolverDevApplicabilityFwd_NONE; - -TEST_P(GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_FP16, - ConvHipImplicitGemm3DGroupFwdXdlops) -{ - this->RunTest(miopen::solver::conv::ConvHipImplicitGemm3DGroupFwdXdlops{}); -}; - -TEST_P(GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_BFP16, - ConvHipImplicitGemm3DGroupFwdXdlops) -{ - this->RunTest(miopen::solver::conv::ConvHipImplicitGemm3DGroupFwdXdlops{}); -}; - -TEST_P(GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_FP32, - ConvHipImplicitGemm3DGroupFwdXdlops) -{ - this->RunTest(miopen::solver::conv::ConvHipImplicitGemm3DGroupFwdXdlops{}); -}; -TEST_P(GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_TF32, - ConvHipImplicitGemm3DGroupFwdXdlops) -{ - this->RunTest(miopen::solver::conv::ConvHipImplicitGemm3DGroupFwdXdlops{}); -}; -TEST_P(CPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlopsDevApplicability_NONE, - ConvHipImplicitGemm3DGroupFwdXdlops) -{ - this->RunTest(miopen::solver::conv::ConvHipImplicitGemm3DGroupFwdXdlops{}); -}; - -// Smoke tests -INSTANTIATE_TEST_SUITE_P(Smoke, - GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_FP16, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoImplicitGEMM), - testing::ValuesIn(GetConvSmokeTestCases(miopenHalf)))); - -INSTANTIATE_TEST_SUITE_P( - Smoke, - GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_BFP16, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoImplicitGEMM), - testing::ValuesIn(GetConvSmokeTestCases(miopenBFloat16)))); - -INSTANTIATE_TEST_SUITE_P(Smoke, - GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_FP32, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoImplicitGEMM), - testing::ValuesIn(GetConvSmokeTestCases(miopenFloat)))); - -INSTANTIATE_TEST_SUITE_P(Smoke, - GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_TF32, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoImplicitGEMM), - testing::ValuesIn(GetConvSmokeTestCases("TF32")))); - -// Full tests -INSTANTIATE_TEST_SUITE_P(Full, - GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_FP16, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoImplicitGEMM), - testing::ValuesIn(GetConvFullTestCases(miopenHalf)))); - -INSTANTIATE_TEST_SUITE_P(Full, - GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_BFP16, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoImplicitGEMM), - testing::ValuesIn(GetConvFullTestCases(miopenBFloat16)))); - -INSTANTIATE_TEST_SUITE_P(Full, - GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_FP32, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoImplicitGEMM), - testing::ValuesIn(GetConvFullTestCases(miopenFloat)))); - -INSTANTIATE_TEST_SUITE_P(Full, - GPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlops_TF32, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoImplicitGEMM), - testing::ValuesIn(GetConvFullTestCases("TF32")))); - -// Device applicability test -INSTANTIATE_TEST_SUITE_P(Smoke, - CPU_UnitTestConvSolverHipImplicitGemm3DGroupFwdXdlopsDevApplicability_NONE, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(GetConvSmokeTestCases(miopenHalf)[0]))); diff --git a/projects/miopen/test/utils/gtest_formating_checks.py b/projects/miopen/test/utils/gtest_formating_checks.py index c87140c9b3c8..2a3892ac8910 100644 --- a/projects/miopen/test/utils/gtest_formating_checks.py +++ b/projects/miopen/test/utils/gtest_formating_checks.py @@ -46,7 +46,7 @@ # Valid enums and Regex for validation VALID_HW_TYPES = {"CPU", "GPU"} -VALID_DATATYPES = {"FP8", "FP16", "FP32", "TF32", "FP64", "BFP16", "BFP8", "I64", "I32", "I16", "I8", "NONE"} +VALID_DATATYPES = {"FP8", "FP16", "FP32", "FP64", "BFP16", "BFP8", "I64", "I32", "I16", "I8", "NONE"} # Our suite (or fixture) naming convention: must start with CPU or GPU, followed by one or more alphanum groups, and end with a valid datatype. TESTSUITE_REGEX = re.compile( r"^(CPU|GPU)_[A-Za-z0-9]+(?:_[A-Za-z0-9]+)*_(" + "|".join(VALID_DATATYPES) + r")$"