Remove support for ROCm < 5.6.0

driver/driver.hpp

@@ -26,7 +26,7 @@
 #ifndef GUARD_MIOPEN_DRIVER_HPP
 #define GUARD_MIOPEN_DRIVER_HPP
 
-#if !defined(_WIN32) && (HIP_PACKAGE_VERSION_FLAT >= 5006000000ULL)
+#if !defined(_WIN32)
 #include <half/half.hpp>
 #else
 #include <half.hpp>

driver/reduce_driver.hpp

@@ -43,7 +43,7 @@
 #include <string>
 #include <cassert>
 #include <type_traits>
-#if !defined(_WIN32) && (HIP_PACKAGE_VERSION_FLAT >= 5006000000ULL)
+#if !defined(_WIN32)
 #include <half/half.hpp>
 #else
 #include <half.hpp>

src/CMakeLists.txt

@@ -413,7 +413,6 @@ if( MIOPEN_BACKEND MATCHES "OpenCL" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN
         ${GPU_REFERENCE_KERNEL_ASM}
         ${GPU_BATCHED_TRANSPOSE_KERNEL_HIP}
         ${GPU_GENERAL_TENSOR_REORDER_KERNEL_HIP_SOURCE}
-        kernels/detect_llvm_amdgcn_buffer_atomic_fadd_f32_float.cpp
         kernels/MIOpenCheckNumerics.cpp
         kernels/MIOpenBatchNormActivBwdPerAct.cl
         kernels/MIOpenBatchNormActivBwdSpatial.cl

src/comgr.cpp

@@ -34,11 +34,10 @@
 #include <miopen/gcn_asm_utils.hpp>
 #include <miopen/kernel.hpp>
 #include <miopen/logger.hpp>
-#include <miopen/rocm_features.hpp>
 #include <miopen/solver/implicitgemm_util.hpp>
 #include <miopen/stringutils.hpp>
 
-#if !defined(_WIN32) && (HIP_PACKAGE_VERSION_FLAT >= 5004000000ULL)
+#if !defined(_WIN32)
 #include <amd_comgr/amd_comgr.h>
 #else
 #include <amd_comgr.h>
@@ -105,29 +104,13 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP)
 #define COMGR_SUPPORTS_PCH (COMGR_VERSION >= 1008000)
 
 #if COMGR_SUPPORTS_PCH
-/// 3.9 reports that HIP PCH is supported, but in fact it is not.
-#define WORKAROUND_SWDEV_257056_PCH_INCORRECTLY_REPORTED 1
 
 #if defined(__HIP_HAS_GET_PCH) && __HIP_HAS_GET_PCH
 #define HIP_SUPPORTS_PCH 1
 #else
 #define HIP_SUPPORTS_PCH 0
 #endif
 
-#if WORKAROUND_SWDEV_257056_PCH_INCORRECTLY_REPORTED
-#if HIP_SUPPORTS_PCH && (HIP_PACKAGE_VERSION_FLAT <= 3009999999ULL)
-#undef HIP_SUPPORTS_PCH
-#define HIP_SUPPORTS_PCH 0
-#endif
-#endif
-
-// '__hipGetPCH' is not available in [4.4, 5.0). See SWDEV-308265.
-#if HIP_SUPPORTS_PCH && (HIP_PACKAGE_VERSION_FLAT >= 4004000000ULL) && \
-    (HIP_PACKAGE_VERSION_FLAT < 5000000000ULL)
-#undef HIP_SUPPORTS_PCH
-#define HIP_SUPPORTS_PCH 0
-#endif
-
 #endif // COMGR_SUPPORTS_PCH
 
 #define PCH_IS_SUPPORTED (COMGR_SUPPORTS_PCH && HIP_SUPPORTS_PCH)
@@ -212,7 +195,7 @@ namespace ocl {
 #error "Wrong OCL_STANDARD"
 #endif
 
-static void AddCompilerOptions(OptionList& list, const miopen::TargetProperties& target)
+static void AddCompilerOptions(OptionList& list)
 {
     list.push_back("-cl-kernel-arg-info");
 #if 0 // For experimients.
@@ -233,20 +216,6 @@ static void AddCompilerOptions(OptionList& list, const miopen::TargetProperties&
         list.push_back("-mcumode");
     }
     list.push_back("-O3");
-
-#if ROCM_FEATURE_TARGETID_OFF
-    // It seems like these options are used only in codegen.
-    // However it seems ok to pass these to compiler.
-    if(target.Sramecc())
-    {
-        if(*target.Sramecc())
-            list.push_back("-msram-ecc");
-        else
-            list.push_back("-mno-sram-ecc");
-    }
-#else
-    std::ignore = target;
-#endif
     list.push_back("-mllvm");
     list.push_back("-amdgpu-internalize-symbols");
 }
@@ -345,18 +314,12 @@ static void RemoveLinkOptionsUnwanted(OptionList& list)
 /// \todo Get list of supported isa names from comgr and select.
 static std::string GetIsaName(const miopen::TargetProperties& target, const bool isHlcBuild)
 {
-#if ROCM_FEATURE_TARGETID_OFF
-    std::ignore                  = isHlcBuild;
-    const char* const ecc_suffix = (target.Sramecc() && *target.Sramecc()) ? "+sram-ecc" : "";
-    return {"amdgcn-amd-amdhsa--" + target.Name() + ecc_suffix};
-#else
     const LcOptionTargetStrings lots(target);
 #if WORKAROUND_ISSUE_1257
     if(isHlcBuild)
         return {"amdgcn-amd-amdhsa--" + lots.device + lots.xnack};
 #endif
     return {"amdgcn-amd-amdhsa--" + lots.targetId};
-#endif
 }
 
 } // namespace lc
@@ -819,10 +782,8 @@ void BuildHip(const std::string& name,
                        + " " + GetDebugCompilerOptionsInsert() //
                        + " " + MIOPEN_STRINGIZE(HIP_COMPILER_FLAGS) +
                        (" -DHIP_PACKAGE_VERSION_FLAT=") + std::to_string(HIP_PACKAGE_VERSION_FLAT);
-#if ROCM_FEATURE_LLVM_AMDGCN_BUFFER_ATOMIC_FADD_F32_RETURNS_FLOAT
             if(miopen::solver::support_amd_buffer_atomic_fadd(target.Name()))
                 raw += " -DCK_AMD_BUFFER_ATOMIC_FADD_RETURNS_FLOAT=1";
-#endif
             auto optCompile = miopen::SplitSpaceSeparated(raw, compiler::lc::GetOptionsNoSplit());
             compiler::lc::hip::RemoveCompilerOptionsUnwanted(optCompile);
             action.SetOptionList(optCompile);
@@ -835,10 +796,8 @@ void BuildHip(const std::string& name,
                        + " " + GetDebugCompilerOptionsInsert() //
                        + " " + MIOPEN_STRINGIZE(HIP_COMPILER_FLAGS) +
                        (" -DHIP_PACKAGE_VERSION_FLAT=") + std::to_string(HIP_PACKAGE_VERSION_FLAT);
-#if ROCM_FEATURE_LLVM_AMDGCN_BUFFER_ATOMIC_FADD_F32_RETURNS_FLOAT
             if(miopen::solver::support_amd_buffer_atomic_fadd(target.Name()))
                 raw += " -DCK_AMD_BUFFER_ATOMIC_FADD_RETURNS_FLOAT=1";
-#endif
 #if PCH_IS_SUPPORTED
             if(compiler::lc::hip::IsPchEnabled())
             {
@@ -928,7 +887,7 @@ void BuildOcl(const std::string& name,
 
         auto optCompile = miopen::SplitSpaceSeparated(options);
         compiler::lc::ocl::RemoveOptionsUnwanted(optCompile);
-        compiler::lc::ocl::AddCompilerOptions(optCompile, target);
+        compiler::lc::ocl::AddCompilerOptions(optCompile);
         action.SetOptionList(optCompile);
 
         const Dataset addedPch;
@@ -1008,10 +967,8 @@ void BuildAsm(const std::string& name,
         SetIsaName(action, target);
         action.SetLogging(true);
         auto optAsm = miopen::SplitSpaceSeparated(options);
-#if ROCM_FEATURE_ASM_REQUIRES_NO_XNACK_OPTION
         if(target.Xnack() && !*target.Xnack())
             optAsm.emplace_back("-mno-xnack");
-#endif
         compiler::lc::gcnasm::RemoveOptionsUnwanted(optAsm);
         action.SetOptionList(optAsm);
 
@@ -1296,10 +1253,8 @@ void BuildHip(const std::string& name,
         opts.push_back("-D__HIP_PLATFORM_HCC__=1"); // Workaround?
 #endif
         opts.push_back("-D__HIP_PLATFORM_AMD__=1"); // Workaround?
-#if ROCM_FEATURE_LLVM_AMDGCN_BUFFER_ATOMIC_FADD_F32_RETURNS_FLOAT
         if(miopen::solver::support_amd_buffer_atomic_fadd(target.Name()))
             opts.push_back("-DCK_AMD_BUFFER_ATOMIC_FADD_RETURNS_FLOAT=1");
-#endif
         opts.push_back("-DHIP_PACKAGE_VERSION_FLAT=" + std::to_string(HIP_PACKAGE_VERSION_FLAT));
         opts.push_back("-DMIOPEN_DONT_USE_HIP_RUNTIME_HEADERS");
 #if HIP_PACKAGE_VERSION_FLAT < 6001024000ULL

src/composable_kernel/composable_kernel/include/utility/reduction_operator.hpp

@@ -224,22 +224,6 @@ struct unary_abs<T, false>
     __device__ inline constexpr T operator()(T a) const { return abs(a); };
 };
 
-// We know for sure that 4.0 has __habs(), but 3.0 does not have it.
-// Let's assume that __habs() exists since 3.5.
-#if HIP_PACKAGE_VERSION_FLAT < 3005000000
-inline __device__ __half __habs(__half x)
-{
-    union
-    {
-        __half half;
-        unsigned short u16;
-    } val;
-    val.half = x;
-    val.u16  = val.u16 & 0x7fff;
-    return val.half;
-}
-#endif
-
 template <bool hasDividing>
 struct unary_abs<half_t, hasDividing>
 {

src/composable_kernel/host/driver_offline/src/conv_bwd_driver_offline.cpp

@@ -3,7 +3,7 @@
 #include <initializer_list>
 #include <cstdlib>
 #include <stdlib.h>
-#if !defined(_WIN32) && (HIP_PACKAGE_VERSION_FLAT >= 5006000000ULL)
+#if !defined(_WIN32)
 #include <half/half.hpp>
 #else
 #include <half.hpp>

src/composable_kernel/host/driver_offline/src/conv_fwd_driver_offline.cpp

@@ -3,7 +3,7 @@
 #include <initializer_list>
 #include <cstdlib>
 #include <stdlib.h>
-#if !defined(_WIN32) && (HIP_PACKAGE_VERSION_FLAT >= 5006000000ULL)
+#if !defined(_WIN32)
 #include <half/half.hpp>
 #else
 #include <half.hpp>

src/fusion.cpp

@@ -45,7 +45,7 @@
 #include <ios>
 #include <algorithm>
 #include <string>
-#if !defined(_WIN32) && (HIP_PACKAGE_VERSION_FLAT >= 5006000000ULL)
+#if !defined(_WIN32)
 #include <half/half.hpp>
 #else
 #include <half.hpp>

src/gemm_v2.cpp

@@ -40,7 +40,7 @@
 #pragma clang diagnostic ignored "-Wunused-macros"
 #define ROCBLAS_BETA_FEATURES_API 1
 #pragma clang diagnostic pop
-#if !defined(_WIN32) && (HIP_PACKAGE_VERSION_FLAT >= 5006000000ULL)
+#if !defined(_WIN32)
 #include <half/half.hpp>
 #else
 #include <half.hpp>

src/hip/handlehip.cpp

@@ -34,7 +34,6 @@
 #include <miopen/invoker.hpp>
 #include <miopen/kernel_cache.hpp>
 #include <miopen/logger.hpp>
-#include <miopen/rocm_features.hpp>
 #include <miopen/stringutils.hpp>
 #include <miopen/target_properties.hpp>
 #include <miopen/timer.hpp>
@@ -57,24 +56,17 @@
 #include <mutex>
 #include <shared_mutex>
 
-#define MIOPEN_WORKAROUND_ROCM_COMPILER_SUPPORT_ISSUE_30 \
-    (MIOPEN_USE_COMGR && BUILD_SHARED_LIBS && (HIP_PACKAGE_VERSION_FLAT < 4003000000ULL))
-
+/// GFX906 and GFX103X are deprecated since 5.7 RC.
 /// hipMemGetInfo constantly fails on gfx906/900 and Navi21.
 /// Brute-force W/A: return fixed values.
-#define WORKAROUND_FAULTY_HIPMEMGETINFO_VEGA_NAVI2X (ROCM_FEATURE_DEPRECATED_VEGA_NAVI2X)
+#define WORKAROUND_FAULTY_HIPMEMGETINFO_VEGA_NAVI2X (HIP_PACKAGE_VERSION_FLAT >= 5007000000ULL)
 
 MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEVICE_CU)
 
 namespace miopen {
 
 namespace {
 
-#if MIOPEN_WORKAROUND_ROCM_COMPILER_SUPPORT_ISSUE_30
-void toCallHipInit() __attribute__((constructor(1000)));
-void toCallHipInit() { hipInit(0); }
-#endif
-
 hipError_t hip_mem_get_info_wrapper(std::size_t* const free, std::size_t* const total)
 {
 #if WORKAROUND_FAULTY_HIPMEMGETINFO_VEGA_NAVI2X
@@ -226,11 +218,7 @@ struct HandleImpl
     {
         hipDeviceProp_t props{};
         hipGetDeviceProperties(&props, device);
-#if ROCM_FEATURE_HIP_GCNARCHNAME_RETURNS_CODENAME
-        const std::string name("gfx" + std::to_string(props.gcnArch));
-#else
         const std::string name(props.gcnArchName);
-#endif
         MIOPEN_LOG_NQI("Raw device name: " << name);
         return name; // NOLINT (performance-no-automatic-move)
     }

src/hip/hip_build_utils.cpp

@@ -30,7 +30,6 @@
 #include <miopen/exec_utils.hpp>
 #include <miopen/logger.hpp>
 #include <miopen/env.hpp>
-#include <miopen/rocm_features.hpp>
 #include <miopen/solver/implicitgemm_util.hpp>
 #include <miopen/target_properties.hpp>
 #include <boost/optional.hpp>
@@ -89,21 +88,13 @@ static boost::filesystem::path HipBuildImpl(boost::optional<TmpDir>& tmp_dir,
     }
 #endif
 
-#if HIP_PACKAGE_VERSION_FLAT < 4001000000ULL
-    params += " --cuda-gpu-arch=" + lots.device;
-#else
     params += " --cuda-gpu-arch=" + lots.device + lots.xnack;
-#endif
     params += " --cuda-device-only";
     params += " -c";
     params += " -O3 ";
     params += " -Wno-unused-command-line-argument -I. ";
     params += MIOPEN_STRINGIZE(HIP_COMPILER_FLAGS);
 
-#if HIP_PACKAGE_VERSION_FLAT < 4004000000ULL
-    params += " -mllvm --amdgpu-spill-vgpr-to-agpr=0";
-#endif
-
 #if MIOPEN_BUILD_DEV
     if(miopen::IsEnabled(ENV(MIOPEN_DEBUG_HIP_VERBOSE)))
     {
@@ -137,18 +128,9 @@ static boost::filesystem::path HipBuildImpl(boost::optional<TmpDir>& tmp_dir,
     // Unbundling is not required for HIP runtime && hip-clang
     tmp_dir->Execute(MIOPEN_OFFLOADBUNDLER_BIN,
                      "--type=o "
-#if(HIP_PACKAGE_VERSION_FLAT >= 4001021072ULL && HIP_PACKAGE_VERSION_FLAT < 4002000000ULL) || \
-    HIP_PACKAGE_VERSION_FLAT >= 4002021072ULL
-                     "--targets=hipv4-amdgcn-amd-amdhsa-"
-#else
-                     "--targets=hip-amdgcn-amd-amdhsa-"
-#endif
-#if HIP_PACKAGE_VERSION_FLAT < 4001000000ULL
-                         + lots.device
-#else
-                         + (std::string{'-'} + lots.device + lots.xnack)
-#endif
-                         + " --inputs=" + bin_file.string() + " --outputs=" + bin_file.string() +
+                     "--targets=hipv4-amdgcn-amd-amdhsa-" +
+                         (std::string{'-'} + lots.device + lots.xnack) +
+                         " --inputs=" + bin_file.string() + " --outputs=" + bin_file.string() +
                          ".hsaco --unbundle");
 
     auto hsaco = std::find_if(boost::filesystem::directory_iterator{tmp_dir->path},
@@ -169,58 +151,14 @@ static boost::filesystem::path HipBuildImpl(boost::optional<TmpDir>& tmp_dir,
 #endif
 }
 
-#ifndef ROCM_FEATURE_LLVM_AMDGCN_BUFFER_ATOMIC_FADD_F32_RETURNS_FLOAT
-static bool
-HipBuildTest(const std::string& program_name, std::string params, const TargetProperties& target)
-{
-    boost::optional<miopen::TmpDir> dir(program_name);
-    std::string source = miopen::GetKernelSrc(program_name);
-    try
-    {
-        std::ignore = HipBuildImpl(dir, program_name, source, params, target, true);
-    }
-    catch(...)
-    {
-        return false;
-    }
-    return true;
-}
-
-static bool DetectIfBufferAtomicFaddReturnsFloatImpl(const TargetProperties& target)
-{
-    const std::string program_name("detect_llvm_amdgcn_buffer_atomic_fadd_f32_float.cpp");
-    std::string params;
-
-    if(HipBuildTest(program_name, params, target))
-    {
-        MIOPEN_LOG_NQI("Yes");
-        return true;
-    }
-    MIOPEN_LOG_NQI("No");
-    return false;
-}
-
-static bool DetectIfBufferAtomicFaddReturnsFloat(const TargetProperties& target)
-{
-    static const bool once = DetectIfBufferAtomicFaddReturnsFloatImpl(target);
-    return once;
-}
-#endif
-
 boost::filesystem::path HipBuild(boost::optional<TmpDir>& tmp_dir,
                                  const std::string& filename,
                                  std::string src,
                                  std::string params,
                                  const TargetProperties& target)
 {
-#ifndef ROCM_FEATURE_LLVM_AMDGCN_BUFFER_ATOMIC_FADD_F32_RETURNS_FLOAT
-    if(miopen::solver::support_amd_buffer_atomic_fadd(target.Name()))
-        if(DetectIfBufferAtomicFaddReturnsFloat(target))
-            params += " -DCK_AMD_BUFFER_ATOMIC_FADD_RETURNS_FLOAT=1";
-#elif ROCM_FEATURE_LLVM_AMDGCN_BUFFER_ATOMIC_FADD_F32_RETURNS_FLOAT
     if(miopen::solver::support_amd_buffer_atomic_fadd(target.Name()))
         params += " -DCK_AMD_BUFFER_ATOMIC_FADD_RETURNS_FLOAT=1";
-#endif
     return HipBuildImpl(tmp_dir, filename, src, params, target, false);
 }