diff --git a/Jenkinsfile b/Jenkinsfile index d60f9262c2..0c1e55efdb 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -789,6 +789,26 @@ pipeline { buildHipClangJobAndReboot(setup_flags: Bf16_flags + Full_test, build_install: "true", gpu_arch: "gfx90a:xnack-") } } + stage('Fp32 OpenCL All gfx908') { + when { + beforeAgent true + expression { params.TARGET_GFX908 && params.DATATYPE_FP32 } + } + agent{ label rocmnode("vega") } + steps{ + buildHipClangJobAndReboot(compiler: 'g++', setup_flags: Full_test, gpu_arch: "gfx908") + } + } + stage('Fp32 OpenCL Install All gfx90a') { + when { + beforeAgent true + expression { params.TARGET_GFX90A && params.DATATYPE_FP32 } + } + agent{ label rocmnode("vega") } + steps{ + buildHipClangJobAndReboot(compiler: 'g++', setup_flags: Full_test, build_install: "true", gpu_arch: "gfx90a:xnack-") + } + } } } diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 8669863f22..fa2ea010be 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -109,6 +109,7 @@ set( MIOpen_Source solver/gemm.cpp solver/gemm_bwd.cpp solver/gemm_wrw.cpp + solver/gemm_common.cpp dropout.cpp dropout_api.cpp readonlyramdb.cpp @@ -192,6 +193,7 @@ set( MIOpen_Source include/miopen/batched_transpose_sol.hpp include/miopen/magic_div.hpp include/miopen/util_sol.hpp + include/miopen/solver/gemm_common.hpp md_graph.cpp mdg_expr.cpp conv/invokers/gcn_asm_1x1u.cpp diff --git a/src/include/miopen/conv/asm_implicit_gemm.hpp b/src/include/miopen/conv/asm_implicit_gemm.hpp index 1d5ca670aa..fe506e1299 100644 --- a/src/include/miopen/conv/asm_implicit_gemm.hpp +++ b/src/include/miopen/conv/asm_implicit_gemm.hpp @@ -25,12 +25,17 @@ *******************************************************************************/ #ifndef CK_ASM_IMPLICITGEMM_HPP_ #define CK_ASM_IMPLICITGEMM_HPP_ + +#include + #include #include #include #include #include +#define WORKAROUND_ISSUE_1317 (MIOPEN_BACKEND_OPENCL) + namespace miopen { namespace solver { diff --git a/src/include/miopen/rocm_features.hpp b/src/include/miopen/rocm_features.hpp index dc4917eba9..cedcbf3eaf 100644 --- a/src/include/miopen/rocm_features.hpp +++ b/src/include/miopen/rocm_features.hpp @@ -55,12 +55,9 @@ /// --input 1, 4, 4, 161, 700 --weights 1, 4, 3, 11, 11 /// --pads_strides_dilations 3 3 3 2 2 2 4 4 4 --trans_output_pads 0 0 0 /// -/// W/A is in effect only when MIOpenGEMM is used (OCL BE) abd includes: -/// - Disabling GEMM for failing configs. -/// - Adding Naive Solvers. Naive solvers are inteded for use as backup for -/// Immediate Mode Fallback when GEMM is disabled. -/// - Note: When MIOpenGEMM is not in use, Naive Solvers are disabled. This minimizes -/// impact of the W/A to the HIP backend. +/// W/A is in effect only when MIOpenGEMM is used (OCL BE) and disables +/// GEMM for the failing configs. When this happens, Naive solvers +/// are used as backup on the Immediate Mode Fallback path. #define WORKAROUND_MIOPENGEMM_SINCE_ROCM41 \ (MIOPEN_USE_MIOPENGEMM && (HIP_PACKAGE_VERSION_FLAT >= 4001000000ULL)) diff --git a/src/include/miopen/solver/gemm_common.hpp b/src/include/miopen/solver/gemm_common.hpp new file mode 100644 index 0000000000..50a22538d2 --- /dev/null +++ b/src/include/miopen/solver/gemm_common.hpp @@ -0,0 +1,44 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2021 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. + * + *******************************************************************************/ + +#ifndef GUARD_SOLVER_GEMM_COMMON_HPP_ +#define GUARD_SOLVER_GEMM_COMMON_HPP_ + +#include + +namespace miopen { +namespace conv { +namespace solver { +namespace gemm { + +bool IsWorkaroundIssue1315(const miopen::ExecutionContext& ctx); + +} // namespace gemm +} // namespace solver +} // namespace conv +} // namespace miopen + +#endif diff --git a/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp index 0dfc549a95..3d316db3fd 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp @@ -867,6 +867,11 @@ ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::Search(const ConvolutionContext& ctx bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::IsApplicable(const ConvolutionContext& ctx) const { +#if WORKAROUND_ISSUE_1317 + if(ctx.IsLayoutDefault()) + if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS_NHWC{})) + return false; +#endif if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS_NHWC{})) return false; diff --git a/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp index 4bc9d48658..61b4f30bc8 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp @@ -792,6 +792,11 @@ ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::GetWorkspaceSize(const ConvolutionCo bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsApplicable(const ConvolutionContext& ctx) const { +#if WORKAROUND_ISSUE_1317 + if(ctx.IsLayoutDefault()) + if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC{})) + return false; +#endif if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC{})) return false; diff --git a/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp index 864852e020..7e151e313a 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp @@ -795,6 +795,11 @@ ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::Search(const ConvolutionContext& ctx bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::IsApplicable(const ConvolutionContext& ctx) const { +#if WORKAROUND_ISSUE_1317 + if(ctx.IsLayoutDefault()) + if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC{})) + return false; +#endif if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC{})) return false; diff --git a/src/solver/gemm.cpp b/src/solver/gemm.cpp index 04f7f0f5e1..6cdad1605c 100644 --- a/src/solver/gemm.cpp +++ b/src/solver/gemm.cpp @@ -32,6 +32,7 @@ #include #include #include +#include #include #include #include @@ -77,10 +78,12 @@ static inline bool IsAnyBufferFp16(const TensorDescriptor& xDesc, } #endif -bool GemmFwdBase::IsApplicable(const ExecutionContext&, +bool GemmFwdBase::IsApplicable(const ExecutionContext& ctx, const conv::ProblemDescription& problem) const { #if MIOPEN_USE_GEMM + if(conv::solver::gemm::IsWorkaroundIssue1315(ctx)) + return false; const auto& xDesc = problem.GetIn(); const auto& wDesc = problem.GetWeights(); const auto& yDesc = problem.GetOut(); @@ -88,6 +91,7 @@ bool GemmFwdBase::IsApplicable(const ExecutionContext&, !(IsAnyBufferBF16(xDesc, yDesc, wDesc) && !IsBf16Supported) && !(IsAnyBufferFp16(xDesc, yDesc, wDesc) && !IsFp16Supported); #else + std::ignore = ctx; std::ignore = problem; return false; #endif diff --git a/src/solver/gemm_bwd.cpp b/src/solver/gemm_bwd.cpp index c29be13349..d030196620 100644 --- a/src/solver/gemm_bwd.cpp +++ b/src/solver/gemm_bwd.cpp @@ -31,6 +31,7 @@ #include #include #include +#include #include #include #include @@ -91,10 +92,12 @@ SlowdownFactor(int n_oper, const double oper_factor, const double multiple_oper_ return 1.0; } -bool GemmBwdBase::IsApplicable(const ExecutionContext&, +bool GemmBwdBase::IsApplicable(const ExecutionContext& ctx, const conv::ProblemDescription& problem) const { #if MIOPEN_USE_GEMM + if(conv::solver::gemm::IsWorkaroundIssue1315(ctx)) + return false; const auto& dyDesc = problem.GetIn(); const auto& wDesc = problem.GetWeights(); const auto& dxDesc = problem.GetOut(); @@ -102,6 +105,7 @@ bool GemmBwdBase::IsApplicable(const ExecutionContext&, !(IsAnyBufferBF16(dxDesc, dyDesc, wDesc) && !IsBf16Supported) && !(IsAnyBufferFp16(dxDesc, dyDesc, wDesc) && !IsFp16Supported); #else + std::ignore = ctx; std::ignore = problem; return false; #endif diff --git a/src/solver/gemm_common.cpp b/src/solver/gemm_common.cpp new file mode 100644 index 0000000000..cbfc6f002d --- /dev/null +++ b/src/solver/gemm_common.cpp @@ -0,0 +1,59 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2021 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 +#include + +#include // std::ignore + +/// This W/A disables all GEMM convolution solvers for xDLOPs +/// targets when MIOpenGEMM is used (OCL BE). More info at +/// https://github.com/ROCmSoftwarePlatform/MIOpen/issues/1315. +/// +/// W/A affects ROCm releases starting from 4.5 and also +/// pre-5.0 Mainline HIP builds, e.g. 9148. +#define WORKAROUND_ISSUE_1315 (MIOPEN_USE_MIOPENGEMM && (HIP_PACKAGE_VERSION_FLAT >= 4004000000ULL)) + +namespace miopen { +namespace conv { +namespace solver { +namespace gemm { + +bool IsWorkaroundIssue1315(const miopen::ExecutionContext& ctx) +{ +#if WORKAROUND_ISSUE_1315 + const auto device = ctx.GetStream().GetTargetProperties().Name(); + return (device == "gfx908") || (device == "gfx90a") || (device == "gfx940"); +#else + std::ignore = ctx; + return false; +#endif +} + +} // namespace gemm +} // namespace solver +} // namespace conv +} // namespace miopen diff --git a/src/solver/gemm_wrw.cpp b/src/solver/gemm_wrw.cpp index 3cd80eadce..0bed5ae5d6 100644 --- a/src/solver/gemm_wrw.cpp +++ b/src/solver/gemm_wrw.cpp @@ -3,6 +3,7 @@ #include #include #include +#include #include #include @@ -60,10 +61,12 @@ SlowdownFactor(int n_oper, const double oper_factor, const double multiple_oper_ return 1.0; } -bool GemmWrwBase::IsApplicable(const ExecutionContext&, +bool GemmWrwBase::IsApplicable(const ExecutionContext& ctx, const conv::ProblemDescription& problem) const { #if MIOPEN_USE_GEMM + if(conv::solver::gemm::IsWorkaroundIssue1315(ctx)) + return false; const auto& dyDesc = problem.GetIn(); const auto& dwDesc = problem.GetWeights(); const auto& xDesc = problem.GetOut(); @@ -72,6 +75,7 @@ bool GemmWrwBase::IsApplicable(const ExecutionContext&, !(IsAnyBufferBF16(xDesc, dyDesc, dwDesc) && !IsBF16PathValid) && !(IsAnyBufferFp16(xDesc, dyDesc, dwDesc) && !IsFp16Supported); #else + std::ignore = ctx; std::ignore = problem; return false; #endif