From 36ab7799d7e0eb641eeb28834326fc74bc44309a Mon Sep 17 00:00:00 2001 From: Artem Tamazov Date: Mon, 22 Jul 2024 08:32:38 +0300 Subject: [PATCH] [COMgr] Prepare for COMgr 3.x (#3107) * amdcomgr-3(01) [cmake] Do not allow different values of MIOPEN_USE_COMGR and MIOPEN_USE_HIPRTC. Remove leftover of support of ROCm older than 5.0. * amdcomgr-3(02) [importance_normal] Fix precompiled binary cache miss for Winograd Fury. Resolves https://github.com/ROCm/MIOpen/pull/2778#discussion_r1670854903 * amdcomgr-3(03) [winograd fury] Add comment. * amdcomgr-3(04) [comgr] Removed comgr::BuildHip() * amdcomgr-3(05) [comgr] Added support for AMD COMgr 3.0. Removed support for AMD COMgr older than 1.7. * amdcomgr-3(06) [comgr] Removed support of unused enum members from to_string(). * amdcomgr-3(08) [comgr] Fix tidy error --------- Co-authored-by: Evgenii Averin <86725875+averinevg@users.noreply.github.com> --- CMakeLists.txt | 20 +- src/comgr.cpp | 336 ++----------------------- src/hipoc/hipoc_program.cpp | 8 +- src/include/miopen/comgr.hpp | 6 - src/solver/conv/conv_wino_fury_RxS.cpp | 9 +- 5 files changed, 42 insertions(+), 337 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index a91bd87733..c78430e195 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -37,6 +37,19 @@ macro(set_var_to_condition var) endif() endmacro() +macro(set_if_bools_are_different var in1 in2) + set(${var} FALSE) + if(${in1}) + if(NOT ${in2}) + set(${var} TRUE) + endif() + else() + if(${in2}) + set(${var} TRUE) + endif() + endif() +endmacro() + get_property(MIOPEN_GENERATOR_IS_MULTI_CONFIG GLOBAL PROPERTY GENERATOR_IS_MULTI_CONFIG) # This has to be initialized before the project() command appears @@ -252,9 +265,14 @@ set(MIOPEN_hip_VERSION ${MIOPEN_hip_VERSION_MAJOR}.${MIOPEN_hip_VERSION_MINOR}.$ # Do not enable HIPRTC by default for older ROCm versions in order to avoid # build time errors, because HIPRTC is a relatively new component. -set_var_to_condition(MIOPEN_USE_HIPRTC_DEFAULT (MIOPEN_USE_COMGR AND (MIOPEN_hip_VERSION VERSION_GREATER_EQUAL 5))) +set_var_to_condition(MIOPEN_USE_HIPRTC_DEFAULT MIOPEN_USE_COMGR) option(MIOPEN_USE_HIPRTC "Use HIPRTC to build HIP kernels instead of COMGR" ${MIOPEN_USE_HIPRTC_DEFAULT}) +set_if_bools_are_different(MIOPEN_CONFIGURATION_ERROR_COMGR_HIPRTC MIOPEN_USE_COMGR MIOPEN_USE_HIPRTC) +if(MIOPEN_CONFIGURATION_ERROR_COMGR_HIPRTC) + message(FATAL_ERROR "MIOPEN_USE_COMGR (${MIOPEN_USE_COMGR}) and MIOPEN_USE_HIPRTC (${MIOPEN_USE_HIPRTC}) should be set to the same value") +endif() + # Do not append system include directories to HIP compiler flags when HIPRTC is used set_var_to_condition(MIOPEN_HIP_COMPILER_USE_SYSTEM_INCLUDE_DIRECTORIES_DEFAULT (NOT (MIOPEN_USE_HIPRTC AND (MIOPEN_hip_VERSION VERSION_GREATER_EQUAL 6.1.40091)))) diff --git a/src/comgr.cpp b/src/comgr.cpp index f64c6377d9..aa53b71bb5 100644 --- a/src/comgr.cpp +++ b/src/comgr.cpp @@ -70,11 +70,6 @@ MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_COMGR_LOG_OPTIONS) /// you would like to log onto console. MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_COMGR_LOG_SOURCE_TEXT) -/// \todo Temporary for debugging: -MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_DEBUG_COMGR_COMPILER_OPTIONS_INSERT) -/// \todo Temporary for debugging: -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMGR_HIP_BUILD_FATBIN) - /// \todo see issue #1222, PR #1316 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_SRAM_EDC_DISABLED) @@ -99,16 +94,18 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP) ((MIOPEN_AMD_COMGR_VERSION_MAJOR * 1000 + MIOPEN_AMD_COMGR_VERSION_MINOR) * 1000 + \ MIOPEN_AMD_COMGR_VERSION_PATCH) +#if COMGR_VERSION < 1007000 +#error "AMD COMgr older than 1.7.0 is not supported" +#endif + #define COMGR_SUPPORTS_PCH (COMGR_VERSION >= 1008000) #if COMGR_SUPPORTS_PCH - #if defined(__HIP_HAS_GET_PCH) && __HIP_HAS_GET_PCH #define HIP_SUPPORTS_PCH 1 #else #define HIP_SUPPORTS_PCH 0 #endif - #endif // COMGR_SUPPORTS_PCH #define PCH_IS_SUPPORTED (COMGR_SUPPORTS_PCH && HIP_SUPPORTS_PCH) @@ -118,10 +115,6 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP) /// have wavesize != 64 (currently gfx10 with default build settings). #define WORKAROUND_ISSUE_1431 PCH_IS_SUPPORTED -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMGR_HIP_PCH_ENFORCE) - -#define COMPILER_LC 1 - #define EC_BASE(comgrcall, info, action) \ do \ { \ @@ -159,7 +152,6 @@ using OptionList = std::vector; /// (minimal compiler abstraction layer). namespace compiler { -#if COMPILER_LC namespace lc { static auto GetOptionsNoSplit() @@ -183,8 +175,6 @@ static void RemoveOptionsUnwanted(OptionList& list) namespace ocl { -#define OCL_COMPILE_SOURCE_WITH_DEVICE_LIBS (COMGR_VERSION >= 1007000) - #define OCL_EARLY_INLINE 1 #define OCL_STANDARD 200 @@ -235,83 +225,6 @@ static void RemoveOptionsUnwanted(OptionList& list) } // namespace ocl -namespace hip { - -#if PCH_IS_SUPPORTED -static bool IsPchEnabled() { return !env::disabled(MIOPEN_DEBUG_COMGR_HIP_PCH_ENFORCE); } -#endif - -static std::string GetPchEnableStatus() -{ -#if PCH_IS_SUPPORTED - auto rv = std::string{IsPchEnabled() ? "1" : "0"}; - if(env::disabled(MIOPEN_DEBUG_COMGR_HIP_PCH_ENFORCE)) - return rv += " (enforced)"; - return rv; -#else - return "0 (not supported)"; -#endif -} - -static bool IsLinkerOption(const std::string& option) -{ - return miopen::StartsWith(option, "-L") || miopen::StartsWith(option, "-Wl,") || - option == "-ldl" || option == "-lm" || option == "--hip-link"; -} - -static void RemoveCommonOptionsUnwanted(OptionList& list) -{ - list.erase( - remove_if( - list.begin(), - list.end(), - [&](const auto& option) { // clang-format off - return miopen::StartsWith(option, "-mcpu=") - || (option == "-hc") - || (option == "-x hip") || (option == "-xhip") - || (option == "--hip-link") - // The following matches current "-lclang_rt.builtins-x86_64" (4.5) as weel as - // upcoming ".../libclang_rt.builtins-x86_64.a" and even future things like - // "...x86_64.../libclang_rt.builtins.a" etc. - || ((option.find("clang_rt.builtins") != std::string::npos) - && (option.find("x86_64") != std::string::npos)) - || miopen::StartsWith(option, "--hip-device-lib-path="); // clang-format on - }), - list.end()); -} - -static void AddCompilerOptions(const OptionList& list) // `const` is for clang-tidy. -{ - // Nothing to do here yet, but let's keep the placeholder for now. - std::ignore = list; -} - -static void RemoveCompilerOptionsUnwanted(OptionList& list) -{ - RemoveCommonOptionsUnwanted(list); - list.erase(remove_if(list.begin(), - list.end(), - [&](const auto& option) { // clang-format off - return (!env::enabled(MIOPEN_DEBUG_COMGR_HIP_BUILD_FATBIN) - && (IsLinkerOption(option))); // clang-format on - }), - list.end()); -} - -static void RemoveLinkOptionsUnwanted(OptionList& list) -{ - RemoveCommonOptionsUnwanted(list); - list.erase(remove_if(list.begin(), - list.end(), - [&](const auto& option) { // clang-format off - return miopen::StartsWith(option, "-D") - || miopen::StartsWith(option, "-isystem"); // clang-format on - }), - list.end()); -} - -} // namespace hip - /// \todo Get list of supported isa names from comgr and select. static std::string GetIsaName(const miopen::TargetProperties& target, const bool isHlcBuild) { @@ -325,7 +238,6 @@ static std::string GetIsaName(const miopen::TargetProperties& target, const bool } // namespace lc #undef OCL_EARLY_INLINE -#endif // COMPILER_LC } // namespace compiler @@ -338,6 +250,12 @@ static inline std::string to_string(const std::string& v) { return {v}; } static inline std::string to_string(const bool& v) { return v ? "true" : "false"; } static inline auto to_string(const std::size_t& v) { return std::to_string(v); } +/// Convert amd_comgr enum members to strings. +/// +/// \note We need support only for the enum members used in this file. +/// Let's skip unused members in order to simplify maintenance +/// of code between different COMgr versions. +/// /// \todo Request comgr to expose this stuff via API. static std::string to_string(const amd_comgr_language_t val) { @@ -347,7 +265,6 @@ static std::string to_string(const amd_comgr_language_t val) AMD_COMGR_LANGUAGE_NONE, AMD_COMGR_LANGUAGE_OPENCL_1_2, AMD_COMGR_LANGUAGE_OPENCL_2_0, - AMD_COMGR_LANGUAGE_HC, AMD_COMGR_LANGUAGE_HIP); return oss.str(); } @@ -360,58 +277,21 @@ static std::string to_string(const amd_comgr_data_kind_t val) AMD_COMGR_DATA_KIND_UNDEF, AMD_COMGR_DATA_KIND_SOURCE, AMD_COMGR_DATA_KIND_INCLUDE, - AMD_COMGR_DATA_KIND_PRECOMPILED_HEADER, - AMD_COMGR_DATA_KIND_DIAGNOSTIC, AMD_COMGR_DATA_KIND_LOG, - AMD_COMGR_DATA_KIND_BC, - AMD_COMGR_DATA_KIND_RELOCATABLE, - AMD_COMGR_DATA_KIND_EXECUTABLE, - AMD_COMGR_DATA_KIND_BYTES, - AMD_COMGR_DATA_KIND_FATBIN); + AMD_COMGR_DATA_KIND_EXECUTABLE); return oss.str(); } static std::string to_string(const amd_comgr_action_kind_t val) { std::ostringstream oss; -#if COMGR_VERSION >= 1007000 MIOPEN_LOG_ENUM(oss, val, - AMD_COMGR_ACTION_SOURCE_TO_PREPROCESSOR, AMD_COMGR_ACTION_ADD_PRECOMPILED_HEADERS, - AMD_COMGR_ACTION_COMPILE_SOURCE_TO_BC, - AMD_COMGR_ACTION_ADD_DEVICE_LIBRARIES, - AMD_COMGR_ACTION_LINK_BC_TO_BC, - AMD_COMGR_ACTION_OPTIMIZE_BC_TO_BC, AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE, - AMD_COMGR_ACTION_CODEGEN_BC_TO_ASSEMBLY, - AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_RELOCATABLE, AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE, AMD_COMGR_ACTION_ASSEMBLE_SOURCE_TO_RELOCATABLE, - AMD_COMGR_ACTION_DISASSEMBLE_RELOCATABLE_TO_SOURCE, - AMD_COMGR_ACTION_DISASSEMBLE_EXECUTABLE_TO_SOURCE, - AMD_COMGR_ACTION_DISASSEMBLE_BYTES_TO_SOURCE, - AMD_COMGR_ACTION_COMPILE_SOURCE_TO_FATBIN, AMD_COMGR_ACTION_COMPILE_SOURCE_WITH_DEVICE_LIBS_TO_BC); -#else - MIOPEN_LOG_ENUM(oss, - val, - AMD_COMGR_ACTION_SOURCE_TO_PREPROCESSOR, - AMD_COMGR_ACTION_ADD_PRECOMPILED_HEADERS, - AMD_COMGR_ACTION_COMPILE_SOURCE_TO_BC, - AMD_COMGR_ACTION_ADD_DEVICE_LIBRARIES, - AMD_COMGR_ACTION_LINK_BC_TO_BC, - AMD_COMGR_ACTION_OPTIMIZE_BC_TO_BC, - AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE, - AMD_COMGR_ACTION_CODEGEN_BC_TO_ASSEMBLY, - AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_RELOCATABLE, - AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE, - AMD_COMGR_ACTION_ASSEMBLE_SOURCE_TO_RELOCATABLE, - AMD_COMGR_ACTION_DISASSEMBLE_RELOCATABLE_TO_SOURCE, - AMD_COMGR_ACTION_DISASSEMBLE_EXECUTABLE_TO_SOURCE, - AMD_COMGR_ACTION_DISASSEMBLE_BYTES_TO_SOURCE, - AMD_COMGR_ACTION_COMPILE_SOURCE_TO_FATBIN); -#endif return oss.str(); } @@ -420,8 +300,7 @@ static bool PrintVersionImpl() std::size_t major = 0; std::size_t minor = 0; (void)amd_comgr_get_version(&major, &minor); - MIOPEN_LOG_NQI("COMgr v." << major << '.' << minor << '.' << MIOPEN_AMD_COMGR_VERSION_PATCH - << ", USE_HIP_PCH: " << compiler::lc::hip::GetPchEnableStatus()); + MIOPEN_LOG_NQI("COMgr v." << major << '.' << minor << '.' << MIOPEN_AMD_COMGR_VERSION_PATCH); return true; } @@ -538,12 +417,6 @@ class Data : ComgrOwner { ECI_THROW(amd_comgr_set_data(handle, bytes.size(), bytes.data()), bytes.size()); } -#if PCH_IS_SUPPORTED - void SetFromBuffer(const char* const buffer, const size_t size) const - { - ECI_THROW(amd_comgr_set_data(handle, size, buffer), size); - } -#endif private: std::size_t GetSize() const @@ -599,21 +472,6 @@ class Dataset : ComgrOwner MIOPEN_LOG_I(text); } } -#if PCH_IS_SUPPORTED - void AddDataHipPch(const char* const content, const size_t size) const - { - const char name[] = "hip.pch"; - const Data d(AMD_COMGR_DATA_KIND_PRECOMPILED_HEADER); - if(env::enabled(MIOPEN_DEBUG_COMGR_LOG_SOURCE_NAMES)) - { - MIOPEN_LOG_I(name << ' ' << size - << " bytes, ptr = " << static_cast(content)); - } - d.SetName(name); - d.SetFromBuffer(content, size); - AddData(d); - } -#endif size_t GetDataCount(const amd_comgr_data_kind_t kind) const { std::size_t count = 0; @@ -728,142 +586,14 @@ static void SetIsaName(const ActionInfo& action, action.SetIsaName(isaName); } -static std::string GetDebugCompilerOptionsInsert() -{ - return env::value(MIOPEN_DEBUG_COMGR_COMPILER_OPTIONS_INSERT); -} - +#if WORKAROUND_ISSUE_1431 static inline bool IsWave64Enforced(const OptionList& opts) { return std::any_of( opts.begin(), opts.end(), [](const std::string& s) { return s == "-mwavefrontsize64"; }); } - -void BuildHip(const std::string& name, - std::string_view text, - const std::string& options, - const miopen::TargetProperties& target, - std::vector& binary) -{ - PrintVersion(); - try - { - const Dataset inputs; - inputs.AddData(name, text, AMD_COMGR_DATA_KIND_SOURCE); - - // For OCL and ASM sources, we do insert contents of include - // files directly into the source text during library build phase by means - // of the addkernels tool. We don't do that for HIP sources, and, therefore - // have to export include files prior compilation. - // Note that we do not need any "subdirs" in the include "pathnames" so far. - for(const auto& inc : GetKernelIncList()) - inputs.AddData(inc.get().string(), GetKernelInc(inc), AMD_COMGR_DATA_KIND_INCLUDE); - -#if PCH_IS_SUPPORTED - if(compiler::lc::hip::IsPchEnabled()) - { - const char* pch = nullptr; - unsigned int pch_size = 0; - __hipGetPCH(&pch, &pch_size); - inputs.AddDataHipPch(pch, pch_size); - } #endif - const ActionInfo action; - action.SetLanguage(AMD_COMGR_LANGUAGE_HIP); - SetIsaName(action, target, true); - action.SetLogging(true); - - const Dataset exe; - if(env::enabled(MIOPEN_DEBUG_COMGR_HIP_BUILD_FATBIN)) - { - auto raw = options // - + " " + GetDebugCompilerOptionsInsert() // - + " " + MIOPEN_STRINGIZE(HIP_COMPILER_FLAGS) + - (" -DHIP_PACKAGE_VERSION_FLAT=") + std::to_string(HIP_PACKAGE_VERSION_FLAT); - if(miopen::solver::support_amd_buffer_atomic_fadd(target.Name())) - raw += " -DCK_AMD_BUFFER_ATOMIC_FADD_RETURNS_FLOAT=1"; - auto optCompile = miopen::SplitSpaceSeparated(raw, compiler::lc::GetOptionsNoSplit()); - compiler::lc::hip::RemoveCompilerOptionsUnwanted(optCompile); - action.SetOptionList(optCompile); - action.Do(AMD_COMGR_ACTION_COMPILE_SOURCE_TO_FATBIN, inputs, exe); - } - else - { - auto raw = std::string(" -O3 ") // Without this, fails in lld. - + options // - + " " + GetDebugCompilerOptionsInsert() // - + " " + MIOPEN_STRINGIZE(HIP_COMPILER_FLAGS) + - (" -DHIP_PACKAGE_VERSION_FLAT=") + std::to_string(HIP_PACKAGE_VERSION_FLAT); - if(miopen::solver::support_amd_buffer_atomic_fadd(target.Name())) - raw += " -DCK_AMD_BUFFER_ATOMIC_FADD_RETURNS_FLOAT=1"; -#if PCH_IS_SUPPORTED - if(compiler::lc::hip::IsPchEnabled()) - { - raw += " -nogpuinc -DMIOPEN_DONT_USE_HIP_RUNTIME_HEADERS"; - } -#endif - auto optCompile = miopen::SplitSpaceSeparated(raw, compiler::lc::GetOptionsNoSplit()); - auto optLink = optCompile; - compiler::lc::hip::RemoveCompilerOptionsUnwanted(optCompile); - compiler::lc::hip::AddCompilerOptions(optCompile); -#if WORKAROUND_ISSUE_1431 - if(compiler::lc::hip::IsPchEnabled()) - { - if((StartsWith(target.Name(), "gfx10") || StartsWith(target.Name(), "gfx11")) && - !IsWave64Enforced(optCompile)) - optCompile.emplace_back("-DWORKAROUND_ISSUE_1431=1"); - } -#endif - action.SetOptionList(optCompile); - const Dataset compiledBc; - action.Do(AMD_COMGR_ACTION_COMPILE_SOURCE_TO_BC, inputs, compiledBc); - - OptionList addDevLibs; - // Use device libs for wavefrontsize64 for non-gfx10 targets - // or when enforced via option. - if(!(StartsWith(target.Name(), "gfx10") || StartsWith(target.Name(), "gfx11")) || - IsWave64Enforced(optCompile)) - { - addDevLibs.push_back("wavefrontsize64"); - } - addDevLibs.push_back("daz_opt"); // Assume that it's ok to flush denormals to zero. - addDevLibs.push_back("finite_only"); // No need to handle INF correcly. - addDevLibs.push_back("unsafe_math"); // Prefer speed over correctness for FP math. - action.SetOptionList(addDevLibs); - const Dataset withDevLibs; - action.Do(AMD_COMGR_ACTION_ADD_DEVICE_LIBRARIES, compiledBc, withDevLibs); - - compiler::lc::hip::RemoveLinkOptionsUnwanted(optLink); - action.SetOptionList(optLink); - const Dataset linkedBc; - action.Do(AMD_COMGR_ACTION_LINK_BC_TO_BC, withDevLibs, linkedBc); - - OptionList codegenBcToRel; - codegenBcToRel.push_back("-O3"); // Nothing more is required at this step. - action.SetOptionList(codegenBcToRel); - const Dataset relocatable; - action.Do(AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE, linkedBc, relocatable); - - action.SetOptionList(OptionList()); - action.Do(AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE, relocatable, exe); - } - - if(exe.GetDataCount(AMD_COMGR_DATA_KIND_EXECUTABLE) < 1) - throw ComgrError{AMD_COMGR_STATUS_ERROR, true, "Executable binary not found"}; - // Assume that the first exec data contains the binary we need. - const auto data = exe.GetData(AMD_COMGR_DATA_KIND_EXECUTABLE, 0); - data.GetBytes(binary); - } - catch(ComgrError& ex) - { - binary.resize(0); // Necessary when "get binary" fails. - MIOPEN_LOG_E("comgr status = " << GetStatusText(ex)); - if(!ex.text.empty()) - MIOPEN_LOG_W(ex.text); - } -} - void BuildOcl(const std::string& name, std::string_view text, const std::string& options, @@ -891,41 +621,8 @@ void BuildOcl(const std::string& name, const Dataset addedPch; action.Do(AMD_COMGR_ACTION_ADD_PRECOMPILED_HEADERS, inputs, addedPch); -#if OCL_COMPILE_SOURCE_WITH_DEVICE_LIBS const Dataset linkedBc; action.Do(AMD_COMGR_ACTION_COMPILE_SOURCE_WITH_DEVICE_LIBS_TO_BC, addedPch, linkedBc); -#else - const Dataset compiledBc; - action.Do(AMD_COMGR_ACTION_COMPILE_SOURCE_TO_BC, addedPch, compiledBc); - - OptionList optLink; - // Use device libs for wavefrontsize64 for non-gfx10 targets - // or when enforced via option. - if(!(StartsWith(target.Name(), "gfx10") || StartsWith(target.Name(), "gfx11")) || - IsWave64Enforced(optCompile)) - { - optLink.push_back("wavefrontsize64"); - } - for(const auto& opt : optCompile) - { - if(opt == "-cl-fp32-correctly-rounded-divide-sqrt") - optLink.push_back("correctly_rounded_sqrt"); - else if(opt == "-cl-denorms-are-zero") - optLink.push_back("daz_opt"); - else if(opt == "-cl-finite-math-only" || opt == "-cl-fast-relaxed-math") - optLink.push_back("finite_only"); - else if(opt == "-cl-unsafe-math-optimizations" || opt == "-cl-fast-relaxed-math") - optLink.push_back("unsafe_math"); - else - { - } // nop - } - action.SetOptionList(optLink); - const Dataset addedDevLibs; - action.Do(AMD_COMGR_ACTION_ADD_DEVICE_LIBRARIES, compiledBc, addedDevLibs); - const Dataset linkedBc; - action.Do(AMD_COMGR_ACTION_LINK_BC_TO_BC, addedDevLibs, linkedBc); -#endif action.SetOptionList(optCompile); const Dataset relocatable; @@ -1015,7 +712,6 @@ using OptionList = std::vector; /// Compiler implementation-specific functionality namespace compiler { -#if COMPILER_LC namespace lc { static inline void RemoveOptionsUnwanted(OptionList& list) @@ -1027,7 +723,6 @@ static inline void RemoveOptionsUnwanted(OptionList& list) } } // namespace lc -#endif // COMPILER_LC } // namespace compiler @@ -1158,6 +853,11 @@ class HiprtcProgram : src_name(src_name_), src_text(src_text_) { LogInputFile(src_name, src_text); + // For OCL and ASM sources, we do insert contents of include + // files directly into the source text during library build phase by means + // of the addkernels tool. We don't do that for HIP sources, and, therefore + // have to export include files prior compilation. + // Note that we do not need any "subdirs" in the include "pathnames" so far. const auto inc_names = miopen::GetKernelIncList(); include_names.reserve(inc_names.size()); for(const auto& inc_name : inc_names) diff --git a/src/hipoc/hipoc_program.cpp b/src/hipoc/hipoc_program.cpp index 3fb2221607..ec090455e8 100644 --- a/src/hipoc/hipoc_program.cpp +++ b/src/hipoc/hipoc_program.cpp @@ -60,7 +60,6 @@ MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_OPENCL_ENFORCE_CODE_OBJECT_VERSION) MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_DEVICE_ARCH) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP) -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_USE_HIPRTC) #if MIOPEN_USE_COMGR #define MIOPEN_WORKAROUND_ROCM_COMPILER_SUPPORT_ISSUE_27 1 @@ -276,12 +275,7 @@ void HIPOCProgramImpl::BuildCodeObjectInMemory(const std::string& params, #endif if(filename.extension() == ".cpp") { -#if MIOPEN_USE_HIPRTC - if(!env::disabled(MIOPEN_DEBUG_USE_HIPRTC)) - hiprtc::BuildHip(filename.string(), src, params, target, binary); - else -#endif // MIOPEN_USE_HIPRTC - comgr::BuildHip(filename.string(), src, params, target, binary); + hiprtc::BuildHip(filename.string(), src, params, target, binary); } else if(filename.extension() == ".s") { diff --git a/src/include/miopen/comgr.hpp b/src/include/miopen/comgr.hpp index f89182a36b..ccb3c58549 100644 --- a/src/include/miopen/comgr.hpp +++ b/src/include/miopen/comgr.hpp @@ -36,12 +36,6 @@ namespace miopen { namespace comgr { -void BuildHip(const std::string& name, - std::string_view text, - const std::string& options, - const miopen::TargetProperties& target, - std::vector& binary); - void BuildOcl(const std::string& name, std::string_view text, const std::string& options, diff --git a/src/solver/conv/conv_wino_fury_RxS.cpp b/src/solver/conv/conv_wino_fury_RxS.cpp index 80c8646aba..e80bd7fc04 100644 --- a/src/solver/conv/conv_wino_fury_RxS.cpp +++ b/src/solver/conv/conv_wino_fury_RxS.cpp @@ -32,9 +32,7 @@ #include #include #include -#if !MIOPEN_USE_COMGR #include -#endif #include #include @@ -375,12 +373,13 @@ ConvWinoFuryRxSCommon::GetSolution(const ExecutionContext& // KernelInfo KernelInfo kernel; -#if !MIOPEN_USE_COMGR + /// Kernel doesn't need ROCM_METADATA_VERSION, but AmdgcnAssemble() + /// uses it to find out required CO version (hack). + /// \todo Delete when COv2 support is removed. KernelBuildParameters options{ - {"ROCM_METADATA_VERSION", 5}, // For AmdgcnAssemble(...) + {"ROCM_METADATA_VERSION", 5}, }; kernel.comp_options = options.GenerateFor(kbp::GcnAsm{}); -#endif kernel.comp_options += std::string(" -mcumode -mwavefrontsize64"); kernel.l_wk.push_back(wg_size);