From dea91c954b75c137d47bf4db89531783664e56e5 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Fri, 17 Dec 2021 15:14:12 -0800 Subject: [PATCH] [MIOpen2.14.0] fix issues in release --- Jenkinsfile | 59 ------------------- src/comgr.cpp | 44 -------------- .../utility/static_kernel_ck_utils_type.hpp | 43 ++++++++++++++ 3 files changed, 43 insertions(+), 103 deletions(-) create mode 100644 src/kernels/static_composable_kernel/include/utility/static_kernel_ck_utils_type.hpp diff --git a/Jenkinsfile b/Jenkinsfile index b36f1d2308..6d2b69285a 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -1,5 +1,4 @@ def rocmnode(name) { -<<<<<<< HEAD return 'rocmtest && miopen && ' + name } @@ -277,42 +276,12 @@ pipeline { Tensile_build_env = "MIOPEN_DEBUG_HIP_KERNELS=0 " Tensile_setup = " -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF" } - parameters { - booleanParam( - name: "STATIC_CHECKS", - defaultValue: true, - description: "") - booleanParam( - name: "SMOKE_TESTS", - defaultValue: true, - description: "") - booleanParam( - name: "FULL_TESTS", - defaultValue: true, - description: "") - booleanParam( - name: "MIOPENTENSILE", - defaultValue: false, - description: "") - booleanParam( - name: "PACKAGES", - defaultValue: true, - description: "") - } stages{ stage("Static checks"){ when { expression { params.STATIC_CHECKS && !params.DISABLE_ALL_STAGES } } parallel{ -<<<<<<< HEAD - stage('Clang Tidy') { - agent{ label rocmnode("nogpu") } -||||||| merged common ancestors - stage('Clang Tidy') { - agent{ label rocmnode("rocmtest") } -======= stage('Hip Tidy') { agent{ label rocmnode("nogpu") } ->>>>>>> release/rocm-rel-4.3 environment{ setup_cmd = "CXX='/opt/rocm/llvm/bin/clang++' cmake -DMIOPEN_BACKEND=HIP -DBUILD_DEV=On .. " build_cmd = "make -j\$(nproc) -k analyze" @@ -368,7 +337,6 @@ pipeline { parallel{ stage('Fp32 OpenCL Debug + Codecov') { agent{ label rocmnode("vega") } ->>>>>>> release/rocm-rel-4.3 steps{ buildHipClangJobAndReboot(compiler: 'g++', build_type: 'debug', config_targets: Smoke_targets, codecov: true) } @@ -425,7 +393,6 @@ pipeline { } stage('Fp32 HipNoGPU Debug') { agent{ label rocmnode("nogpu") } ->>>>>>> release/rocm-rel-4.3 environment{ HipNoGPU_flags = "-DMIOPEN_BACKEND=HIPNOGPU -DMIOPEN_INSTALL_CXX_HEADERS=On" build_cmd = "make -j\$(nproc)" @@ -522,16 +489,6 @@ pipeline { } } } -<<<<<<< HEAD - - // Run fp16, bfp16, and int8 quick tests - stage("Fast low precision"){ - when { expression { params.SMOKE_TESTS } } -||||||| merged common ancestors - - // Run fp16, bfp16, and int8 quick tests - stage("Fast low precision"){ -======= stage("Smoke Fp16/Bf16/Int8"){ when { expression { params.SMOKE_FP16_BF16_INT8 && !params.DISABLE_ALL_STAGES } } environment{ @@ -885,30 +842,14 @@ pipeline { stage("Packages"){ when { expression { params.PACKAGES && !params.DISABLE_ALL_STAGES } } parallel { -<<<<<<< HEAD - stage('OpenCL Release Package') { - agent{ label rocmnode("nogpu") } -||||||| merged common ancestors - stage('OpenCL Release Package') { - agent{ label rocmnode("rocmtest") } -======= stage('OpenCL Package') { agent{ label rocmnode("nogpu") } ->>>>>>> release/rocm-rel-4.3 steps{ buildHipClangJobAndReboot(compiler: 'g++', package_build: "true", gpu_arch: "gfx900;gfx906;gfx908;gfx90a") } } -<<<<<<< HEAD - stage("HIP Release Package /opt/rocm"){ - agent{ label rocmnode("nogpu") } -||||||| merged common ancestors - stage("HIP Release Package /opt/rocm"){ - agent{ label rocmnode("rocmtest") } -======= stage("HIP Package /opt/rocm"){ agent{ label rocmnode("nogpu") } ->>>>>>> release/rocm-rel-4.3 steps{ buildHipClangJobAndReboot( package_build: "true", prefixpath: '/opt/rocm', gpu_arch: "gfx900;gfx906;gfx908;gfx90a") } diff --git a/src/comgr.cpp b/src/comgr.cpp index 032ff1151b..ac567ded04 100644 --- a/src/comgr.cpp +++ b/src/comgr.cpp @@ -202,20 +202,6 @@ static void AddCompilerOptions(OptionList& list, const miopen::TargetProperties& #if ROCM_FEATURE_TARGETID_OFF // It seems like these options are used only in codegen. // However it seems ok to pass these to compiler. -<<<<<<< HEAD - if(target.Sramecc()) - list.push_back("-msram-ecc"); - else - list.push_back("-mno-sram-ecc"); -#else - std::ignore = target; -#endif -||||||| merged common ancestors - if(IsEnabledFeatureSramEcc(target.Name())) - list.push_back("-msram-ecc"); - else - list.push_back("-mno-sram-ecc"); -======= if(target.Sramecc()) { if(*target.Sramecc()) @@ -226,7 +212,6 @@ static void AddCompilerOptions(OptionList& list, const miopen::TargetProperties& #else std::ignore = target; #endif ->>>>>>> release/rocm-rel-4.3 list.push_back("-mllvm"); list.push_back("-amdgpu-internalize-symbols"); } @@ -323,18 +308,6 @@ static void RemoveLinkOptionsUnwanted(OptionList& list) /// \todo Get list of supported isa names from comgr and select. static std::string GetIsaName(const miopen::TargetProperties& target) { -<<<<<<< HEAD -#if ROCM_FEATURE_TARGETID_OFF - const char* const ecc_suffix = target.Sramecc() ? "+sram-ecc" : ""; - return {"amdgcn-amd-amdhsa--" + target.Name() + ecc_suffix}; -#else - const LcOptionTargetStrings lots(target); - return {"amdgcn-amd-amdhsa--" + lots.targetId}; -#endif -||||||| merged common ancestors - const char* const ecc_suffix = IsEnabledFeatureSramEcc(device) ? "+sram-ecc" : ""; - return {"amdgcn-amd-amdhsa--" + device + ecc_suffix}; -======= #if ROCM_FEATURE_TARGETID_OFF const char* const ecc_suffix = (target.Sramecc() && *target.Sramecc()) ? "+sram-ecc" : ""; return {"amdgcn-amd-amdhsa--" + target.Name() + ecc_suffix}; @@ -342,7 +315,6 @@ static std::string GetIsaName(const miopen::TargetProperties& target) const LcOptionTargetStrings lots(target); return {"amdgcn-amd-amdhsa--" + lots.targetId}; #endif ->>>>>>> release/rocm-rel-4.3 } } // namespace lc @@ -783,18 +755,10 @@ void BuildHip(const std::string& name, + " " + GetDebugCompilerOptionsInsert() // + " " + MIOPEN_STRINGIZE(HIP_COMPILER_FLAGS) + (" -DHIP_PACKAGE_VERSION_FLAT=") + std::to_string(HIP_PACKAGE_VERSION_FLAT); -<<<<<<< HEAD -#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 -||||||| merged common ancestors -======= #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 ->>>>>>> release/rocm-rel-4.3 auto optCompile = miopen::SplitSpaceSeparated(raw, compiler::lc::GetOptionsNoSplit()); compiler::lc::hip::RemoveCompilerOptionsUnwanted(optCompile); action.SetOptionList(optCompile); @@ -807,18 +771,10 @@ void BuildHip(const std::string& name, + " " + GetDebugCompilerOptionsInsert() // + " " + MIOPEN_STRINGIZE(HIP_COMPILER_FLAGS) + (" -DHIP_PACKAGE_VERSION_FLAT=") + std::to_string(HIP_PACKAGE_VERSION_FLAT); -<<<<<<< HEAD -#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 -||||||| merged common ancestors -======= #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 ->>>>>>> release/rocm-rel-4.3 #if COMGR_SUPPORTS_PCH if(compiler::lc::hip::IsPchEnabled()) { diff --git a/src/kernels/static_composable_kernel/include/utility/static_kernel_ck_utils_type.hpp b/src/kernels/static_composable_kernel/include/utility/static_kernel_ck_utils_type.hpp new file mode 100644 index 0000000000..6220c20bac --- /dev/null +++ b/src/kernels/static_composable_kernel/include/utility/static_kernel_ck_utils_type.hpp @@ -0,0 +1,43 @@ +#ifndef CK_UTILS_TYPE_HPP +#define CK_UTILS_TYPE_HPP + +#include "static_kernel_integral_constant.hpp" + +namespace ck { + +template +struct Sequence; + +template +struct is_same : public integral_constant +{ +}; + +template +struct is_same : public integral_constant +{ +}; + +template +struct is_static : integral_constant +{ +}; + +template +struct is_static> : integral_constant +{ +}; + +template +struct is_static> : integral_constant +{ +}; + +template +using remove_reference_t = typename std::remove_reference::type; + +template +using remove_cv_t = typename std::remove_cv::type; + +} // namespace ck +#endif